fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2026-05-06 05:08:08 +02:00
Code Issues Projects Releases Packages Wiki Activity Actions 1

radeonsi: move si_shader_llvm_build.c content into si_shader_llvm.c

Browse source 
Reviewed-by: Timothy Arceri <tarceri@itsqueeze.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3421>
This commit is contained in:
Marek Olšák 2020-01-15 18:06:02 -05:00 • committed by Marge Bot
parent cd5b99c541
commit ab33ba987a
5 changed files with 153 additions and 183 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

1
src/gallium/drivers/radeonsi/Makefile.sources
View file

@ -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 \

1
src/gallium/drivers/radeonsi/meson.build
View file

@ -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',

75
src/gallium/drivers/radeonsi/si_shader_internal.h
View file

@ -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);

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

@ -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);
}

143
src/gallium/drivers/radeonsi/si_shader_llvm_build.c
View file

@ -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);
}