2 * Copyright 2016 Advanced Micro Devices, Inc.
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:
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
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.
25 #include "si_shader_internal.h"
27 #include "ac_llvm_util.h"
28 #include "util/u_memory.h"
30 struct si_llvm_diagnostics
{
31 struct pipe_debug_callback
*debug
;
35 static void si_diagnostic_handler(LLVMDiagnosticInfoRef di
, void *context
)
37 struct si_llvm_diagnostics
*diag
= (struct si_llvm_diagnostics
*)context
;
38 LLVMDiagnosticSeverity severity
= LLVMGetDiagInfoSeverity(di
);
39 const char *severity_str
= NULL
;
43 severity_str
= "error";
46 severity_str
= "warning";
54 char *description
= LLVMGetDiagInfoDescription(di
);
56 pipe_debug_message(diag
->debug
, SHADER_INFO
,
57 "LLVM diagnostic (%s): %s", severity_str
, description
);
59 if (severity
== LLVMDSError
) {
61 fprintf(stderr
,"LLVM triggered Diagnostic Handler: %s\n", description
);
64 LLVMDisposeMessage(description
);
68 * Compile an LLVM module to machine code.
70 * @returns 0 for success, 1 for failure
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
)
77 struct ac_compiler_passes
*passes
= compiler
->passes
;
80 passes
= compiler
->passes_wave32
;
81 else if (less_optimized
&& compiler
->low_opt_passes
)
82 passes
= compiler
->low_opt_passes
;
84 struct si_llvm_diagnostics diag
;
85 LLVMContextRef llvm_ctx
;
90 /* Setup Diagnostic Handler*/
91 llvm_ctx
= LLVMGetModuleContext(M
);
93 LLVMContextSetDiagnosticHandler(llvm_ctx
, si_diagnostic_handler
, &diag
);
96 if (!ac_compile_module_to_elf(passes
, M
, (char **)&binary
->elf_buffer
,
100 if (diag
.retval
!= 0)
101 pipe_debug_message(debug
, SHADER_INFO
, "LLVM compile failed");
105 void si_shader_binary_clean(struct si_shader_binary
*binary
)
107 free((void *)binary
->elf_buffer
);
108 binary
->elf_buffer
= NULL
;
110 free(binary
->llvm_ir_string
);
111 binary
->llvm_ir_string
= NULL
;
114 void si_llvm_context_init(struct si_shader_context
*ctx
,
115 struct si_screen
*sscreen
,
116 struct ac_llvm_compiler
*compiler
,
119 /* Initialize the gallivm object:
120 * We are only using the module, context, and builder fields of this struct.
121 * This should be enough for us to be able to pass our gallivm struct to the
122 * helper functions in the gallivm module.
124 memset(ctx
, 0, sizeof(*ctx
));
125 ctx
->screen
= sscreen
;
126 ctx
->compiler
= compiler
;
128 ac_llvm_context_init(&ctx
->ac
, compiler
, sscreen
->info
.chip_class
,
129 sscreen
->info
.family
,
130 AC_FLOAT_MODE_NO_SIGNED_ZEROS_FP_MATH
,
133 ctx
->voidt
= LLVMVoidTypeInContext(ctx
->ac
.context
);
134 ctx
->i1
= LLVMInt1TypeInContext(ctx
->ac
.context
);
135 ctx
->i8
= LLVMInt8TypeInContext(ctx
->ac
.context
);
136 ctx
->i32
= LLVMInt32TypeInContext(ctx
->ac
.context
);
137 ctx
->i64
= LLVMInt64TypeInContext(ctx
->ac
.context
);
138 ctx
->i128
= LLVMIntTypeInContext(ctx
->ac
.context
, 128);
139 ctx
->f32
= LLVMFloatTypeInContext(ctx
->ac
.context
);
140 ctx
->v2i32
= LLVMVectorType(ctx
->i32
, 2);
141 ctx
->v4i32
= LLVMVectorType(ctx
->i32
, 4);
142 ctx
->v4f32
= LLVMVectorType(ctx
->f32
, 4);
143 ctx
->v8i32
= LLVMVectorType(ctx
->i32
, 8);
145 ctx
->i32_0
= LLVMConstInt(ctx
->i32
, 0, 0);
146 ctx
->i32_1
= LLVMConstInt(ctx
->i32
, 1, 0);
147 ctx
->i1false
= LLVMConstInt(ctx
->i1
, 0, 0);
148 ctx
->i1true
= LLVMConstInt(ctx
->i1
, 1, 0);
151 /* Set the context to a certain shader. Can be called repeatedly
152 * to change the shader. */
153 void si_llvm_context_set_ir(struct si_shader_context
*ctx
,
154 struct si_shader
*shader
)
156 struct si_shader_selector
*sel
= shader
->selector
;
157 const struct si_shader_info
*info
= &sel
->info
;
159 ctx
->shader
= shader
;
160 ctx
->type
= sel
->type
;
162 ctx
->num_const_buffers
= util_last_bit(info
->const_buffers_declared
);
163 ctx
->num_shader_buffers
= util_last_bit(info
->shader_buffers_declared
);
165 ctx
->num_samplers
= util_last_bit(info
->samplers_declared
);
166 ctx
->num_images
= util_last_bit(info
->images_declared
);
169 void si_llvm_create_func(struct si_shader_context
*ctx
, const char *name
,
170 LLVMTypeRef
*return_types
, unsigned num_return_elems
,
171 unsigned max_workgroup_size
)
173 LLVMTypeRef ret_type
;
174 enum ac_llvm_calling_convention call_conv
;
175 enum pipe_shader_type real_shader_type
;
177 if (num_return_elems
)
178 ret_type
= LLVMStructTypeInContext(ctx
->ac
.context
,
180 num_return_elems
, true);
182 ret_type
= ctx
->voidt
;
184 real_shader_type
= ctx
->type
;
186 /* LS is merged into HS (TCS), and ES is merged into GS. */
187 if (ctx
->screen
->info
.chip_class
>= GFX9
) {
188 if (ctx
->shader
->key
.as_ls
)
189 real_shader_type
= PIPE_SHADER_TESS_CTRL
;
190 else if (ctx
->shader
->key
.as_es
|| ctx
->shader
->key
.as_ngg
)
191 real_shader_type
= PIPE_SHADER_GEOMETRY
;
194 switch (real_shader_type
) {
195 case PIPE_SHADER_VERTEX
:
196 case PIPE_SHADER_TESS_EVAL
:
197 call_conv
= AC_LLVM_AMDGPU_VS
;
199 case PIPE_SHADER_TESS_CTRL
:
200 call_conv
= AC_LLVM_AMDGPU_HS
;
202 case PIPE_SHADER_GEOMETRY
:
203 call_conv
= AC_LLVM_AMDGPU_GS
;
205 case PIPE_SHADER_FRAGMENT
:
206 call_conv
= AC_LLVM_AMDGPU_PS
;
208 case PIPE_SHADER_COMPUTE
:
209 call_conv
= AC_LLVM_AMDGPU_CS
;
212 unreachable("Unhandle shader type");
215 /* Setup the function */
216 ctx
->return_type
= ret_type
;
217 ctx
->main_fn
= ac_build_main(&ctx
->args
, &ctx
->ac
, call_conv
, name
,
218 ret_type
, ctx
->ac
.module
);
219 ctx
->return_value
= LLVMGetUndef(ctx
->return_type
);
221 if (ctx
->screen
->info
.address32_hi
) {
222 ac_llvm_add_target_dep_function_attr(ctx
->main_fn
,
223 "amdgpu-32bit-address-high-bits",
224 ctx
->screen
->info
.address32_hi
);
227 LLVMAddTargetDependentFunctionAttr(ctx
->main_fn
,
228 "no-signed-zeros-fp-math",
231 ac_llvm_set_workgroup_size(ctx
->main_fn
, max_workgroup_size
);
234 void si_llvm_optimize_module(struct si_shader_context
*ctx
)
236 /* Dump LLVM IR before any optimization passes */
237 if (ctx
->screen
->debug_flags
& DBG(PREOPT_IR
) &&
238 si_can_dump_shader(ctx
->screen
, ctx
->type
))
239 LLVMDumpModule(ctx
->ac
.module
);
242 LLVMRunPassManager(ctx
->compiler
->passmgr
, ctx
->ac
.module
);
243 LLVMDisposeBuilder(ctx
->ac
.builder
);
246 void si_llvm_dispose(struct si_shader_context
*ctx
)
248 LLVMDisposeModule(ctx
->ac
.module
);
249 LLVMContextDispose(ctx
->ac
.context
);
250 ac_llvm_context_dispose(&ctx
->ac
);