From: Nicolai Hähnle <[email protected]>
v2: update for LLVMValueRefs in ac_shader_abi
---
src/amd/common/ac_nir_to_llvm.c | 38 ++++++++++++++++++--------------------
1 file changed, 18 insertions(+), 20 deletions(-)
diff --git a/src/amd/common/ac_nir_to_llvm.c b/src/amd/common/ac_nir_to_llvm.c
index e72747a..40d9cc2 100644
--- a/src/amd/common/ac_nir_to_llvm.c
+++ b/src/amd/common/ac_nir_to_llvm.c
@@ -23,20 +23,21 @@
#include "ac_nir_to_llvm.h"
#include "ac_llvm_build.h"
#include "ac_llvm_util.h"
#include "ac_binary.h"
#include "sid.h"
#include "nir/nir.h"
#include "../vulkan/radv_descriptor_set.h"
#include "util/bitscan.h"
#include <llvm-c/Transforms/Scalar.h>
+#include "ac_shader_abi.h"
#include "ac_shader_info.h"
#include "ac_exp_param.h"
enum radeon_llvm_calling_convention {
RADEON_LLVM_AMDGPU_VS = 87,
RADEON_LLVM_AMDGPU_GS = 88,
RADEON_LLVM_AMDGPU_PS = 89,
RADEON_LLVM_AMDGPU_CS = 90,
};
@@ -50,45 +51,42 @@ enum desc_type {
DESC_IMAGE,
DESC_FMASK,
DESC_SAMPLER,
DESC_BUFFER,
};
struct nir_to_llvm_context {
struct ac_llvm_context ac;
const struct ac_nir_compiler_options *options;
struct ac_shader_variant_info *shader_info;
+ struct ac_shader_abi abi;
+
unsigned max_workgroup_size;
LLVMContextRef context;
LLVMModuleRef module;
LLVMBuilderRef builder;
LLVMValueRef main_function;
struct hash_table *defs;
struct hash_table *phis;
LLVMValueRef descriptor_sets[AC_UD_MAX_SETS];
LLVMValueRef ring_offsets;
LLVMValueRef push_constants;
LLVMValueRef num_work_groups;
LLVMValueRef workgroup_ids;
LLVMValueRef local_invocation_ids;
LLVMValueRef tg_size;
LLVMValueRef vertex_buffers;
- LLVMValueRef base_vertex;
- LLVMValueRef start_instance;
- LLVMValueRef draw_index;
- LLVMValueRef vertex_id;
LLVMValueRef rel_auto_id;
LLVMValueRef vs_prim_id;
- LLVMValueRef instance_id;
LLVMValueRef ls_out_layout;
LLVMValueRef es2gs_offset;
LLVMValueRef tcs_offchip_layout;
LLVMValueRef tcs_out_offsets;
LLVMValueRef tcs_out_layout;
LLVMValueRef tcs_in_layout;
LLVMValueRef oc_lds;
LLVMValueRef tess_factor_offset;
LLVMValueRef tcs_patch_id;
@@ -737,34 +735,34 @@ static void create_function(struct nir_to_llvm_context
*ctx)
if (ctx->shader_info->info.cs.grid_components_used)
add_user_sgpr_argument(&args, LLVMVectorType(ctx->i32,
ctx->shader_info->info.cs.grid_components_used), &ctx->num_work_groups); /*
grid size */
add_sgpr_argument(&args, LLVMVectorType(ctx->i32, 3),
&ctx->workgroup_ids);
add_sgpr_argument(&args, ctx->i32, &ctx->tg_size);
add_vgpr_argument(&args, LLVMVectorType(ctx->i32, 3),
&ctx->local_invocation_ids);
break;
case MESA_SHADER_VERTEX:
if (!ctx->is_gs_copy_shader) {
if (ctx->shader_info->info.vs.has_vertex_buffers)
add_user_sgpr_argument(&args,
const_array(ctx->v4i32, 16), &ctx->vertex_buffers); /* vertex buffers */
- add_user_sgpr_argument(&args, ctx->i32,
&ctx->base_vertex); // base vertex
- add_user_sgpr_argument(&args, ctx->i32,
&ctx->start_instance);// start instance
+ add_user_sgpr_argument(&args, ctx->i32,
&ctx->abi.base_vertex); // base vertex
+ add_user_sgpr_argument(&args, ctx->i32,
&ctx->abi.start_instance);// start instance
if (ctx->shader_info->info.vs.needs_draw_id)
- add_user_sgpr_argument(&args, ctx->i32,
&ctx->draw_index); // draw id
+ add_user_sgpr_argument(&args, ctx->i32,
&ctx->abi.draw_id); // draw id
}
if (ctx->options->key.vs.as_es)
add_sgpr_argument(&args, ctx->i32, &ctx->es2gs_offset);
// es2gs offset
else if (ctx->options->key.vs.as_ls)
add_user_sgpr_argument(&args, ctx->i32,
&ctx->ls_out_layout); // ls out layout
- add_vgpr_argument(&args, ctx->i32, &ctx->vertex_id); // vertex
id
+ add_vgpr_argument(&args, ctx->i32, &ctx->abi.vertex_id); //
vertex id
if (!ctx->is_gs_copy_shader) {
add_vgpr_argument(&args, ctx->i32, &ctx->rel_auto_id);
// rel auto id
add_vgpr_argument(&args, ctx->i32, &ctx->vs_prim_id);
// vs prim id
- add_vgpr_argument(&args, ctx->i32, &ctx->instance_id);
// instance id
+ add_vgpr_argument(&args, ctx->i32,
&ctx->abi.instance_id); // instance id
}
break;
case MESA_SHADER_TESS_CTRL:
add_user_sgpr_argument(&args, ctx->i32,
&ctx->tcs_offchip_layout); // tcs offchip layout
add_user_sgpr_argument(&args, ctx->i32, &ctx->tcs_out_offsets);
// tcs out offsets
add_user_sgpr_argument(&args, ctx->i32, &ctx->tcs_out_layout);
// tcs out layout
add_user_sgpr_argument(&args, ctx->i32, &ctx->tcs_in_layout);
// tcs in layout
add_sgpr_argument(&args, ctx->i32, &ctx->oc_lds); // param oc
lds
add_sgpr_argument(&args, ctx->i32, &ctx->tess_factor_offset);
// tess factor offset
add_vgpr_argument(&args, ctx->i32, &ctx->tcs_patch_id); //
patch id
@@ -3915,36 +3913,36 @@ static void visit_intrinsic(struct nir_to_llvm_context
*ctx,
nir_intrinsic_instr *instr)
{
LLVMValueRef result = NULL;
switch (instr->intrinsic) {
case nir_intrinsic_load_work_group_id: {
result = ctx->workgroup_ids;
break;
}
case nir_intrinsic_load_base_vertex: {
- result = ctx->base_vertex;
+ result = ctx->abi.base_vertex;
break;
}
case nir_intrinsic_load_vertex_id_zero_base: {
- result = ctx->vertex_id;
+ result = ctx->abi.vertex_id;
break;
}
case nir_intrinsic_load_local_invocation_id: {
result = ctx->local_invocation_ids;
break;
}
case nir_intrinsic_load_base_instance:
- result = ctx->start_instance;
+ result = ctx->abi.start_instance;
break;
case nir_intrinsic_load_draw_id:
- result = ctx->draw_index;
+ result = ctx->abi.draw_id;
break;
case nir_intrinsic_load_invocation_id:
if (ctx->stage == MESA_SHADER_TESS_CTRL)
result = unpack_param(ctx, ctx->tcs_rel_ids, 8, 5);
else
result = ctx->gs_invocation_id;
break;
case nir_intrinsic_load_primitive_id:
if (ctx->stage == MESA_SHADER_GEOMETRY) {
ctx->shader_info->gs.uses_prim_id = true;
@@ -3966,21 +3964,21 @@ static void visit_intrinsic(struct nir_to_llvm_context
*ctx,
ctx->shader_info->fs.force_persample = true;
result = load_sample_pos(ctx);
break;
case nir_intrinsic_load_sample_mask_in:
result = ctx->sample_coverage;
break;
case nir_intrinsic_load_front_face:
result = ctx->front_face;
break;
case nir_intrinsic_load_instance_id:
- result = ctx->instance_id;
+ result = ctx->abi.instance_id;
ctx->shader_info->vs.vgpr_comp_cnt = MAX2(3,
ctx->shader_info->vs.vgpr_comp_cnt);
break;
case nir_intrinsic_load_num_work_groups:
result = ctx->num_work_groups;
break;
case nir_intrinsic_load_local_invocation_index:
result = visit_load_local_invocation_index(ctx);
break;
case nir_intrinsic_load_push_constant:
@@ -4754,27 +4752,27 @@ handle_vs_input_decl(struct nir_to_llvm_context *ctx,
LLVMValueRef t_list;
LLVMValueRef input;
LLVMValueRef buffer_index;
int index = variable->data.location - VERT_ATTRIB_GENERIC0;
int idx = variable->data.location;
unsigned attrib_count = glsl_count_attribute_slots(variable->type,
true);
variable->data.driver_location = idx * 4;
if (ctx->options->key.vs.instance_rate_inputs & (1u << index)) {
- buffer_index = LLVMBuildAdd(ctx->builder, ctx->instance_id,
- ctx->start_instance, "");
+ buffer_index = LLVMBuildAdd(ctx->builder, ctx->abi.instance_id,
+ ctx->abi.start_instance, "");
ctx->shader_info->vs.vgpr_comp_cnt = MAX2(3,
ctx->shader_info->vs.vgpr_comp_cnt);
} else
- buffer_index = LLVMBuildAdd(ctx->builder, ctx->vertex_id,
- ctx->base_vertex, "");
+ buffer_index = LLVMBuildAdd(ctx->builder, ctx->abi.vertex_id,
+ ctx->abi.base_vertex, "");
for (unsigned i = 0; i < attrib_count; ++i, ++idx) {
t_offset = LLVMConstInt(ctx->i32, index + i, false);
t_list = ac_build_indexed_load_const(&ctx->ac, t_list_ptr,
t_offset);
input = ac_build_buffer_load_format(&ctx->ac, t_list,
buffer_index,
LLVMConstInt(ctx->i32, 0,
false),
true);
@@ -6240,21 +6238,21 @@ void ac_compile_nir_shader(LLVMTargetMachineRef tm,
default:
break;
}
}
static void
ac_gs_copy_shader_emit(struct nir_to_llvm_context *ctx)
{
LLVMValueRef args[9];
args[0] = ctx->gsvs_ring;
- args[1] = LLVMBuildMul(ctx->builder, ctx->vertex_id,
LLVMConstInt(ctx->i32, 4, false), "");
+ args[1] = LLVMBuildMul(ctx->builder, ctx->abi.vertex_id,
LLVMConstInt(ctx->i32, 4, false), "");
args[3] = ctx->i32zero;
args[4] = ctx->i32one; /* OFFEN */
args[5] = ctx->i32zero; /* IDXEN */
args[6] = ctx->i32one; /* GLC */
args[7] = ctx->i32one; /* SLC */
args[8] = ctx->i32zero; /* TFE */
int idx = 0;
for (unsigned i = 0; i < RADEON_LLVM_MAX_OUTPUTS; ++i) {
--
2.9.3
_______________________________________________
mesa-dev mailing list
[email protected]
https://lists.freedesktop.org/mailman/listinfo/mesa-dev
