diff --git a/src/compiler/spirv/spirv_to_nir.c b/src/compiler/spirv/spirv_to_nir.c
index 5cb7691506e..9f2e91b5a3b 100644
--- a/src/compiler/spirv/spirv_to_nir.c
+++ b/src/compiler/spirv/spirv_to_nir.c
@@ -3086,6 +3086,14 @@ vtn_handle_texture(struct vtn_builder *b, SpvOp opcode,
       is_shadow && glsl_get_components(ret_type->type) == 1;
    instr->component = gather_component;
 
+   /* If SpvCapabilityImageGatherBiasLodAMD is enabled, texture gather without an explicit LOD
+    * has an implicit one (instead of using level 0).
+    */
+   if (texop == nir_texop_tg4 && b->image_gather_bias_lod &&
+       !(operands & SpvImageOperandsLodMask)) {
+      instr->is_gather_implicit_lod = true;
+   }
+
    /* The Vulkan spec says:
     *
     *    "If an instruction loads from or stores to a resource (including
@@ -4797,6 +4805,7 @@ vtn_handle_preamble_instruction(struct vtn_builder *b, SpvOp opcode,
 
       case SpvCapabilityImageGatherBiasLodAMD:
          spv_check_supported(amd_image_gather_bias_lod, cap);
+         b->image_gather_bias_lod = true;
          break;
 
       case SpvCapabilityAtomicFloat16AddEXT:
diff --git a/src/compiler/spirv/vtn_private.h b/src/compiler/spirv/vtn_private.h
index bd65a60d9bd..72d81fd0b89 100644
--- a/src/compiler/spirv/vtn_private.h
+++ b/src/compiler/spirv/vtn_private.h
@@ -738,6 +738,7 @@ struct vtn_builder {
    struct vtn_value *entry_point;
    struct vtn_value *workgroup_size_builtin;
    bool variable_pointers;
+   bool image_gather_bias_lod;
 
    uint32_t *interface_ids;
    size_t interface_ids_count;