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