spirv,nir: add support for BaryCoord{NoPersp}KHR builtins

This introduces new intrinsics nir_intrinsic_load_barycentric_coord_xxx
with 3-components instead of expanding the existing ones that are
supposed to interpolate input varyings, while BaryCoord is a sysval
on most hardware.

Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/23254>

src/compiler/nir/nir_divergence_analysis.c

@@ -478,6 +478,11 @@ visit_intrinsic(nir_shader *shader, nir_intrinsic_instr *instr)
    case nir_intrinsic_load_barycentric_model:
    case nir_intrinsic_load_barycentric_at_sample:
    case nir_intrinsic_load_barycentric_at_offset:
+   case nir_intrinsic_load_barycentric_coord_pixel:
+   case nir_intrinsic_load_barycentric_coord_centroid:
+   case nir_intrinsic_load_barycentric_coord_sample:
+   case nir_intrinsic_load_barycentric_coord_at_sample:
+   case nir_intrinsic_load_barycentric_coord_at_offset:
    case nir_intrinsic_interp_deref_at_offset:
    case nir_intrinsic_interp_deref_at_sample:
    case nir_intrinsic_interp_deref_at_centroid:

src/compiler/nir/nir_gather_info.c

@@ -700,6 +700,19 @@ gather_intrinsic_info(nir_intrinsic_instr *instr, nir_shader *shader,
          shader->info.fs.uses_sample_qualifier = true;
       break;
 
+   case nir_intrinsic_load_barycentric_coord_pixel:
+   case nir_intrinsic_load_barycentric_coord_centroid:
+   case nir_intrinsic_load_barycentric_coord_sample:
+   case nir_intrinsic_load_barycentric_coord_at_offset:
+   case nir_intrinsic_load_barycentric_coord_at_sample:
+      if (nir_intrinsic_interp_mode(instr) == INTERP_MODE_SMOOTH ||
+          nir_intrinsic_interp_mode(instr) == INTERP_MODE_NONE) {
+         BITSET_SET(shader->info.system_values_read, SYSTEM_VALUE_BARYCENTRIC_PERSP_COORD);
+      } else if (nir_intrinsic_interp_mode(instr) == INTERP_MODE_NOPERSPECTIVE) {
+         BITSET_SET(shader->info.system_values_read, SYSTEM_VALUE_BARYCENTRIC_LINEAR_COORD);
+      }
+      break;
+
    case nir_intrinsic_quad_broadcast:
    case nir_intrinsic_quad_swap_horizontal:
    case nir_intrinsic_quad_swap_vertical:

src/compiler/nir/nir_intrinsics.py

@@ -934,6 +934,10 @@ system_value("user_data_amd", 4)
 #
 # The vec2 value produced by these intrinsics is intended for use as the
 # barycoord source of a load_interpolated_input intrinsic.
+#
+# The vec3 variants are intended to be used for input barycentric coordinates
+# which are system values on most hardware, compared to the vec2 variants which
+# interpolates input varyings.
 
 def barycentric(name, dst_comp, src_comp=[]):
     intrinsic("load_barycentric_" + name, src_comp=src_comp, dest_comp=dst_comp,
@@ -941,13 +945,18 @@ def barycentric(name, dst_comp, src_comp=[]):
 
 # no sources.
 barycentric("pixel", 2)
+barycentric("coord_pixel", 3)
 barycentric("centroid", 2)
+barycentric("coord_centroid", 3)
 barycentric("sample", 2)
+barycentric("coord_sample", 3)
 barycentric("model", 3)
 # src[] = { sample_id }.
 barycentric("at_sample", 2, [1])
+barycentric("coord_at_sample", 3, [1])
 # src[] = { offset.xy }.
 barycentric("at_offset", 2, [2])
+barycentric("coord_at_offset", 3, [2])
 
 # Load sample position:
 #

src/compiler/nir/nir_lower_system_values.c

@@ -115,6 +115,39 @@ lower_system_value_instr(nir_builder *b, nir_instr *instr, void *_state)
    case nir_intrinsic_load_workgroup_size:
       return sanitize_32bit_sysval(b, intrin);
 
+   case nir_intrinsic_interp_deref_at_centroid:
+   case nir_intrinsic_interp_deref_at_sample:
+   case nir_intrinsic_interp_deref_at_offset: {
+      nir_deref_instr *deref = nir_src_as_deref(intrin->src[0]);
+      if (!nir_deref_mode_is(deref, nir_var_system_value))
+         return NULL;
+
+      nir_variable *var = deref->var;
+      enum glsl_interp_mode interp_mode;
+
+      if (var->data.location == SYSTEM_VALUE_BARYCENTRIC_PERSP_COORD) {
+         interp_mode = INTERP_MODE_SMOOTH;
+      } else {
+         assert(var->data.location == SYSTEM_VALUE_BARYCENTRIC_LINEAR_COORD);
+         interp_mode = INTERP_MODE_NOPERSPECTIVE;
+      }
+
+      switch (intrin->intrinsic) {
+      case nir_intrinsic_interp_deref_at_centroid:
+         return nir_load_barycentric_coord_centroid(b, 32, .interp_mode = interp_mode);
+      case nir_intrinsic_interp_deref_at_sample:
+         assert(intrin->src[1].is_ssa);
+         return nir_load_barycentric_coord_at_sample(b, 32, intrin->src[1].ssa,
+                                                     .interp_mode = interp_mode);
+      case nir_intrinsic_interp_deref_at_offset:
+         assert(intrin->src[1].is_ssa);
+         return nir_load_barycentric_coord_at_offset(b, 32, intrin->src[1].ssa,
+                                                     .interp_mode = interp_mode);
+      default:
+         unreachable("Bogus interpolateAt() intrinsic.");
+      }
+   }
+
    case nir_intrinsic_load_deref: {
       nir_deref_instr *deref = nir_src_as_deref(intrin->src[0]);
       if (!nir_deref_mode_is(deref, nir_var_system_value))
@@ -215,6 +248,26 @@ lower_system_value_instr(nir_builder *b, nir_instr *instr, void *_state)
          return nir_load_barycentric(b, nir_intrinsic_load_barycentric_model,
                                      INTERP_MODE_NONE);
 
+      case SYSTEM_VALUE_BARYCENTRIC_LINEAR_COORD:
+      case SYSTEM_VALUE_BARYCENTRIC_PERSP_COORD: {
+         enum glsl_interp_mode interp_mode;
+
+         if (var->data.location == SYSTEM_VALUE_BARYCENTRIC_PERSP_COORD) {
+            interp_mode = INTERP_MODE_SMOOTH;
+         } else {
+            assert(var->data.location == SYSTEM_VALUE_BARYCENTRIC_LINEAR_COORD);
+            interp_mode = INTERP_MODE_NOPERSPECTIVE;
+         }
+
+         if (var->data.sample) {
+            return nir_load_barycentric_coord_sample(b, 32, .interp_mode = interp_mode);
+         } else if (var->data.centroid) {
+            return nir_load_barycentric_coord_centroid(b, 32, .interp_mode = interp_mode);
+         } else {
+            return nir_load_barycentric_coord_pixel(b, 32, .interp_mode = interp_mode);
+         }
+      }
+
       case SYSTEM_VALUE_HELPER_INVOCATION: {
          /* When demote operation is used, reading the HelperInvocation
           * needs to use Volatile memory access semantics to provide the

src/compiler/shader_enums.h

@@ -825,6 +825,14 @@ typedef enum
    SYSTEM_VALUE_BARYCENTRIC_LINEAR_SAMPLE,
    SYSTEM_VALUE_BARYCENTRIC_PULL_MODEL,
 
+   /**
+    * \name VK_KHR_fragment_shader_barycentric
+    */
+   /*@{*/
+   SYSTEM_VALUE_BARYCENTRIC_PERSP_COORD,
+   SYSTEM_VALUE_BARYCENTRIC_LINEAR_COORD,
+   /*@}*/
+
    /**
     * \name Ray tracing shader system values
     */

src/compiler/spirv/vtn_variables.c

@@ -1186,6 +1186,14 @@ vtn_get_builtin_location(struct vtn_builder *b,
       *location = SYSTEM_VALUE_RAY_TRIANGLE_VERTEX_POSITIONS;
       set_mode_system_value(b, mode);
       break;
+   case SpvBuiltInBaryCoordKHR:
+      *location = SYSTEM_VALUE_BARYCENTRIC_PERSP_COORD;
+      set_mode_system_value(b, mode);
+      break;
+   case SpvBuiltInBaryCoordNoPerspKHR:
+      *location = SYSTEM_VALUE_BARYCENTRIC_LINEAR_COORD;
+      set_mode_system_value(b, mode);
+      break;
 
    default:
       vtn_fail("Unsupported builtin: %s (%u)",