1 /*
2 * Copyright 2014 Advanced Micro Devices, Inc.
3 *
4 * SPDX-License-Identifier: MIT
5 */
6 /* based on pieces from si_pipe.c and radeon_llvm_emit.c */
7 #include "ac_llvm_util.h"
8
9 #include "ac_llvm_build.h"
10 #include "c11/threads.h"
11 #include "util/bitscan.h"
12 #include "util/u_math.h"
13 #include <llvm-c/Core.h>
14 #include <llvm-c/Support.h>
15
16 #include <assert.h>
17 #include <stdio.h>
18 #include <string.h>
19
ac_init_llvm_target(void)20 static void ac_init_llvm_target(void)
21 {
22 LLVMInitializeAMDGPUTargetInfo();
23 LLVMInitializeAMDGPUTarget();
24 LLVMInitializeAMDGPUTargetMC();
25 LLVMInitializeAMDGPUAsmPrinter();
26
27 /* For inline assembly. */
28 LLVMInitializeAMDGPUAsmParser();
29
30 /* For ACO disassembly. */
31 LLVMInitializeAMDGPUDisassembler();
32
33 const char *argv[] = {
34 /* error messages prefix */
35 "mesa",
36 "-amdgpu-atomic-optimizations=true",
37 /* image_msaa_load currently doesn't work with LLVM + GFX12 */
38 "-amdgpu-enable-image-intrinsic-optimizer=false",
39 };
40
41 ac_reset_llvm_all_options_occurrences();
42 LLVMParseCommandLineOptions(ARRAY_SIZE(argv), argv, NULL);
43
44 ac_llvm_run_atexit_for_destructors();
45 }
46
ac_init_shared_llvm_once(void)47 PUBLIC void ac_init_shared_llvm_once(void)
48 {
49 static once_flag ac_init_llvm_target_once_flag = ONCE_FLAG_INIT;
50 call_once(&ac_init_llvm_target_once_flag, ac_init_llvm_target);
51 }
52
53 #if !LLVM_IS_SHARED
54 static once_flag ac_init_static_llvm_target_once_flag = ONCE_FLAG_INIT;
ac_init_static_llvm_once(void)55 static void ac_init_static_llvm_once(void)
56 {
57 call_once(&ac_init_static_llvm_target_once_flag, ac_init_llvm_target);
58 }
59 #endif
60
ac_init_llvm_once(void)61 void ac_init_llvm_once(void)
62 {
63 #if LLVM_IS_SHARED
64 ac_init_shared_llvm_once();
65 #else
66 ac_init_static_llvm_once();
67 #endif
68 }
69
ac_get_llvm_target(const char * triple)70 LLVMTargetRef ac_get_llvm_target(const char *triple)
71 {
72 LLVMTargetRef target = NULL;
73 char *err_message = NULL;
74
75 if (LLVMGetTargetFromTriple(triple, &target, &err_message)) {
76 fprintf(stderr, "Cannot find target for triple %s ", triple);
77 if (err_message) {
78 fprintf(stderr, "%s\n", err_message);
79 }
80 LLVMDisposeMessage(err_message);
81 return NULL;
82 }
83 return target;
84 }
85
ac_create_target_machine(enum radeon_family family,enum ac_target_machine_options tm_options,LLVMCodeGenOptLevel level,const char ** out_triple)86 static LLVMTargetMachineRef ac_create_target_machine(enum radeon_family family,
87 enum ac_target_machine_options tm_options,
88 LLVMCodeGenOptLevel level,
89 const char **out_triple)
90 {
91 assert(family >= CHIP_TAHITI);
92 const char *triple = (tm_options & AC_TM_SUPPORTS_SPILL) ? "amdgcn-mesa-mesa3d" : "amdgcn--";
93 LLVMTargetRef target = ac_get_llvm_target(triple);
94 const char *name = ac_get_llvm_processor_name(family);
95
96 LLVMTargetMachineRef tm =
97 LLVMCreateTargetMachine(target, triple, name, "", level,
98 LLVMRelocDefault, LLVMCodeModelDefault);
99
100 if (!ac_is_llvm_processor_supported(tm, name)) {
101 LLVMDisposeTargetMachine(tm);
102 fprintf(stderr, "amd: LLVM doesn't support %s, bailing out...\n", name);
103 return NULL;
104 }
105
106 if (out_triple)
107 *out_triple = triple;
108
109 return tm;
110 }
111
ac_get_llvm_attribute(LLVMContextRef ctx,const char * str)112 LLVMAttributeRef ac_get_llvm_attribute(LLVMContextRef ctx, const char *str)
113 {
114 return LLVMCreateEnumAttribute(ctx, LLVMGetEnumAttributeKindForName(str, strlen(str)), 0);
115 }
116
ac_add_function_attr(LLVMContextRef ctx,LLVMValueRef function,int attr_idx,const char * attr)117 void ac_add_function_attr(LLVMContextRef ctx, LLVMValueRef function, int attr_idx,
118 const char *attr)
119 {
120 assert(LLVMIsAFunction(function));
121 LLVMAddAttributeAtIndex(function, attr_idx, ac_get_llvm_attribute(ctx, attr));
122 }
123
ac_dump_module(LLVMModuleRef module)124 void ac_dump_module(LLVMModuleRef module)
125 {
126 char *str = LLVMPrintModuleToString(module);
127 fprintf(stderr, "%s", str);
128 LLVMDisposeMessage(str);
129 }
130
ac_llvm_add_target_dep_function_attr(LLVMValueRef F,const char * name,unsigned value)131 void ac_llvm_add_target_dep_function_attr(LLVMValueRef F, const char *name, unsigned value)
132 {
133 char str[16];
134
135 snprintf(str, sizeof(str), "0x%x", value);
136 LLVMAddTargetDependentFunctionAttr(F, name, str);
137 }
138
ac_llvm_set_workgroup_size(LLVMValueRef F,unsigned size)139 void ac_llvm_set_workgroup_size(LLVMValueRef F, unsigned size)
140 {
141 if (!size)
142 return;
143
144 char str[32];
145 snprintf(str, sizeof(str), "%u,%u", size, size);
146 LLVMAddTargetDependentFunctionAttr(F, "amdgpu-flat-work-group-size", str);
147 }
148
ac_llvm_set_target_features(LLVMValueRef F,struct ac_llvm_context * ctx,bool wgp_mode)149 void ac_llvm_set_target_features(LLVMValueRef F, struct ac_llvm_context *ctx, bool wgp_mode)
150 {
151 char features[2048];
152
153 snprintf(features, sizeof(features), "+DumpCode%s%s%s",
154 /* GFX9 has broken VGPR indexing, so always promote alloca to scratch. */
155 ctx->gfx_level == GFX9 ? ",-promote-alloca" : "",
156 /* Wave32 is the default. */
157 ctx->gfx_level >= GFX10 && ctx->wave_size == 64 ?
158 ",+wavefrontsize64,-wavefrontsize32" : "",
159 ctx->gfx_level >= GFX10 && !wgp_mode ? ",+cumode" : "");
160
161 LLVMAddTargetDependentFunctionAttr(F, "target-features", features);
162 }
163
ac_init_llvm_compiler(struct ac_llvm_compiler * compiler,enum radeon_family family,enum ac_target_machine_options tm_options)164 bool ac_init_llvm_compiler(struct ac_llvm_compiler *compiler, enum radeon_family family,
165 enum ac_target_machine_options tm_options)
166 {
167 const char *triple;
168 memset(compiler, 0, sizeof(*compiler));
169
170 compiler->tm = ac_create_target_machine(family, tm_options, LLVMCodeGenLevelDefault, &triple);
171 if (!compiler->tm)
172 return false;
173
174 if (tm_options & AC_TM_CREATE_LOW_OPT) {
175 compiler->low_opt_tm =
176 ac_create_target_machine(family, tm_options, LLVMCodeGenLevelLess, NULL);
177 if (!compiler->low_opt_tm)
178 goto fail;
179 }
180
181 compiler->target_library_info = ac_create_target_library_info(triple);
182 if (!compiler->target_library_info)
183 goto fail;
184
185 compiler->passmgr =
186 ac_create_passmgr(compiler->target_library_info, tm_options & AC_TM_CHECK_IR);
187 if (!compiler->passmgr)
188 goto fail;
189
190 return true;
191 fail:
192 ac_destroy_llvm_compiler(compiler);
193 return false;
194 }
195
ac_destroy_llvm_compiler(struct ac_llvm_compiler * compiler)196 void ac_destroy_llvm_compiler(struct ac_llvm_compiler *compiler)
197 {
198 ac_destroy_llvm_passes(compiler->passes);
199 ac_destroy_llvm_passes(compiler->low_opt_passes);
200
201 if (compiler->passmgr)
202 LLVMDisposePassManager(compiler->passmgr);
203 if (compiler->target_library_info)
204 ac_dispose_target_library_info(compiler->target_library_info);
205 if (compiler->low_opt_tm)
206 LLVMDisposeTargetMachine(compiler->low_opt_tm);
207 if (compiler->tm)
208 LLVMDisposeTargetMachine(compiler->tm);
209 }
210