src/gallium/drivers/radeonsi/si_shader_llvm.c

   1 /*
   2  * Copyright 2016 Advanced Micro Devices, Inc.
   3  * All Rights Reserved.
   4  *
   5  * Permission is hereby granted, free of charge, to any person obtaining a
   6  * copy of this software and associated documentation files (the "Software"),
   7  * to deal in the Software without restriction, including without limitation
   8  * on the rights to use, copy, modify, merge, publish, distribute, sub
   9  * license, and/or sell copies of the Software, and to permit persons to whom
  10  * the Software is furnished to do so, subject to the following conditions:
  11  *
  12  * The above copyright notice and this permission notice (including the next
  13  * paragraph) shall be included in all copies or substantial portions of the
  14  * Software.
  15  *
  16  * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
  17  * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
  18  * FITNESS FOR A PARTICULAR PURPOSE AND NON-INFRINGEMENT. IN NO EVENT SHALL
  19  * THE AUTHOR(S) AND/OR THEIR SUPPLIERS BE LIABLE FOR ANY CLAIM,
  20  * DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR
  21  * OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE
  22  * USE OR OTHER DEALINGS IN THE SOFTWARE.
  23  */
  24
  25 #include "si_shader_internal.h"
  26 #include "si_pipe.h"
  27 #include "ac_llvm_util.h"
  28 #include "util/u_memory.h"
  29
  30 struct si_llvm_diagnostics {
  31         struct pipe_debug_callback *debug;
  32         unsigned retval;
  33 };
  34
  35 static void si_diagnostic_handler(LLVMDiagnosticInfoRef di, void *context)
  36 {
  37         struct si_llvm_diagnostics *diag = (struct si_llvm_diagnostics *)context;
  38         LLVMDiagnosticSeverity severity = LLVMGetDiagInfoSeverity(di);
  39         const char *severity_str = NULL;
  40
  41         switch (severity) {
  42         case LLVMDSError:
  43                 severity_str = "error";
  44                 break;
  45         case LLVMDSWarning:
  46                 severity_str = "warning";
  47                 break;
  48         case LLVMDSRemark:
  49         case LLVMDSNote:
  50         default:
  51                 return;
  52         }
  53
  54         char *description = LLVMGetDiagInfoDescription(di);
  55
  56         pipe_debug_message(diag->debug, SHADER_INFO,
  57                            "LLVM diagnostic (%s): %s", severity_str, description);
  58
  59         if (severity == LLVMDSError) {
  60                 diag->retval = 1;
  61                 fprintf(stderr,"LLVM triggered Diagnostic Handler: %s\n", description);
  62         }
  63
  64         LLVMDisposeMessage(description);
  65 }
  66
  67 /**
  68  * Compile an LLVM module to machine code.
  69  *
  70  * @returns 0 for success, 1 for failure
  71  */
  72 unsigned si_llvm_compile(LLVMModuleRef M, struct si_shader_binary *binary,
  73                          struct ac_llvm_compiler *compiler,
  74                          struct pipe_debug_callback *debug,
  75                          bool less_optimized, unsigned wave_size)
  76 {
  77         struct ac_compiler_passes *passes = compiler->passes;
  78
  79         if (wave_size == 32)
  80                 passes = compiler->passes_wave32;
  81         else if (less_optimized && compiler->low_opt_passes)
  82                 passes = compiler->low_opt_passes;
  83
  84         struct si_llvm_diagnostics diag;
  85         LLVMContextRef llvm_ctx;
  86
  87         diag.debug = debug;
  88         diag.retval = 0;
  89
  90         /* Setup Diagnostic Handler*/
  91         llvm_ctx = LLVMGetModuleContext(M);
  92
  93         LLVMContextSetDiagnosticHandler(llvm_ctx, si_diagnostic_handler, &diag);
  94
  95         /* Compile IR. */
  96         if (!ac_compile_module_to_elf(passes, M, (char **)&binary->elf_buffer,
  97                                       &binary->elf_size))
  98                 diag.retval = 1;
  99
 100         if (diag.retval != 0)
 101                 pipe_debug_message(debug, SHADER_INFO, "LLVM compile failed");
 102         return diag.retval;
 103 }
 104
 105 void si_shader_binary_clean(struct si_shader_binary *binary)
 106 {
 107         free((void *)binary->elf_buffer);
 108         binary->elf_buffer = NULL;
 109
 110         free(binary->llvm_ir_string);
 111         binary->llvm_ir_string = NULL;
 112 }
 113
 114 void si_llvm_context_init(struct si_shader_context *ctx,
 115                           struct si_screen *sscreen,
 116                           struct ac_llvm_compiler *compiler,
 117                           unsigned wave_size,
 118                           unsigned ballot_mask_bits)
 119 {
 120         /* Initialize the gallivm object:
 121          * We are only using the module, context, and builder fields of this struct.
 122          * This should be enough for us to be able to pass our gallivm struct to the
 123          * helper functions in the gallivm module.
 124          */
 125         memset(ctx, 0, sizeof(*ctx));
 126         ctx->screen = sscreen;
 127         ctx->compiler = compiler;
 128
 129         ac_llvm_context_init(&ctx->ac, compiler, sscreen->info.chip_class,
 130                              sscreen->info.family,
 131                              AC_FLOAT_MODE_NO_SIGNED_ZEROS_FP_MATH,
 132                              wave_size, ballot_mask_bits);
 133
 134         ctx->voidt = LLVMVoidTypeInContext(ctx->ac.context);
 135         ctx->i1 = LLVMInt1TypeInContext(ctx->ac.context);
 136         ctx->i8 = LLVMInt8TypeInContext(ctx->ac.context);
 137         ctx->i32 = LLVMInt32TypeInContext(ctx->ac.context);
 138         ctx->i64 = LLVMInt64TypeInContext(ctx->ac.context);
 139         ctx->i128 = LLVMIntTypeInContext(ctx->ac.context, 128);
 140         ctx->f32 = LLVMFloatTypeInContext(ctx->ac.context);
 141         ctx->v2i32 = LLVMVectorType(ctx->i32, 2);
 142         ctx->v4i32 = LLVMVectorType(ctx->i32, 4);
 143         ctx->v4f32 = LLVMVectorType(ctx->f32, 4);
 144         ctx->v8i32 = LLVMVectorType(ctx->i32, 8);
 145
 146         ctx->i32_0 = LLVMConstInt(ctx->i32, 0, 0);
 147         ctx->i32_1 = LLVMConstInt(ctx->i32, 1, 0);
 148         ctx->i1false = LLVMConstInt(ctx->i1, 0, 0);
 149         ctx->i1true = LLVMConstInt(ctx->i1, 1, 0);
 150 }
 151
 152 /* Set the context to a certain shader. Can be called repeatedly
 153  * to change the shader. */
 154 void si_llvm_context_set_ir(struct si_shader_context *ctx,
 155                             struct si_shader *shader)
 156 {
 157         struct si_shader_selector *sel = shader->selector;
 158         const struct si_shader_info *info = &sel->info;
 159
 160         ctx->shader = shader;
 161         ctx->type = sel->type;
 162
 163         ctx->num_const_buffers = util_last_bit(info->const_buffers_declared);
 164         ctx->num_shader_buffers = util_last_bit(info->shader_buffers_declared);
 165
 166         ctx->num_samplers = util_last_bit(info->samplers_declared);
 167         ctx->num_images = util_last_bit(info->images_declared);
 168 }
 169
 170 void si_llvm_create_func(struct si_shader_context *ctx, const char *name,
 171                          LLVMTypeRef *return_types, unsigned num_return_elems,
 172                          unsigned max_workgroup_size)
 173 {
 174         LLVMTypeRef ret_type;
 175         enum ac_llvm_calling_convention call_conv;
 176         enum pipe_shader_type real_shader_type;
 177
 178         if (num_return_elems)
 179                 ret_type = LLVMStructTypeInContext(ctx->ac.context,
 180                                                    return_types,
 181                                                    num_return_elems, true);
 182         else
 183                 ret_type = ctx->voidt;
 184
 185         real_shader_type = ctx->type;
 186
 187         /* LS is merged into HS (TCS), and ES is merged into GS. */
 188         if (ctx->screen->info.chip_class >= GFX9) {
 189                 if (ctx->shader->key.as_ls)
 190                         real_shader_type = PIPE_SHADER_TESS_CTRL;
 191                 else if (ctx->shader->key.as_es || ctx->shader->key.as_ngg)
 192                         real_shader_type = PIPE_SHADER_GEOMETRY;
 193         }
 194
 195         switch (real_shader_type) {
 196         case PIPE_SHADER_VERTEX:
 197         case PIPE_SHADER_TESS_EVAL:
 198                 call_conv = AC_LLVM_AMDGPU_VS;
 199                 break;
 200         case PIPE_SHADER_TESS_CTRL:
 201                 call_conv = AC_LLVM_AMDGPU_HS;
 202                 break;
 203         case PIPE_SHADER_GEOMETRY:
 204                 call_conv = AC_LLVM_AMDGPU_GS;
 205                 break;
 206         case PIPE_SHADER_FRAGMENT:
 207                 call_conv = AC_LLVM_AMDGPU_PS;
 208                 break;
 209         case PIPE_SHADER_COMPUTE:
 210                 call_conv = AC_LLVM_AMDGPU_CS;
 211                 break;
 212         default:
 213                 unreachable("Unhandle shader type");
 214         }
 215
 216         /* Setup the function */
 217         ctx->return_type = ret_type;
 218         ctx->main_fn = ac_build_main(&ctx->args, &ctx->ac, call_conv, name,
 219                                      ret_type, ctx->ac.module);
 220         ctx->return_value = LLVMGetUndef(ctx->return_type);
 221
 222         if (ctx->screen->info.address32_hi) {
 223                 ac_llvm_add_target_dep_function_attr(ctx->main_fn,
 224                                                      "amdgpu-32bit-address-high-bits",
 225                                                      ctx->screen->info.address32_hi);
 226         }
 227
 228         LLVMAddTargetDependentFunctionAttr(ctx->main_fn,
 229                                            "no-signed-zeros-fp-math",
 230                                            "true");
 231
 232         ac_llvm_set_workgroup_size(ctx->main_fn, max_workgroup_size);
 233 }
 234
 235 void si_llvm_optimize_module(struct si_shader_context *ctx)
 236 {
 237         /* Dump LLVM IR before any optimization passes */
 238         if (ctx->screen->debug_flags & DBG(PREOPT_IR) &&
 239             si_can_dump_shader(ctx->screen, ctx->type))
 240                 LLVMDumpModule(ctx->ac.module);
 241
 242         /* Run the pass */
 243         LLVMRunPassManager(ctx->compiler->passmgr, ctx->ac.module);
 244         LLVMDisposeBuilder(ctx->ac.builder);
 245 }
 246
 247 void si_llvm_dispose(struct si_shader_context *ctx)
 248 {
 249         LLVMDisposeModule(ctx->ac.module);
 250         LLVMContextDispose(ctx->ac.context);
 251         ac_llvm_context_dispose(&ctx->ac);
 252 }