ClangFormat: apply to source, most of intern
[blender.git] / intern / cycles / app / cycles_cubin_cc.cpp
1 /*
2  * Copyright 2017 Blender Foundation
3  *
4  * Licensed under the Apache License, Version 2.0 (the "License");
5  * you may not use this file except in compliance with the License.
6  * You may obtain a copy of the License at
7  *
8  * http://www.apache.org/licenses/LICENSE-2.0
9  *
10  * Unless required by applicable law or agreed to in writing, software
11  * distributed under the License is distributed on an "AS IS" BASIS,
12  * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13  * See the License for the specific language governing permissions and
14  * limitations under the License.
15  */
16
17 #include <stdio.h>
18 #include <stdint.h>
19
20 #include <string>
21 #include <vector>
22
23 #include <OpenImageIO/argparse.h>
24 #include <OpenImageIO/filesystem.h>
25
26 #include "cuew.h"
27
28 #ifdef _MSC_VER
29 #  include <Windows.h>
30 #endif
31
32 using std::string;
33 using std::vector;
34
35 namespace std {
36 template<typename T> std::string to_string(const T &n)
37 {
38   std::ostringstream s;
39   s << n;
40   return s.str();
41 }
42 }  // namespace std
43
44 class CompilationSettings {
45  public:
46   CompilationSettings() : target_arch(0), bits(64), verbose(false), fast_math(false)
47   {
48   }
49
50   string cuda_toolkit_dir;
51   string input_file;
52   string output_file;
53   string ptx_file;
54   vector<string> defines;
55   vector<string> includes;
56   int target_arch;
57   int bits;
58   bool verbose;
59   bool fast_math;
60 };
61
62 static bool compile_cuda(CompilationSettings &settings)
63 {
64   const char *headers[] = {"stdlib.h", "float.h", "math.h", "stdio.h"};
65   const char *header_content[] = {"\n", "\n", "\n", "\n"};
66
67   printf("Building %s\n", settings.input_file.c_str());
68
69   string code;
70   if (!OIIO::Filesystem::read_text_file(settings.input_file, code)) {
71     fprintf(stderr, "Error: unable to read %s\n", settings.input_file.c_str());
72     return false;
73   }
74
75   vector<string> options;
76   for (size_t i = 0; i < settings.includes.size(); i++) {
77     options.push_back("-I" + settings.includes[i]);
78   }
79
80   for (size_t i = 0; i < settings.defines.size(); i++) {
81     options.push_back("-D" + settings.defines[i]);
82   }
83   options.push_back("-D__KERNEL_CUDA_VERSION__=" + std::to_string(cuewNvrtcVersion()));
84   options.push_back("-arch=compute_" + std::to_string(settings.target_arch));
85   options.push_back("--device-as-default-execution-space");
86   if (settings.fast_math)
87     options.push_back("--use_fast_math");
88
89   nvrtcProgram prog;
90   nvrtcResult result = nvrtcCreateProgram(&prog,
91                                           code.c_str(),                      // buffer
92                                           NULL,                              // name
93                                           sizeof(headers) / sizeof(void *),  // numHeaders
94                                           header_content,                    // headers
95                                           headers);                          // includeNames
96
97   if (result != NVRTC_SUCCESS) {
98     fprintf(stderr, "Error: nvrtcCreateProgram failed (%d)\n\n", (int)result);
99     return false;
100   }
101
102   /* Tranfer options to a classic C array. */
103   vector<const char *> opts(options.size());
104   for (size_t i = 0; i < options.size(); i++) {
105     opts[i] = options[i].c_str();
106   }
107
108   result = nvrtcCompileProgram(prog, options.size(), &opts[0]);
109
110   if (result != NVRTC_SUCCESS) {
111     fprintf(stderr, "Error: nvrtcCompileProgram failed (%d)\n\n", (int)result);
112
113     size_t log_size;
114     nvrtcGetProgramLogSize(prog, &log_size);
115
116     vector<char> log(log_size);
117     nvrtcGetProgramLog(prog, &log[0]);
118     fprintf(stderr, "%s\n", &log[0]);
119
120     return false;
121   }
122
123   /* Retrieve the ptx code. */
124   size_t ptx_size;
125   result = nvrtcGetPTXSize(prog, &ptx_size);
126   if (result != NVRTC_SUCCESS) {
127     fprintf(stderr, "Error: nvrtcGetPTXSize failed (%d)\n\n", (int)result);
128     return false;
129   }
130
131   vector<char> ptx_code(ptx_size);
132   result = nvrtcGetPTX(prog, &ptx_code[0]);
133   if (result != NVRTC_SUCCESS) {
134     fprintf(stderr, "Error: nvrtcGetPTX failed (%d)\n\n", (int)result);
135     return false;
136   }
137
138   /* Write a file in the temp folder with the ptx code. */
139   settings.ptx_file = OIIO::Filesystem::temp_directory_path() + "/" +
140                       OIIO::Filesystem::unique_path();
141   FILE *f = fopen(settings.ptx_file.c_str(), "wb");
142   fwrite(&ptx_code[0], 1, ptx_size, f);
143   fclose(f);
144
145   return true;
146 }
147
148 static bool link_ptxas(CompilationSettings &settings)
149 {
150   string cudapath = "";
151   if (settings.cuda_toolkit_dir.size())
152     cudapath = settings.cuda_toolkit_dir + "/bin/";
153
154   string ptx = "\"" + cudapath + "ptxas\" " + settings.ptx_file + " -o " + settings.output_file +
155                " --gpu-name sm_" + std::to_string(settings.target_arch) + " -m" +
156                std::to_string(settings.bits);
157
158   if (settings.verbose) {
159     ptx += " --verbose";
160     printf("%s\n", ptx.c_str());
161   }
162
163   int pxresult = system(ptx.c_str());
164   if (pxresult) {
165     fprintf(stderr, "Error: ptxas failed (%d)\n\n", pxresult);
166     return false;
167   }
168
169   if (!OIIO::Filesystem::remove(settings.ptx_file)) {
170     fprintf(stderr, "Error: removing %s\n\n", settings.ptx_file.c_str());
171   }
172
173   return true;
174 }
175
176 static bool init(CompilationSettings &settings)
177 {
178 #ifdef _MSC_VER
179   if (settings.cuda_toolkit_dir.size()) {
180     SetDllDirectory((settings.cuda_toolkit_dir + "/bin").c_str());
181   }
182 #else
183   (void)settings;
184 #endif
185
186   int cuewresult = cuewInit(CUEW_INIT_NVRTC);
187   if (cuewresult != CUEW_SUCCESS) {
188     fprintf(stderr, "Error: cuew init fialed (0x%d)\n\n", cuewresult);
189     return false;
190   }
191
192   if (cuewNvrtcVersion() < 80) {
193     fprintf(stderr, "Error: only cuda 8 and higher is supported, %d\n\n", cuewCompilerVersion());
194     return false;
195   }
196
197   if (!nvrtcCreateProgram) {
198     fprintf(stderr, "Error: nvrtcCreateProgram not resolved\n");
199     return false;
200   }
201
202   if (!nvrtcCompileProgram) {
203     fprintf(stderr, "Error: nvrtcCompileProgram not resolved\n");
204     return false;
205   }
206
207   if (!nvrtcGetProgramLogSize) {
208     fprintf(stderr, "Error: nvrtcGetProgramLogSize not resolved\n");
209     return false;
210   }
211
212   if (!nvrtcGetProgramLog) {
213     fprintf(stderr, "Error: nvrtcGetProgramLog not resolved\n");
214     return false;
215   }
216
217   if (!nvrtcGetPTXSize) {
218     fprintf(stderr, "Error: nvrtcGetPTXSize not resolved\n");
219     return false;
220   }
221
222   if (!nvrtcGetPTX) {
223     fprintf(stderr, "Error: nvrtcGetPTX not resolved\n");
224     return false;
225   }
226
227   return true;
228 }
229
230 static bool parse_parameters(int argc, const char **argv, CompilationSettings &settings)
231 {
232   OIIO::ArgParse ap;
233   ap.options("Usage: cycles_cubin_cc [options]",
234              "-target %d",
235              &settings.target_arch,
236              "target shader model",
237              "-m %d",
238              &settings.bits,
239              "Cuda architecture bits",
240              "-i %s",
241              &settings.input_file,
242              "Input source filename",
243              "-o %s",
244              &settings.output_file,
245              "Output cubin filename",
246              "-I %L",
247              &settings.includes,
248              "Add additional includepath",
249              "-D %L",
250              &settings.defines,
251              "Add additional defines",
252              "-v",
253              &settings.verbose,
254              "Use verbose logging",
255              "--use_fast_math",
256              &settings.fast_math,
257              "Use fast math",
258              "-cuda-toolkit-dir %s",
259              &settings.cuda_toolkit_dir,
260              "path to the cuda toolkit binary directory",
261              NULL);
262
263   if (ap.parse(argc, argv) < 0) {
264     fprintf(stderr, "%s\n", ap.geterror().c_str());
265     ap.usage();
266     return false;
267   }
268
269   if (!settings.output_file.size()) {
270     fprintf(stderr, "Error: Output file not set(-o), required\n\n");
271     return false;
272   }
273
274   if (!settings.input_file.size()) {
275     fprintf(stderr, "Error: Input file not set(-i, required\n\n");
276     return false;
277   }
278
279   if (!settings.target_arch) {
280     fprintf(stderr, "Error: target shader model not set (-target), required\n\n");
281     return false;
282   }
283
284   return true;
285 }
286
287 int main(int argc, const char **argv)
288 {
289   CompilationSettings settings;
290
291   if (!parse_parameters(argc, argv, settings)) {
292     fprintf(stderr, "Error: invalid parameters, exiting\n");
293     exit(EXIT_FAILURE);
294   }
295
296   if (!init(settings)) {
297     fprintf(stderr, "Error: initialization error, exiting\n");
298     exit(EXIT_FAILURE);
299   }
300
301   if (!compile_cuda(settings)) {
302     fprintf(stderr, "Error: compilation error, exiting\n");
303     exit(EXIT_FAILURE);
304   }
305
306   if (!link_ptxas(settings)) {
307     exit(EXIT_FAILURE);
308   }
309
310   return 0;
311 }