Fix T54105: random walk SSS missing in branched indirect paths.
[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 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         {
164                 ptx += " --verbose";
165                 printf("%s\n", ptx.c_str());
166         }
167         
168         int pxresult = system(ptx.c_str());
169         if(pxresult) {
170                 fprintf(stderr, "Error: ptxas failed (%x)\n\n", pxresult);
171                 return false;
172         }
173
174         if(!OIIO::Filesystem::remove(settings.ptx_file)) {
175                 fprintf(stderr, "Error: removing %s\n\n", settings.ptx_file.c_str());
176         }
177
178         return true;
179 }
180
181 bool init(CompilationSettings &settings)
182 {
183 #ifdef _MSC_VER
184         if(settings.cuda_toolkit_dir.size()) {
185                 SetDllDirectory((settings.cuda_toolkit_dir + "/bin").c_str());
186         }
187 #endif
188
189         int cuewresult = cuewInit(CUEW_INIT_NVRTC);
190         if(cuewresult != CUEW_SUCCESS) {
191                 fprintf(stderr, "Error: cuew init fialed (0x%x)\n\n", cuewresult);
192                 return false;
193         }
194
195         if(cuewNvrtcVersion() < 80) {
196                 fprintf(stderr, "Error: only cuda 8 and higher is supported, %d\n\n", cuewCompilerVersion());
197                 return false;
198         }
199
200         if(!nvrtcCreateProgram) {
201                 fprintf(stderr, "Error: nvrtcCreateProgram not resolved\n");
202                 return false;
203         }
204
205         if(!nvrtcCompileProgram) {
206                 fprintf(stderr, "Error: nvrtcCompileProgram not resolved\n");
207                 return false;
208         }
209
210         if(!nvrtcGetProgramLogSize) {
211                 fprintf(stderr, "Error: nvrtcGetProgramLogSize not resolved\n");
212                 return false;
213         }
214
215         if(!nvrtcGetProgramLog) {
216                 fprintf(stderr, "Error: nvrtcGetProgramLog not resolved\n");
217                 return false;
218         }
219
220         if(!nvrtcGetPTXSize) {
221                 fprintf(stderr, "Error: nvrtcGetPTXSize not resolved\n");
222                 return false;
223         }
224
225         if(!nvrtcGetPTX) {
226                 fprintf(stderr, "Error: nvrtcGetPTX not resolved\n");
227                 return false;
228         }
229
230         return true;
231 }
232
233 bool parse_parameters(int argc, const char **argv, CompilationSettings &settings)
234 {
235         OIIO::ArgParse ap;
236         ap.options("Usage: cycles_cubin_cc [options]",
237                 "-target %d", &settings.target_arch, "target shader model",
238                 "-m %d", &settings.bits, "Cuda architecture bits",
239                 "-i %s", &settings.input_file, "Input source filename",
240                 "-o %s", &settings.output_file, "Output cubin filename",
241                 "-I %L", &settings.includes, "Add additional includepath",
242                 "-D %L", &settings.defines, "Add additional defines",
243                 "-v", &settings.verbose, "Use verbose logging",
244                 "--use_fast_math", &settings.fast_math, "Use fast math",
245                 "-cuda-toolkit-dir %s", &settings.cuda_toolkit_dir, "path to the cuda toolkit binary directory",
246                 NULL);
247
248         if(ap.parse(argc, argv) < 0) {
249                 fprintf(stderr, "%s\n", ap.geterror().c_str());
250                 ap.usage();
251                 return false;
252         }
253
254         if(!settings.output_file.size()) {
255                 fprintf(stderr, "Error: Output file not set(-o), required\n\n");
256                 return false;
257         }
258
259         if(!settings.input_file.size()) {
260                 fprintf(stderr, "Error: Input file not set(-i, required\n\n");
261                 return false;
262         }
263
264         if(!settings.target_arch) {
265                 fprintf(stderr, "Error: target shader model not set (-target), required\n\n");
266                 return false;
267         }
268
269         return true;
270 }
271
272 int main(int argc, const char **argv)
273 {
274         CompilationSettings settings;
275
276         if(!parse_parameters(argc, argv, settings)) {
277                 fprintf(stderr, "Error: invalid parameters, exiting\n");
278                 exit(EXIT_FAILURE);
279         }
280
281         if(!init(settings)) {
282                 fprintf(stderr, "Error: initialization error, exiting\n");
283                 exit(EXIT_FAILURE);
284         }
285
286         if(!compile_cuda(settings)) {
287                 fprintf(stderr, "Error: compilation error, exiting\n");
288                 exit(EXIT_FAILURE);
289         }
290
291         if(!link_ptxas(settings)) {
292                 exit(EXIT_FAILURE);
293         }
294
295         return 0;
296 }