diff --git a/source/blender/draw/engines/eevee_next/shaders/eevee_light_culling_sort_comp.glsl b/source/blender/draw/engines/eevee_next/shaders/eevee_light_culling_sort_comp.glsl
index e98b170cd4c..27f6a63cdd3 100644
--- a/source/blender/draw/engines/eevee_next/shaders/eevee_light_culling_sort_comp.glsl
+++ b/source/blender/draw/engines/eevee_next/shaders/eevee_light_culling_sort_comp.glsl
@@ -11,6 +11,11 @@ shared float zdists_cache[gl_WorkGroupSize.x];
 
 void main()
 {
+  /* Early exit if no lights are present to prevent out of bounds buffer read. */
+  if (light_cull_buf.visible_count == 0) {
+    return;
+  }
+
   uint src_index = gl_GlobalInvocationID.x;
   bool valid_thread = true;
 
diff --git a/source/blender/draw/engines/eevee_next/shaders/eevee_shadow_page_mask_comp.glsl b/source/blender/draw/engines/eevee_next/shaders/eevee_shadow_page_mask_comp.glsl
index 14d91f451b7..03b7509caca 100644
--- a/source/blender/draw/engines/eevee_next/shaders/eevee_shadow_page_mask_comp.glsl
+++ b/source/blender/draw/engines/eevee_next/shaders/eevee_shadow_page_mask_comp.glsl
@@ -40,7 +40,7 @@ void main()
           tile.is_used = true;
         }
         /* Reset count for next level. */
-        usage_grid[tile_co.y][tile_co.x] = 0u;
+        usage_grid[tile_co.y / 2][tile_co.x / 2] = 0u;
       }
 
       barrier();
diff --git a/source/blender/draw/engines/eevee_next/shaders/eevee_shadow_tilemap_init_comp.glsl b/source/blender/draw/engines/eevee_next/shaders/eevee_shadow_tilemap_init_comp.glsl
index 318efcd2084..e6cda65818f 100644
--- a/source/blender/draw/engines/eevee_next/shaders/eevee_shadow_tilemap_init_comp.glsl
+++ b/source/blender/draw/engines/eevee_next/shaders/eevee_shadow_tilemap_init_comp.glsl
@@ -57,7 +57,8 @@ void main()
 
   ivec2 tile_co = ivec2(gl_GlobalInvocationID.xy);
   ivec2 tile_shifted = tile_co + tilemap.grid_shift;
-  ivec2 tile_wrapped = ivec2(tile_shifted % SHADOW_TILEMAP_RES);
+  /* Ensure value is shifted into positive range to avoid modulo on negative. */
+  ivec2 tile_wrapped = ivec2((ivec2(SHADOW_TILEMAP_RES) + tile_shifted) % SHADOW_TILEMAP_RES);
 
   /* If this tile was shifted in and contains old information, update it.
    * Note that cubemap always shift all tiles on update. */
@@ -72,7 +73,7 @@ void main()
   uint lod_size = uint(SHADOW_TILEMAP_RES);
   for (int lod = 0; lod <= lod_max; lod++, lod_size >>= 1u) {
     bool thread_active = all(lessThan(tile_co, ivec2(lod_size)));
-    ShadowTileDataPacked tile;
+    ShadowTileDataPacked tile = 0;
     int tile_load = shadow_tile_offset(tile_wrapped, tilemap.tiles_index, lod);
     if (thread_active) {
       tile = init_tile_data(tiles_buf[tile_load], do_update);
diff --git a/source/blender/gpu/shaders/metal/mtl_shader_defines.msl b/source/blender/gpu/shaders/metal/mtl_shader_defines.msl
index 355ad423e5e..0329de7f561 100644
--- a/source/blender/gpu/shaders/metal/mtl_shader_defines.msl
+++ b/source/blender/gpu/shaders/metal/mtl_shader_defines.msl
@@ -66,7 +66,7 @@ using uvec4 = uint4;
 
 /* Compute decorators. */
 #define TG threadgroup
-#define barrier() threadgroup_barrier(mem_flags::mem_threadgroup)
+#define barrier() threadgroup_barrier(mem_flags::mem_threadgroup | mem_flags::mem_device | mem_flags::mem_texture)
 
 #ifdef MTL_USE_WORKGROUP_SIZE
 /* Compute workgroup size. */