Mesa (master): radeonsi: move more LLVM functions into si_shader_llvm.c
GitLab Mirror
gitlab-mirror at kemper.freedesktop.org
Thu Jan 23 19:32:51 UTC 2020
Module: Mesa
Branch: master
Commit: bd19d144a10f81946ca7d4180cb990d71bc6f0e2
URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=bd19d144a10f81946ca7d4180cb990d71bc6f0e2
Author: Marek Olšák <marek.olsak at amd.com>
Date: Wed Jan 15 18:41:06 2020 -0500
radeonsi: move more LLVM functions into si_shader_llvm.c
Reviewed-by: Timothy Arceri <tarceri at itsqueeze.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3421>
---
src/gallium/drivers/radeonsi/si_shader.c | 388 +---------------------
src/gallium/drivers/radeonsi/si_shader_internal.h | 24 +-
src/gallium/drivers/radeonsi/si_shader_llvm.c | 382 ++++++++++++++++++++-
3 files changed, 398 insertions(+), 396 deletions(-)
diff --git a/src/gallium/drivers/radeonsi/si_shader.c b/src/gallium/drivers/radeonsi/si_shader.c
index 0fc616b4865..b3739533a9c 100644
--- a/src/gallium/drivers/radeonsi/si_shader.c
+++ b/src/gallium/drivers/radeonsi/si_shader.c
@@ -44,7 +44,7 @@ static const char scratch_rsrc_dword1_symbol[] =
static void si_dump_shader_key(const struct si_shader *shader, FILE *f);
/** Whether the shader runs as a combination of multiple API shaders */
-static bool is_multi_part_shader(struct si_shader_context *ctx)
+bool si_is_multi_part_shader(struct si_shader_context *ctx)
{
if (ctx->screen->info.chip_class <= GFX8)
return false;
@@ -58,7 +58,7 @@ static bool is_multi_part_shader(struct si_shader_context *ctx)
/** Whether the shader runs on a merged HW stage (LSHS or ESGS) */
bool si_is_merged_shader(struct si_shader_context *ctx)
{
- return ctx->shader->key.as_ngg || is_multi_part_shader(ctx);
+ return ctx->shader->key.as_ngg || si_is_multi_part_shader(ctx);
}
/**
@@ -145,105 +145,6 @@ unsigned si_shader_io_get_unique_index(unsigned semantic_name, unsigned index,
}
}
-/**
- * Get the value of a shader input parameter and extract a bitfield.
- */
-static LLVMValueRef unpack_llvm_param(struct si_shader_context *ctx,
- LLVMValueRef value, unsigned rshift,
- unsigned bitwidth)
-{
- if (LLVMGetTypeKind(LLVMTypeOf(value)) == LLVMFloatTypeKind)
- value = ac_to_integer(&ctx->ac, value);
-
- if (rshift)
- value = LLVMBuildLShr(ctx->ac.builder, value,
- LLVMConstInt(ctx->ac.i32, rshift, 0), "");
-
- if (rshift + bitwidth < 32) {
- unsigned mask = (1 << bitwidth) - 1;
- value = LLVMBuildAnd(ctx->ac.builder, value,
- LLVMConstInt(ctx->ac.i32, mask, 0), "");
- }
-
- return value;
-}
-
-LLVMValueRef si_unpack_param(struct si_shader_context *ctx,
- struct ac_arg param, unsigned rshift,
- unsigned bitwidth)
-{
- LLVMValueRef value = ac_get_arg(&ctx->ac, param);
-
- return unpack_llvm_param(ctx, value, rshift, bitwidth);
-}
-
-LLVMValueRef si_get_primitive_id(struct si_shader_context *ctx,
- unsigned swizzle)
-{
- if (swizzle > 0)
- return ctx->ac.i32_0;
-
- switch (ctx->type) {
- case PIPE_SHADER_VERTEX:
- return ac_get_arg(&ctx->ac, ctx->vs_prim_id);
- case PIPE_SHADER_TESS_CTRL:
- return ac_get_arg(&ctx->ac, ctx->args.tcs_patch_id);
- case PIPE_SHADER_TESS_EVAL:
- return ac_get_arg(&ctx->ac, ctx->args.tes_patch_id);
- case PIPE_SHADER_GEOMETRY:
- return ac_get_arg(&ctx->ac, ctx->args.gs_prim_id);
- default:
- assert(0);
- return ctx->ac.i32_0;
- }
-}
-
-static LLVMValueRef get_block_size(struct ac_shader_abi *abi)
-{
- struct si_shader_context *ctx = si_shader_context_from_abi(abi);
-
- LLVMValueRef values[3];
- LLVMValueRef result;
- unsigned i;
- unsigned *properties = ctx->shader->selector->info.properties;
-
- if (properties[TGSI_PROPERTY_CS_FIXED_BLOCK_WIDTH] != 0) {
- unsigned sizes[3] = {
- properties[TGSI_PROPERTY_CS_FIXED_BLOCK_WIDTH],
- properties[TGSI_PROPERTY_CS_FIXED_BLOCK_HEIGHT],
- properties[TGSI_PROPERTY_CS_FIXED_BLOCK_DEPTH]
- };
-
- for (i = 0; i < 3; ++i)
- values[i] = LLVMConstInt(ctx->ac.i32, sizes[i], 0);
-
- result = ac_build_gather_values(&ctx->ac, values, 3);
- } else {
- result = ac_get_arg(&ctx->ac, ctx->block_size);
- }
-
- return result;
-}
-
-void si_declare_compute_memory(struct si_shader_context *ctx)
-{
- struct si_shader_selector *sel = ctx->shader->selector;
- unsigned lds_size = sel->info.properties[TGSI_PROPERTY_CS_LOCAL_SIZE];
-
- LLVMTypeRef i8p = LLVMPointerType(ctx->ac.i8, AC_ADDR_SPACE_LDS);
- LLVMValueRef var;
-
- assert(!ctx->ac.lds);
-
- var = LLVMAddGlobalInAddressSpace(ctx->ac.module,
- LLVMArrayType(ctx->ac.i8, lds_size),
- "compute_lds",
- AC_ADDR_SPACE_LDS);
- LLVMSetAlignment(var, 64 * 1024);
-
- ctx->ac.lds = LLVMBuildBitCast(ctx->ac.builder, var, i8p, "");
-}
-
static void si_dump_streamout(struct pipe_stream_output_info *so)
{
unsigned i;
@@ -291,7 +192,7 @@ static void declare_streamout_params(struct si_shader_context *ctx,
}
}
-static unsigned si_get_max_workgroup_size(const struct si_shader *shader)
+unsigned si_get_max_workgroup_size(const struct si_shader *shader)
{
switch (shader->selector->type) {
case PIPE_SHADER_VERTEX:
@@ -1531,7 +1432,7 @@ static bool si_build_main_function(struct si_shader_context *ctx,
si_llvm_init_ps_callbacks(ctx);
break;
case PIPE_SHADER_COMPUTE:
- ctx->abi.load_local_group_size = get_block_size;
+ ctx->abi.load_local_group_size = si_llvm_get_block_size;
break;
default:
assert(!"Unsupported shader type");
@@ -1776,287 +1677,6 @@ static void si_get_vs_prolog_key(const struct si_shader_info *info,
shader_out->info.uses_instanceid = true;
}
-/**
- * Given a list of shader part functions, build a wrapper function that
- * runs them in sequence to form a monolithic shader.
- */
-void si_build_wrapper_function(struct si_shader_context *ctx, LLVMValueRef *parts,
- unsigned num_parts, unsigned main_part,
- unsigned next_shader_first_part)
-{
- LLVMBuilderRef builder = ctx->ac.builder;
- /* PS epilog has one arg per color component; gfx9 merged shader
- * prologs need to forward 40 SGPRs.
- */
- LLVMValueRef initial[AC_MAX_ARGS], out[AC_MAX_ARGS];
- LLVMTypeRef function_type;
- unsigned num_first_params;
- unsigned num_out, initial_num_out;
- ASSERTED unsigned num_out_sgpr; /* used in debug checks */
- ASSERTED unsigned initial_num_out_sgpr; /* used in debug checks */
- unsigned num_sgprs, num_vgprs;
- unsigned gprs;
-
- memset(&ctx->args, 0, sizeof(ctx->args));
-
- for (unsigned i = 0; i < num_parts; ++i) {
- ac_add_function_attr(ctx->ac.context, parts[i], -1,
- AC_FUNC_ATTR_ALWAYSINLINE);
- LLVMSetLinkage(parts[i], LLVMPrivateLinkage);
- }
-
- /* The parameters of the wrapper function correspond to those of the
- * first part in terms of SGPRs and VGPRs, but we use the types of the
- * main part to get the right types. This is relevant for the
- * dereferenceable attribute on descriptor table pointers.
- */
- num_sgprs = 0;
- num_vgprs = 0;
-
- function_type = LLVMGetElementType(LLVMTypeOf(parts[0]));
- num_first_params = LLVMCountParamTypes(function_type);
-
- for (unsigned i = 0; i < num_first_params; ++i) {
- LLVMValueRef param = LLVMGetParam(parts[0], i);
-
- if (ac_is_sgpr_param(param)) {
- assert(num_vgprs == 0);
- num_sgprs += ac_get_type_size(LLVMTypeOf(param)) / 4;
- } else {
- num_vgprs += ac_get_type_size(LLVMTypeOf(param)) / 4;
- }
- }
-
- gprs = 0;
- while (gprs < num_sgprs + num_vgprs) {
- LLVMValueRef param = LLVMGetParam(parts[main_part], ctx->args.arg_count);
- LLVMTypeRef type = LLVMTypeOf(param);
- unsigned size = ac_get_type_size(type) / 4;
-
- /* This is going to get casted anyways, so we don't have to
- * have the exact same type. But we do have to preserve the
- * pointer-ness so that LLVM knows about it.
- */
- enum ac_arg_type arg_type = AC_ARG_INT;
- if (LLVMGetTypeKind(type) == LLVMPointerTypeKind) {
- type = LLVMGetElementType(type);
-
- if (LLVMGetTypeKind(type) == LLVMVectorTypeKind) {
- if (LLVMGetVectorSize(type) == 4)
- arg_type = AC_ARG_CONST_DESC_PTR;
- else if (LLVMGetVectorSize(type) == 8)
- arg_type = AC_ARG_CONST_IMAGE_PTR;
- else
- assert(0);
- } else if (type == ctx->ac.f32) {
- arg_type = AC_ARG_CONST_FLOAT_PTR;
- } else {
- assert(0);
- }
- }
-
- ac_add_arg(&ctx->args, gprs < num_sgprs ? AC_ARG_SGPR : AC_ARG_VGPR,
- size, arg_type, NULL);
-
- assert(ac_is_sgpr_param(param) == (gprs < num_sgprs));
- assert(gprs + size <= num_sgprs + num_vgprs &&
- (gprs >= num_sgprs || gprs + size <= num_sgprs));
-
- gprs += size;
- }
-
- /* Prepare the return type. */
- unsigned num_returns = 0;
- LLVMTypeRef returns[AC_MAX_ARGS], last_func_type, return_type;
-
- last_func_type = LLVMGetElementType(LLVMTypeOf(parts[num_parts - 1]));
- return_type = LLVMGetReturnType(last_func_type);
-
- switch (LLVMGetTypeKind(return_type)) {
- case LLVMStructTypeKind:
- num_returns = LLVMCountStructElementTypes(return_type);
- assert(num_returns <= ARRAY_SIZE(returns));
- LLVMGetStructElementTypes(return_type, returns);
- break;
- case LLVMVoidTypeKind:
- break;
- default:
- unreachable("unexpected type");
- }
-
- si_llvm_create_func(ctx, "wrapper", returns, num_returns,
- si_get_max_workgroup_size(ctx->shader));
-
- if (si_is_merged_shader(ctx))
- ac_init_exec_full_mask(&ctx->ac);
-
- /* Record the arguments of the function as if they were an output of
- * a previous part.
- */
- num_out = 0;
- num_out_sgpr = 0;
-
- for (unsigned i = 0; i < ctx->args.arg_count; ++i) {
- LLVMValueRef param = LLVMGetParam(ctx->main_fn, i);
- LLVMTypeRef param_type = LLVMTypeOf(param);
- LLVMTypeRef out_type = ctx->args.args[i].file == AC_ARG_SGPR ? ctx->ac.i32 : ctx->ac.f32;
- unsigned size = ac_get_type_size(param_type) / 4;
-
- if (size == 1) {
- if (LLVMGetTypeKind(param_type) == LLVMPointerTypeKind) {
- param = LLVMBuildPtrToInt(builder, param, ctx->ac.i32, "");
- param_type = ctx->ac.i32;
- }
-
- if (param_type != out_type)
- param = LLVMBuildBitCast(builder, param, out_type, "");
- out[num_out++] = param;
- } else {
- LLVMTypeRef vector_type = LLVMVectorType(out_type, size);
-
- if (LLVMGetTypeKind(param_type) == LLVMPointerTypeKind) {
- param = LLVMBuildPtrToInt(builder, param, ctx->ac.i64, "");
- param_type = ctx->ac.i64;
- }
-
- if (param_type != vector_type)
- param = LLVMBuildBitCast(builder, param, vector_type, "");
-
- for (unsigned j = 0; j < size; ++j)
- out[num_out++] = LLVMBuildExtractElement(
- builder, param, LLVMConstInt(ctx->ac.i32, j, 0), "");
- }
-
- if (ctx->args.args[i].file == AC_ARG_SGPR)
- num_out_sgpr = num_out;
- }
-
- memcpy(initial, out, sizeof(out));
- initial_num_out = num_out;
- initial_num_out_sgpr = num_out_sgpr;
-
- /* Now chain the parts. */
- LLVMValueRef ret = NULL;
- for (unsigned part = 0; part < num_parts; ++part) {
- LLVMValueRef in[AC_MAX_ARGS];
- LLVMTypeRef ret_type;
- unsigned out_idx = 0;
- unsigned num_params = LLVMCountParams(parts[part]);
-
- /* Merged shaders are executed conditionally depending
- * on the number of enabled threads passed in the input SGPRs. */
- if (is_multi_part_shader(ctx) && part == 0) {
- LLVMValueRef ena, count = initial[3];
-
- count = LLVMBuildAnd(builder, count,
- LLVMConstInt(ctx->ac.i32, 0x7f, 0), "");
- ena = LLVMBuildICmp(builder, LLVMIntULT,
- ac_get_thread_id(&ctx->ac), count, "");
- ac_build_ifcc(&ctx->ac, ena, 6506);
- }
-
- /* Derive arguments for the next part from outputs of the
- * previous one.
- */
- for (unsigned param_idx = 0; param_idx < num_params; ++param_idx) {
- LLVMValueRef param;
- LLVMTypeRef param_type;
- bool is_sgpr;
- unsigned param_size;
- LLVMValueRef arg = NULL;
-
- param = LLVMGetParam(parts[part], param_idx);
- param_type = LLVMTypeOf(param);
- param_size = ac_get_type_size(param_type) / 4;
- is_sgpr = ac_is_sgpr_param(param);
-
- if (is_sgpr) {
- ac_add_function_attr(ctx->ac.context, parts[part],
- param_idx + 1, AC_FUNC_ATTR_INREG);
- } else if (out_idx < num_out_sgpr) {
- /* Skip returned SGPRs the current part doesn't
- * declare on the input. */
- out_idx = num_out_sgpr;
- }
-
- assert(out_idx + param_size <= (is_sgpr ? num_out_sgpr : num_out));
-
- if (param_size == 1)
- arg = out[out_idx];
- else
- arg = ac_build_gather_values(&ctx->ac, &out[out_idx], param_size);
-
- if (LLVMTypeOf(arg) != param_type) {
- if (LLVMGetTypeKind(param_type) == LLVMPointerTypeKind) {
- if (LLVMGetPointerAddressSpace(param_type) ==
- AC_ADDR_SPACE_CONST_32BIT) {
- arg = LLVMBuildBitCast(builder, arg, ctx->ac.i32, "");
- arg = LLVMBuildIntToPtr(builder, arg, param_type, "");
- } else {
- arg = LLVMBuildBitCast(builder, arg, ctx->ac.i64, "");
- arg = LLVMBuildIntToPtr(builder, arg, param_type, "");
- }
- } else {
- arg = LLVMBuildBitCast(builder, arg, param_type, "");
- }
- }
-
- in[param_idx] = arg;
- out_idx += param_size;
- }
-
- ret = ac_build_call(&ctx->ac, parts[part], in, num_params);
-
- if (is_multi_part_shader(ctx) &&
- part + 1 == next_shader_first_part) {
- ac_build_endif(&ctx->ac, 6506);
-
- /* The second half of the merged shader should use
- * the inputs from the toplevel (wrapper) function,
- * not the return value from the last call.
- *
- * That's because the last call was executed condi-
- * tionally, so we can't consume it in the main
- * block.
- */
- memcpy(out, initial, sizeof(initial));
- num_out = initial_num_out;
- num_out_sgpr = initial_num_out_sgpr;
- continue;
- }
-
- /* Extract the returned GPRs. */
- ret_type = LLVMTypeOf(ret);
- num_out = 0;
- num_out_sgpr = 0;
-
- if (LLVMGetTypeKind(ret_type) != LLVMVoidTypeKind) {
- assert(LLVMGetTypeKind(ret_type) == LLVMStructTypeKind);
-
- unsigned ret_size = LLVMCountStructElementTypes(ret_type);
-
- for (unsigned i = 0; i < ret_size; ++i) {
- LLVMValueRef val =
- LLVMBuildExtractValue(builder, ret, i, "");
-
- assert(num_out < ARRAY_SIZE(out));
- out[num_out++] = val;
-
- if (LLVMTypeOf(val) == ctx->ac.i32) {
- assert(num_out_sgpr + 1 == num_out);
- num_out_sgpr = num_out;
- }
- }
- }
- }
-
- /* Return the value from the last part. */
- if (LLVMGetTypeKind(LLVMTypeOf(ret)) == LLVMVoidTypeKind)
- LLVMBuildRetVoid(builder);
- else
- LLVMBuildRet(builder, ret);
-}
-
static bool si_should_optimize_less(struct ac_llvm_compiler *compiler,
struct si_shader_selector *sel)
{
diff --git a/src/gallium/drivers/radeonsi/si_shader_internal.h b/src/gallium/drivers/radeonsi/si_shader_internal.h
index d9b84ab1302..542466ee205 100644
--- a/src/gallium/drivers/radeonsi/si_shader_internal.h
+++ b/src/gallium/drivers/radeonsi/si_shader_internal.h
@@ -196,23 +196,14 @@ si_shader_context_from_abi(struct ac_shader_abi *abi)
return container_of(abi, ctx, abi);
}
+bool si_is_multi_part_shader(struct si_shader_context *ctx);
bool si_is_merged_shader(struct si_shader_context *ctx);
-void si_declare_compute_memory(struct si_shader_context *ctx);
-LLVMValueRef si_get_primitive_id(struct si_shader_context *ctx,
- unsigned swizzle);
void si_add_arg_checked(struct ac_shader_args *args,
enum ac_arg_regfile file,
unsigned registers, enum ac_arg_type type,
struct ac_arg *arg,
unsigned idx);
-bool si_nir_build_llvm(struct si_shader_context *ctx, struct nir_shader *nir);
-
-LLVMValueRef si_unpack_param(struct si_shader_context *ctx,
- struct ac_arg param, unsigned rshift,
- unsigned bitwidth);
-void si_build_wrapper_function(struct si_shader_context *ctx, LLVMValueRef *parts,
- unsigned num_parts, unsigned main_part,
- unsigned next_shader_first_part);
+unsigned si_get_max_workgroup_size(const struct si_shader *shader);
bool si_need_ps_prolog(const union si_shader_part_key *key);
void si_get_ps_prolog_key(struct si_shader *shader,
union si_shader_part_key *key,
@@ -276,6 +267,17 @@ void si_llvm_emit_barrier(struct si_shader_context *ctx);
void si_llvm_declare_esgs_ring(struct si_shader_context *ctx);
void si_init_exec_from_input(struct si_shader_context *ctx, struct ac_arg param,
unsigned bitoffset);
+LLVMValueRef si_unpack_param(struct si_shader_context *ctx,
+ struct ac_arg param, unsigned rshift,
+ unsigned bitwidth);
+LLVMValueRef si_get_primitive_id(struct si_shader_context *ctx,
+ unsigned swizzle);
+LLVMValueRef si_llvm_get_block_size(struct ac_shader_abi *abi);
+void si_llvm_declare_compute_memory(struct si_shader_context *ctx);
+bool si_nir_build_llvm(struct si_shader_context *ctx, struct nir_shader *nir);
+void si_build_wrapper_function(struct si_shader_context *ctx, LLVMValueRef *parts,
+ unsigned num_parts, unsigned main_part,
+ unsigned next_shader_first_part);
/* si_shader_llvm_gs.c */
LLVMValueRef si_is_es_thread(struct si_shader_context *ctx);
diff --git a/src/gallium/drivers/radeonsi/si_shader_llvm.c b/src/gallium/drivers/radeonsi/si_shader_llvm.c
index d7336ea6d87..4ddcbccfac0 100644
--- a/src/gallium/drivers/radeonsi/si_shader_llvm.c
+++ b/src/gallium/drivers/radeonsi/si_shader_llvm.c
@@ -365,6 +365,105 @@ void si_init_exec_from_input(struct si_shader_context *ctx, struct ac_arg param,
ctx->ac.voidt, args, 2, AC_FUNC_ATTR_CONVERGENT);
}
+/**
+ * Get the value of a shader input parameter and extract a bitfield.
+ */
+static LLVMValueRef unpack_llvm_param(struct si_shader_context *ctx,
+ LLVMValueRef value, unsigned rshift,
+ unsigned bitwidth)
+{
+ if (LLVMGetTypeKind(LLVMTypeOf(value)) == LLVMFloatTypeKind)
+ value = ac_to_integer(&ctx->ac, value);
+
+ if (rshift)
+ value = LLVMBuildLShr(ctx->ac.builder, value,
+ LLVMConstInt(ctx->ac.i32, rshift, 0), "");
+
+ if (rshift + bitwidth < 32) {
+ unsigned mask = (1 << bitwidth) - 1;
+ value = LLVMBuildAnd(ctx->ac.builder, value,
+ LLVMConstInt(ctx->ac.i32, mask, 0), "");
+ }
+
+ return value;
+}
+
+LLVMValueRef si_unpack_param(struct si_shader_context *ctx,
+ struct ac_arg param, unsigned rshift,
+ unsigned bitwidth)
+{
+ LLVMValueRef value = ac_get_arg(&ctx->ac, param);
+
+ return unpack_llvm_param(ctx, value, rshift, bitwidth);
+}
+
+LLVMValueRef si_get_primitive_id(struct si_shader_context *ctx,
+ unsigned swizzle)
+{
+ if (swizzle > 0)
+ return ctx->ac.i32_0;
+
+ switch (ctx->type) {
+ case PIPE_SHADER_VERTEX:
+ return ac_get_arg(&ctx->ac, ctx->vs_prim_id);
+ case PIPE_SHADER_TESS_CTRL:
+ return ac_get_arg(&ctx->ac, ctx->args.tcs_patch_id);
+ case PIPE_SHADER_TESS_EVAL:
+ return ac_get_arg(&ctx->ac, ctx->args.tes_patch_id);
+ case PIPE_SHADER_GEOMETRY:
+ return ac_get_arg(&ctx->ac, ctx->args.gs_prim_id);
+ default:
+ assert(0);
+ return ctx->ac.i32_0;
+ }
+}
+
+LLVMValueRef si_llvm_get_block_size(struct ac_shader_abi *abi)
+{
+ struct si_shader_context *ctx = si_shader_context_from_abi(abi);
+
+ LLVMValueRef values[3];
+ LLVMValueRef result;
+ unsigned i;
+ unsigned *properties = ctx->shader->selector->info.properties;
+
+ if (properties[TGSI_PROPERTY_CS_FIXED_BLOCK_WIDTH] != 0) {
+ unsigned sizes[3] = {
+ properties[TGSI_PROPERTY_CS_FIXED_BLOCK_WIDTH],
+ properties[TGSI_PROPERTY_CS_FIXED_BLOCK_HEIGHT],
+ properties[TGSI_PROPERTY_CS_FIXED_BLOCK_DEPTH]
+ };
+
+ for (i = 0; i < 3; ++i)
+ values[i] = LLVMConstInt(ctx->ac.i32, sizes[i], 0);
+
+ result = ac_build_gather_values(&ctx->ac, values, 3);
+ } else {
+ result = ac_get_arg(&ctx->ac, ctx->block_size);
+ }
+
+ return result;
+}
+
+void si_llvm_declare_compute_memory(struct si_shader_context *ctx)
+{
+ struct si_shader_selector *sel = ctx->shader->selector;
+ unsigned lds_size = sel->info.properties[TGSI_PROPERTY_CS_LOCAL_SIZE];
+
+ LLVMTypeRef i8p = LLVMPointerType(ctx->ac.i8, AC_ADDR_SPACE_LDS);
+ LLVMValueRef var;
+
+ assert(!ctx->ac.lds);
+
+ var = LLVMAddGlobalInAddressSpace(ctx->ac.module,
+ LLVMArrayType(ctx->ac.i8, lds_size),
+ "compute_lds",
+ AC_ADDR_SPACE_LDS);
+ LLVMSetAlignment(var, 64 * 1024);
+
+ ctx->ac.lds = LLVMBuildBitCast(ctx->ac.builder, var, i8p, "");
+}
+
bool si_nir_build_llvm(struct si_shader_context *ctx, struct nir_shader *nir)
{
if (nir->info.stage == MESA_SHADER_VERTEX) {
@@ -417,9 +516,290 @@ bool si_nir_build_llvm(struct si_shader_context *ctx, struct nir_shader *nir)
if (ctx->shader->selector->info.properties[TGSI_PROPERTY_CS_LOCAL_SIZE]) {
assert(gl_shader_stage_is_compute(nir->info.stage));
- si_declare_compute_memory(ctx);
+ si_llvm_declare_compute_memory(ctx);
}
ac_nir_translate(&ctx->ac, &ctx->abi, &ctx->args, nir);
return true;
}
+
+/**
+ * Given a list of shader part functions, build a wrapper function that
+ * runs them in sequence to form a monolithic shader.
+ */
+void si_build_wrapper_function(struct si_shader_context *ctx, LLVMValueRef *parts,
+ unsigned num_parts, unsigned main_part,
+ unsigned next_shader_first_part)
+{
+ LLVMBuilderRef builder = ctx->ac.builder;
+ /* PS epilog has one arg per color component; gfx9 merged shader
+ * prologs need to forward 40 SGPRs.
+ */
+ LLVMValueRef initial[AC_MAX_ARGS], out[AC_MAX_ARGS];
+ LLVMTypeRef function_type;
+ unsigned num_first_params;
+ unsigned num_out, initial_num_out;
+ ASSERTED unsigned num_out_sgpr; /* used in debug checks */
+ ASSERTED unsigned initial_num_out_sgpr; /* used in debug checks */
+ unsigned num_sgprs, num_vgprs;
+ unsigned gprs;
+
+ memset(&ctx->args, 0, sizeof(ctx->args));
+
+ for (unsigned i = 0; i < num_parts; ++i) {
+ ac_add_function_attr(ctx->ac.context, parts[i], -1,
+ AC_FUNC_ATTR_ALWAYSINLINE);
+ LLVMSetLinkage(parts[i], LLVMPrivateLinkage);
+ }
+
+ /* The parameters of the wrapper function correspond to those of the
+ * first part in terms of SGPRs and VGPRs, but we use the types of the
+ * main part to get the right types. This is relevant for the
+ * dereferenceable attribute on descriptor table pointers.
+ */
+ num_sgprs = 0;
+ num_vgprs = 0;
+
+ function_type = LLVMGetElementType(LLVMTypeOf(parts[0]));
+ num_first_params = LLVMCountParamTypes(function_type);
+
+ for (unsigned i = 0; i < num_first_params; ++i) {
+ LLVMValueRef param = LLVMGetParam(parts[0], i);
+
+ if (ac_is_sgpr_param(param)) {
+ assert(num_vgprs == 0);
+ num_sgprs += ac_get_type_size(LLVMTypeOf(param)) / 4;
+ } else {
+ num_vgprs += ac_get_type_size(LLVMTypeOf(param)) / 4;
+ }
+ }
+
+ gprs = 0;
+ while (gprs < num_sgprs + num_vgprs) {
+ LLVMValueRef param = LLVMGetParam(parts[main_part], ctx->args.arg_count);
+ LLVMTypeRef type = LLVMTypeOf(param);
+ unsigned size = ac_get_type_size(type) / 4;
+
+ /* This is going to get casted anyways, so we don't have to
+ * have the exact same type. But we do have to preserve the
+ * pointer-ness so that LLVM knows about it.
+ */
+ enum ac_arg_type arg_type = AC_ARG_INT;
+ if (LLVMGetTypeKind(type) == LLVMPointerTypeKind) {
+ type = LLVMGetElementType(type);
+
+ if (LLVMGetTypeKind(type) == LLVMVectorTypeKind) {
+ if (LLVMGetVectorSize(type) == 4)
+ arg_type = AC_ARG_CONST_DESC_PTR;
+ else if (LLVMGetVectorSize(type) == 8)
+ arg_type = AC_ARG_CONST_IMAGE_PTR;
+ else
+ assert(0);
+ } else if (type == ctx->ac.f32) {
+ arg_type = AC_ARG_CONST_FLOAT_PTR;
+ } else {
+ assert(0);
+ }
+ }
+
+ ac_add_arg(&ctx->args, gprs < num_sgprs ? AC_ARG_SGPR : AC_ARG_VGPR,
+ size, arg_type, NULL);
+
+ assert(ac_is_sgpr_param(param) == (gprs < num_sgprs));
+ assert(gprs + size <= num_sgprs + num_vgprs &&
+ (gprs >= num_sgprs || gprs + size <= num_sgprs));
+
+ gprs += size;
+ }
+
+ /* Prepare the return type. */
+ unsigned num_returns = 0;
+ LLVMTypeRef returns[AC_MAX_ARGS], last_func_type, return_type;
+
+ last_func_type = LLVMGetElementType(LLVMTypeOf(parts[num_parts - 1]));
+ return_type = LLVMGetReturnType(last_func_type);
+
+ switch (LLVMGetTypeKind(return_type)) {
+ case LLVMStructTypeKind:
+ num_returns = LLVMCountStructElementTypes(return_type);
+ assert(num_returns <= ARRAY_SIZE(returns));
+ LLVMGetStructElementTypes(return_type, returns);
+ break;
+ case LLVMVoidTypeKind:
+ break;
+ default:
+ unreachable("unexpected type");
+ }
+
+ si_llvm_create_func(ctx, "wrapper", returns, num_returns,
+ si_get_max_workgroup_size(ctx->shader));
+
+ if (si_is_merged_shader(ctx))
+ ac_init_exec_full_mask(&ctx->ac);
+
+ /* Record the arguments of the function as if they were an output of
+ * a previous part.
+ */
+ num_out = 0;
+ num_out_sgpr = 0;
+
+ for (unsigned i = 0; i < ctx->args.arg_count; ++i) {
+ LLVMValueRef param = LLVMGetParam(ctx->main_fn, i);
+ LLVMTypeRef param_type = LLVMTypeOf(param);
+ LLVMTypeRef out_type = ctx->args.args[i].file == AC_ARG_SGPR ? ctx->ac.i32 : ctx->ac.f32;
+ unsigned size = ac_get_type_size(param_type) / 4;
+
+ if (size == 1) {
+ if (LLVMGetTypeKind(param_type) == LLVMPointerTypeKind) {
+ param = LLVMBuildPtrToInt(builder, param, ctx->ac.i32, "");
+ param_type = ctx->ac.i32;
+ }
+
+ if (param_type != out_type)
+ param = LLVMBuildBitCast(builder, param, out_type, "");
+ out[num_out++] = param;
+ } else {
+ LLVMTypeRef vector_type = LLVMVectorType(out_type, size);
+
+ if (LLVMGetTypeKind(param_type) == LLVMPointerTypeKind) {
+ param = LLVMBuildPtrToInt(builder, param, ctx->ac.i64, "");
+ param_type = ctx->ac.i64;
+ }
+
+ if (param_type != vector_type)
+ param = LLVMBuildBitCast(builder, param, vector_type, "");
+
+ for (unsigned j = 0; j < size; ++j)
+ out[num_out++] = LLVMBuildExtractElement(
+ builder, param, LLVMConstInt(ctx->ac.i32, j, 0), "");
+ }
+
+ if (ctx->args.args[i].file == AC_ARG_SGPR)
+ num_out_sgpr = num_out;
+ }
+
+ memcpy(initial, out, sizeof(out));
+ initial_num_out = num_out;
+ initial_num_out_sgpr = num_out_sgpr;
+
+ /* Now chain the parts. */
+ LLVMValueRef ret = NULL;
+ for (unsigned part = 0; part < num_parts; ++part) {
+ LLVMValueRef in[AC_MAX_ARGS];
+ LLVMTypeRef ret_type;
+ unsigned out_idx = 0;
+ unsigned num_params = LLVMCountParams(parts[part]);
+
+ /* Merged shaders are executed conditionally depending
+ * on the number of enabled threads passed in the input SGPRs. */
+ if (si_is_multi_part_shader(ctx) && part == 0) {
+ LLVMValueRef ena, count = initial[3];
+
+ count = LLVMBuildAnd(builder, count,
+ LLVMConstInt(ctx->ac.i32, 0x7f, 0), "");
+ ena = LLVMBuildICmp(builder, LLVMIntULT,
+ ac_get_thread_id(&ctx->ac), count, "");
+ ac_build_ifcc(&ctx->ac, ena, 6506);
+ }
+
+ /* Derive arguments for the next part from outputs of the
+ * previous one.
+ */
+ for (unsigned param_idx = 0; param_idx < num_params; ++param_idx) {
+ LLVMValueRef param;
+ LLVMTypeRef param_type;
+ bool is_sgpr;
+ unsigned param_size;
+ LLVMValueRef arg = NULL;
+
+ param = LLVMGetParam(parts[part], param_idx);
+ param_type = LLVMTypeOf(param);
+ param_size = ac_get_type_size(param_type) / 4;
+ is_sgpr = ac_is_sgpr_param(param);
+
+ if (is_sgpr) {
+ ac_add_function_attr(ctx->ac.context, parts[part],
+ param_idx + 1, AC_FUNC_ATTR_INREG);
+ } else if (out_idx < num_out_sgpr) {
+ /* Skip returned SGPRs the current part doesn't
+ * declare on the input. */
+ out_idx = num_out_sgpr;
+ }
+
+ assert(out_idx + param_size <= (is_sgpr ? num_out_sgpr : num_out));
+
+ if (param_size == 1)
+ arg = out[out_idx];
+ else
+ arg = ac_build_gather_values(&ctx->ac, &out[out_idx], param_size);
+
+ if (LLVMTypeOf(arg) != param_type) {
+ if (LLVMGetTypeKind(param_type) == LLVMPointerTypeKind) {
+ if (LLVMGetPointerAddressSpace(param_type) ==
+ AC_ADDR_SPACE_CONST_32BIT) {
+ arg = LLVMBuildBitCast(builder, arg, ctx->ac.i32, "");
+ arg = LLVMBuildIntToPtr(builder, arg, param_type, "");
+ } else {
+ arg = LLVMBuildBitCast(builder, arg, ctx->ac.i64, "");
+ arg = LLVMBuildIntToPtr(builder, arg, param_type, "");
+ }
+ } else {
+ arg = LLVMBuildBitCast(builder, arg, param_type, "");
+ }
+ }
+
+ in[param_idx] = arg;
+ out_idx += param_size;
+ }
+
+ ret = ac_build_call(&ctx->ac, parts[part], in, num_params);
+
+ if (si_is_multi_part_shader(ctx) &&
+ part + 1 == next_shader_first_part) {
+ ac_build_endif(&ctx->ac, 6506);
+
+ /* The second half of the merged shader should use
+ * the inputs from the toplevel (wrapper) function,
+ * not the return value from the last call.
+ *
+ * That's because the last call was executed condi-
+ * tionally, so we can't consume it in the main
+ * block.
+ */
+ memcpy(out, initial, sizeof(initial));
+ num_out = initial_num_out;
+ num_out_sgpr = initial_num_out_sgpr;
+ continue;
+ }
+
+ /* Extract the returned GPRs. */
+ ret_type = LLVMTypeOf(ret);
+ num_out = 0;
+ num_out_sgpr = 0;
+
+ if (LLVMGetTypeKind(ret_type) != LLVMVoidTypeKind) {
+ assert(LLVMGetTypeKind(ret_type) == LLVMStructTypeKind);
+
+ unsigned ret_size = LLVMCountStructElementTypes(ret_type);
+
+ for (unsigned i = 0; i < ret_size; ++i) {
+ LLVMValueRef val =
+ LLVMBuildExtractValue(builder, ret, i, "");
+
+ assert(num_out < ARRAY_SIZE(out));
+ out[num_out++] = val;
+
+ if (LLVMTypeOf(val) == ctx->ac.i32) {
+ assert(num_out_sgpr + 1 == num_out);
+ num_out_sgpr = num_out;
+ }
+ }
+ }
+ }
+
+ /* Return the value from the last part. */
+ if (LLVMGetTypeKind(LLVMTypeOf(ret)) == LLVMVoidTypeKind)
+ LLVMBuildRetVoid(builder);
+ else
+ LLVMBuildRet(builder, ret);
+}
More information about the mesa-commit
mailing list