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"
30 #include "tgsi/tgsi_from_mesa.h"
31 #include "util/u_memory.h"
33 struct si_llvm_diagnostics
{
34 struct pipe_debug_callback
*debug
;
38 static void si_diagnostic_handler(LLVMDiagnosticInfoRef di
, void *context
)
40 struct si_llvm_diagnostics
*diag
= (struct si_llvm_diagnostics
*)context
;
41 LLVMDiagnosticSeverity severity
= LLVMGetDiagInfoSeverity(di
);
42 const char *severity_str
= NULL
;
46 severity_str
= "error";
49 severity_str
= "warning";
57 char *description
= LLVMGetDiagInfoDescription(di
);
59 pipe_debug_message(diag
->debug
, SHADER_INFO
,
60 "LLVM diagnostic (%s): %s", severity_str
, description
);
62 if (severity
== LLVMDSError
) {
64 fprintf(stderr
,"LLVM triggered Diagnostic Handler: %s\n", description
);
67 LLVMDisposeMessage(description
);
70 int si_compile_llvm(struct si_screen
*sscreen
,
71 struct si_shader_binary
*binary
,
72 struct ac_shader_config
*conf
,
73 struct ac_llvm_compiler
*compiler
,
74 struct ac_llvm_context
*ac
,
75 struct pipe_debug_callback
*debug
,
76 enum pipe_shader_type shader_type
,
80 unsigned count
= p_atomic_inc_return(&sscreen
->num_compilations
);
82 if (si_can_dump_shader(sscreen
, shader_type
)) {
83 fprintf(stderr
, "radeonsi: Compiling shader %d\n", count
);
85 if (!(sscreen
->debug_flags
& (DBG(NO_IR
) | DBG(PREOPT_IR
)))) {
86 fprintf(stderr
, "%s LLVM IR:\n\n", name
);
87 ac_dump_module(ac
->module
);
88 fprintf(stderr
, "\n");
92 if (sscreen
->record_llvm_ir
) {
93 char *ir
= LLVMPrintModuleToString(ac
->module
);
94 binary
->llvm_ir_string
= strdup(ir
);
95 LLVMDisposeMessage(ir
);
98 if (!si_replace_shader(count
, binary
)) {
99 struct ac_compiler_passes
*passes
= compiler
->passes
;
101 if (ac
->wave_size
== 32)
102 passes
= compiler
->passes_wave32
;
103 else if (less_optimized
&& compiler
->low_opt_passes
)
104 passes
= compiler
->low_opt_passes
;
106 struct si_llvm_diagnostics diag
= {debug
};
107 LLVMContextSetDiagnosticHandler(ac
->context
, si_diagnostic_handler
, &diag
);
109 if (!ac_compile_module_to_elf(passes
, ac
->module
,
110 (char **)&binary
->elf_buffer
,
114 if (diag
.retval
!= 0) {
115 pipe_debug_message(debug
, SHADER_INFO
, "LLVM compilation failed");
120 struct ac_rtld_binary rtld
;
121 if (!ac_rtld_open(&rtld
, (struct ac_rtld_open_info
){
122 .info
= &sscreen
->info
,
123 .shader_type
= tgsi_processor_to_shader_stage(shader_type
),
124 .wave_size
= ac
->wave_size
,
126 .elf_ptrs
= &binary
->elf_buffer
,
127 .elf_sizes
= &binary
->elf_size
}))
130 bool ok
= ac_rtld_read_config(&rtld
, conf
);
131 ac_rtld_close(&rtld
);
135 /* Enable 64-bit and 16-bit denormals, because there is no performance
138 * If denormals are enabled, all floating-point output modifiers are
141 * Don't enable denormals for 32-bit floats, because:
142 * - Floating-point output modifiers would be ignored by the hw.
143 * - Some opcodes don't support denormals, such as v_mad_f32. We would
144 * have to stop using those.
145 * - GFX6 & GFX7 would be very slow.
147 conf
->float_mode
|= V_00B028_FP_64_DENORMS
;
152 void si_shader_binary_clean(struct si_shader_binary
*binary
)
154 free((void *)binary
->elf_buffer
);
155 binary
->elf_buffer
= NULL
;
157 free(binary
->llvm_ir_string
);
158 binary
->llvm_ir_string
= NULL
;
161 void si_llvm_context_init(struct si_shader_context
*ctx
,
162 struct si_screen
*sscreen
,
163 struct ac_llvm_compiler
*compiler
,
166 memset(ctx
, 0, sizeof(*ctx
));
167 ctx
->screen
= sscreen
;
168 ctx
->compiler
= compiler
;
170 ac_llvm_context_init(&ctx
->ac
, compiler
, sscreen
->info
.chip_class
,
171 sscreen
->info
.family
,
172 AC_FLOAT_MODE_NO_SIGNED_ZEROS_FP_MATH
,
176 /* Set the context to a certain shader. Can be called repeatedly
177 * to change the shader. */
178 void si_llvm_context_set_ir(struct si_shader_context
*ctx
,
179 struct si_shader
*shader
)
181 struct si_shader_selector
*sel
= shader
->selector
;
182 const struct si_shader_info
*info
= &sel
->info
;
184 ctx
->shader
= shader
;
185 ctx
->type
= sel
->type
;
187 ctx
->num_const_buffers
= util_last_bit(info
->const_buffers_declared
);
188 ctx
->num_shader_buffers
= util_last_bit(info
->shader_buffers_declared
);
190 ctx
->num_samplers
= util_last_bit(info
->samplers_declared
);
191 ctx
->num_images
= util_last_bit(info
->images_declared
);
194 void si_llvm_create_func(struct si_shader_context
*ctx
, const char *name
,
195 LLVMTypeRef
*return_types
, unsigned num_return_elems
,
196 unsigned max_workgroup_size
)
198 LLVMTypeRef ret_type
;
199 enum ac_llvm_calling_convention call_conv
;
200 enum pipe_shader_type real_shader_type
;
202 if (num_return_elems
)
203 ret_type
= LLVMStructTypeInContext(ctx
->ac
.context
,
205 num_return_elems
, true);
207 ret_type
= ctx
->ac
.voidt
;
209 real_shader_type
= ctx
->type
;
211 /* LS is merged into HS (TCS), and ES is merged into GS. */
212 if (ctx
->screen
->info
.chip_class
>= GFX9
) {
213 if (ctx
->shader
->key
.as_ls
)
214 real_shader_type
= PIPE_SHADER_TESS_CTRL
;
215 else if (ctx
->shader
->key
.as_es
|| ctx
->shader
->key
.as_ngg
)
216 real_shader_type
= PIPE_SHADER_GEOMETRY
;
219 switch (real_shader_type
) {
220 case PIPE_SHADER_VERTEX
:
221 case PIPE_SHADER_TESS_EVAL
:
222 call_conv
= AC_LLVM_AMDGPU_VS
;
224 case PIPE_SHADER_TESS_CTRL
:
225 call_conv
= AC_LLVM_AMDGPU_HS
;
227 case PIPE_SHADER_GEOMETRY
:
228 call_conv
= AC_LLVM_AMDGPU_GS
;
230 case PIPE_SHADER_FRAGMENT
:
231 call_conv
= AC_LLVM_AMDGPU_PS
;
233 case PIPE_SHADER_COMPUTE
:
234 call_conv
= AC_LLVM_AMDGPU_CS
;
237 unreachable("Unhandle shader type");
240 /* Setup the function */
241 ctx
->return_type
= ret_type
;
242 ctx
->main_fn
= ac_build_main(&ctx
->args
, &ctx
->ac
, call_conv
, name
,
243 ret_type
, ctx
->ac
.module
);
244 ctx
->return_value
= LLVMGetUndef(ctx
->return_type
);
246 if (ctx
->screen
->info
.address32_hi
) {
247 ac_llvm_add_target_dep_function_attr(ctx
->main_fn
,
248 "amdgpu-32bit-address-high-bits",
249 ctx
->screen
->info
.address32_hi
);
252 LLVMAddTargetDependentFunctionAttr(ctx
->main_fn
,
253 "no-signed-zeros-fp-math",
256 ac_llvm_set_workgroup_size(ctx
->main_fn
, max_workgroup_size
);
259 void si_llvm_optimize_module(struct si_shader_context
*ctx
)
261 /* Dump LLVM IR before any optimization passes */
262 if (ctx
->screen
->debug_flags
& DBG(PREOPT_IR
) &&
263 si_can_dump_shader(ctx
->screen
, ctx
->type
))
264 LLVMDumpModule(ctx
->ac
.module
);
267 LLVMRunPassManager(ctx
->compiler
->passmgr
, ctx
->ac
.module
);
268 LLVMDisposeBuilder(ctx
->ac
.builder
);
271 void si_llvm_dispose(struct si_shader_context
*ctx
)
273 LLVMDisposeModule(ctx
->ac
.module
);
274 LLVMContextDispose(ctx
->ac
.context
);
275 ac_llvm_context_dispose(&ctx
->ac
);