From: Nicolai Hähnle <nicolai.haeh...@amd.com>
---
src/amd/common/ac_nir_to_llvm.c | 74 +++++++++++++++++++++++------------------
1 file changed, 41 insertions(+), 33 deletions(-)
diff --git a/src/amd/common/ac_nir_to_llvm.c b/src/amd/common/ac_nir_to_llvm.c
index 68865bd..e65f167 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;
@@ -258,65 +256,70 @@ struct arg_info {
LLVMValueRef *assign[MAX_ARGS];
unsigned array_params_mask;
uint8_t count;
uint8_t user_sgpr_count;
uint8_t sgpr_count;
uint8_t num_user_sgprs_used;
uint8_t num_sgprs_used;
uint8_t num_vgprs_used;
};
-static inline void
+static inline int
add_argument(struct arg_info *info,
LLVMTypeRef type, LLVMValueRef *param_ptr)
{
- assert(info->count < MAX_ARGS);
- info->assign[info->count] = param_ptr;
- info->types[info->count] = type;
+ int idx = info->count;
+ assert(idx < MAX_ARGS);
+ info->assign[idx] = param_ptr;
+ info->types[idx] = type;
info->count++;
+ return idx;
}
-static inline void
+static inline int
add_sgpr_argument(struct arg_info *info,
LLVMTypeRef type, LLVMValueRef *param_ptr)
{
- add_argument(info, type, param_ptr);
+ int idx = add_argument(info, type, param_ptr);
info->num_sgprs_used += llvm_get_type_size(type) / 4;
info->sgpr_count++;
+ return idx;
}
-static inline void
+static inline int
add_user_sgpr_argument(struct arg_info *info,
LLVMTypeRef type,
LLVMValueRef *param_ptr)
{
- add_sgpr_argument(info, type, param_ptr);
+ int idx = add_sgpr_argument(info, type, param_ptr);
info->num_user_sgprs_used += llvm_get_type_size(type) / 4;
info->user_sgpr_count++;
+ return idx;
}
-static inline void
+static inline int
add_vgpr_argument(struct arg_info *info,
LLVMTypeRef type,
LLVMValueRef *param_ptr)
{
- add_argument(info, type, param_ptr);
+ int idx = add_argument(info, type, param_ptr);
info->num_vgprs_used += llvm_get_type_size(type) / 4;
+ return idx;
}
-static inline void
+static inline int
add_user_sgpr_array_argument(struct arg_info *info,
LLVMTypeRef type,
LLVMValueRef *param_ptr)
{
info->array_params_mask |= (1 << info->count);
- add_user_sgpr_argument(info, type, param_ptr);
+ return add_user_sgpr_argument(info, type, param_ptr);
}
static void assign_arguments(LLVMValueRef main_function,
struct arg_info *info)
{
unsigned i;
for (i = 0; i < info->count; i++) {
if (info->assign[i])
*info->assign[i] = LLVMGetParam(main_function, i);
}
@@ -739,34 +742,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->v16i8, 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
+ ctx->abi.param_base_vertex =
add_user_sgpr_argument(&args, ctx->i32, NULL); // base vertex
+ ctx->abi.param_start_instance =
add_user_sgpr_argument(&args, ctx->i32, NULL);// start instance
if (ctx->shader_info->info.vs.needs_draw_id)
- add_user_sgpr_argument(&args, ctx->i32,
&ctx->draw_index); // draw id
+ ctx->abi.param_draw_id =
add_user_sgpr_argument(&args, ctx->i32, NULL); // 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
+ ctx->abi.param_vertex_id = add_vgpr_argument(&args, ctx->i32,
NULL); // 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
+ ctx->abi.param_instance_id = add_vgpr_argument(&args,
ctx->i32, NULL); // 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
@@ -3836,36 +3839,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 = LLVMGetParam(ctx->main_function,
ctx->abi.param_base_vertex);
break;
}
case nir_intrinsic_load_vertex_id_zero_base: {
- result = ctx->vertex_id;
+ result = LLVMGetParam(ctx->main_function,
ctx->abi.param_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 = LLVMGetParam(ctx->main_function,
ctx->abi.param_start_instance);
break;
case nir_intrinsic_load_draw_id:
- result = ctx->draw_index;
+ result = LLVMGetParam(ctx->main_function,
ctx->abi.param_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)
result = ctx->gs_prim_id;
@@ -3884,21 +3887,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 = LLVMGetParam(ctx->main_function,
ctx->abi.param_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:
@@ -4672,27 +4675,30 @@ 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,
+ LLVMGetParam(ctx->main_function,
ctx->abi.param_instance_id),
+ LLVMGetParam(ctx->main_function,
ctx->abi.param_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, "");
+ } else {
+ buffer_index = LLVMBuildAdd(ctx->builder,
+ LLVMGetParam(ctx->main_function,
ctx->abi.param_vertex_id),
+ LLVMGetParam(ctx->main_function,
ctx->abi.param_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);
@@ -6140,21 +6146,23 @@ 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,
+ LLVMGetParam(ctx->main_function,
ctx->abi.param_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
mesa-dev@lists.freedesktop.org
https://lists.freedesktop.org/mailman/listinfo/mesa-dev
