#include "si_shader_internal.h"
#include "si_pipe.h"
-#include "ac_llvm_util.h"
+#include "ac_rtld.h"
+#include "sid.h"
+
+#include "tgsi/tgsi_from_mesa.h"
#include "util/u_memory.h"
struct si_llvm_diagnostics {
LLVMDisposeMessage(description);
}
-/**
- * Compile an LLVM module to machine code.
- *
- * @returns 0 for success, 1 for failure
- */
-unsigned si_llvm_compile(LLVMModuleRef M, struct si_shader_binary *binary,
- struct ac_llvm_compiler *compiler,
- struct pipe_debug_callback *debug,
- bool less_optimized, unsigned wave_size)
+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)
{
- struct ac_compiler_passes *passes = compiler->passes;
+ unsigned count = p_atomic_inc_return(&sscreen->num_compilations);
- if (wave_size == 32)
- passes = compiler->passes_wave32;
- else if (less_optimized && compiler->low_opt_passes)
- passes = compiler->low_opt_passes;
+ if (si_can_dump_shader(sscreen, shader_type)) {
+ fprintf(stderr, "radeonsi: Compiling shader %d\n", count);
- struct si_llvm_diagnostics diag;
- LLVMContextRef llvm_ctx;
+ if (!(sscreen->debug_flags & (DBG(NO_IR) | DBG(PREOPT_IR)))) {
+ fprintf(stderr, "%s LLVM IR:\n\n", name);
+ ac_dump_module(ac->module);
+ fprintf(stderr, "\n");
+ }
+ }
- diag.debug = debug;
- diag.retval = 0;
+ if (sscreen->record_llvm_ir) {
+ char *ir = LLVMPrintModuleToString(ac->module);
+ binary->llvm_ir_string = strdup(ir);
+ LLVMDisposeMessage(ir);
+ }
- /* Setup Diagnostic Handler*/
- llvm_ctx = LLVMGetModuleContext(M);
+ if (!si_replace_shader(count, binary)) {
+ struct ac_compiler_passes *passes = compiler->passes;
- LLVMContextSetDiagnosticHandler(llvm_ctx, si_diagnostic_handler, &diag);
+ if (ac->wave_size == 32)
+ passes = compiler->passes_wave32;
+ else if (less_optimized && compiler->low_opt_passes)
+ passes = compiler->low_opt_passes;
- /* Compile IR. */
- if (!ac_compile_module_to_elf(passes, M, (char **)&binary->elf_buffer,
- &binary->elf_size))
- diag.retval = 1;
+ struct si_llvm_diagnostics diag = {debug};
+ LLVMContextSetDiagnosticHandler(ac->context, si_diagnostic_handler, &diag);
- if (diag.retval != 0)
- pipe_debug_message(debug, SHADER_INFO, "LLVM compile failed");
- return diag.retval;
-}
+ if (!ac_compile_module_to_elf(passes, ac->module,
+ (char **)&binary->elf_buffer,
+ &binary->elf_size))
+ diag.retval = 1;
-void si_shader_binary_clean(struct si_shader_binary *binary)
-{
- free((void *)binary->elf_buffer);
- binary->elf_buffer = NULL;
+ if (diag.retval != 0) {
+ pipe_debug_message(debug, SHADER_INFO, "LLVM compilation failed");
+ return diag.retval;
+ }
+ }
- free(binary->llvm_ir_string);
- binary->llvm_ir_string = NULL;
+ struct ac_rtld_binary rtld;
+ if (!ac_rtld_open(&rtld, (struct ac_rtld_open_info){
+ .info = &sscreen->info,
+ .shader_type = tgsi_processor_to_shader_stage(shader_type),
+ .wave_size = ac->wave_size,
+ .num_parts = 1,
+ .elf_ptrs = &binary->elf_buffer,
+ .elf_sizes = &binary->elf_size }))
+ return -1;
+
+ bool ok = ac_rtld_read_config(&rtld, conf);
+ ac_rtld_close(&rtld);
+ if (!ok)
+ return -1;
+
+ /* Enable 64-bit and 16-bit denormals, because there is no performance
+ * cost.
+ *
+ * If denormals are enabled, all floating-point output modifiers are
+ * ignored.
+ *
+ * Don't enable denormals for 32-bit floats, because:
+ * - Floating-point output modifiers would be ignored by the hw.
+ * - Some opcodes don't support denormals, such as v_mad_f32. We would
+ * have to stop using those.
+ * - GFX6 & GFX7 would be very slow.
+ */
+ conf->float_mode |= V_00B028_FP_64_DENORMS;
+
+ return 0;
}
void si_llvm_context_init(struct si_shader_context *ctx,
struct si_screen *sscreen,
struct ac_llvm_compiler *compiler,
- unsigned wave_size,
- unsigned ballot_mask_bits)
+ unsigned wave_size)
{
- /* Initialize the gallivm object:
- * We are only using the module, context, and builder fields of this struct.
- * This should be enough for us to be able to pass our gallivm struct to the
- * helper functions in the gallivm module.
- */
memset(ctx, 0, sizeof(*ctx));
ctx->screen = sscreen;
ctx->compiler = compiler;
ac_llvm_context_init(&ctx->ac, compiler, sscreen->info.chip_class,
sscreen->info.family,
AC_FLOAT_MODE_NO_SIGNED_ZEROS_FP_MATH,
- wave_size, ballot_mask_bits);
-
- ctx->voidt = LLVMVoidTypeInContext(ctx->ac.context);
- ctx->i1 = LLVMInt1TypeInContext(ctx->ac.context);
- ctx->i8 = LLVMInt8TypeInContext(ctx->ac.context);
- ctx->i32 = LLVMInt32TypeInContext(ctx->ac.context);
- ctx->i64 = LLVMInt64TypeInContext(ctx->ac.context);
- ctx->i128 = LLVMIntTypeInContext(ctx->ac.context, 128);
- ctx->f32 = LLVMFloatTypeInContext(ctx->ac.context);
- ctx->v2i32 = LLVMVectorType(ctx->i32, 2);
- ctx->v4i32 = LLVMVectorType(ctx->i32, 4);
- ctx->v4f32 = LLVMVectorType(ctx->f32, 4);
- ctx->v8i32 = LLVMVectorType(ctx->i32, 8);
-
- ctx->i32_0 = LLVMConstInt(ctx->i32, 0, 0);
- ctx->i32_1 = LLVMConstInt(ctx->i32, 1, 0);
- ctx->i1false = LLVMConstInt(ctx->i1, 0, 0);
- ctx->i1true = LLVMConstInt(ctx->i1, 1, 0);
+ wave_size, 64);
}
-/* Set the context to a certain TGSI shader. Can be called repeatedly
- * to change the shader. */
-void si_llvm_context_set_ir(struct si_shader_context *ctx,
- struct si_shader *shader)
-{
- struct si_shader_selector *sel = shader->selector;
- const struct tgsi_shader_info *info = &sel->info;
-
- ctx->shader = shader;
- ctx->type = sel->type;
-
- ctx->num_const_buffers = util_last_bit(info->const_buffers_declared);
- ctx->num_shader_buffers = util_last_bit(info->shader_buffers_declared);
-
- ctx->num_samplers = util_last_bit(info->samplers_declared);
- ctx->num_images = util_last_bit(info->images_declared);
-}
-
-void si_llvm_create_func(struct si_shader_context *ctx,
- const char *name,
- LLVMTypeRef *return_types, unsigned num_return_elems)
+void si_llvm_create_func(struct si_shader_context *ctx, const char *name,
+ LLVMTypeRef *return_types, unsigned num_return_elems,
+ unsigned max_workgroup_size)
{
LLVMTypeRef ret_type;
enum ac_llvm_calling_convention call_conv;
return_types,
num_return_elems, true);
else
- ret_type = ctx->voidt;
+ ret_type = ctx->ac.voidt;
real_shader_type = ctx->type;
ctx->return_type = ret_type;
ctx->main_fn = ac_build_main(&ctx->args, &ctx->ac, call_conv, name,
ret_type, ctx->ac.module);
+ ctx->return_value = LLVMGetUndef(ctx->return_type);
+
+ if (ctx->screen->info.address32_hi) {
+ ac_llvm_add_target_dep_function_attr(ctx->main_fn,
+ "amdgpu-32bit-address-high-bits",
+ ctx->screen->info.address32_hi);
+ }
+
+ LLVMAddTargetDependentFunctionAttr(ctx->main_fn,
+ "no-signed-zeros-fp-math",
+ "true");
+
+ ac_llvm_set_workgroup_size(ctx->main_fn, max_workgroup_size);
}
void si_llvm_optimize_module(struct si_shader_context *ctx)
projects / mesa.git / blobdiff