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 };
38
39 ac_reset_llvm_all_options_occurrences();
40 LLVMParseCommandLineOptions(ARRAY_SIZE(argv), argv, NULL);
41
42 ac_llvm_run_atexit_for_destructors();
43 }
44
ac_init_shared_llvm_once(void)45 PUBLIC void ac_init_shared_llvm_once(void)
46 {
47 static once_flag ac_init_llvm_target_once_flag = ONCE_FLAG_INIT;
48 call_once(&ac_init_llvm_target_once_flag, ac_init_llvm_target);
49 }
50
51 #if !LLVM_IS_SHARED
52 static once_flag ac_init_static_llvm_target_once_flag = ONCE_FLAG_INIT;
ac_init_static_llvm_once(void)53 static void ac_init_static_llvm_once(void)
54 {
55 call_once(&ac_init_static_llvm_target_once_flag, ac_init_llvm_target);
56 }
57 #endif
58
ac_init_llvm_once(void)59 void ac_init_llvm_once(void)
60 {
61 #if LLVM_IS_SHARED
62 ac_init_shared_llvm_once();
63 #else
64 ac_init_static_llvm_once();
65 #endif
66 }
67
ac_get_llvm_target(const char * triple)68 LLVMTargetRef ac_get_llvm_target(const char *triple)
69 {
70 LLVMTargetRef target = NULL;
71 char *err_message = NULL;
72
73 if (LLVMGetTargetFromTriple(triple, &target, &err_message)) {
74 fprintf(stderr, "Cannot find target for triple %s ", triple);
75 if (err_message) {
76 fprintf(stderr, "%s\n", err_message);
77 }
78 LLVMDisposeMessage(err_message);
79 return NULL;
80 }
81 return target;
82 }
83
ac_create_target_machine(enum radeon_family family,enum ac_target_machine_options tm_options,LLVMCodeGenOptLevel level,const char ** out_triple)84 static LLVMTargetMachineRef ac_create_target_machine(enum radeon_family family,
85 enum ac_target_machine_options tm_options,
86 LLVMCodeGenOptLevel level,
87 const char **out_triple)
88 {
89 assert(family >= CHIP_TAHITI);
90 const char *triple = (tm_options & AC_TM_SUPPORTS_SPILL) ? "amdgcn-mesa-mesa3d" : "amdgcn--";
91 LLVMTargetRef target = ac_get_llvm_target(triple);
92 const char *name = ac_get_llvm_processor_name(family);
93
94 LLVMTargetMachineRef tm =
95 LLVMCreateTargetMachine(target, triple, name, "", level,
96 LLVMRelocDefault, LLVMCodeModelDefault);
97
98 if (!ac_is_llvm_processor_supported(tm, name)) {
99 LLVMDisposeTargetMachine(tm);
100 fprintf(stderr, "amd: LLVM doesn't support %s, bailing out...\n", name);
101 return NULL;
102 }
103
104 if (out_triple)
105 *out_triple = triple;
106
107 return tm;
108 }
109
ac_get_llvm_attribute(LLVMContextRef ctx,const char * str)110 LLVMAttributeRef ac_get_llvm_attribute(LLVMContextRef ctx, const char *str)
111 {
112 return LLVMCreateEnumAttribute(ctx, LLVMGetEnumAttributeKindForName(str, strlen(str)), 0);
113 }
114
ac_add_function_attr(LLVMContextRef ctx,LLVMValueRef function,int attr_idx,const char * attr)115 void ac_add_function_attr(LLVMContextRef ctx, LLVMValueRef function, int attr_idx,
116 const char *attr)
117 {
118 assert(LLVMIsAFunction(function));
119 LLVMAddAttributeAtIndex(function, attr_idx, ac_get_llvm_attribute(ctx, attr));
120 }
121
ac_dump_module(LLVMModuleRef module)122 void ac_dump_module(LLVMModuleRef module)
123 {
124 char *str = LLVMPrintModuleToString(module);
125 fprintf(stderr, "%s", str);
126 LLVMDisposeMessage(str);
127 }
128
ac_llvm_add_target_dep_function_attr(LLVMValueRef F,const char * name,unsigned value)129 void ac_llvm_add_target_dep_function_attr(LLVMValueRef F, const char *name, unsigned value)
130 {
131 char str[16];
132
133 snprintf(str, sizeof(str), "0x%x", value);
134 LLVMAddTargetDependentFunctionAttr(F, name, str);
135 }
136
ac_llvm_set_workgroup_size(LLVMValueRef F,unsigned size)137 void ac_llvm_set_workgroup_size(LLVMValueRef F, unsigned size)
138 {
139 if (!size)
140 return;
141
142 char str[32];
143 snprintf(str, sizeof(str), "%u,%u", size, size);
144 LLVMAddTargetDependentFunctionAttr(F, "amdgpu-flat-work-group-size", str);
145 }
146
ac_llvm_set_target_features(LLVMValueRef F,struct ac_llvm_context * ctx,bool wgp_mode)147 void ac_llvm_set_target_features(LLVMValueRef F, struct ac_llvm_context *ctx, bool wgp_mode)
148 {
149 char features[2048];
150
151 snprintf(features, sizeof(features), "+DumpCode%s%s%s",
152 /* GFX9 has broken VGPR indexing, so always promote alloca to scratch. */
153 ctx->gfx_level == GFX9 ? ",-promote-alloca" : "",
154 /* Wave32 is the default. */
155 ctx->gfx_level >= GFX10 && ctx->wave_size == 64 ?
156 ",+wavefrontsize64,-wavefrontsize32" : "",
157 ctx->gfx_level >= GFX10 && !wgp_mode ? ",+cumode" : "");
158
159 LLVMAddTargetDependentFunctionAttr(F, "target-features", features);
160 }
161
ac_init_llvm_compiler(struct ac_llvm_compiler * compiler,enum radeon_family family,enum ac_target_machine_options tm_options)162 bool ac_init_llvm_compiler(struct ac_llvm_compiler *compiler, enum radeon_family family,
163 enum ac_target_machine_options tm_options)
164 {
165 const char *triple;
166 memset(compiler, 0, sizeof(*compiler));
167
168 compiler->tm = ac_create_target_machine(family, tm_options, LLVMCodeGenLevelDefault, &triple);
169 if (!compiler->tm)
170 return false;
171
172 if (tm_options & AC_TM_CREATE_LOW_OPT) {
173 compiler->low_opt_tm =
174 ac_create_target_machine(family, tm_options, LLVMCodeGenLevelLess, NULL);
175 if (!compiler->low_opt_tm)
176 goto fail;
177 }
178
179 compiler->target_library_info = ac_create_target_library_info(triple);
180 if (!compiler->target_library_info)
181 goto fail;
182
183 compiler->passmgr =
184 ac_create_passmgr(compiler->target_library_info, tm_options & AC_TM_CHECK_IR);
185 if (!compiler->passmgr)
186 goto fail;
187
188 return true;
189 fail:
190 ac_destroy_llvm_compiler(compiler);
191 return false;
192 }
193
ac_destroy_llvm_compiler(struct ac_llvm_compiler * compiler)194 void ac_destroy_llvm_compiler(struct ac_llvm_compiler *compiler)
195 {
196 ac_destroy_llvm_passes(compiler->passes);
197 ac_destroy_llvm_passes(compiler->low_opt_passes);
198
199 if (compiler->passmgr)
200 LLVMDisposePassManager(compiler->passmgr);
201 if (compiler->target_library_info)
202 ac_dispose_target_library_info(compiler->target_library_info);
203 if (compiler->low_opt_tm)
204 LLVMDisposeTargetMachine(compiler->low_opt_tm);
205 if (compiler->tm)
206 LLVMDisposeTargetMachine(compiler->tm);
207 }
208