microsoft/compiler: Handle load_sample_pos_at_id

Reviewed-by: Sil Vilerino <sivileri@microsoft.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/14624>

src/microsoft/compiler/dxil_function.c

@@ -74,6 +74,8 @@ static struct  predefined_func_descr predefined_funcs[] = {
 {"dx.op.legacyF32ToF16", "i", "if", DXIL_ATTR_KIND_READ_ONLY},
 {"dx.op.makeDouble", "g", "iii", DXIL_ATTR_KIND_READ_NONE},
 {"dx.op.splitDouble", "G", "ig", DXIL_ATTR_KIND_READ_NONE},
+{"dx.op.texture2DMSGetSamplePosition", "S", "i@i", DXIL_ATTR_KIND_READ_ONLY},
+{"dx.op.renderTargetGetSamplePosition", "S", "ii", DXIL_ATTR_KIND_READ_ONLY},
 };
 
 struct func_descr {
@@ -175,6 +177,7 @@ get_type_from_string(struct dxil_module *mod, const char *param_descr,
    case DXIL_FUNC_PARAM_FROM_OVERLOAD:  return dxil_get_overload_type(mod, overload);
    case DXIL_FUNC_PARAM_RESRET: return dxil_module_get_resret_type(mod, overload);
    case DXIL_FUNC_PARAM_DIM: return dxil_module_get_dimret_type(mod);
+   case DXIL_FUNC_PARAM_SAMPLE_POS: return dxil_module_get_samplepos_type(mod);
    case DXIL_FUNC_PARAM_CBUF_RET: return dxil_module_get_cbuf_ret_type(mod, overload);
    case DXIL_FUNC_PARAM_SPLIT_DOUBLE: return dxil_module_get_split_double_ret_type(mod);
    case DXIL_FUNC_PARAM_POINTER: {

src/microsoft/compiler/dxil_function.h

@@ -41,6 +41,7 @@
 #define DXIL_FUNC_PARAM_CBUF_RET 'B'
 #define DXIL_FUNC_PARAM_DIM 'D'
 #define DXIL_FUNC_PARAM_SPLIT_DOUBLE 'G'
+#define DXIL_FUNC_PARAM_SAMPLE_POS 'S'
 
 #include "dxil_module.h"
 #include "util/rb_tree.h"

src/microsoft/compiler/dxil_module.c

@@ -878,6 +878,17 @@ dxil_module_get_dimret_type(struct dxil_module *m)
    return dxil_module_get_struct_type(m, "dx.types.Dimensions", dimret, 4);
 }
 
+const struct dxil_type *
+dxil_module_get_samplepos_type(struct dxil_module *m)
+{
+   const struct dxil_type *float_type = dxil_module_get_float_type(m, 32);
+
+   const struct dxil_type *samplepos[] =
+      { float_type, float_type };
+
+   return dxil_module_get_struct_type(m, "dx.types.SamplePos", samplepos, 2);
+}
+
 const struct dxil_type *
 dxil_module_add_function_type(struct dxil_module *m,
                               const struct dxil_type *ret_type,

src/microsoft/compiler/dxil_module.h

@@ -277,6 +277,9 @@ dxil_module_get_resret_type(struct dxil_module *m, enum overload_type overload);
 const struct dxil_type *
 dxil_module_get_dimret_type(struct dxil_module *m);
 
+const struct dxil_type *
+dxil_module_get_samplepos_type(struct dxil_module *m);
+
 const struct dxil_type *
 dxil_module_get_struct_type(struct dxil_module *m,
                             const char *name,

src/microsoft/compiler/nir_to_dxil.c

@@ -242,6 +242,10 @@ enum dxil_intr {
    DXIL_INTR_TEXTURE_SIZE = 72,
    DXIL_INTR_TEXTURE_GATHER = 73,
 
+   DXIL_INTR_TEXTURE2DMS_GET_SAMPLE_POSITION = 75,
+   DXIL_INTR_RENDER_TARGET_GET_SAMPLE_POSITION = 76,
+   DXIL_INTR_RENDER_TARGET_GET_SAMPLE_COUNT = 77,
+
    DXIL_INTR_ATOMIC_BINOP = 78,
    DXIL_INTR_ATOMIC_CMPXCHG = 79,
    DXIL_INTR_BARRIER = 80,
@@ -3591,6 +3595,38 @@ emit_load_vulkan_descriptor(struct ntd_context *ctx, nir_intrinsic_instr *intr)
    return true;
 }
 
+static bool
+emit_load_sample_pos_from_id(struct ntd_context *ctx, nir_intrinsic_instr *intr)
+{
+   const struct dxil_func *func = dxil_get_function(&ctx->mod, "dx.op.renderTargetGetSamplePosition", DXIL_NONE);
+   if (!func)
+      return false;
+
+   const struct dxil_value *opcode = dxil_module_get_int32_const(&ctx->mod, DXIL_INTR_RENDER_TARGET_GET_SAMPLE_POSITION);
+   if (!opcode)
+      return false;
+
+   const struct dxil_value *args[] = {
+      opcode,
+      get_src(ctx, &intr->src[0], 0, nir_type_uint32),
+   };
+   if (!args[1])
+      return false;
+
+   const struct dxil_value *v = dxil_emit_call(&ctx->mod, func, args, ARRAY_SIZE(args));
+   if (!v)
+      return false;
+
+   for (unsigned i = 0; i < 2; ++i) {
+      /* GL coords go from 0 -> 1, D3D from -0.5 -> 0.5 */
+      const struct dxil_value *coord = dxil_emit_binop(&ctx->mod, DXIL_BINOP_ADD,
+         dxil_emit_extractval(&ctx->mod, v, i),
+         dxil_module_get_float_const(&ctx->mod, 0.5f), 0);
+      store_dest(ctx, &intr->dest, i, coord, nir_type_float32);
+   }
+   return true;
+}
+
 static bool
 emit_intrinsic(struct ntd_context *ctx, nir_intrinsic_instr *intr)
 {
@@ -3745,6 +3781,9 @@ emit_intrinsic(struct ntd_context *ctx, nir_intrinsic_instr *intr)
    case nir_intrinsic_load_vulkan_descriptor:
       return emit_load_vulkan_descriptor(ctx, intr);
 
+   case nir_intrinsic_load_sample_pos_from_id:
+      return emit_load_sample_pos_from_id(ctx, intr);
+
    case nir_intrinsic_load_num_workgroups:
    case nir_intrinsic_load_workgroup_size:
    default: