radv/llvm: use the ring_offsets shader arg

Besides being nicer, this also fixes load_sample_positions_amd with LLVM.

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>

src/amd/vulkan/radv_nir_lower_abi.c

@@ -34,7 +34,6 @@ typedef struct {
    const struct radv_shader_args *args;
    const struct radv_shader_info *info;
    const struct radv_pipeline_key *pl_key;
-   bool use_llvm;
    uint32_t address32_hi;
 } lower_abi_state;
 
@@ -83,22 +82,12 @@ lower_abi_instr(nir_builder *b, nir_instr *instr, void *state)
 
    switch (intrin->intrinsic) {
    case nir_intrinsic_load_ring_tess_factors_amd:
-      if (s->use_llvm) {
-         progress = false;
-         break;
-      }
-
       replacement = load_ring(b, RING_HS_TESS_FACTOR, s);
       break;
    case nir_intrinsic_load_ring_tess_factors_offset_amd:
       replacement = ac_nir_load_arg(b, &s->args->ac, s->args->ac.tcs_factor_offset);
       break;
    case nir_intrinsic_load_ring_tess_offchip_amd:
-      if (s->use_llvm) {
-         progress = false;
-         break;
-      }
-
       replacement = load_ring(b, RING_HS_TESS_OFFCHIP, s);
       break;
    case nir_intrinsic_load_ring_tess_offchip_offset_amd:
@@ -117,19 +106,9 @@ lower_abi_instr(nir_builder *b, nir_instr *instr, void *state)
       }
       break;
    case nir_intrinsic_load_ring_esgs_amd:
-      if (s->use_llvm) {
-         progress = false;
-         break;
-      }
-
       replacement = load_ring(b, stage == MESA_SHADER_GEOMETRY ? RING_ESGS_GS : RING_ESGS_VS, s);
       break;
    case nir_intrinsic_load_ring_gsvs_amd:
-      if (s->use_llvm) {
-         progress = false;
-         break;
-      }
-
       replacement = load_ring(b, RING_GSVS_VS, s);
       break;
    case nir_intrinsic_load_ring_es2gs_offset_amd:
@@ -137,11 +116,6 @@ lower_abi_instr(nir_builder *b, nir_instr *instr, void *state)
       break;
 
    case nir_intrinsic_load_ring_attr_amd:
-      if (s->use_llvm) {
-         progress = false;
-         break;
-      }
-
       replacement = load_ring(b, RING_PS_ATTR, s);
 
       nir_ssa_def *dword1 = nir_channel(b, replacement, 1);
@@ -446,14 +420,13 @@ lower_abi_instr(nir_builder *b, nir_instr *instr, void *state)
 void
 radv_nir_lower_abi(nir_shader *shader, enum amd_gfx_level gfx_level,
                    const struct radv_shader_info *info, const struct radv_shader_args *args,
-                   const struct radv_pipeline_key *pl_key, bool use_llvm, uint32_t address32_hi)
+                   const struct radv_pipeline_key *pl_key, uint32_t address32_hi)
 {
    lower_abi_state state = {
       .gfx_level = gfx_level,
       .info = info,
       .args = args,
       .pl_key = pl_key,
-      .use_llvm = use_llvm,
       .address32_hi = address32_hi,
    };
 

src/amd/vulkan/radv_nir_to_llvm.c

@@ -56,17 +56,12 @@ struct radv_shader_context {
 
    LLVMValueRef descriptor_sets[MAX_SETS];
 
-   LLVMValueRef ring_offsets;
-
    LLVMValueRef vs_rel_patch_id;
 
    LLVMValueRef gs_wave_id;
 
    LLVMValueRef esgs_ring;
    LLVMValueRef gsvs_ring[4];
-   LLVMValueRef hs_ring_tess_offchip;
-   LLVMValueRef hs_ring_tess_factor;
-   LLVMValueRef attr_ring;
 
    uint64_t output_mask;
 };
@@ -174,12 +169,6 @@ create_function(struct radv_shader_context *ctx, gl_shader_stage stage, bool has
                            get_llvm_calling_convention(ctx->main_function.value, stage),
                            ctx->max_workgroup_size, ctx->options);
 
-   ctx->ring_offsets = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.implicit.buffer.ptr",
-                                          LLVMPointerType(ctx->ac.i8, AC_ADDR_SPACE_CONST), NULL, 0,
-                                          AC_FUNC_ATTR_READNONE);
-   ctx->ring_offsets = LLVMBuildBitCast(ctx->ac.builder, ctx->ring_offsets,
-                                        ac_array_in_const_addr_space(ctx->ac.v4i32), "");
-
    load_descriptor_sets(ctx);
 
    if (stage == MESA_SHADER_TESS_CTRL ||
@@ -1199,16 +1188,6 @@ ac_setup_rings(struct radv_shader_context *ctx)
 {
    struct ac_llvm_pointer ring_offsets = { .t = ctx->ac.v4i32, .v = ctx->ring_offsets };
 
-   if (ctx->options->gfx_level <= GFX8 &&
-       (ctx->stage == MESA_SHADER_GEOMETRY ||
-        (ctx->stage == MESA_SHADER_VERTEX && ctx->shader_info->vs.as_es) ||
-        (ctx->stage == MESA_SHADER_TESS_EVAL && ctx->shader_info->tes.as_es))) {
-      unsigned ring = ctx->stage == MESA_SHADER_GEOMETRY ? RING_ESGS_GS : RING_ESGS_VS;
-      LLVMValueRef offset = LLVMConstInt(ctx->ac.i32, ring, false);
-
-      ctx->esgs_ring = ac_build_load_to_sgpr(&ctx->ac, ring_offsets, offset);
-   }
-
    if (ctx->args->is_gs_copy_shader) {
       ctx->gsvs_ring[0] = ac_build_load_to_sgpr(&ctx->ac, ring_offsets, LLVMConstInt(ctx->ac.i32, RING_GSVS_VS, false));
    }
@@ -1265,26 +1244,6 @@ ac_setup_rings(struct radv_shader_context *ctx)
          ctx->gsvs_ring[stream] = ring;
       }
    }
-
-   if (ctx->stage == MESA_SHADER_TESS_CTRL || ctx->stage == MESA_SHADER_TESS_EVAL) {
-      ctx->hs_ring_tess_offchip = ac_build_load_to_sgpr(
-         &ctx->ac, ring_offsets, LLVMConstInt(ctx->ac.i32, RING_HS_TESS_OFFCHIP, false));
-      ctx->hs_ring_tess_factor = ac_build_load_to_sgpr(
-         &ctx->ac, ring_offsets, LLVMConstInt(ctx->ac.i32, RING_HS_TESS_FACTOR, false));
-   }
-
-   if (ctx->options->gfx_level >= GFX11 &&
-       ((ctx->stage == MESA_SHADER_VERTEX && !ctx->shader_info->vs.as_es && !ctx->shader_info->vs.as_ls) ||
-        (ctx->stage == MESA_SHADER_TESS_EVAL && !ctx->shader_info->tes.as_es) ||
-        (ctx->stage == MESA_SHADER_GEOMETRY))) {
-      ctx->attr_ring = ac_build_load_to_sgpr(&ctx->ac, ring_offsets,
-                                             LLVMConstInt(ctx->ac.i32, RING_PS_ATTR, false));
-
-      LLVMValueRef tmp = LLVMBuildExtractElement(ctx->ac.builder, ctx->attr_ring, ctx->ac.i32_1, "");
-      uint32_t stride = S_008F04_STRIDE(16 * ctx->shader_info->outinfo.param_exports);
-      tmp = LLVMBuildOr(ctx->ac.builder, tmp, LLVMConstInt(ctx->ac.i32, stride, false), "");
-      ctx->attr_ring = LLVMBuildInsertElement(ctx->ac.builder, ctx->attr_ring, tmp, ctx->ac.i32_1, "");
-   }
 }
 
 /* Fixup the HW not emitting the TCS regs if there are no HS threads. */
@@ -1324,35 +1283,20 @@ prepare_gs_input_vgprs(struct radv_shader_context *ctx, bool merged)
 static void
 declare_esgs_ring(struct radv_shader_context *ctx)
 {
-   if (ctx->esgs_ring)
-      return;
-
    assert(!LLVMGetNamedGlobal(ctx->ac.module, "esgs_ring"));
 
-   ctx->esgs_ring = LLVMAddGlobalInAddressSpace(ctx->ac.module, LLVMArrayType(ctx->ac.i32, 0),
-                                                "esgs_ring", AC_ADDR_SPACE_LDS);
-   LLVMSetLinkage(ctx->esgs_ring, LLVMExternalLinkage);
-   LLVMSetAlignment(ctx->esgs_ring, 64 * 1024);
+   LLVMValueRef esgs_ring = LLVMAddGlobalInAddressSpace(ctx->ac.module, LLVMArrayType(ctx->ac.i32, 0),
+                                                        "esgs_ring", AC_ADDR_SPACE_LDS);
+   LLVMSetLinkage(esgs_ring, LLVMExternalLinkage);
+   LLVMSetAlignment(esgs_ring, 64 * 1024);
 }
 
 static LLVMValueRef radv_intrinsic_load(struct ac_shader_abi *abi, nir_intrinsic_op op)
 {
-   struct radv_shader_context *ctx = radv_shader_context_from_abi(abi);
-
    switch (op) {
    case nir_intrinsic_load_base_vertex:
    case nir_intrinsic_load_first_vertex:
       return radv_load_base_vertex(abi, op == nir_intrinsic_load_base_vertex);
-   case nir_intrinsic_load_ring_tess_factors_amd:
-      return ctx->hs_ring_tess_factor;
-   case nir_intrinsic_load_ring_tess_offchip_amd:
-      return ctx->hs_ring_tess_offchip;
-   case nir_intrinsic_load_ring_esgs_amd:
-      return ctx->esgs_ring;
-   case nir_intrinsic_load_ring_attr_amd:
-      return ctx->attr_ring;
-   case nir_intrinsic_load_ring_gsvs_amd:
-      return ctx->gsvs_ring[0];
    default:
       return NULL;
    }

src/amd/vulkan/radv_pipeline.c

@@ -3896,7 +3896,6 @@ radv_postprocess_nir(struct radv_pipeline *pipeline,
 
    NIR_PASS(_, stage->nir, ac_nir_lower_global_access);
    NIR_PASS_V(stage->nir, radv_nir_lower_abi, gfx_level, &stage->info, &stage->args, pipeline_key,
-              radv_use_llvm_for_stage(device, stage->stage),
               device->physical_device->rad_info.address32_hi);
    radv_optimize_nir_algebraic(
       stage->nir, io_to_mem || lowered_ngg || stage->stage == MESA_SHADER_COMPUTE ||

src/amd/vulkan/radv_shader.h

@@ -547,8 +547,7 @@ nir_shader *radv_shader_spirv_to_nir(struct radv_device *device,
 
 void radv_nir_lower_abi(nir_shader *shader, enum amd_gfx_level gfx_level,
                         const struct radv_shader_info *info, const struct radv_shader_args *args,
-                        const struct radv_pipeline_key *pl_key, bool use_llvm,
-                        uint32_t address32_hi);
+                        const struct radv_pipeline_key *pl_key, uint32_t address32_hi);
 
 void radv_init_shader_arenas(struct radv_device *device);
 void radv_destroy_shader_arenas(struct radv_device *device);

src/amd/vulkan/radv_shader_args.c

@@ -578,9 +578,7 @@ radv_declare_shader_args(enum amd_gfx_level gfx_level, const struct radv_pipelin
    allocate_user_sgprs(gfx_level, info, args, stage, has_previous_stage, previous_stage,
                        needs_view_index, has_ngg_query, key, &user_sgpr_info);
 
-   if (args->explicit_scratch_args) {
-      ac_add_arg(&args->ac, AC_ARG_SGPR, 2, AC_ARG_CONST_DESC_PTR, &args->ac.ring_offsets);
-   }
+   ac_add_arg(&args->ac, AC_ARG_SGPR, 2, AC_ARG_CONST_DESC_PTR, &args->ac.ring_offsets);
    if (stage == MESA_SHADER_TASK) {
       ac_add_arg(&args->ac, AC_ARG_SGPR, 2, AC_ARG_CONST_DESC_PTR, &args->task_ring_offsets);
    }