fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2026-01-10 10:20:20 +01:00
Code Issues Projects Releases Packages Wiki Activity Actions 2

ac/llvm: don't declare LDS as an array for HS & GS & CS, use IntToPtr(0)

Browse source 
We don't need all this stuff anymore.

Reviewed-by: Pierre-Eric Pelloux-Prayer <pierre-eric.pelloux-prayer@amd.com>
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/35529>
This commit is contained in:
Marek Olšák 2025-06-11 16:34:05 -04:00 • committed by Marge Bot
parent 5ded4f3c7d
commit f6aecfb886
5 changed files with 12 additions and 93 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

11
src/amd/llvm/ac_llvm_build.c
View file

@ -2243,17 +2243,6 @@ void ac_init_exec_full_mask(struct ac_llvm_context *ctx)
ac_build_intrinsic(ctx, "llvm.amdgcn.init.exec", ctx->voidt, &full_mask, 1, 0);
}
void ac_declare_lds_as_pointer(struct ac_llvm_context *ctx)
{
unsigned lds_size = ctx->gfx_level >= GFX7 ? 65536 : 32768;
LLVMTypeRef type = LLVMArrayType(ctx->i32, lds_size / 4);
ctx->lds = (struct ac_llvm_pointer) {
.value = LLVMBuildIntToPtr(ctx->builder, ctx->i32_0,
LLVMPointerType(type, AC_ADDR_SPACE_LDS), "lds"),
.pointee_type = type
};
}
LLVMValueRef ac_find_lsb(struct ac_llvm_context *ctx, LLVMTypeRef dst_type, LLVMValueRef src0)
{
unsigned src0_bitsize = ac_get_elem_bits(ctx, LLVMTypeOf(src0));

4
src/amd/llvm/ac_llvm_build.h
View file

@ -142,8 +142,6 @@ struct ac_llvm_context {
bool exports_color_null;
bool exports_mrtz;
struct ac_llvm_pointer lds;
LLVMValueRef ring_offsets;
int ring_offsets_index;
};
@ -408,8 +406,6 @@ LLVMValueRef ac_build_sudot_4x8(struct ac_llvm_context *ctx, LLVMValueRef s0, LL
void ac_init_exec_full_mask(struct ac_llvm_context *ctx);
void ac_declare_lds_as_pointer(struct ac_llvm_context *ac);
LLVMValueRef ac_find_lsb(struct ac_llvm_context *ctx, LLVMTypeRef dst_type, LLVMValueRef src0);
LLVMTypeRef ac_array_in_const_addr_space(LLVMTypeRef elem_type);

33
src/amd/llvm/ac_nir_to_llvm.c
View file

@ -31,6 +31,7 @@ struct ac_nir_context {
struct ac_llvm_pointer scratch;
struct ac_llvm_pointer constant_data;
struct ac_llvm_pointer lds;
struct hash_table *defs;
struct hash_table *phis;
@ -57,10 +58,20 @@ static LLVMValueRef get_src(struct ac_nir_context *nir, nir_src src)
static LLVMValueRef get_memory_ptr(struct ac_nir_context *ctx, nir_src src, unsigned c_off)
{
if (!ctx->lds.value) {
unsigned lds_size = ctx->ac.gfx_level >= GFX7 ? 65536 : 32768;
LLVMTypeRef type = LLVMArrayType(ctx->ac.i32, lds_size / 4);
ctx->lds = (struct ac_llvm_pointer) {
.value = LLVMBuildIntToPtr(ctx->ac.builder, ctx->ac.i32_0,
LLVMPointerType(type, AC_ADDR_SPACE_LDS), "lds"),
.pointee_type = type,
};
}
LLVMValueRef ptr = get_src(ctx, src);
ptr = LLVMBuildAdd(ctx->ac.builder, ptr, LLVMConstInt(ctx->ac.i32, c_off, 0), "");
/* LDS is used here as a i8 pointer. */
return LLVMBuildGEP2(ctx->ac.builder, ctx->ac.i8, ctx->ac.lds.value, &ptr, 1, "");
return LLVMBuildGEP2(ctx->ac.builder, ctx->ac.i8, ctx->lds.value, &ptr, 1, "");
}
static LLVMBasicBlockRef get_block(struct ac_nir_context *nir, const struct nir_block *b)
@ -4024,23 +4035,6 @@ static void setup_constant_data(struct ac_nir_context *ctx, struct nir_shader *s
};
}
static void setup_shared(struct ac_nir_context *ctx, struct nir_shader *nir)
{
if (ctx->ac.lds.value)
return;
LLVMTypeRef type = LLVMArrayType(ctx->ac.i8, nir->info.shared_size);
LLVMValueRef lds =
LLVMAddGlobalInAddressSpace(ctx->ac.module, type, "compute_lds", AC_ADDR_SPACE_LDS);
LLVMSetAlignment(lds, 64 * 1024);
ctx->ac.lds = (struct ac_llvm_pointer) {
.value = lds,
.pointee_type = type
};
}
static void setup_gds(struct ac_nir_context *ctx, nir_function_impl *impl)
{
bool has_gds_atomic = false;
@ -4099,9 +4093,6 @@ bool ac_nir_translate(struct ac_llvm_context *ac, struct ac_shader_abi *abi,
setup_constant_data(&ctx, nir);
setup_gds(&ctx, func->impl);
if (gl_shader_stage_is_compute(nir->info.stage))
setup_shared(&ctx, nir);
if ((ret = visit_cf_list(&ctx, &func->impl->body)))
phi_post_pass(&ctx);

7
src/amd/vulkan/radv_nir_to_llvm.c
View file

@ -108,13 +108,6 @@ create_function(struct radv_shader_context *ctx, gl_shader_stage stage, bool has
ctx->main_function = create_llvm_function(&ctx->ac, ctx->ac.module, ctx->ac.builder, &ctx->args->ac,
get_llvm_calling_convention(ctx->main_function.value, stage),
ctx->max_workgroup_size, ctx->options);
if (stage == MESA_SHADER_TESS_CTRL || (stage == MESA_SHADER_VERTEX && ctx->shader_info->vs.as_ls) ||
ctx->shader_info->is_ngg ||
/* GFX9 has the ESGS ring buffer in LDS. */
(stage == MESA_SHADER_GEOMETRY && has_previous_stage)) {
ac_declare_lds_as_pointer(&ctx->ac);
}
}
static LLVMValueRef

50
src/gallium/drivers/radeonsi/si_shader_llvm.c
View file

@ -196,23 +196,6 @@ static void si_llvm_create_main_func(struct si_shader_context *ctx)
ac_llvm_add_target_dep_function_attr(
ctx->main_fn.value, "InitialPSInputAddr", SI_SPI_PS_INPUT_ADDR_FOR_PROLOG);
}
if (ctx->stage <= MESA_SHADER_GEOMETRY &&
(shader->key.ge.as_ls || ctx->stage == MESA_SHADER_TESS_CTRL ||
shader->key.ge.as_es || shader->key.ge.as_ngg)) {
/* The LSHS size is not known until draw time, so we append it
* at the end of whatever LDS use there may be in the rest of
* the shader (currently none, unless LLVM decides to do its
* own LDS-based lowering).
*/
ctx->ac.lds = (struct ac_llvm_pointer) {
.value = LLVMAddGlobalInAddressSpace(ctx->ac.module, LLVMArrayType(ctx->ac.i32, 0),
"__lds_end", AC_ADDR_SPACE_LDS),
.pointee_type = LLVMArrayType(ctx->ac.i32, 0)
};
LLVMSetAlignment(ctx->ac.lds.value, 256);
}
}
static void si_llvm_optimize_module(struct si_shader_context *ctx)
@ -314,27 +297,6 @@ LLVMValueRef si_unpack_param(struct si_shader_context *ctx, struct ac_arg param,
return unpack_llvm_param(ctx, value, rshift, bitwidth);
}
static void si_llvm_declare_compute_memory(struct si_shader_context *ctx)
{
struct si_shader_selector *sel = ctx->shader->selector;
unsigned lds_size = sel->info.base.shared_size;
LLVMTypeRef i8p = LLVMPointerType(ctx->ac.i8, AC_ADDR_SPACE_LDS);
LLVMValueRef var;
assert(!ctx->ac.lds.value);
LLVMTypeRef type = LLVMArrayType(ctx->ac.i8, lds_size);
var = LLVMAddGlobalInAddressSpace(ctx->ac.module, type,
"compute_lds", AC_ADDR_SPACE_LDS);
LLVMSetAlignment(var, 64 * 1024);
ctx->ac.lds = (struct ac_llvm_pointer) {
.value = LLVMBuildBitCast(ctx->ac.builder, var, i8p, ""),
.pointee_type = type,
};
}
/**
* Given two parts (LS/HS or ES/GS) of a merged shader, build a wrapper function that
* runs them in sequence to form a monolithic shader.
@ -481,16 +443,10 @@ static bool si_llvm_translate_nir(struct si_shader_context *ctx, struct si_shade
si_llvm_create_main_func(ctx);
switch (ctx->stage) {
case MESA_SHADER_VERTEX:
break;
case MESA_SHADER_TESS_CTRL:
si_llvm_init_tcs_callbacks(ctx);
break;
case MESA_SHADER_GEOMETRY:
break;
case MESA_SHADER_FRAGMENT: {
ctx->abi.kill_ps_if_inf_interp =
ctx->screen->options.no_infinite_interp &&
@ -500,12 +456,6 @@ static bool si_llvm_translate_nir(struct si_shader_context *ctx, struct si_shade
break;
}
case MESA_SHADER_COMPUTE:
case MESA_SHADER_KERNEL:
if (ctx->shader->selector->info.base.shared_size)
si_llvm_declare_compute_memory(ctx);
break;
default:
break;
}