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 "util/u_memory.h"
29 struct si_llvm_diagnostics
{
30 struct pipe_debug_callback
*debug
;
34 static void si_diagnostic_handler(LLVMDiagnosticInfoRef di
, void *context
)
36 struct si_llvm_diagnostics
*diag
= (struct si_llvm_diagnostics
*)context
;
37 LLVMDiagnosticSeverity severity
= LLVMGetDiagInfoSeverity(di
);
38 const char *severity_str
= NULL
;
42 severity_str
= "error";
45 severity_str
= "warning";
53 char *description
= LLVMGetDiagInfoDescription(di
);
55 pipe_debug_message(diag
->debug
, SHADER_INFO
,
56 "LLVM diagnostic (%s): %s", severity_str
, description
);
58 if (severity
== LLVMDSError
) {
60 fprintf(stderr
,"LLVM triggered Diagnostic Handler: %s\n", description
);
63 LLVMDisposeMessage(description
);
67 * Compile an LLVM module to machine code.
69 * @returns 0 for success, 1 for failure
71 unsigned si_llvm_compile(LLVMModuleRef M
, struct si_shader_binary
*binary
,
72 struct ac_llvm_compiler
*compiler
,
73 struct pipe_debug_callback
*debug
,
74 bool less_optimized
, unsigned wave_size
)
76 struct ac_compiler_passes
*passes
= compiler
->passes
;
79 passes
= compiler
->passes_wave32
;
80 else if (less_optimized
&& compiler
->low_opt_passes
)
81 passes
= compiler
->low_opt_passes
;
83 struct si_llvm_diagnostics diag
;
84 LLVMContextRef llvm_ctx
;
89 /* Setup Diagnostic Handler*/
90 llvm_ctx
= LLVMGetModuleContext(M
);
92 LLVMContextSetDiagnosticHandler(llvm_ctx
, si_diagnostic_handler
, &diag
);
95 if (!ac_compile_module_to_elf(passes
, M
, (char **)&binary
->elf_buffer
,
100 pipe_debug_message(debug
, SHADER_INFO
, "LLVM compile failed");
104 void si_shader_binary_clean(struct si_shader_binary
*binary
)
106 free((void *)binary
->elf_buffer
);
107 binary
->elf_buffer
= NULL
;
109 free(binary
->llvm_ir_string
);
110 binary
->llvm_ir_string
= NULL
;
113 void si_llvm_context_init(struct si_shader_context
*ctx
,
114 struct si_screen
*sscreen
,
115 struct ac_llvm_compiler
*compiler
,
118 /* Initialize the gallivm object:
119 * We are only using the module, context, and builder fields of this struct.
120 * This should be enough for us to be able to pass our gallivm struct to the
121 * helper functions in the gallivm module.
123 memset(ctx
, 0, sizeof(*ctx
));
124 ctx
->screen
= sscreen
;
125 ctx
->compiler
= compiler
;
127 ac_llvm_context_init(&ctx
->ac
, compiler
, sscreen
->info
.chip_class
,
128 sscreen
->info
.family
,
129 AC_FLOAT_MODE_NO_SIGNED_ZEROS_FP_MATH
,
132 ctx
->voidt
= LLVMVoidTypeInContext(ctx
->ac
.context
);
133 ctx
->i1
= LLVMInt1TypeInContext(ctx
->ac
.context
);
134 ctx
->i8
= LLVMInt8TypeInContext(ctx
->ac
.context
);
135 ctx
->i32
= LLVMInt32TypeInContext(ctx
->ac
.context
);
136 ctx
->i64
= LLVMInt64TypeInContext(ctx
->ac
.context
);
137 ctx
->i128
= LLVMIntTypeInContext(ctx
->ac
.context
, 128);
138 ctx
->f32
= LLVMFloatTypeInContext(ctx
->ac
.context
);
139 ctx
->v2i32
= LLVMVectorType(ctx
->i32
, 2);
140 ctx
->v4i32
= LLVMVectorType(ctx
->i32
, 4);
141 ctx
->v4f32
= LLVMVectorType(ctx
->f32
, 4);
142 ctx
->v8i32
= LLVMVectorType(ctx
->i32
, 8);
144 ctx
->i32_0
= LLVMConstInt(ctx
->i32
, 0, 0);
145 ctx
->i32_1
= LLVMConstInt(ctx
->i32
, 1, 0);
146 ctx
->i1false
= LLVMConstInt(ctx
->i1
, 0, 0);
147 ctx
->i1true
= LLVMConstInt(ctx
->i1
, 1, 0);
150 /* Set the context to a certain shader. Can be called repeatedly
151 * to change the shader. */
152 void si_llvm_context_set_ir(struct si_shader_context
*ctx
,
153 struct si_shader
*shader
)
155 struct si_shader_selector
*sel
= shader
->selector
;
156 const struct si_shader_info
*info
= &sel
->info
;
158 ctx
->shader
= shader
;
159 ctx
->type
= sel
->type
;
161 ctx
->num_const_buffers
= util_last_bit(info
->const_buffers_declared
);
162 ctx
->num_shader_buffers
= util_last_bit(info
->shader_buffers_declared
);
164 ctx
->num_samplers
= util_last_bit(info
->samplers_declared
);
165 ctx
->num_images
= util_last_bit(info
->images_declared
);
168 void si_llvm_create_func(struct si_shader_context
*ctx
, const char *name
,
169 LLVMTypeRef
*return_types
, unsigned num_return_elems
,
170 unsigned max_workgroup_size
)
172 LLVMTypeRef ret_type
;
173 enum ac_llvm_calling_convention call_conv
;
174 enum pipe_shader_type real_shader_type
;
176 if (num_return_elems
)
177 ret_type
= LLVMStructTypeInContext(ctx
->ac
.context
,
179 num_return_elems
, true);
181 ret_type
= ctx
->voidt
;
183 real_shader_type
= ctx
->type
;
185 /* LS is merged into HS (TCS), and ES is merged into GS. */
186 if (ctx
->screen
->info
.chip_class
>= GFX9
) {
187 if (ctx
->shader
->key
.as_ls
)
188 real_shader_type
= PIPE_SHADER_TESS_CTRL
;
189 else if (ctx
->shader
->key
.as_es
|| ctx
->shader
->key
.as_ngg
)
190 real_shader_type
= PIPE_SHADER_GEOMETRY
;
193 switch (real_shader_type
) {
194 case PIPE_SHADER_VERTEX
:
195 case PIPE_SHADER_TESS_EVAL
:
196 call_conv
= AC_LLVM_AMDGPU_VS
;
198 case PIPE_SHADER_TESS_CTRL
:
199 call_conv
= AC_LLVM_AMDGPU_HS
;
201 case PIPE_SHADER_GEOMETRY
:
202 call_conv
= AC_LLVM_AMDGPU_GS
;
204 case PIPE_SHADER_FRAGMENT
:
205 call_conv
= AC_LLVM_AMDGPU_PS
;
207 case PIPE_SHADER_COMPUTE
:
208 call_conv
= AC_LLVM_AMDGPU_CS
;
211 unreachable("Unhandle shader type");
214 /* Setup the function */
215 ctx
->return_type
= ret_type
;
216 ctx
->main_fn
= ac_build_main(&ctx
->args
, &ctx
->ac
, call_conv
, name
,
217 ret_type
, ctx
->ac
.module
);
218 ctx
->return_value
= LLVMGetUndef(ctx
->return_type
);
220 if (ctx
->screen
->info
.address32_hi
) {
221 ac_llvm_add_target_dep_function_attr(ctx
->main_fn
,
222 "amdgpu-32bit-address-high-bits",
223 ctx
->screen
->info
.address32_hi
);
226 LLVMAddTargetDependentFunctionAttr(ctx
->main_fn
,
227 "no-signed-zeros-fp-math",
230 ac_llvm_set_workgroup_size(ctx
->main_fn
, max_workgroup_size
);
233 void si_llvm_optimize_module(struct si_shader_context
*ctx
)
235 /* Dump LLVM IR before any optimization passes */
236 if (ctx
->screen
->debug_flags
& DBG(PREOPT_IR
) &&
237 si_can_dump_shader(ctx
->screen
, ctx
->type
))
238 LLVMDumpModule(ctx
->ac
.module
);
241 LLVMRunPassManager(ctx
->compiler
->passmgr
, ctx
->ac
.module
);
242 LLVMDisposeBuilder(ctx
->ac
.builder
);
245 void si_llvm_dispose(struct si_shader_context
*ctx
)
247 LLVMDisposeModule(ctx
->ac
.module
);
248 LLVMContextDispose(ctx
->ac
.context
);
249 ac_llvm_context_dispose(&ctx
->ac
);