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 \
'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',
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,
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);
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);
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);
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);
+}
+++ /dev/null
-/*
- * 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);
-}
projects / mesa.git / commitdiff