From ab33ba987a622e1d222654d77e811b168f499917 Mon Sep 17 00:00:00 2001 From: =?utf8?q?Marek=20Ol=C5=A1=C3=A1k?= Date: Wed, 15 Jan 2020 18:06:02 -0500 Subject: [PATCH] radeonsi: move si_shader_llvm_build.c content into si_shader_llvm.c Reviewed-by: Timothy Arceri Part-of: --- src/gallium/drivers/radeonsi/Makefile.sources | 1 - src/gallium/drivers/radeonsi/meson.build | 1 - .../drivers/radeonsi/si_shader_internal.h | 75 +++++---- src/gallium/drivers/radeonsi/si_shader_llvm.c | 116 ++++++++++++++ .../drivers/radeonsi/si_shader_llvm_build.c | 143 ------------------ 5 files changed, 153 insertions(+), 183 deletions(-) delete mode 100644 src/gallium/drivers/radeonsi/si_shader_llvm_build.c diff --git a/src/gallium/drivers/radeonsi/Makefile.sources b/src/gallium/drivers/radeonsi/Makefile.sources index bc4f9bc2166..68278186278 100644 --- a/src/gallium/drivers/radeonsi/Makefile.sources +++ b/src/gallium/drivers/radeonsi/Makefile.sources @@ -36,7 +36,6 @@ C_SOURCES := \ si_shader.h \ si_shader_internal.h \ si_shader_llvm.c \ - si_shader_llvm_build.c \ si_shader_llvm_gs.c \ si_shader_llvm_ps.c \ si_shader_llvm_resources.c \ diff --git a/src/gallium/drivers/radeonsi/meson.build b/src/gallium/drivers/radeonsi/meson.build index 16e313e37c1..357869eb94c 100644 --- a/src/gallium/drivers/radeonsi/meson.build +++ b/src/gallium/drivers/radeonsi/meson.build @@ -51,7 +51,6 @@ files_libradeonsi = files( 'si_shader.h', 'si_shader_internal.h', 'si_shader_llvm.c', - 'si_shader_llvm_build.c', 'si_shader_llvm_gs.c', 'si_shader_llvm_ps.c', 'si_shader_llvm_resources.c', diff --git a/src/gallium/drivers/radeonsi/si_shader_internal.h b/src/gallium/drivers/radeonsi/si_shader_internal.h index e0f71b4635e..af88bad47ed 100644 --- a/src/gallium/drivers/radeonsi/si_shader_internal.h +++ b/src/gallium/drivers/radeonsi/si_shader_internal.h @@ -196,18 +196,6 @@ si_shader_context_from_abi(struct ac_shader_abi *abi) return container_of(abi, ctx, abi); } -void si_llvm_context_init(struct si_shader_context *ctx, - struct si_screen *sscreen, - struct ac_llvm_compiler *compiler, - unsigned wave_size); -void si_llvm_create_func(struct si_shader_context *ctx, const char *name, - LLVMTypeRef *return_types, unsigned num_return_elems, - unsigned max_workgroup_size); - -void si_llvm_dispose(struct si_shader_context *ctx); - -void si_llvm_optimize_module(struct si_shader_context *ctx); - LLVMValueRef si_nir_load_input_tes(struct ac_shader_abi *abi, LLVMTypeRef type, LLVMValueRef vertex_index, @@ -222,17 +210,6 @@ LLVMValueRef si_nir_load_input_tes(struct ac_shader_abi *abi, bool load_input); bool si_is_merged_shader(struct si_shader_context *ctx); LLVMValueRef si_get_sample_id(struct si_shader_context *ctx); -LLVMValueRef si_buffer_load_const(struct si_shader_context *ctx, - LLVMValueRef resource, LLVMValueRef offset); -void si_llvm_build_ret(struct si_shader_context *ctx, LLVMValueRef ret); -LLVMValueRef si_prolog_get_rw_buffers(struct si_shader_context *ctx); -LLVMValueRef si_build_gather_64bit(struct si_shader_context *ctx, - LLVMTypeRef type, LLVMValueRef val1, - LLVMValueRef val2); -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); void si_declare_compute_memory(struct si_shader_context *ctx); LLVMValueRef si_get_primitive_id(struct si_shader_context *ctx, unsigned swizzle); @@ -255,21 +232,6 @@ void si_get_ps_prolog_key(struct si_shader *shader, bool separate_prolog); void si_get_ps_epilog_key(struct si_shader *shader, union si_shader_part_key *key); -LLVMValueRef si_insert_input_ret(struct si_shader_context *ctx, LLVMValueRef ret, - struct ac_arg param, unsigned return_index); -LLVMValueRef si_insert_input_ret_float(struct si_shader_context *ctx, LLVMValueRef ret, - struct ac_arg param, unsigned return_index); -LLVMValueRef si_insert_input_ptr(struct si_shader_context *ctx, LLVMValueRef ret, - struct ac_arg param, unsigned return_index); -int si_compile_llvm(struct si_screen *sscreen, - struct si_shader_binary *binary, - struct ac_shader_config *conf, - struct ac_llvm_compiler *compiler, - struct ac_llvm_context *ac, - struct pipe_debug_callback *debug, - enum pipe_shader_type shader_type, - const char *name, - bool less_optimized); void si_fix_resource_usage(struct si_screen *sscreen, struct si_shader *shader); void si_create_function(struct si_shader_context *ctx, bool ngg_cull_shader); @@ -291,6 +253,43 @@ void gfx10_ngg_gs_emit_prologue(struct si_shader_context *ctx); void gfx10_ngg_gs_emit_epilogue(struct si_shader_context *ctx); void gfx10_ngg_calculate_subgroup_info(struct si_shader *shader); +/* si_shader_llvm.c */ +int si_compile_llvm(struct si_screen *sscreen, + struct si_shader_binary *binary, + struct ac_shader_config *conf, + struct ac_llvm_compiler *compiler, + struct ac_llvm_context *ac, + struct pipe_debug_callback *debug, + enum pipe_shader_type shader_type, + const char *name, + bool less_optimized); +void si_llvm_context_init(struct si_shader_context *ctx, + struct si_screen *sscreen, + struct ac_llvm_compiler *compiler, + unsigned wave_size); +void si_llvm_create_func(struct si_shader_context *ctx, const char *name, + LLVMTypeRef *return_types, unsigned num_return_elems, + unsigned max_workgroup_size); +void si_llvm_optimize_module(struct si_shader_context *ctx); +void si_llvm_dispose(struct si_shader_context *ctx); +LLVMValueRef si_buffer_load_const(struct si_shader_context *ctx, + LLVMValueRef resource, LLVMValueRef offset); +void si_llvm_build_ret(struct si_shader_context *ctx, LLVMValueRef ret); +LLVMValueRef si_insert_input_ret(struct si_shader_context *ctx, LLVMValueRef ret, + struct ac_arg param, unsigned return_index); +LLVMValueRef si_insert_input_ret_float(struct si_shader_context *ctx, LLVMValueRef ret, + struct ac_arg param, unsigned return_index); +LLVMValueRef si_insert_input_ptr(struct si_shader_context *ctx, LLVMValueRef ret, + struct ac_arg param, unsigned return_index); +LLVMValueRef si_prolog_get_rw_buffers(struct si_shader_context *ctx); +LLVMValueRef si_build_gather_64bit(struct si_shader_context *ctx, + LLVMTypeRef type, LLVMValueRef val1, + LLVMValueRef val2); +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); + /* si_shader_llvm_gs.c */ LLVMValueRef si_is_es_thread(struct si_shader_context *ctx); LLVMValueRef si_is_gs_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 0ea102624e9..47497b96216 100644 --- a/src/gallium/drivers/radeonsi/si_shader_llvm.c +++ b/src/gallium/drivers/radeonsi/si_shader_llvm.c @@ -247,3 +247,119 @@ void si_llvm_dispose(struct si_shader_context *ctx) LLVMContextDispose(ctx->ac.context); ac_llvm_context_dispose(&ctx->ac); } + +/** + * Load a dword from a constant buffer. + */ +LLVMValueRef si_buffer_load_const(struct si_shader_context *ctx, + LLVMValueRef resource, LLVMValueRef offset) +{ + return ac_build_buffer_load(&ctx->ac, resource, 1, NULL, offset, NULL, + 0, 0, true, true); +} + +void si_llvm_build_ret(struct si_shader_context *ctx, LLVMValueRef ret) +{ + if (LLVMGetTypeKind(LLVMTypeOf(ret)) == LLVMVoidTypeKind) + LLVMBuildRetVoid(ctx->ac.builder); + else + LLVMBuildRet(ctx->ac.builder, ret); +} + +LLVMValueRef si_insert_input_ret(struct si_shader_context *ctx, LLVMValueRef ret, + struct ac_arg param, unsigned return_index) +{ + return LLVMBuildInsertValue(ctx->ac.builder, ret, + ac_get_arg(&ctx->ac, param), + return_index, ""); +} + +LLVMValueRef si_insert_input_ret_float(struct si_shader_context *ctx, LLVMValueRef ret, + struct ac_arg param, unsigned return_index) +{ + LLVMBuilderRef builder = ctx->ac.builder; + LLVMValueRef p = ac_get_arg(&ctx->ac, param); + + return LLVMBuildInsertValue(builder, ret, + ac_to_float(&ctx->ac, p), + return_index, ""); +} + +LLVMValueRef si_insert_input_ptr(struct si_shader_context *ctx, LLVMValueRef ret, + struct ac_arg param, unsigned return_index) +{ + LLVMBuilderRef builder = ctx->ac.builder; + LLVMValueRef ptr = ac_get_arg(&ctx->ac, param); + ptr = LLVMBuildPtrToInt(builder, ptr, ctx->ac.i32, ""); + return LLVMBuildInsertValue(builder, ret, ptr, return_index, ""); +} + +LLVMValueRef si_prolog_get_rw_buffers(struct si_shader_context *ctx) +{ + LLVMValueRef ptr[2], list; + bool merged_shader = si_is_merged_shader(ctx); + + ptr[0] = LLVMGetParam(ctx->main_fn, (merged_shader ? 8 : 0) + SI_SGPR_RW_BUFFERS); + list = LLVMBuildIntToPtr(ctx->ac.builder, ptr[0], + ac_array_in_const32_addr_space(ctx->ac.v4i32), ""); + return list; +} + +LLVMValueRef si_build_gather_64bit(struct si_shader_context *ctx, + LLVMTypeRef type, LLVMValueRef val1, + LLVMValueRef val2) +{ + LLVMValueRef values[2] = { + ac_to_integer(&ctx->ac, val1), + ac_to_integer(&ctx->ac, val2), + }; + LLVMValueRef result = ac_build_gather_values(&ctx->ac, values, 2); + return LLVMBuildBitCast(ctx->ac.builder, result, type, ""); +} + +void si_llvm_emit_barrier(struct si_shader_context *ctx) +{ + /* GFX6 only (thanks to a hw bug workaround): + * The real barrier instruction isn’t needed, because an entire patch + * always fits into a single wave. + */ + if (ctx->screen->info.chip_class == GFX6 && + ctx->type == PIPE_SHADER_TESS_CTRL) { + ac_build_waitcnt(&ctx->ac, AC_WAIT_LGKM | AC_WAIT_VLOAD | AC_WAIT_VSTORE); + return; + } + + ac_build_s_barrier(&ctx->ac); +} + +/* Ensure that the esgs ring is declared. + * + * We declare it with 64KB alignment as a hint that the + * pointer value will always be 0. + */ +void si_llvm_declare_esgs_ring(struct si_shader_context *ctx) +{ + if (ctx->esgs_ring) + return; + + assert(!LLVMGetNamedGlobal(ctx->ac.module, "esgs_ring")); + + ctx->esgs_ring = LLVMAddGlobalInAddressSpace( + ctx->ac.module, LLVMArrayType(ctx->ac.i32, 0), + "esgs_ring", + AC_ADDR_SPACE_LDS); + LLVMSetLinkage(ctx->esgs_ring, LLVMExternalLinkage); + LLVMSetAlignment(ctx->esgs_ring, 64 * 1024); +} + +void si_init_exec_from_input(struct si_shader_context *ctx, struct ac_arg param, + unsigned bitoffset) +{ + LLVMValueRef args[] = { + ac_get_arg(&ctx->ac, param), + LLVMConstInt(ctx->ac.i32, bitoffset, 0), + }; + ac_build_intrinsic(&ctx->ac, + "llvm.amdgcn.init.exec.from.input", + ctx->ac.voidt, args, 2, AC_FUNC_ATTR_CONVERGENT); +} diff --git a/src/gallium/drivers/radeonsi/si_shader_llvm_build.c b/src/gallium/drivers/radeonsi/si_shader_llvm_build.c deleted file mode 100644 index 829b9a2fb33..00000000000 --- a/src/gallium/drivers/radeonsi/si_shader_llvm_build.c +++ /dev/null @@ -1,143 +0,0 @@ -/* - * Copyright 2017 Advanced Micro Devices, Inc. - * All Rights Reserved. - * - * 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 - * on the rights to use, copy, modify, merge, publish, distribute, sub - * license, 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 NON-INFRINGEMENT. IN NO EVENT SHALL - * THE AUTHOR(S) AND/OR THEIR SUPPLIERS 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 "si_shader_internal.h" -#include "si_pipe.h" -#include "sid.h" - -/** - * Load a dword from a constant buffer. - */ -LLVMValueRef si_buffer_load_const(struct si_shader_context *ctx, - LLVMValueRef resource, LLVMValueRef offset) -{ - return ac_build_buffer_load(&ctx->ac, resource, 1, NULL, offset, NULL, - 0, 0, true, true); -} - -void si_llvm_build_ret(struct si_shader_context *ctx, LLVMValueRef ret) -{ - if (LLVMGetTypeKind(LLVMTypeOf(ret)) == LLVMVoidTypeKind) - LLVMBuildRetVoid(ctx->ac.builder); - else - LLVMBuildRet(ctx->ac.builder, ret); -} - -LLVMValueRef si_insert_input_ret(struct si_shader_context *ctx, LLVMValueRef ret, - struct ac_arg param, unsigned return_index) -{ - return LLVMBuildInsertValue(ctx->ac.builder, ret, - ac_get_arg(&ctx->ac, param), - return_index, ""); -} - -LLVMValueRef si_insert_input_ret_float(struct si_shader_context *ctx, LLVMValueRef ret, - struct ac_arg param, unsigned return_index) -{ - LLVMBuilderRef builder = ctx->ac.builder; - LLVMValueRef p = ac_get_arg(&ctx->ac, param); - - return LLVMBuildInsertValue(builder, ret, - ac_to_float(&ctx->ac, p), - return_index, ""); -} - -LLVMValueRef si_insert_input_ptr(struct si_shader_context *ctx, LLVMValueRef ret, - struct ac_arg param, unsigned return_index) -{ - LLVMBuilderRef builder = ctx->ac.builder; - LLVMValueRef ptr = ac_get_arg(&ctx->ac, param); - ptr = LLVMBuildPtrToInt(builder, ptr, ctx->ac.i32, ""); - return LLVMBuildInsertValue(builder, ret, ptr, return_index, ""); -} - -LLVMValueRef si_prolog_get_rw_buffers(struct si_shader_context *ctx) -{ - LLVMValueRef ptr[2], list; - bool merged_shader = si_is_merged_shader(ctx); - - ptr[0] = LLVMGetParam(ctx->main_fn, (merged_shader ? 8 : 0) + SI_SGPR_RW_BUFFERS); - list = LLVMBuildIntToPtr(ctx->ac.builder, ptr[0], - ac_array_in_const32_addr_space(ctx->ac.v4i32), ""); - return list; -} - -LLVMValueRef si_build_gather_64bit(struct si_shader_context *ctx, - LLVMTypeRef type, LLVMValueRef val1, - LLVMValueRef val2) -{ - LLVMValueRef values[2] = { - ac_to_integer(&ctx->ac, val1), - ac_to_integer(&ctx->ac, val2), - }; - LLVMValueRef result = ac_build_gather_values(&ctx->ac, values, 2); - return LLVMBuildBitCast(ctx->ac.builder, result, type, ""); -} - -void si_llvm_emit_barrier(struct si_shader_context *ctx) -{ - /* GFX6 only (thanks to a hw bug workaround): - * The real barrier instruction isn’t needed, because an entire patch - * always fits into a single wave. - */ - if (ctx->screen->info.chip_class == GFX6 && - ctx->type == PIPE_SHADER_TESS_CTRL) { - ac_build_waitcnt(&ctx->ac, AC_WAIT_LGKM | AC_WAIT_VLOAD | AC_WAIT_VSTORE); - return; - } - - ac_build_s_barrier(&ctx->ac); -} - -/* Ensure that the esgs ring is declared. - * - * We declare it with 64KB alignment as a hint that the - * pointer value will always be 0. - */ -void si_llvm_declare_esgs_ring(struct si_shader_context *ctx) -{ - if (ctx->esgs_ring) - return; - - assert(!LLVMGetNamedGlobal(ctx->ac.module, "esgs_ring")); - - ctx->esgs_ring = LLVMAddGlobalInAddressSpace( - ctx->ac.module, LLVMArrayType(ctx->ac.i32, 0), - "esgs_ring", - AC_ADDR_SPACE_LDS); - LLVMSetLinkage(ctx->esgs_ring, LLVMExternalLinkage); - LLVMSetAlignment(ctx->esgs_ring, 64 * 1024); -} - -void si_init_exec_from_input(struct si_shader_context *ctx, struct ac_arg param, - unsigned bitoffset) -{ - LLVMValueRef args[] = { - ac_get_arg(&ctx->ac, param), - LLVMConstInt(ctx->ac.i32, bitoffset, 0), - }; - ac_build_intrinsic(&ctx->ac, - "llvm.amdgcn.init.exec.from.input", - ctx->ac.voidt, args, 2, AC_FUNC_ATTR_CONVERGENT); -} -- 2.30.2