2 * Copyright 2014 Advanced Micro Devices, Inc.
4 * Permission is hereby granted, free of charge, to any person obtaining a
5 * copy of this software and associated documentation files (the
6 * "Software"), to deal in the Software without restriction, including
7 * without limitation the rights to use, copy, modify, merge, publish,
8 * distribute, sub license, and/or sell copies of the Software, and to
9 * permit persons to whom the Software is furnished to do so, subject to
10 * the following conditions:
12 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
13 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
14 * FITNESS FOR A PARTICULAR PURPOSE AND NON-INFRINGEMENT. IN NO EVENT SHALL
15 * THE COPYRIGHT HOLDERS, AUTHORS AND/OR ITS SUPPLIERS BE LIABLE FOR ANY CLAIM,
16 * DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR
17 * OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE
18 * USE OR OTHER DEALINGS IN THE SOFTWARE.
20 * The above copyright notice and this permission notice (including the
21 * next paragraph) shall be included in all copies or substantial portions
25 /* based on pieces from si_pipe.c and radeon_llvm_emit.c */
26 #include "ac_llvm_util.h"
28 #include "ac_llvm_build.h"
29 #include "c11/threads.h"
30 #include "gallivm/lp_bld_misc.h"
31 #include "util/bitscan.h"
32 #include "util/u_math.h"
33 #include <llvm-c/Core.h>
34 #include <llvm-c/Support.h>
35 #include <llvm-c/Transforms/IPO.h>
36 #include <llvm-c/Transforms/Scalar.h>
37 #include <llvm-c/Transforms/Utils.h>
43 static void ac_init_llvm_target()
45 LLVMInitializeAMDGPUTargetInfo();
46 LLVMInitializeAMDGPUTarget();
47 LLVMInitializeAMDGPUTargetMC();
48 LLVMInitializeAMDGPUAsmPrinter();
50 /* For inline assembly. */
51 LLVMInitializeAMDGPUAsmParser();
53 /* For ACO disassembly. */
54 LLVMInitializeAMDGPUDisassembler();
56 /* Workaround for bug in llvm 4.0 that causes image intrinsics
58 * https://reviews.llvm.org/D26348
60 * "mesa" is the prefix for error messages.
62 * -global-isel-abort=2 is a no-op unless global isel has been enabled.
63 * This option tells the backend to fall-back to SelectionDAG and print
64 * a diagnostic message if global isel fails.
66 const char *argv
[] = {
68 "-simplifycfg-sink-common=false",
69 "-global-isel-abort=2",
70 #if LLVM_VERSION_MAJOR >= 10
71 /* Atomic optimizations require LLVM 10.0 for gfx10 support. */
72 "-amdgpu-atomic-optimizations=true",
74 #if LLVM_VERSION_MAJOR >= 11
75 /* This was disabled by default in: https://reviews.llvm.org/D77228 */
76 "-structurizecfg-skip-uniform-regions",
79 LLVMParseCommandLineOptions(ARRAY_SIZE(argv
), argv
, NULL
);
82 PUBLIC
void ac_init_shared_llvm_once(void)
84 static once_flag ac_init_llvm_target_once_flag
= ONCE_FLAG_INIT
;
85 call_once(&ac_init_llvm_target_once_flag
, ac_init_llvm_target
);
89 static once_flag ac_init_static_llvm_target_once_flag
= ONCE_FLAG_INIT
;
90 static void ac_init_static_llvm_once(void)
92 call_once(&ac_init_static_llvm_target_once_flag
, ac_init_llvm_target
);
96 void ac_init_llvm_once(void)
99 ac_init_shared_llvm_once();
101 ac_init_static_llvm_once();
105 static LLVMTargetRef
ac_get_llvm_target(const char *triple
)
107 LLVMTargetRef target
= NULL
;
108 char *err_message
= NULL
;
110 if (LLVMGetTargetFromTriple(triple
, &target
, &err_message
)) {
111 fprintf(stderr
, "Cannot find target for triple %s ", triple
);
113 fprintf(stderr
, "%s\n", err_message
);
115 LLVMDisposeMessage(err_message
);
121 const char *ac_get_llvm_processor_name(enum radeon_family family
)
177 case CHIP_SIENNA_CICHLID
:
178 case CHIP_NAVY_FLOUNDER
:
185 static LLVMTargetMachineRef
ac_create_target_machine(enum radeon_family family
,
186 enum ac_target_machine_options tm_options
,
187 LLVMCodeGenOptLevel level
,
188 const char **out_triple
)
190 assert(family
>= CHIP_TAHITI
);
192 const char *triple
= (tm_options
& AC_TM_SUPPORTS_SPILL
) ? "amdgcn-mesa-mesa3d" : "amdgcn--";
193 LLVMTargetRef target
= ac_get_llvm_target(triple
);
195 snprintf(features
, sizeof(features
), "+DumpCode%s%s%s%s%s",
196 LLVM_VERSION_MAJOR
>= 11 ? "" : ",-fp32-denormals,+fp64-denormals",
197 family
>= CHIP_NAVI10
&& !(tm_options
& AC_TM_WAVE32
)
198 ? ",+wavefrontsize64,-wavefrontsize32"
200 family
<= CHIP_NAVI14
&& tm_options
& AC_TM_FORCE_ENABLE_XNACK
? ",+xnack" : "",
201 family
<= CHIP_NAVI14
&& tm_options
& AC_TM_FORCE_DISABLE_XNACK
? ",-xnack" : "",
202 tm_options
& AC_TM_PROMOTE_ALLOCA_TO_SCRATCH
? ",-promote-alloca" : "");
204 LLVMTargetMachineRef tm
=
205 LLVMCreateTargetMachine(target
, triple
, ac_get_llvm_processor_name(family
), features
, level
,
206 LLVMRelocDefault
, LLVMCodeModelDefault
);
209 *out_triple
= triple
;
210 if (tm_options
& AC_TM_ENABLE_GLOBAL_ISEL
)
211 ac_enable_global_isel(tm
);
215 static LLVMPassManagerRef
ac_create_passmgr(LLVMTargetLibraryInfoRef target_library_info
,
218 LLVMPassManagerRef passmgr
= LLVMCreatePassManager();
222 if (target_library_info
)
223 LLVMAddTargetLibraryInfo(target_library_info
, passmgr
);
226 LLVMAddVerifierPass(passmgr
);
227 LLVMAddAlwaysInlinerPass(passmgr
);
228 /* Normally, the pass manager runs all passes on one function before
229 * moving onto another. Adding a barrier no-op pass forces the pass
230 * manager to run the inliner on all functions first, which makes sure
231 * that the following passes are only run on the remaining non-inline
232 * function, so it removes useless work done on dead inline functions.
234 ac_llvm_add_barrier_noop_pass(passmgr
);
235 /* This pass should eliminate all the load and store instructions. */
236 LLVMAddPromoteMemoryToRegisterPass(passmgr
);
237 LLVMAddScalarReplAggregatesPass(passmgr
);
238 LLVMAddLICMPass(passmgr
);
239 LLVMAddAggressiveDCEPass(passmgr
);
240 LLVMAddCFGSimplificationPass(passmgr
);
241 /* This is recommended by the instruction combining pass. */
242 LLVMAddEarlyCSEMemSSAPass(passmgr
);
243 LLVMAddInstructionCombiningPass(passmgr
);
247 static const char *attr_to_str(enum ac_func_attr attr
)
250 case AC_FUNC_ATTR_ALWAYSINLINE
:
251 return "alwaysinline";
252 case AC_FUNC_ATTR_INREG
:
254 case AC_FUNC_ATTR_NOALIAS
:
256 case AC_FUNC_ATTR_NOUNWIND
:
258 case AC_FUNC_ATTR_READNONE
:
260 case AC_FUNC_ATTR_READONLY
:
262 case AC_FUNC_ATTR_WRITEONLY
:
264 case AC_FUNC_ATTR_INACCESSIBLE_MEM_ONLY
:
265 return "inaccessiblememonly";
266 case AC_FUNC_ATTR_CONVERGENT
:
269 fprintf(stderr
, "Unhandled function attribute: %x\n", attr
);
274 void ac_add_function_attr(LLVMContextRef ctx
, LLVMValueRef function
, int attr_idx
,
275 enum ac_func_attr attr
)
277 const char *attr_name
= attr_to_str(attr
);
278 unsigned kind_id
= LLVMGetEnumAttributeKindForName(attr_name
, strlen(attr_name
));
279 LLVMAttributeRef llvm_attr
= LLVMCreateEnumAttribute(ctx
, kind_id
, 0);
281 if (LLVMIsAFunction(function
))
282 LLVMAddAttributeAtIndex(function
, attr_idx
, llvm_attr
);
284 LLVMAddCallSiteAttribute(function
, attr_idx
, llvm_attr
);
287 void ac_add_func_attributes(LLVMContextRef ctx
, LLVMValueRef function
, unsigned attrib_mask
)
289 attrib_mask
|= AC_FUNC_ATTR_NOUNWIND
;
290 attrib_mask
&= ~AC_FUNC_ATTR_LEGACY
;
292 while (attrib_mask
) {
293 enum ac_func_attr attr
= 1u << u_bit_scan(&attrib_mask
);
294 ac_add_function_attr(ctx
, function
, -1, attr
);
298 void ac_dump_module(LLVMModuleRef module
)
300 char *str
= LLVMPrintModuleToString(module
);
301 fprintf(stderr
, "%s", str
);
302 LLVMDisposeMessage(str
);
305 void ac_llvm_add_target_dep_function_attr(LLVMValueRef F
, const char *name
, unsigned value
)
309 snprintf(str
, sizeof(str
), "0x%x", value
);
310 LLVMAddTargetDependentFunctionAttr(F
, name
, str
);
313 void ac_llvm_set_workgroup_size(LLVMValueRef F
, unsigned size
)
319 snprintf(str
, sizeof(str
), "%u,%u", size
, size
);
320 LLVMAddTargetDependentFunctionAttr(F
, "amdgpu-flat-work-group-size", str
);
323 unsigned ac_count_scratch_private_memory(LLVMValueRef function
)
325 unsigned private_mem_vgprs
= 0;
327 /* Process all LLVM instructions. */
328 LLVMBasicBlockRef bb
= LLVMGetFirstBasicBlock(function
);
330 LLVMValueRef next
= LLVMGetFirstInstruction(bb
);
333 LLVMValueRef inst
= next
;
334 next
= LLVMGetNextInstruction(next
);
336 if (LLVMGetInstructionOpcode(inst
) != LLVMAlloca
)
339 LLVMTypeRef type
= LLVMGetElementType(LLVMTypeOf(inst
));
340 /* No idea why LLVM aligns allocas to 4 elements. */
341 unsigned alignment
= LLVMGetAlignment(inst
);
342 unsigned dw_size
= align(ac_get_type_size(type
) / 4, alignment
);
343 private_mem_vgprs
+= dw_size
;
345 bb
= LLVMGetNextBasicBlock(bb
);
348 return private_mem_vgprs
;
351 bool ac_init_llvm_compiler(struct ac_llvm_compiler
*compiler
, enum radeon_family family
,
352 enum ac_target_machine_options tm_options
)
355 memset(compiler
, 0, sizeof(*compiler
));
357 compiler
->tm
= ac_create_target_machine(family
, tm_options
, LLVMCodeGenLevelDefault
, &triple
);
361 if (tm_options
& AC_TM_CREATE_LOW_OPT
) {
362 compiler
->low_opt_tm
=
363 ac_create_target_machine(family
, tm_options
, LLVMCodeGenLevelLess
, NULL
);
364 if (!compiler
->low_opt_tm
)
368 if (family
>= CHIP_NAVI10
) {
369 assert(!(tm_options
& AC_TM_CREATE_LOW_OPT
));
370 compiler
->tm_wave32
=
371 ac_create_target_machine(family
, tm_options
| AC_TM_WAVE32
, LLVMCodeGenLevelDefault
, NULL
);
372 if (!compiler
->tm_wave32
)
376 compiler
->target_library_info
= ac_create_target_library_info(triple
);
377 if (!compiler
->target_library_info
)
381 ac_create_passmgr(compiler
->target_library_info
, tm_options
& AC_TM_CHECK_IR
);
382 if (!compiler
->passmgr
)
387 ac_destroy_llvm_compiler(compiler
);
391 void ac_destroy_llvm_compiler(struct ac_llvm_compiler
*compiler
)
393 ac_destroy_llvm_passes(compiler
->passes
);
394 ac_destroy_llvm_passes(compiler
->passes_wave32
);
395 ac_destroy_llvm_passes(compiler
->low_opt_passes
);
397 if (compiler
->passmgr
)
398 LLVMDisposePassManager(compiler
->passmgr
);
399 if (compiler
->target_library_info
)
400 ac_dispose_target_library_info(compiler
->target_library_info
);
401 if (compiler
->low_opt_tm
)
402 LLVMDisposeTargetMachine(compiler
->low_opt_tm
);
404 LLVMDisposeTargetMachine(compiler
->tm
);
405 if (compiler
->tm_wave32
)
406 LLVMDisposeTargetMachine(compiler
->tm_wave32
);