diff --git a/src/microsoft/compiler/dxil_function.c b/src/microsoft/compiler/dxil_function.c
index a0121855df6..d8e0a331ced 100644
--- a/src/microsoft/compiler/dxil_function.c
+++ b/src/microsoft/compiler/dxil_function.c
@@ -90,6 +90,7 @@ static struct  predefined_func_descr predefined_funcs[] = {
 {"dx.op.loadOutputControlPoint", "O", "iiici", DXIL_ATTR_KIND_READ_NONE},
 {"dx.op.createHandleFromBinding", "@", "i#ib", DXIL_ATTR_KIND_READ_NONE},
 {"dx.op.annotateHandle", "@", "i@P", DXIL_ATTR_KIND_READ_NONE},
+{"dx.op.isHelperLane", "b", "i", DXIL_ATTR_KIND_READ_ONLY},
 };
 
 struct func_descr {
diff --git a/src/microsoft/compiler/nir_to_dxil.c b/src/microsoft/compiler/nir_to_dxil.c
index 53e917841d5..bfd77e108d9 100644
--- a/src/microsoft/compiler/nir_to_dxil.c
+++ b/src/microsoft/compiler/nir_to_dxil.c
@@ -330,6 +330,8 @@ enum dxil_intr {
 
    DXIL_INTR_ANNOTATE_HANDLE = 216,
    DXIL_INTR_CREATE_HANDLE_FROM_BINDING = 217,
+
+   DXIL_INTR_IS_HELPER_LANE = 221,
 };
 
 enum dxil_atomic_op {
@@ -4542,6 +4544,10 @@ emit_intrinsic(struct ntd_context *ctx, nir_intrinsic_instr *intr)
    case nir_intrinsic_load_sample_pos_from_id:
       return emit_load_sample_pos_from_id(ctx, intr);
 
+   case nir_intrinsic_load_helper_invocation:
+      return emit_load_unary_external_function(
+         ctx, intr, "dx.op.isHelperLane", DXIL_INTR_IS_HELPER_LANE);
+
    case nir_intrinsic_load_num_workgroups:
    case nir_intrinsic_load_workgroup_size:
    default:
diff --git a/src/microsoft/spirv_to_dxil/spirv_to_dxil.c b/src/microsoft/spirv_to_dxil/spirv_to_dxil.c
index 9876c10939e..994b8499f8d 100644
--- a/src/microsoft/spirv_to_dxil/spirv_to_dxil.c
+++ b/src/microsoft/spirv_to_dxil/spirv_to_dxil.c
@@ -119,10 +119,17 @@ spirv_to_dxil(const uint32_t *words, size_t word_count,
 
    glsl_type_singleton_init_or_ref();
 
+   struct nir_to_dxil_options opts = {
+      .environment = DXIL_ENVIRONMENT_VULKAN,
+      .shader_model_max = SHADER_MODEL_6_2,
+      .validator_version_max = DXIL_VALIDATOR_1_4,
+   };
+
    struct nir_shader_compiler_options nir_options = *dxil_get_nir_compiler_options();
    // We will manually handle base_vertex when vertex_id and instance_id have
    // have been already converted to zero-base.
    nir_options.lower_base_vertex = !conf->zero_based_vertex_instance_id;
+   nir_options.lower_helper_invocation = opts.shader_model_max < SHADER_MODEL_6_6;
 
    nir_shader *nir = spirv_to_nir(
       words, word_count, (struct nir_spirv_specialization *)specializations,
@@ -144,12 +151,6 @@ spirv_to_dxil(const uint32_t *words, size_t word_count,
    if (dgb_opts->dump_nir)
       nir_print_shader(nir, stderr);
 
-   struct nir_to_dxil_options opts = {
-      .environment = DXIL_ENVIRONMENT_VULKAN,
-      .shader_model_max = SHADER_MODEL_6_2,
-      .validator_version_max = DXIL_VALIDATOR_1_4,
-   };
-
    struct dxil_logger logger_inner = {.priv = logger->priv,
                                       .log = logger->log};