From c2ec23ab846bf049aad25fd0cd0696f9e18ea260 Mon Sep 17 00:00:00 2001 From: Samuel Pitoiset Date: Thu, 12 May 2022 15:50:04 +0200 Subject: [PATCH] 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 Reviewed-by: Rhys Perry Part-of: --- src/compiler/nir/nir_divergence_analysis.c | 5 ++ src/compiler/nir/nir_gather_info.c | 13 ++++++ src/compiler/nir/nir_intrinsics.py | 9 ++++ src/compiler/nir/nir_lower_system_values.c | 53 ++++++++++++++++++++++ src/compiler/shader_enums.h | 8 ++++ src/compiler/spirv/vtn_variables.c | 8 ++++ 6 files changed, 96 insertions(+) diff --git a/src/compiler/nir/nir_divergence_analysis.c b/src/compiler/nir/nir_divergence_analysis.c index 59886020458..1799ce80f1a 100644 --- a/src/compiler/nir/nir_divergence_analysis.c +++ b/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: diff --git a/src/compiler/nir/nir_gather_info.c b/src/compiler/nir/nir_gather_info.c index c89ea09df90..3efba5794bc 100644 --- a/src/compiler/nir/nir_gather_info.c +++ b/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: diff --git a/src/compiler/nir/nir_intrinsics.py b/src/compiler/nir/nir_intrinsics.py index 21ba72c211e..861fb5fba55 100644 --- a/src/compiler/nir/nir_intrinsics.py +++ b/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: # diff --git a/src/compiler/nir/nir_lower_system_values.c b/src/compiler/nir/nir_lower_system_values.c index 3165b0c1cf4..d91366c82d0 100644 --- a/src/compiler/nir/nir_lower_system_values.c +++ b/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 diff --git a/src/compiler/shader_enums.h b/src/compiler/shader_enums.h index e7ebd120ef2..17dd94bac3c 100644 --- a/src/compiler/shader_enums.h +++ b/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 */ diff --git a/src/compiler/spirv/vtn_variables.c b/src/compiler/spirv/vtn_variables.c index 55bc4b66a2d..5ffd4bd1b36 100644 --- a/src/compiler/spirv/vtn_variables.c +++ b/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)",