Merge branch 'master' into blender2.8
[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>
37         std::string to_string(const T &n) {
38                 std::ostringstream s;
39                 s << n;
40                 return s.str();
41         }
42 }
43
44 class CompilationSettings
45 {
46 public:
47         CompilationSettings()
48         : target_arch(0),
49           bits(64),
50           verbose(false),
51           fast_math(false)
52         {}
53
54         string cuda_toolkit_dir;
55         string input_file;
56         string output_file;
57         string ptx_file;
58         vector<string> defines;
59         vector<string> includes;
60         int target_arch;
61         int bits;
62         bool verbose;
63         bool fast_math;
64 };
65
66 bool compile_cuda(CompilationSettings &settings)
67 {
68         const char* headers[] = {"stdlib.h" , "float.h", "math.h", "stdio.h"};
69         const char* header_content[] = {"\n", "\n", "\n", "\n"};
70
71         printf("Building %s\n", settings.input_file.c_str());
72
73         string code;
74         if(!OIIO::Filesystem::read_text_file(settings.input_file, code)) {
75                 fprintf(stderr, "Error: unable to read %s\n", settings.input_file.c_str());
76                 return false;
77         }
78
79         vector<string> options;
80         for(size_t i = 0; i < settings.includes.size(); i++) {
81                 options.push_back("-I" + settings.includes[i]);
82         }
83
84         for(size_t i = 0; i < settings.defines.size(); i++) {
85                 options.push_back("-D" + settings.defines[i]);
86         }
87         options.push_back("-D__KERNEL_CUDA_VERSION__=" + std::to_string(cuewNvrtcVersion()));
88         options.push_back("-arch=compute_" + std::to_string(settings.target_arch));
89         options.push_back("--device-as-default-execution-space");
90         if(settings.fast_math)
91                 options.push_back("--use_fast_math");
92
93         nvrtcProgram prog;
94         nvrtcResult result = nvrtcCreateProgram(&prog,
95                 code.c_str(),                    // buffer
96                 NULL,                            // name
97                 sizeof(headers) / sizeof(void*), // numHeaders
98                 header_content,                  // headers
99                 headers);                        // includeNames
100
101         if(result != NVRTC_SUCCESS) {
102                 fprintf(stderr, "Error: nvrtcCreateProgram failed (%x)\n\n", result);
103                 return false;
104         }
105
106         /* Tranfer options to a classic C array. */
107         vector<const char*> opts(options.size());
108         for(size_t i = 0; i < options.size(); i++) {
109                 opts[i] = options[i].c_str();
110         }
111
112         result = nvrtcCompileProgram(prog, options.size(), &opts[0]);
113
114         if(result != NVRTC_SUCCESS) {
115                 fprintf(stderr, "Error: nvrtcCompileProgram failed (%x)\n\n", result);
116
117                 size_t log_size;
118                 nvrtcGetProgramLogSize(prog, &log_size);
119
120                 vector<char> log(log_size);
121                 nvrtcGetProgramLog(prog, &log[0]);
122                 fprintf(stderr, "%s\n", &log[0]);
123
124                 return false;
125         }
126
127         /* Retrieve the ptx code. */
128         size_t ptx_size;
129         result = nvrtcGetPTXSize(prog, &ptx_size);
130         if(result != NVRTC_SUCCESS) {
131                 fprintf(stderr, "Error: nvrtcGetPTXSize failed (%x)\n\n", result);
132                 return false;
133         }
134
135         vector<char> ptx_code(ptx_size);
136         result = nvrtcGetPTX(prog, &ptx_code[0]);
137         if(result != NVRTC_SUCCESS) {
138                 fprintf(stderr, "Error: nvrtcGetPTX failed (%x)\n\n", result);
139                 return false;
140         }
141
142         /* Write a file in the temp folder with the ptx code. */
143         settings.ptx_file = OIIO::Filesystem::temp_directory_path() + "/" + OIIO::Filesystem::unique_path();
144         FILE * f= fopen(settings.ptx_file.c_str(), "wb");
145         fwrite(&ptx_code[0], 1, ptx_size, f);
146         fclose(f);
147
148         return true;
149 }
150
151 bool link_ptxas(CompilationSettings &settings)
152 {
153         string cudapath = "";
154         if(settings.cuda_toolkit_dir.size())
155                 cudapath = settings.cuda_toolkit_dir + "/bin/";
156
157         string ptx = "\"" +cudapath + "ptxas\" " + settings.ptx_file +
158                                         " -o " + settings.output_file +
159                                         " --gpu-name sm_" + std::to_string(settings.target_arch) +
160                                         " -m" + std::to_string(settings.bits);
161
162         if(settings.verbose) {
163                 ptx += " --verbose";
164                 printf("%s\n", ptx.c_str());
165         }
166
167         int pxresult = system(ptx.c_str());
168         if(pxresult) {
169                 fprintf(stderr, "Error: ptxas failed (%x)\n\n", pxresult);
170                 return false;
171         }
172
173         if(!OIIO::Filesystem::remove(settings.ptx_file)) {
174                 fprintf(stderr, "Error: removing %s\n\n", settings.ptx_file.c_str());
175         }
176
177         return true;
178 }
179
180 bool init(CompilationSettings &settings)
181 {
182 #ifdef _MSC_VER
183         if(settings.cuda_toolkit_dir.size()) {
184                 SetDllDirectory((settings.cuda_toolkit_dir + "/bin").c_str());
185         }
186 #endif
187
188         int cuewresult = cuewInit(CUEW_INIT_NVRTC);
189         if(cuewresult != CUEW_SUCCESS) {
190                 fprintf(stderr, "Error: cuew init fialed (0x%x)\n\n", cuewresult);
191                 return false;
192         }
193
194         if(cuewNvrtcVersion() < 80) {
195                 fprintf(stderr, "Error: only cuda 8 and higher is supported, %d\n\n", cuewCompilerVersion());
196                 return false;
197         }
198
199         if(!nvrtcCreateProgram) {
200                 fprintf(stderr, "Error: nvrtcCreateProgram not resolved\n");
201                 return false;
202         }
203
204         if(!nvrtcCompileProgram) {
205                 fprintf(stderr, "Error: nvrtcCompileProgram not resolved\n");
206                 return false;
207         }
208
209         if(!nvrtcGetProgramLogSize) {
210                 fprintf(stderr, "Error: nvrtcGetProgramLogSize not resolved\n");
211                 return false;
212         }
213
214         if(!nvrtcGetProgramLog) {
215                 fprintf(stderr, "Error: nvrtcGetProgramLog not resolved\n");
216                 return false;
217         }
218
219         if(!nvrtcGetPTXSize) {
220                 fprintf(stderr, "Error: nvrtcGetPTXSize not resolved\n");
221                 return false;
222         }
223
224         if(!nvrtcGetPTX) {
225                 fprintf(stderr, "Error: nvrtcGetPTX not resolved\n");
226                 return false;
227         }
228
229         return true;
230 }
231
232 bool parse_parameters(int argc, const char **argv, CompilationSettings &settings)
233 {
234         OIIO::ArgParse ap;
235         ap.options("Usage: cycles_cubin_cc [options]",
236                 "-target %d", &settings.target_arch, "target shader model",
237                 "-m %d", &settings.bits, "Cuda architecture bits",
238                 "-i %s", &settings.input_file, "Input source filename",
239                 "-o %s", &settings.output_file, "Output cubin filename",
240                 "-I %L", &settings.includes, "Add additional includepath",
241                 "-D %L", &settings.defines, "Add additional defines",
242                 "-v", &settings.verbose, "Use verbose logging",
243                 "--use_fast_math", &settings.fast_math, "Use fast math",
244                 "-cuda-toolkit-dir %s", &settings.cuda_toolkit_dir, "path to the cuda toolkit binary directory",
245                 NULL);
246
247         if(ap.parse(argc, argv) < 0) {
248                 fprintf(stderr, "%s\n", ap.geterror().c_str());
249                 ap.usage();
250                 return false;
251         }
252
253         if(!settings.output_file.size()) {
254                 fprintf(stderr, "Error: Output file not set(-o), required\n\n");
255                 return false;
256         }
257
258         if(!settings.input_file.size()) {
259                 fprintf(stderr, "Error: Input file not set(-i, required\n\n");
260                 return false;
261         }
262
263         if(!settings.target_arch) {
264                 fprintf(stderr, "Error: target shader model not set (-target), required\n\n");
265                 return false;
266         }
267
268         return true;
269 }
270
271 int main(int argc, const char **argv)
272 {
273         CompilationSettings settings;
274
275         if(!parse_parameters(argc, argv, settings)) {
276                 fprintf(stderr, "Error: invalid parameters, exiting\n");
277                 exit(EXIT_FAILURE);
278         }
279
280         if(!init(settings)) {
281                 fprintf(stderr, "Error: initialization error, exiting\n");
282                 exit(EXIT_FAILURE);
283         }
284
285         if(!compile_cuda(settings)) {
286                 fprintf(stderr, "Error: compilation error, exiting\n");
287                 exit(EXIT_FAILURE);
288         }
289
290         if(!link_ptxas(settings)) {
291                 exit(EXIT_FAILURE);
292         }
293
294         return 0;
295 }