From 607720151f8bf9e8fcc483c7e007ad7b39a6fc58 Mon Sep 17 00:00:00 2001
From: Jesse Natalie <jenatali@microsoft.com>
Date: Mon, 25 Mar 2024 07:36:00 -0700
Subject: [PATCH] microsoft/compiler: Handle base vertex/instance sysvals as
 DXIL intrinsics

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/28383>
---
 src/microsoft/compiler/dxil_function.c |  2 ++
 src/microsoft/compiler/nir_to_dxil.c   | 12 ++++++++++++
 2 files changed, 14 insertions(+)

diff --git a/src/microsoft/compiler/dxil_function.c b/src/microsoft/compiler/dxil_function.c
index c68e1683078..483e85be61f 100644
--- a/src/microsoft/compiler/dxil_function.c
+++ b/src/microsoft/compiler/dxil_function.c
@@ -113,6 +113,8 @@ static struct  predefined_func_descr predefined_funcs[] = {
 {"dx.op.quadReadLaneAt", "O", "iOi", DXIL_ATTR_KIND_NO_UNWIND},
 {"dx.op.quadOp", "O", "iOc", DXIL_ATTR_KIND_NO_UNWIND},
 {"dx.op.dot4AddPacked", "i", "iiii", DXIL_ATTR_KIND_READ_NONE},
+{"dx.op.startVertexLocation", "i", "i", DXIL_ATTR_KIND_READ_NONE},
+{"dx.op.startInstanceLocation", "i", "i", DXIL_ATTR_KIND_READ_NONE},
 };
 
 struct func_descr {
diff --git a/src/microsoft/compiler/nir_to_dxil.c b/src/microsoft/compiler/nir_to_dxil.c
index 3b983bf6476..d4b1845f509 100644
--- a/src/microsoft/compiler/nir_to_dxil.c
+++ b/src/microsoft/compiler/nir_to_dxil.c
@@ -390,6 +390,9 @@ enum dxil_intr {
    DXIL_INTR_SAMPLE_CMP_LEVEL = 224,
    DXIL_INTR_SAMPLE_CMP_GRAD = 254,
    DXIL_INTR_SAMPLE_CMP_BIAS = 255,
+
+   DXIL_INTR_START_VERTEX_LOCATION = 256,
+   DXIL_INTR_START_INSTANCE_LOCATION = 257,
 };
 
 enum dxil_atomic_op {
@@ -4939,6 +4942,15 @@ emit_intrinsic(struct ntd_context *ctx, nir_intrinsic_instr *intr)
    case nir_intrinsic_exclusive_scan:
       return emit_reduce(ctx, intr);
 
+   case nir_intrinsic_load_first_vertex:
+      ctx->mod.feats.extended_command_info = true;
+      return emit_load_unary_external_function(ctx, intr, "dx.op.startVertexLocation",
+                                               DXIL_INTR_START_VERTEX_LOCATION, nir_type_int);
+   case nir_intrinsic_load_base_instance:
+      ctx->mod.feats.extended_command_info = true;
+      return emit_load_unary_external_function(ctx, intr, "dx.op.startInstanceLocation",
+                                               DXIL_INTR_START_INSTANCE_LOCATION, nir_type_int);
+
    case nir_intrinsic_load_num_workgroups:
    case nir_intrinsic_load_workgroup_size:
    default: