KonstantinSeurer
/
mesa
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

ac/nir, radv, radeonsi: Switch to using ac_shader_args 

Reviewed-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Acked-by: Marek Olšák <marek.olsak@amd.com>
This commit is contained in:
Connor Abbott 2019-11-11 12:50:12 +01:00
parent 9885af3bdf
commit 3b143369a5
13 changed files with 1546 additions and 1542 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

95
src/amd/llvm/ac_nir_to_llvm.c
View File

@ -38,6 +38,7 @@
struct ac_nir_context {
struct ac_llvm_context ac;
struct ac_shader_abi *abi;
const struct ac_shader_args *args;
gl_shader_stage stage;
shader_info *info;
@ -1435,16 +1436,22 @@ static LLVMValueRef visit_load_push_constant(struct ac_nir_context *ctx,
offset += LLVMConstIntGetZExtValue(src0);
offset /= 4;
offset -= ctx->abi->base_inline_push_consts;
offset -= ctx->args->base_inline_push_consts;
if (offset + count <= ctx->abi->num_inline_push_consts) {
unsigned num_inline_push_consts = ctx->args->num_inline_push_consts;
if (offset + count <= num_inline_push_consts) {
LLVMValueRef push_constants[num_inline_push_consts];
for (unsigned i = 0; i < num_inline_push_consts; i++)
push_constants[i] = ac_get_arg(&ctx->ac,
ctx->args->inline_push_consts[i]);
return ac_build_gather_values(&ctx->ac,
ctx->abi->inline_push_consts + offset,
push_constants + offset,
count);
}
}
ptr = LLVMBuildGEP(ctx->ac.builder, ctx->abi->push_constants, &addr, 1, "");
ptr = LLVMBuildGEP(ctx->ac.builder,
ac_get_arg(&ctx->ac, ctx->args->push_constants), &addr, 1, "");
if (instr->dest.ssa.bit_size == 8) {
unsigned load_dwords = instr->dest.ssa.num_components > 1 ? 2 : 1;
@ -2902,7 +2909,8 @@ visit_load_local_invocation_index(struct ac_nir_context *ctx)
{
LLVMValueRef result;
LLVMValueRef thread_id = ac_get_thread_id(&ctx->ac);
result = LLVMBuildAnd(ctx->ac.builder, ctx->abi->tg_size,
result = LLVMBuildAnd(ctx->ac.builder,
ac_get_arg(&ctx->ac, ctx->args->tg_size),
LLVMConstInt(ctx->ac.i32, 0xfc0, false), "");
if (ctx->ac.wave_size == 32)
@ -2917,7 +2925,8 @@ visit_load_subgroup_id(struct ac_nir_context *ctx)
{
if (ctx->stage == MESA_SHADER_COMPUTE) {
LLVMValueRef result;
result = LLVMBuildAnd(ctx->ac.builder, ctx->abi->tg_size,
result = LLVMBuildAnd(ctx->ac.builder,
ac_get_arg(&ctx->ac, ctx->args->tg_size),
LLVMConstInt(ctx->ac.i32, 0xfc0, false), "");
return LLVMBuildLShr(ctx->ac.builder, result, LLVMConstInt(ctx->ac.i32, 6, false), "");
} else {
@ -2929,7 +2938,8 @@ static LLVMValueRef
visit_load_num_subgroups(struct ac_nir_context *ctx)
{
if (ctx->stage == MESA_SHADER_COMPUTE) {
return LLVMBuildAnd(ctx->ac.builder, ctx->abi->tg_size,
return LLVMBuildAnd(ctx->ac.builder,
ac_get_arg(&ctx->ac, ctx->args->tg_size),
LLVMConstInt(ctx->ac.i32, 0x3f, false), "");
} else {
return LLVMConstInt(ctx->ac.i32, 1, false);
@ -3059,8 +3069,10 @@ static LLVMValueRef load_sample_pos(struct ac_nir_context *ctx)
LLVMValueRef values[2];
LLVMValueRef pos[2];
pos[0] = ac_to_float(&ctx->ac, ctx->abi->frag_pos[0]);
pos[1] = ac_to_float(&ctx->ac, ctx->abi->frag_pos[1]);
pos[0] = ac_to_float(&ctx->ac,
ac_get_arg(&ctx->ac, ctx->args->frag_pos[0]));
pos[1] = ac_to_float(&ctx->ac,
ac_get_arg(&ctx->ac, ctx->args->frag_pos[1]));
values[0] = ac_build_fract(&ctx->ac, pos[0], 32);
values[1] = ac_build_fract(&ctx->ac, pos[1], 32);
@ -3077,19 +3089,19 @@ static LLVMValueRef lookup_interp_param(struct ac_nir_context *ctx,
case INTERP_MODE_SMOOTH:
case INTERP_MODE_NONE:
if (location == INTERP_CENTER)
return ctx->abi->persp_center;
return ac_get_arg(&ctx->ac, ctx->args->persp_center);
else if (location == INTERP_CENTROID)
return ctx->abi->persp_centroid;
else if (location == INTERP_SAMPLE)
return ctx->abi->persp_sample;
return ac_get_arg(&ctx->ac, ctx->args->persp_sample);
break;
case INTERP_MODE_NOPERSPECTIVE:
if (location == INTERP_CENTER)
return ctx->abi->linear_center;
return ac_get_arg(&ctx->ac, ctx->args->linear_center);
else if (location == INTERP_CENTROID)
return ctx->abi->linear_centroid;
else if (location == INTERP_SAMPLE)
return ctx->abi->linear_sample;
return ac_get_arg(&ctx->ac, ctx->args->linear_sample);
break;
}
return NULL;
@ -3203,10 +3215,10 @@ static LLVMValueRef load_interpolated_input(struct ac_nir_context *ctx,
LLVMValueRef llvm_chan = LLVMConstInt(ctx->ac.i32, comp_start + comp, false);
if (bitsize == 16) {
values[comp] = ac_build_fs_interp_f16(&ctx->ac, llvm_chan, attr_number,
ctx->abi->prim_mask, i, j);
ac_get_arg(&ctx->ac, ctx->args->prim_mask), i, j);
} else {
values[comp] = ac_build_fs_interp(&ctx->ac, llvm_chan, attr_number,
ctx->abi->prim_mask, i, j);
ac_get_arg(&ctx->ac, ctx->args->prim_mask), i, j);
}
}
@ -3234,7 +3246,7 @@ static LLVMValueRef load_flat_input(struct ac_nir_context *ctx,
LLVMConstInt(ctx->ac.i32, 2, false),
llvm_chan,
attr_number,
ctx->abi->prim_mask);
ac_get_arg(&ctx->ac, ctx->args->prim_mask));
values[chan] = LLVMBuildBitCast(ctx->ac.builder, values[chan], ctx->ac.i32, "");
values[chan] = LLVMBuildTruncOrBitCast(ctx->ac.builder, values[chan],
bit_size == 16 ? ctx->ac.i16 : ctx->ac.i32, "");
@ -3274,8 +3286,8 @@ static void visit_intrinsic(struct ac_nir_context *ctx,
LLVMValueRef values[3];
for (int i = 0; i < 3; i++) {
values[i] = ctx->abi->workgroup_ids[i] ?
ctx->abi->workgroup_ids[i] : ctx->ac.i32_0;
values[i] = ctx->args->workgroup_ids[i].used ?
ac_get_arg(&ctx->ac, ctx->args->workgroup_ids[i]) : ctx->ac.i32_0;
}
result = ac_build_gather_values(&ctx->ac, values, 3);
@ -3289,51 +3301,56 @@ static void visit_intrinsic(struct ac_nir_context *ctx,
result = ctx->abi->load_local_group_size(ctx->abi);
break;
case nir_intrinsic_load_vertex_id:
result = LLVMBuildAdd(ctx->ac.builder, ctx->abi->vertex_id,
ctx->abi->base_vertex, "");
result = LLVMBuildAdd(ctx->ac.builder,
ac_get_arg(&ctx->ac, ctx->args->vertex_id),
ac_get_arg(&ctx->ac, ctx->args->base_vertex), "");
break;
case nir_intrinsic_load_vertex_id_zero_base: {
result = ctx->abi->vertex_id;
break;
}
case nir_intrinsic_load_local_invocation_id: {
result = ctx->abi->local_invocation_ids;
result = ac_get_arg(&ctx->ac, ctx->args->local_invocation_ids);
break;
}
case nir_intrinsic_load_base_instance:
result = ctx->abi->start_instance;
result = ac_get_arg(&ctx->ac, ctx->args->start_instance);
break;
case nir_intrinsic_load_draw_id:
result = ctx->abi->draw_id;
result = ac_get_arg(&ctx->ac, ctx->args->draw_id);
break;
case nir_intrinsic_load_view_index:
result = ctx->abi->view_index;
result = ac_get_arg(&ctx->ac, ctx->args->view_index);
break;
case nir_intrinsic_load_invocation_id:
if (ctx->stage == MESA_SHADER_TESS_CTRL) {
result = ac_unpack_param(&ctx->ac, ctx->abi->tcs_rel_ids, 8, 5);
result = ac_unpack_param(&ctx->ac,
ac_get_arg(&ctx->ac, ctx->args->tcs_rel_ids),
8, 5);
} else {
if (ctx->ac.chip_class >= GFX10) {
result = LLVMBuildAnd(ctx->ac.builder,
ctx->abi->gs_invocation_id,
ac_get_arg(&ctx->ac, ctx->args->gs_invocation_id),
LLVMConstInt(ctx->ac.i32, 127, 0), "");
} else {
result = ctx->abi->gs_invocation_id;
result = ac_get_arg(&ctx->ac, ctx->args->gs_invocation_id);
}
}
break;
case nir_intrinsic_load_primitive_id:
if (ctx->stage == MESA_SHADER_GEOMETRY) {
result = ctx->abi->gs_prim_id;
result = ac_get_arg(&ctx->ac, ctx->args->gs_prim_id);
} else if (ctx->stage == MESA_SHADER_TESS_CTRL) {
result = ctx->abi->tcs_patch_id;
result = ac_get_arg(&ctx->ac, ctx->args->tcs_patch_id);
} else if (ctx->stage == MESA_SHADER_TESS_EVAL) {
result = ctx->abi->tes_patch_id;
result = ac_get_arg(&ctx->ac, ctx->args->tes_patch_id);
} else
fprintf(stderr, "Unknown primitive id intrinsic: %d", ctx->stage);
break;
case nir_intrinsic_load_sample_id:
result = ac_unpack_param(&ctx->ac, ctx->abi->ancillary, 8, 4);
result = ac_unpack_param(&ctx->ac,
ac_get_arg(&ctx->ac, ctx->args->ancillary),
8, 4);
break;
case nir_intrinsic_load_sample_pos:
result = load_sample_pos(ctx);
@ -3343,10 +3360,11 @@ static void visit_intrinsic(struct ac_nir_context *ctx,
break;
case nir_intrinsic_load_frag_coord: {
LLVMValueRef values[4] = {
ctx->abi->frag_pos[0],
ctx->abi->frag_pos[1],
ctx->abi->frag_pos[2],
ac_build_fdiv(&ctx->ac, ctx->ac.f32_1, ctx->abi->frag_pos[3])
ac_get_arg(&ctx->ac, ctx->args->frag_pos[0]),
ac_get_arg(&ctx->ac, ctx->args->frag_pos[1]),
ac_get_arg(&ctx->ac, ctx->args->frag_pos[2]),
ac_build_fdiv(&ctx->ac, ctx->ac.f32_1,
ac_get_arg(&ctx->ac, ctx->args->frag_pos[3]))
};
result = ac_to_integer(&ctx->ac,
ac_build_gather_values(&ctx->ac, values, 4));
@ -3356,7 +3374,7 @@ static void visit_intrinsic(struct ac_nir_context *ctx,
result = ctx->abi->inputs[ac_llvm_reg_index_soa(VARYING_SLOT_LAYER, 0)];
break;
case nir_intrinsic_load_front_face:
result = ctx->abi->front_face;
result = ac_get_arg(&ctx->ac, ctx->args->front_face);
break;
case nir_intrinsic_load_helper_invocation:
result = ac_build_load_helper_invocation(&ctx->ac);
@ -3375,7 +3393,7 @@ static void visit_intrinsic(struct ac_nir_context *ctx,
result = ctx->abi->instance_id;
break;
case nir_intrinsic_load_num_work_groups:
result = ctx->abi->num_work_groups;
result = ac_get_arg(&ctx->ac, ctx->args->num_work_groups);
break;
case nir_intrinsic_load_local_invocation_index:
result = visit_load_local_invocation_index(ctx);
@ -4714,13 +4732,14 @@ setup_shared(struct ac_nir_context *ctx,
}
void ac_nir_translate(struct ac_llvm_context *ac, struct ac_shader_abi *abi,
struct nir_shader *nir)
const struct ac_shader_args *args, struct nir_shader *nir)
{
struct ac_nir_context ctx = {};
struct nir_function *func;
ctx.ac = *ac;
ctx.abi = abi;
ctx.args = args;
ctx.stage = nir->info.stage;
ctx.info = &nir->info;

3
src/amd/llvm/ac_nir_to_llvm.h
View File

@ -34,6 +34,7 @@ struct nir_shader;
struct nir_variable;
struct ac_llvm_context;
struct ac_shader_abi;
struct ac_shader_args;
/* Interpolation locations */
#define INTERP_CENTER 0
@ -50,7 +51,7 @@ void ac_lower_indirect_derefs(struct nir_shader *nir, enum chip_class);
bool ac_are_tessfactors_def_in_all_invocs(const struct nir_shader *nir);
void ac_nir_translate(struct ac_llvm_context *ac, struct ac_shader_abi *abi,
struct nir_shader *nir);
const struct ac_shader_args *args, struct nir_shader *nir);
void
ac_handle_shader_output_decl(struct ac_llvm_context *ctx,

44
src/amd/llvm/ac_shader_abi.h
View File

@ -25,6 +25,7 @@
#define AC_SHADER_ABI_H
#include <llvm-c/Core.h>
#include <assert.h>
#include "ac_shader_args.h"
#include "compiler/shader_enums.h"
@ -49,47 +50,14 @@ enum ac_descriptor_type {
* radv to share a compiler backend.
*/
struct ac_shader_abi {
LLVMValueRef base_vertex;
LLVMValueRef start_instance;
LLVMValueRef draw_id;
LLVMValueRef outputs[AC_LLVM_MAX_OUTPUTS * 4];
/* These input registers sometimes need to be fixed up. */
LLVMValueRef vertex_id;
LLVMValueRef instance_id;
LLVMValueRef tcs_patch_id;
LLVMValueRef tcs_rel_ids;
LLVMValueRef tes_patch_id;
LLVMValueRef gs_prim_id;
LLVMValueRef gs_invocation_id;
/* PS */
LLVMValueRef frag_pos[4];
LLVMValueRef front_face;
LLVMValueRef ancillary;
LLVMValueRef sample_coverage;
LLVMValueRef prim_mask;
LLVMValueRef color0;
LLVMValueRef color1;
LLVMValueRef persp_centroid, linear_centroid;
LLVMValueRef color0, color1;
LLVMValueRef user_data;
LLVMValueRef persp_sample;
LLVMValueRef persp_center;
LLVMValueRef persp_centroid;
LLVMValueRef linear_sample;
LLVMValueRef linear_center;
LLVMValueRef linear_centroid;
/* CS */
LLVMValueRef local_invocation_ids;
LLVMValueRef num_work_groups;
LLVMValueRef workgroup_ids[3];
LLVMValueRef tg_size;
/* Vulkan only */
LLVMValueRef push_constants;
LLVMValueRef inline_push_consts[AC_MAX_INLINE_PUSH_CONSTS];
unsigned num_inline_push_consts;
unsigned base_inline_push_consts;
LLVMValueRef view_index;
LLVMValueRef outputs[AC_LLVM_MAX_OUTPUTS * 4];
/* For VS and PS: pre-loaded shader inputs.
*

1
src/amd/vulkan/meson.build
View File

@ -101,6 +101,7 @@ libradv_files = files(
'radv_radeon_winsys.h',
'radv_shader.c',
'radv_shader.h',
'radv_shader_args.h',
'radv_shader_helper.h',
'radv_shader_info.c',
'radv_query.c',

1403
src/amd/vulkan/radv_nir_to_llvm.c
View File

File diff suppressed because it is too large Load Diff

76
src/amd/vulkan/radv_shader_args.h Normal file
View File

@ -0,0 +1,76 @@
/*
* Copyright © 2019 Valve Corporation.
*
* Permission is hereby granted, free of charge, to any person obtaining a
* copy of this software and associated documentation files (the "Software"),
* to deal in the Software without restriction, including without limitation
* the rights to use, copy, modify, merge, publish, distribute, sublicense,
* and/or sell copies of the Software, and to permit persons to whom the
* Software is furnished to do so, subject to the following conditions:
*
* The above copyright notice and this permission notice (including the next
* paragraph) shall be included in all copies or substantial portions of the
* Software.
*
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
* THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
* FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
* IN THE SOFTWARE.
*/
#include "ac_shader_args.h"
#include "radv_constants.h"
#include "util/list.h"
#include "amd_family.h"
struct radv_shader_args {
struct ac_shader_args ac;
struct radv_shader_info *shader_info;
const struct radv_nir_compiler_options *options;
struct ac_arg descriptor_sets[MAX_SETS];
struct ac_arg ring_offsets;
struct ac_arg vertex_buffers;
struct ac_arg rel_auto_id;
struct ac_arg vs_prim_id;
struct ac_arg es2gs_offset;
struct ac_arg oc_lds;
struct ac_arg merged_wave_info;
struct ac_arg tess_factor_offset;
struct ac_arg tes_rel_patch_id;
struct ac_arg tes_u;
struct ac_arg tes_v;
/* HW GS */
/* On gfx10:
* - bits 0..10: ordered_wave_id
* - bits 12..20: number of vertices in group
* - bits 22..30: number of primitives in group
*/
struct ac_arg gs_tg_info;
struct ac_arg gs2vs_offset;
struct ac_arg gs_wave_id;
struct ac_arg gs_vtx_offset[6];
/* Streamout */
struct ac_arg streamout_buffers;
struct ac_arg streamout_write_idx;
struct ac_arg streamout_config;
struct ac_arg streamout_offset[4];
bool is_gs_copy_shader;
};
static inline struct radv_shader_args *
radv_shader_args_from_ac(struct ac_shader_args *args)
{
struct radv_shader_args *radv_args = NULL;
return (struct radv_shader_args *) container_of(args, radv_args, ac);
}

47
src/gallium/drivers/radeonsi/gfx10_shader_ngg.c
View File

@ -31,12 +31,12 @@
static LLVMValueRef get_wave_id_in_tg(struct si_shader_context *ctx)
{
return si_unpack_param(ctx, ctx->param_merged_wave_info, 24, 4);
return si_unpack_param(ctx, ctx->merged_wave_info, 24, 4);
}
static LLVMValueRef get_tgsize(struct si_shader_context *ctx)
{
return si_unpack_param(ctx, ctx->param_merged_wave_info, 28, 4);
return si_unpack_param(ctx, ctx->merged_wave_info, 28, 4);
}
static LLVMValueRef get_thread_id_in_tg(struct si_shader_context *ctx)
@ -50,32 +50,22 @@ static LLVMValueRef get_thread_id_in_tg(struct si_shader_context *ctx)
static LLVMValueRef ngg_get_vtx_cnt(struct si_shader_context *ctx)
{
return ac_build_bfe(&ctx->ac, ctx->gs_tg_info,
LLVMConstInt(ctx->ac.i32, 12, false),
LLVMConstInt(ctx->ac.i32, 9, false),
false);
return si_unpack_param(ctx, ctx->gs_tg_info, 12, 9);
}
static LLVMValueRef ngg_get_prim_cnt(struct si_shader_context *ctx)
{
return ac_build_bfe(&ctx->ac, ctx->gs_tg_info,
LLVMConstInt(ctx->ac.i32, 22, false),
LLVMConstInt(ctx->ac.i32, 9, false),
false);
return si_unpack_param(ctx, ctx->gs_tg_info, 22, 9);
}
static LLVMValueRef ngg_get_ordered_id(struct si_shader_context *ctx)
{
return ac_build_bfe(&ctx->ac, ctx->gs_tg_info,
ctx->i32_0,
LLVMConstInt(ctx->ac.i32, 11, false),
false);
return si_unpack_param(ctx, ctx->gs_tg_info, 0, 11);
}
static LLVMValueRef ngg_get_query_buf(struct si_shader_context *ctx)
{
LLVMValueRef buf_ptr = LLVMGetParam(ctx->main_fn,
ctx->param_rw_buffers);
LLVMValueRef buf_ptr = ac_get_arg(&ctx->ac, ctx->rw_buffers);
return ac_build_load_to_sgpr(&ctx->ac, buf_ptr,
LLVMConstInt(ctx->i32, GFX10_GS_QUERY_BUF, false));
@ -212,7 +202,7 @@ static void build_streamout(struct si_shader_context *ctx,
struct tgsi_shader_info *info = &ctx->shader->selector->info;
struct pipe_stream_output_info *so = &ctx->shader->selector->so;
LLVMBuilderRef builder = ctx->ac.builder;
LLVMValueRef buf_ptr = LLVMGetParam(ctx->main_fn, ctx->param_rw_buffers);
LLVMValueRef buf_ptr = ac_get_arg(&ctx->ac, ctx->rw_buffers);
LLVMValueRef tid = get_thread_id_in_tg(ctx);
LLVMValueRef tmp, tmp2;
LLVMValueRef i32_2 = LLVMConstInt(ctx->i32, 2, false);
@ -583,16 +573,16 @@ void gfx10_emit_ngg_epilogue(struct ac_shader_abi *abi,
ac_build_endif(&ctx->ac, ctx->merged_wrap_if_label);
LLVMValueRef prims_in_wave = si_unpack_param(ctx, ctx->param_merged_wave_info, 8, 8);
LLVMValueRef vtx_in_wave = si_unpack_param(ctx, ctx->param_merged_wave_info, 0, 8);
LLVMValueRef prims_in_wave = si_unpack_param(ctx, ctx->merged_wave_info, 8, 8);
LLVMValueRef vtx_in_wave = si_unpack_param(ctx, ctx->merged_wave_info, 0, 8);
LLVMValueRef is_gs_thread = LLVMBuildICmp(builder, LLVMIntULT,
ac_get_thread_id(&ctx->ac), prims_in_wave, "");
LLVMValueRef is_es_thread = LLVMBuildICmp(builder, LLVMIntULT,
ac_get_thread_id(&ctx->ac), vtx_in_wave, "");
LLVMValueRef vtxindex[] = {
si_unpack_param(ctx, ctx->param_gs_vtx01_offset, 0, 16),
si_unpack_param(ctx, ctx->param_gs_vtx01_offset, 16, 16),
si_unpack_param(ctx, ctx->param_gs_vtx23_offset, 0, 16),
si_unpack_param(ctx, ctx->gs_vtx01_offset, 0, 16),
si_unpack_param(ctx, ctx->gs_vtx01_offset, 16, 16),
si_unpack_param(ctx, ctx->gs_vtx23_offset, 0, 16),
};
/* Determine the number of vertices per primitive. */
@ -606,7 +596,7 @@ void gfx10_emit_ngg_epilogue(struct ac_shader_abi *abi,
num_vertices_val = LLVMConstInt(ctx->i32, 3, 0);
} else {
/* Extract OUTPRIM field. */
tmp = si_unpack_param(ctx, ctx->param_vs_state_bits, 2, 2);
tmp = si_unpack_param(ctx, ctx->vs_state_bits, 2, 2);
num_vertices_val = LLVMBuildAdd(builder, tmp, ctx->i32_1, "");
num_vertices = 3; /* TODO: optimize for points & lines */
}
@ -673,14 +663,14 @@ void gfx10_emit_ngg_epilogue(struct ac_shader_abi *abi,
ac_build_ifcc(&ctx->ac, is_gs_thread, 5400);
/* Extract the PROVOKING_VTX_INDEX field. */
LLVMValueRef provoking_vtx_in_prim =
si_unpack_param(ctx, ctx->param_vs_state_bits, 4, 2);
si_unpack_param(ctx, ctx->vs_state_bits, 4, 2);
/* provoking_vtx_index = vtxindex[provoking_vtx_in_prim]; */
LLVMValueRef indices = ac_build_gather_values(&ctx->ac, vtxindex, 3);
LLVMValueRef provoking_vtx_index =
LLVMBuildExtractElement(builder, indices, provoking_vtx_in_prim, "");
LLVMBuildStore(builder, ctx->abi.gs_prim_id,
LLVMBuildStore(builder, ac_get_arg(&ctx->ac, ctx->args.gs_prim_id),
ac_build_gep0(&ctx->ac, ctx->esgs_ring, provoking_vtx_index));
ac_build_endif(&ctx->ac, 5400);
}
@ -690,7 +680,7 @@ void gfx10_emit_ngg_epilogue(struct ac_shader_abi *abi,
/* Update query buffer */
/* TODO: this won't catch 96-bit clear_buffer via transform feedback. */
if (!info->properties[TGSI_PROPERTY_VS_BLIT_SGPRS_AMD]) {
tmp = si_unpack_param(ctx, ctx->param_vs_state_bits, 6, 1);
tmp = si_unpack_param(ctx, ctx->vs_state_bits, 6, 1);
tmp = LLVMBuildTrunc(builder, tmp, ctx->i1, "");
ac_build_ifcc(&ctx->ac, tmp, 5029); /* if (STREAMOUT_QUERY_ENABLED) */
tmp = LLVMBuildICmp(builder, LLVMIntEQ, get_wave_id_in_tg(ctx), ctx->ac.i32_0, "");
@ -752,7 +742,8 @@ void gfx10_emit_ngg_epilogue(struct ac_shader_abi *abi,
continue;
}
tmp = LLVMBuildLShr(builder, ctx->abi.gs_invocation_id,
tmp = LLVMBuildLShr(builder,
ac_get_arg(&ctx->ac, ctx->args.gs_invocation_id),
LLVMConstInt(ctx->ac.i32, 8 + i, false), "");
prim.edgeflag[i] = LLVMBuildTrunc(builder, tmp, ctx->ac.i1, "");
@ -1099,7 +1090,7 @@ void gfx10_ngg_gs_emit_epilogue(struct si_shader_context *ctx)
}
/* Write shader query data. */
tmp = si_unpack_param(ctx, ctx->param_vs_state_bits, 6, 1);
tmp = si_unpack_param(ctx, ctx->vs_state_bits, 6, 1);
tmp = LLVMBuildTrunc(builder, tmp, ctx->i1, "");
ac_build_ifcc(&ctx->ac, tmp, 5109); /* if (STREAMOUT_QUERY_ENABLED) */
unsigned num_query_comps = sel->so.num_outputs ? 8 : 4;

105
src/gallium/drivers/radeonsi/si_compute_prim_discard.c
View File

@ -318,50 +318,51 @@ void si_build_prim_discard_compute_shader(struct si_shader_context *ctx)
ac_add_function_attr(ctx->ac.context, vs, -1, AC_FUNC_ATTR_ALWAYSINLINE);
LLVMSetLinkage(vs, LLVMPrivateLinkage);
LLVMTypeRef const_desc_type;
enum ac_arg_type const_desc_type;
if (ctx->shader->selector->info.const_buffers_declared == 1 &&
ctx->shader->selector->info.shader_buffers_declared == 0)
const_desc_type = ctx->f32;
const_desc_type = AC_ARG_CONST_FLOAT_PTR;
else
const_desc_type = ctx->v4i32;
const_desc_type = AC_ARG_CONST_DESC_PTR;
struct si_function_info fninfo;
si_init_function_info(&fninfo);
memset(&ctx->args, 0, sizeof(ctx->args));
LLVMValueRef index_buffers_and_constants, vertex_counter, vb_desc, const_desc;
LLVMValueRef base_vertex, start_instance, block_id, local_id, ordered_wave_id;
LLVMValueRef restart_index, vp_scale[2], vp_translate[2], smallprim_precision;
LLVMValueRef num_prims_udiv_multiplier, num_prims_udiv_terms, sampler_desc;
LLVMValueRef last_wave_prim_id, vertex_count_addr;
struct ac_arg param_index_buffers_and_constants, param_vertex_counter;
struct ac_arg param_vb_desc, param_const_desc;
struct ac_arg param_base_vertex, param_start_instance;
struct ac_arg param_block_id, param_local_id, param_ordered_wave_id;
struct ac_arg param_restart_index, param_smallprim_precision;
struct ac_arg param_num_prims_udiv_multiplier, param_num_prims_udiv_terms;
struct ac_arg param_sampler_desc, param_last_wave_prim_id, param_vertex_count_addr;
add_arg_assign(&fninfo, ARG_SGPR, ac_array_in_const32_addr_space(ctx->v4i32),
&index_buffers_and_constants);
add_arg_assign(&fninfo, ARG_SGPR, ctx->i32, &vertex_counter);
add_arg_assign(&fninfo, ARG_SGPR, ctx->i32, &last_wave_prim_id);
add_arg_assign(&fninfo, ARG_SGPR, ctx->i32, &vertex_count_addr);
add_arg_assign(&fninfo, ARG_SGPR, ac_array_in_const32_addr_space(ctx->v4i32),
&vb_desc);
add_arg_assign(&fninfo, ARG_SGPR, ac_array_in_const32_addr_space(const_desc_type),
&const_desc);
add_arg_assign(&fninfo, ARG_SGPR, ac_array_in_const32_addr_space(ctx->v8i32),
&sampler_desc);
add_arg_assign(&fninfo, ARG_SGPR, ctx->i32, &base_vertex);
add_arg_assign(&fninfo, ARG_SGPR, ctx->i32, &start_instance);
add_arg_assign(&fninfo, ARG_SGPR, ctx->i32, &num_prims_udiv_multiplier);
add_arg_assign(&fninfo, ARG_SGPR, ctx->i32, &num_prims_udiv_terms);
add_arg_assign(&fninfo, ARG_SGPR, ctx->i32, &restart_index);
add_arg_assign(&fninfo, ARG_SGPR, ctx->f32, &smallprim_precision);
ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_CONST_DESC_PTR,
&param_index_buffers_and_constants);
ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &param_vertex_counter);
ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &param_last_wave_prim_id);
ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &param_vertex_count_addr);
ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_CONST_DESC_PTR,
&param_vb_desc);
ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, const_desc_type,
&param_const_desc);
ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_CONST_IMAGE_PTR,
&param_sampler_desc);
ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &param_base_vertex);
ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &param_start_instance);
ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &param_num_prims_udiv_multiplier);
ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &param_num_prims_udiv_terms);
ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &param_restart_index);
ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_FLOAT, &param_smallprim_precision);
/* Block ID and thread ID inputs. */
add_arg_assign(&fninfo, ARG_SGPR, ctx->i32, &block_id);
ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &param_block_id);
if (VERTEX_COUNTER_GDS_MODE == 2)
add_arg_assign(&fninfo, ARG_SGPR, ctx->i32, &ordered_wave_id);
add_arg_assign(&fninfo, ARG_VGPR, ctx->i32, &local_id);
ac_add_arg(&ctx->args, AC_ARG_SGPR, 1, AC_ARG_INT, &param_ordered_wave_id);
ac_add_arg(&ctx->args, AC_ARG_VGPR, 1, AC_ARG_INT, &param_local_id);
/* Create the compute shader function. */
unsigned old_type = ctx->type;
ctx->type = PIPE_SHADER_COMPUTE;
si_create_function(ctx, "prim_discard_cs", NULL, 0, &fninfo, THREADGROUP_SIZE);
si_create_function(ctx, "prim_discard_cs", NULL, 0, THREADGROUP_SIZE);
ctx->type = old_type;
if (VERTEX_COUNTER_GDS_MODE == 1) {
@ -376,14 +377,14 @@ void si_build_prim_discard_compute_shader(struct si_shader_context *ctx)
vs_params[num_vs_params++] = LLVMGetUndef(LLVMTypeOf(LLVMGetParam(vs, 0))); /* RW_BUFFERS */
vs_params[num_vs_params++] = LLVMGetUndef(LLVMTypeOf(LLVMGetParam(vs, 1))); /* BINDLESS */
vs_params[num_vs_params++] = const_desc;
vs_params[num_vs_params++] = sampler_desc;
vs_params[num_vs_params++] = ac_get_arg(&ctx->ac, param_const_desc);
vs_params[num_vs_params++] = ac_get_arg(&ctx->ac, param_sampler_desc);
vs_params[num_vs_params++] = LLVMConstInt(ctx->i32,
S_VS_STATE_INDEXED(key->opt.cs_indexed), 0);
vs_params[num_vs_params++] = base_vertex;
vs_params[num_vs_params++] = start_instance;
vs_params[num_vs_params++] = ac_get_arg(&ctx->ac, param_base_vertex);
vs_params[num_vs_params++] = ac_get_arg(&ctx->ac, param_start_instance);
vs_params[num_vs_params++] = ctx->i32_0; /* DrawID */
vs_params[num_vs_params++] = vb_desc;
vs_params[num_vs_params++] = ac_get_arg(&ctx->ac, param_vb_desc);
vs_params[(param_vertex_id = num_vs_params++)] = NULL; /* VertexID */
vs_params[(param_instance_id = num_vs_params++)] = NULL; /* InstanceID */
@ -396,6 +397,7 @@ void si_build_prim_discard_compute_shader(struct si_shader_context *ctx)
/* Load descriptors. (load 8 dwords at once) */
LLVMValueRef input_indexbuf, output_indexbuf, tmp, desc[8];
LLVMValueRef index_buffers_and_constants = ac_get_arg(&ctx->ac, param_index_buffers_and_constants);
tmp = LLVMBuildPointerCast(builder, index_buffers_and_constants,
ac_array_in_const32_addr_space(ctx->v8i32), "");
tmp = ac_build_load_to_sgpr(&ctx->ac, tmp, ctx->i32_0);
@ -408,12 +410,17 @@ void si_build_prim_discard_compute_shader(struct si_shader_context *ctx)
/* Compute PrimID and InstanceID. */
LLVMValueRef global_thread_id =
ac_build_imad(&ctx->ac, block_id,
LLVMConstInt(ctx->i32, THREADGROUP_SIZE, 0), local_id);
ac_build_imad(&ctx->ac, ac_get_arg(&ctx->ac, param_block_id),
LLVMConstInt(ctx->i32, THREADGROUP_SIZE, 0),
ac_get_arg(&ctx->ac, param_local_id));
LLVMValueRef prim_id = global_thread_id; /* PrimID within an instance */
LLVMValueRef instance_id = ctx->i32_0;
if (key->opt.cs_instancing) {
LLVMValueRef num_prims_udiv_terms =
ac_get_arg(&ctx->ac, param_num_prims_udiv_terms);
LLVMValueRef num_prims_udiv_multiplier =
ac_get_arg(&ctx->ac, param_num_prims_udiv_multiplier);
/* Unpack num_prims_udiv_terms. */
LLVMValueRef post_shift = LLVMBuildAnd(builder, num_prims_udiv_terms,
LLVMConstInt(ctx->i32, 0x1f, 0), "");
@ -477,6 +484,8 @@ void si_build_prim_discard_compute_shader(struct si_shader_context *ctx)
}
}
LLVMValueRef ordered_wave_id = ac_get_arg(&ctx->ac, param_ordered_wave_id);
/* Extract the ordered wave ID. */
if (VERTEX_COUNTER_GDS_MODE == 2) {
ordered_wave_id = LLVMBuildLShr(builder, ordered_wave_id,
@ -485,7 +494,8 @@ void si_build_prim_discard_compute_shader(struct si_shader_context *ctx)
LLVMConstInt(ctx->i32, 0xfff, 0), "");
}
LLVMValueRef thread_id =
LLVMBuildAnd(builder, local_id, LLVMConstInt(ctx->i32, 63, 0), "");
LLVMBuildAnd(builder, ac_get_arg(&ctx->ac, param_local_id),
LLVMConstInt(ctx->i32, 63, 0), "");
/* Every other triangle in a strip has a reversed vertex order, so we
* need to swap vertices of odd primitives to get the correct primitive
@ -493,6 +503,7 @@ void si_build_prim_discard_compute_shader(struct si_shader_context *ctx)
* restart complicates it, because a strip can start anywhere.
*/
LLVMValueRef prim_restart_accepted = ctx->i1true;
LLVMValueRef vertex_counter = ac_get_arg(&ctx->ac, param_vertex_counter);
if (key->opt.cs_prim_type == PIPE_PRIM_TRIANGLE_STRIP) {
/* Without primitive restart, odd primitives have reversed orientation.
@ -520,7 +531,8 @@ void si_build_prim_discard_compute_shader(struct si_shader_context *ctx)
for (unsigned i = 0; i < 3; i++) {
LLVMValueRef not_reset = LLVMBuildICmp(builder, LLVMIntNE, index[i],
restart_index, "");
ac_get_arg(&ctx->ac, param_restart_index),
"");
if (i == 0)
index0_is_reset = LLVMBuildNot(builder, not_reset, "");
prim_restart_accepted = LLVMBuildAnd(builder, prim_restart_accepted,
@ -680,6 +692,7 @@ void si_build_prim_discard_compute_shader(struct si_shader_context *ctx)
LLVMValueRef vp = ac_build_load_invariant(&ctx->ac, index_buffers_and_constants,
LLVMConstInt(ctx->i32, 2, 0));
vp = LLVMBuildBitCast(builder, vp, ctx->v4f32, "");
LLVMValueRef vp_scale[2], vp_translate[2];
vp_scale[0] = ac_llvm_extract_elem(&ctx->ac, vp, 0);
vp_scale[1] = ac_llvm_extract_elem(&ctx->ac, vp, 1);
vp_translate[0] = ac_llvm_extract_elem(&ctx->ac, vp, 2);
@ -699,7 +712,8 @@ void si_build_prim_discard_compute_shader(struct si_shader_context *ctx)
LLVMValueRef accepted =
ac_cull_triangle(&ctx->ac, pos, prim_restart_accepted,
vp_scale, vp_translate, smallprim_precision,
vp_scale, vp_translate,
ac_get_arg(&ctx->ac, param_smallprim_precision),
&options);
LLVMValueRef accepted_threadmask = ac_get_i1_sgpr_mask(&ctx->ac, accepted);
@ -788,7 +802,8 @@ void si_build_prim_discard_compute_shader(struct si_shader_context *ctx)
if (VERTEX_COUNTER_GDS_MODE == 2) {
ac_build_ifcc(&ctx->ac,
LLVMBuildICmp(builder, LLVMIntEQ, global_thread_id,
last_wave_prim_id, ""), 12606);
ac_get_arg(&ctx->ac, param_last_wave_prim_id), ""),
12606);
LLVMValueRef count = LLVMBuildAdd(builder, start, num_prims_accepted, "");
count = LLVMBuildMul(builder, count,
LLVMConstInt(ctx->i32, vertices_per_prim, 0), "");
@ -798,7 +813,7 @@ void si_build_prim_discard_compute_shader(struct si_shader_context *ctx)
*/
if (ctx->screen->info.chip_class <= GFX8) {
LLVMValueRef desc[] = {
vertex_count_addr,
ac_get_arg(&ctx->ac, param_vertex_count_addr),
LLVMConstInt(ctx->i32,
S_008F04_BASE_ADDRESS_HI(ctx->screen->info.address32_hi), 0),
LLVMConstInt(ctx->i32, 4, 0),
@ -810,7 +825,9 @@ void si_build_prim_discard_compute_shader(struct si_shader_context *ctx)
ctx->i32_0, 0, ac_glc | ac_slc, false);
} else {
LLVMBuildStore(builder, count,
si_expand_32bit_pointer(ctx, vertex_count_addr));
si_expand_32bit_pointer(ctx,
ac_get_arg(&ctx->ac,
param_vertex_count_addr)));
}
ac_build_endif(&ctx->ac, 12606);
} else {

1131
src/gallium/drivers/radeonsi/si_shader.c
View File

File diff suppressed because it is too large Load Diff

106
src/gallium/drivers/radeonsi/si_shader_internal.h
View File

@ -44,22 +44,6 @@ struct pipe_debug_callback;
#define RADEON_LLVM_MAX_SYSTEM_VALUES 11
#define RADEON_LLVM_MAX_ADDRS 16
enum si_arg_regfile {
ARG_SGPR,
ARG_VGPR
};
/**
* Used to collect types and other info about arguments of the LLVM function
* before the function is created.
*/
struct si_function_info {
LLVMTypeRef types[100];
LLVMValueRef *assign[100];
unsigned num_sgpr_params;
unsigned num_params;
};
struct si_shader_output_values {
LLVMValueRef values[4];
unsigned semantic_name;
@ -82,6 +66,7 @@ struct si_shader_context {
unsigned num_images;
unsigned num_samplers;
struct ac_shader_args args;
struct ac_shader_abi abi;
/** This function is responsible for initilizing the inputs array and will be
@ -123,19 +108,26 @@ struct si_shader_context {
LLVMValueRef main_fn;
LLVMTypeRef return_type;
/* Parameter indices for LLVMGetParam. */
int param_rw_buffers;
int param_const_and_shader_buffers;
int param_samplers_and_images;
int param_bindless_samplers_and_images;
struct ac_arg const_and_shader_buffers;
struct ac_arg samplers_and_images;
/* For merged shaders, the per-stage descriptors for the stage other
* than the one we're processing, used to pass them through from the
* first stage to the second.
*/
struct ac_arg other_const_and_shader_buffers;
struct ac_arg other_samplers_and_images;
struct ac_arg rw_buffers;
struct ac_arg bindless_samplers_and_images;
/* Common inputs for merged shaders. */
int param_merged_wave_info;
int param_merged_scratch_offset;
struct ac_arg merged_wave_info;
struct ac_arg merged_scratch_offset;
/* API VS */
int param_vertex_buffers;
int param_rel_auto_id;
int param_vs_prim_id;
int param_vertex_index0;
struct ac_arg vertex_buffers;
struct ac_arg rel_auto_id;
struct ac_arg vs_prim_id;
struct ac_arg vertex_index0;
/* VS states and layout of LS outputs / TCS inputs at the end
* [0] = clamp vertex color
* [1] = indexed
@ -144,12 +136,12 @@ struct si_shader_context {
* [24:31] = stride between vertices in DW = num_inputs * 4
* max = 32*4
*/
int param_vs_state_bits;
int param_vs_blit_inputs;
struct ac_arg vs_state_bits;
struct ac_arg vs_blit_inputs;
/* HW VS */
int param_streamout_config;
int param_streamout_write_index;
int param_streamout_offset[4];
struct ac_arg streamout_config;
struct ac_arg streamout_write_index;
struct ac_arg streamout_offset[4];
/* API TCS & TES */
/* Layout of TCS outputs in the offchip buffer
@ -161,7 +153,7 @@ struct si_shader_context {
* [12:31] = the offset of per patch attributes in the buffer in bytes.
* max = NUM_PATCHES*32*32*16
*/
int param_tcs_offchip_layout;
struct ac_arg tcs_offchip_layout;
/* API TCS */
/* Offsets where TCS outputs and TCS patch outputs live in LDS:
@ -169,41 +161,43 @@ struct si_shader_context {
* [16:31] = TCS output patch0 offset for per-patch / 16
* max = (NUM_PATCHES + 1) * 32*32
*/
int param_tcs_out_lds_offsets;
struct ac_arg tcs_out_lds_offsets;
/* Layout of TCS outputs / TES inputs:
* [0:12] = stride between output patches in DW, num_outputs * num_vertices * 4
* max = 32*32*4 + 32*4
* [13:18] = gl_PatchVerticesIn, max = 32
* [19:31] = high 13 bits of the 32-bit address of tessellation ring buffers
*/
int param_tcs_out_lds_layout;
int param_tcs_offchip_offset;
int param_tcs_factor_offset;
struct ac_arg tcs_out_lds_layout;
struct ac_arg tcs_offchip_offset;
struct ac_arg tcs_factor_offset;
/* API TES */
int param_tes_offchip_addr;
int param_tes_u;
int param_tes_v;
int param_tes_rel_patch_id;
struct ac_arg tes_offchip_addr;
struct ac_arg tes_u;
struct ac_arg tes_v;
struct ac_arg tes_rel_patch_id;
/* HW ES */
int param_es2gs_offset;
struct ac_arg es2gs_offset;
/* HW GS */
/* On gfx10:
* - bits 0..10: ordered_wave_id
* - bits 12..20: number of vertices in group
* - bits 22..30: number of primitives in group
*/
LLVMValueRef gs_tg_info;
struct ac_arg gs_tg_info;
/* API GS */
int param_gs2vs_offset;
int param_gs_wave_id; /* GFX6 */
LLVMValueRef gs_vtx_offset[6]; /* in dwords (GFX6) */
int param_gs_vtx01_offset; /* in dwords (GFX9) */
int param_gs_vtx23_offset; /* in dwords (GFX9) */
int param_gs_vtx45_offset; /* in dwords (GFX9) */
struct ac_arg gs2vs_offset;
struct ac_arg gs_wave_id; /* GFX6 */
struct ac_arg gs_vtx_offset[6]; /* in dwords (GFX6) */
struct ac_arg gs_vtx01_offset; /* in dwords (GFX9) */
struct ac_arg gs_vtx23_offset; /* in dwords (GFX9) */
struct ac_arg gs_vtx45_offset; /* in dwords (GFX9) */
/* PS */
struct ac_arg pos_fixed_pt;
/* CS */
int param_block_size;
int param_cs_user_data;
struct ac_arg block_size;
struct ac_arg cs_user_data;
struct ac_llvm_compiler *compiler;
@ -252,14 +246,9 @@ si_shader_context_from_abi(struct ac_shader_abi *abi)
return container_of(abi, ctx, abi);
}
void si_init_function_info(struct si_function_info *fninfo);
unsigned add_arg_assign(struct si_function_info *fninfo,
enum si_arg_regfile regfile, LLVMTypeRef type,
LLVMValueRef *assign);
void si_create_function(struct si_shader_context *ctx,
const char *name,
LLVMTypeRef *returns, unsigned num_returns,
struct si_function_info *fninfo,
unsigned max_workgroup_size);
unsigned si_llvm_compile(LLVMModuleRef M, struct si_shader_binary *binary,
struct ac_llvm_compiler *compiler,
@ -287,8 +276,7 @@ void si_llvm_context_set_ir(struct si_shader_context *ctx,
void si_llvm_create_func(struct si_shader_context *ctx,
const char *name,
LLVMTypeRef *return_types, unsigned num_return_elems,
LLVMTypeRef *ParamTypes, unsigned ParamCount);
LLVMTypeRef *return_types, unsigned num_return_elems);
void si_llvm_dispose(struct si_shader_context *ctx);
@ -385,7 +373,7 @@ void si_llvm_load_input_fs(
bool si_nir_build_llvm(struct si_shader_context *ctx, struct nir_shader *nir);
LLVMValueRef si_unpack_param(struct si_shader_context *ctx,
unsigned param, unsigned rshift,
struct ac_arg param, unsigned rshift,
unsigned bitwidth);
void gfx10_emit_ngg_epilogue(struct ac_shader_abi *abi,

19
src/gallium/drivers/radeonsi/si_shader_nir.c
View File

@ -1047,19 +1047,19 @@ si_nir_lookup_interp_param(struct ac_shader_abi *abi,
case INTERP_MODE_SMOOTH:
case INTERP_MODE_NONE:
if (location == INTERP_CENTER)
return ctx->abi.persp_center;
return ac_get_arg(&ctx->ac, ctx->args.persp_center);
else if (location == INTERP_CENTROID)
return ctx->abi.persp_centroid;
else if (location == INTERP_SAMPLE)
return ctx->abi.persp_sample;
return ac_get_arg(&ctx->ac, ctx->args.persp_sample);
break;
case INTERP_MODE_NOPERSPECTIVE:
if (location == INTERP_CENTER)
return ctx->abi.linear_center;
return ac_get_arg(&ctx->ac, ctx->args.linear_center);
else if (location == INTERP_CENTROID)
return ctx->abi.linear_centroid;
return ac_get_arg(&ctx->ac, ctx->args.linear_centroid);
else if (location == INTERP_SAMPLE)
return ctx->abi.linear_sample;
return ac_get_arg(&ctx->ac, ctx->args.linear_sample);
break;
default:
assert(!"Unhandled interpolation mode.");
@ -1082,8 +1082,7 @@ si_nir_load_sampler_desc(struct ac_shader_abi *abi,
assert(desc_type <= AC_DESC_BUFFER);
if (bindless) {
LLVMValueRef list =
LLVMGetParam(ctx->main_fn, ctx->param_bindless_samplers_and_images);
LLVMValueRef list = ac_get_arg(&ctx->ac, ctx->bindless_samplers_and_images);
/* dynamic_index is the bindless handle */
if (image) {
@ -1114,7 +1113,7 @@ si_nir_load_sampler_desc(struct ac_shader_abi *abi,
unsigned num_slots = image ? ctx->num_images : ctx->num_samplers;
assert(const_index < num_slots || dynamic_index);
LLVMValueRef list = LLVMGetParam(ctx->main_fn, ctx->param_samplers_and_images);
LLVMValueRef list = ac_get_arg(&ctx->ac, ctx->samplers_and_images);
LLVMValueRef index = LLVMConstInt(ctx->ac.i32, const_index, false);
if (dynamic_index) {
@ -1230,7 +1229,7 @@ bool si_nir_build_llvm(struct si_shader_context *ctx, struct nir_shader *nir)
ctx->shader->key.mono.u.ps.interpolate_at_sample_force_center;
} else if (nir->info.stage == MESA_SHADER_COMPUTE) {
if (nir->info.cs.user_data_components_amd) {
ctx->abi.user_data = LLVMGetParam(ctx->main_fn, ctx->param_cs_user_data);
ctx->abi.user_data = ac_get_arg(&ctx->ac, ctx->cs_user_data);
ctx->abi.user_data = ac_build_expand_to_vec4(&ctx->ac, ctx->abi.user_data,
nir->info.cs.user_data_components_amd);
}
@ -1248,7 +1247,7 @@ bool si_nir_build_llvm(struct si_shader_context *ctx, struct nir_shader *nir)
assert(gl_shader_stage_is_compute(nir->info.stage));
si_declare_compute_memory(ctx);
}
ac_nir_translate(&ctx->ac, &ctx->abi, nir);
ac_nir_translate(&ctx->ac, &ctx->abi, &ctx->args, nir);
return true;
}

19
src/gallium/drivers/radeonsi/si_shader_tgsi_mem.c
View File

@ -224,8 +224,8 @@ image_fetch_rsrc(
/* Bindless descriptors are accessible from a different pair of
* user SGPR indices.
*/
rsrc_ptr = LLVMGetParam(ctx->main_fn,
ctx->param_bindless_samplers_and_images);
rsrc_ptr = ac_get_arg(&ctx->ac,
ctx->bindless_samplers_and_images);
index = lp_build_emit_fetch_src(bld_base, image, TGSI_TYPE_UNSIGNED, 0);
/* Bindless image descriptors use 16-dword slots. */
@ -235,7 +235,7 @@ image_fetch_rsrc(
if (fmask)
index = LLVMBuildAdd(ctx->ac.builder, index, ctx->i32_1, "");
} else {
rsrc_ptr = LLVMGetParam(ctx->main_fn, ctx->param_samplers_and_images);
rsrc_ptr = ac_get_arg(&ctx->ac, ctx->samplers_and_images);
if (!image->Register.Indirect) {
index = LLVMConstInt(ctx->i32, image->Register.Index, 0);
@ -1126,7 +1126,7 @@ static void tex_fetch_ptrs(struct lp_build_tgsi_context *bld_base,
LLVMValueRef *fmask_ptr)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
LLVMValueRef list = LLVMGetParam(ctx->main_fn, ctx->param_samplers_and_images);
LLVMValueRef list = ac_get_arg(&ctx->ac, ctx->samplers_and_images);
const struct tgsi_full_instruction *inst = emit_data->inst;
const struct tgsi_full_src_register *reg;
unsigned target = inst->Texture.Texture;
@ -1152,8 +1152,7 @@ static void tex_fetch_ptrs(struct lp_build_tgsi_context *bld_base,
/* Bindless descriptors are accessible from a different pair of
* user SGPR indices.
*/
list = LLVMGetParam(ctx->main_fn,
ctx->param_bindless_samplers_and_images);
list = ac_get_arg(&ctx->ac, ctx->bindless_samplers_and_images);
index = lp_build_emit_fetch_src(bld_base, reg,
TGSI_TYPE_UNSIGNED, 0);
@ -1735,7 +1734,7 @@ static LLVMValueRef si_llvm_emit_fbfetch(struct si_shader_context *ctx)
/* Load the image descriptor. */
STATIC_ASSERT(SI_PS_IMAGE_COLORBUF0 % 2 == 0);
ptr = LLVMGetParam(ctx->main_fn, ctx->param_rw_buffers);
ptr = ac_get_arg(&ctx->ac, ctx->rw_buffers);
ptr = LLVMBuildPointerCast(ctx->ac.builder, ptr,
ac_array_in_const32_addr_space(ctx->v8i32), "");
image = ac_build_load_to_sgpr(&ctx->ac, ptr,
@ -1743,14 +1742,14 @@ static LLVMValueRef si_llvm_emit_fbfetch(struct si_shader_context *ctx)
unsigned chan = 0;
args.coords[chan++] = si_unpack_param(ctx, SI_PARAM_POS_FIXED_PT, 0, 16);
args.coords[chan++] = si_unpack_param(ctx, ctx->pos_fixed_pt, 0, 16);
if (!ctx->shader->key.mono.u.ps.fbfetch_is_1D)
args.coords[chan++] = si_unpack_param(ctx, SI_PARAM_POS_FIXED_PT, 16, 16);
args.coords[chan++] = si_unpack_param(ctx, ctx->pos_fixed_pt, 16, 16);
/* Get the current render target layer index. */
if (ctx->shader->key.mono.u.ps.fbfetch_layered)
args.coords[chan++] = si_unpack_param(ctx, SI_PARAM_ANCILLARY, 16, 11);
args.coords[chan++] = si_unpack_param(ctx, ctx->args.ancillary, 16, 11);
if (ctx->shader->key.mono.u.ps.fbfetch_msaa)
args.coords[chan++] = si_get_sample_id(ctx);

39
src/gallium/drivers/radeonsi/si_shader_tgsi_setup.c
View File

@ -27,14 +27,6 @@
#include "ac_llvm_util.h"
#include "util/u_memory.h"
enum si_llvm_calling_convention {
RADEON_LLVM_AMDGPU_VS = 87,
RADEON_LLVM_AMDGPU_GS = 88,
RADEON_LLVM_AMDGPU_PS = 89,
RADEON_LLVM_AMDGPU_CS = 90,
RADEON_LLVM_AMDGPU_HS = 93,
};
struct si_llvm_diagnostics {
struct pipe_debug_callback *debug;
unsigned retval;
@ -1091,12 +1083,10 @@ void si_llvm_context_set_ir(struct si_shader_context *ctx,
void si_llvm_create_func(struct si_shader_context *ctx,
const char *name,
LLVMTypeRef *return_types, unsigned num_return_elems,
LLVMTypeRef *ParamTypes, unsigned ParamCount)
LLVMTypeRef *return_types, unsigned num_return_elems)
{
LLVMTypeRef main_fn_type, ret_type;
LLVMBasicBlockRef main_fn_body;
enum si_llvm_calling_convention call_conv;
LLVMTypeRef ret_type;
enum ac_llvm_calling_convention call_conv;
enum pipe_shader_type real_shader_type;
if (num_return_elems)
@ -1106,14 +1096,6 @@ void si_llvm_create_func(struct si_shader_context *ctx,
else
ret_type = ctx->voidt;
/* Setup the function */
ctx->return_type = ret_type;
main_fn_type = LLVMFunctionType(ret_type, ParamTypes, ParamCount, 0);
ctx->main_fn = LLVMAddFunction(ctx->gallivm.module, name, main_fn_type);
main_fn_body = LLVMAppendBasicBlockInContext(ctx->ac.context,
ctx->main_fn, "main_body");
LLVMPositionBuilderAtEnd(ctx->ac.builder, main_fn_body);
real_shader_type = ctx->type;
/* LS is merged into HS (TCS), and ES is merged into GS. */
@ -1127,25 +1109,28 @@ void si_llvm_create_func(struct si_shader_context *ctx,
switch (real_shader_type) {
case PIPE_SHADER_VERTEX:
case PIPE_SHADER_TESS_EVAL:
call_conv = RADEON_LLVM_AMDGPU_VS;
call_conv = AC_LLVM_AMDGPU_VS;
break;
case PIPE_SHADER_TESS_CTRL:
call_conv = RADEON_LLVM_AMDGPU_HS;
call_conv = AC_LLVM_AMDGPU_HS;
break;
case PIPE_SHADER_GEOMETRY:
call_conv = RADEON_LLVM_AMDGPU_GS;
call_conv = AC_LLVM_AMDGPU_GS;
break;
case PIPE_SHADER_FRAGMENT:
call_conv = RADEON_LLVM_AMDGPU_PS;
call_conv = AC_LLVM_AMDGPU_PS;
break;
case PIPE_SHADER_COMPUTE:
call_conv = RADEON_LLVM_AMDGPU_CS;
call_conv = AC_LLVM_AMDGPU_CS;
break;
default:
unreachable("Unhandle shader type");
}
LLVMSetFunctionCallConv(ctx->main_fn, call_conv);
/* Setup the function */
ctx->return_type = ret_type;
ctx->main_fn = ac_build_main(&ctx->args, &ctx->ac, call_conv, name,
ret_type, ctx->gallivm.module);
}
void si_llvm_optimize_module(struct si_shader_context *ctx)