Blender  V3.3
cycles_cubin_cc.cpp
Go to the documentation of this file.
1 /* SPDX-License-Identifier: Apache-2.0
2  * Copyright 2017-2022 Blender Foundation */
3 
4 #include <stdint.h>
5 #include <stdio.h>
6 
7 #include <string>
8 #include <vector>
9 
10 #include <OpenImageIO/argparse.h>
11 #include <OpenImageIO/filesystem.h>
12 
13 #include "cuew.h"
14 
15 #ifdef _MSC_VER
16 # include <Windows.h>
17 #endif
18 
19 using std::string;
20 using std::vector;
21 
22 namespace std {
23 template<typename T> std::string to_string(const T &n)
24 {
25  std::ostringstream s;
26  s << n;
27  return s.str();
28 }
29 } // namespace std
30 
32  public:
34  : target_arch(0), bits(64), verbose(false), fast_math(false), ptx_only(false)
35  {
36  }
37 
39  string input_file;
40  string output_file;
41  string ptx_file;
45  int bits;
46  bool verbose;
47  bool fast_math;
48  bool ptx_only;
49 };
50 
51 static bool compile_cuda(CompilationSettings &settings)
52 {
53  const char *headers[] = {"stdlib.h", "float.h", "math.h", "stdio.h", "stddef.h"};
54  const char *header_content[] = {"\n", "\n", "\n", "\n", "\n"};
55 
56  printf("Building %s\n", settings.input_file.c_str());
57 
58  string code;
59  if (!OIIO::Filesystem::read_text_file(settings.input_file, code)) {
60  fprintf(stderr, "Error: unable to read %s\n", settings.input_file.c_str());
61  return false;
62  }
63 
65  for (size_t i = 0; i < settings.includes.size(); i++) {
66  options.push_back("-I" + settings.includes[i]);
67  }
68 
69  for (size_t i = 0; i < settings.defines.size(); i++) {
70  options.push_back("-D" + settings.defines[i]);
71  }
72  options.push_back("-D__KERNEL_CUDA_VERSION__=" + std::to_string(cuewNvrtcVersion()));
73  options.push_back("-arch=compute_" + std::to_string(settings.target_arch));
74  options.push_back("--device-as-default-execution-space");
75  options.push_back("-DCYCLES_CUBIN_CC");
76  options.push_back("--std=c++11");
77  if (settings.fast_math)
78  options.push_back("--use_fast_math");
79 
80  nvrtcProgram prog;
81  nvrtcResult result = nvrtcCreateProgram(&prog,
82  code.c_str(), // buffer
83  NULL, // name
84  sizeof(headers) / sizeof(void *), // numHeaders
85  header_content, // headers
86  headers); // includeNames
87 
88  if (result != NVRTC_SUCCESS) {
89  fprintf(stderr, "Error: nvrtcCreateProgram failed (%d)\n\n", (int)result);
90  return false;
91  }
92 
93  /* Transfer options to a classic C array. */
94  vector<const char *> opts(options.size());
95  for (size_t i = 0; i < options.size(); i++) {
96  opts[i] = options[i].c_str();
97  }
98 
99  result = nvrtcCompileProgram(prog, options.size(), &opts[0]);
100 
101  if (result != NVRTC_SUCCESS) {
102  fprintf(stderr, "Error: nvrtcCompileProgram failed (%d)\n\n", (int)result);
103 
104  size_t log_size;
105  nvrtcGetProgramLogSize(prog, &log_size);
106 
107  vector<char> log(log_size);
108  nvrtcGetProgramLog(prog, &log[0]);
109  fprintf(stderr, "%s\n", &log[0]);
110 
111  return false;
112  }
113 
114  /* Retrieve the ptx code. */
115  size_t ptx_size;
116  result = nvrtcGetPTXSize(prog, &ptx_size);
117  if (result != NVRTC_SUCCESS) {
118  fprintf(stderr, "Error: nvrtcGetPTXSize failed (%d)\n\n", (int)result);
119  return false;
120  }
121 
122  vector<char> ptx_code(ptx_size);
123  result = nvrtcGetPTX(prog, &ptx_code[0]);
124  if (result != NVRTC_SUCCESS) {
125  fprintf(stderr, "Error: nvrtcGetPTX failed (%d)\n\n", (int)result);
126  return false;
127  }
128  if (settings.ptx_only) {
129  settings.ptx_file = settings.output_file;
130  }
131  else {
132  /* Write a file in the temp folder with the ptx code. */
133  settings.ptx_file = OIIO::Filesystem::temp_directory_path() + "/" +
134  OIIO::Filesystem::unique_path();
135  }
136  FILE *f = fopen(settings.ptx_file.c_str(), "wb");
137  fwrite(&ptx_code[0], 1, ptx_size, f);
138  fclose(f);
139 
140  return true;
141 }
142 
143 static bool link_ptxas(CompilationSettings &settings)
144 {
145  string cudapath = "";
146  if (settings.cuda_toolkit_dir.size())
147  cudapath = settings.cuda_toolkit_dir + "/bin/";
148 
149  string ptx = "\"" + cudapath + "ptxas\" " + settings.ptx_file + " -o " + settings.output_file +
150  " --gpu-name sm_" + std::to_string(settings.target_arch) + " -m" +
151  std::to_string(settings.bits);
152 
153  if (settings.verbose) {
154  ptx += " --verbose";
155  printf("%s\n", ptx.c_str());
156  }
157 
158  int pxresult = system(ptx.c_str());
159  if (pxresult) {
160  fprintf(stderr, "Error: ptxas failed (%d)\n\n", pxresult);
161  return false;
162  }
163 
164  if (!OIIO::Filesystem::remove(settings.ptx_file)) {
165  fprintf(stderr, "Error: removing %s\n\n", settings.ptx_file.c_str());
166  }
167 
168  return true;
169 }
170 
171 static bool init(CompilationSettings &settings)
172 {
173 #ifdef _MSC_VER
174  if (settings.cuda_toolkit_dir.size()) {
175  SetDllDirectory((settings.cuda_toolkit_dir + "/bin").c_str());
176  }
177 #else
178  (void)settings;
179 #endif
180 
181  int cuewresult = cuewInit(CUEW_INIT_NVRTC);
182  if (cuewresult != CUEW_SUCCESS) {
183  fprintf(stderr, "Error: cuew init fialed (0x%d)\n\n", cuewresult);
184  return false;
185  }
186 
187  if (cuewNvrtcVersion() < 80) {
188  fprintf(stderr, "Error: only cuda 8 and higher is supported, %d\n\n", cuewCompilerVersion());
189  return false;
190  }
191 
192  if (!nvrtcCreateProgram) {
193  fprintf(stderr, "Error: nvrtcCreateProgram not resolved\n");
194  return false;
195  }
196 
197  if (!nvrtcCompileProgram) {
198  fprintf(stderr, "Error: nvrtcCompileProgram not resolved\n");
199  return false;
200  }
201 
202  if (!nvrtcGetProgramLogSize) {
203  fprintf(stderr, "Error: nvrtcGetProgramLogSize not resolved\n");
204  return false;
205  }
206 
207  if (!nvrtcGetProgramLog) {
208  fprintf(stderr, "Error: nvrtcGetProgramLog not resolved\n");
209  return false;
210  }
211 
212  if (!nvrtcGetPTXSize) {
213  fprintf(stderr, "Error: nvrtcGetPTXSize not resolved\n");
214  return false;
215  }
216 
217  if (!nvrtcGetPTX) {
218  fprintf(stderr, "Error: nvrtcGetPTX not resolved\n");
219  return false;
220  }
221 
222  return true;
223 }
224 
225 static bool parse_parameters(int argc, const char **argv, CompilationSettings &settings)
226 {
227  OIIO::ArgParse ap;
228  ap.options("Usage: cycles_cubin_cc [options]",
229  "-target %d",
230  &settings.target_arch,
231  "target shader model",
232  "-m %d",
233  &settings.bits,
234  "Cuda architecture bits",
235  "-i %s",
236  &settings.input_file,
237  "Input source filename",
238  "-o %s",
239  &settings.output_file,
240  "Output cubin filename",
241  "-I %L",
242  &settings.includes,
243  "Add additional includepath",
244  "-D %L",
245  &settings.defines,
246  "Add additional defines",
247  "-ptx",
248  &settings.ptx_only,
249  "emit PTX code",
250  "-v",
251  &settings.verbose,
252  "Use verbose logging",
253  "--use_fast_math",
254  &settings.fast_math,
255  "Use fast math",
256  "-cuda-toolkit-dir %s",
257  &settings.cuda_toolkit_dir,
258  "path to the cuda toolkit binary directory",
259  NULL);
260 
261  if (ap.parse(argc, argv) < 0) {
262  fprintf(stderr, "%s\n", ap.geterror().c_str());
263  ap.usage();
264  return false;
265  }
266 
267  if (!settings.output_file.size()) {
268  fprintf(stderr, "Error: Output file not set(-o), required\n\n");
269  return false;
270  }
271 
272  if (!settings.input_file.size()) {
273  fprintf(stderr, "Error: Input file not set(-i, required\n\n");
274  return false;
275  }
276 
277  if (!settings.target_arch) {
278  fprintf(stderr, "Error: target shader model not set (-target), required\n\n");
279  return false;
280  }
281 
282  return true;
283 }
284 
285 int main(int argc, const char **argv)
286 {
287  CompilationSettings settings;
288 
289  if (!parse_parameters(argc, argv, settings)) {
290  fprintf(stderr, "Error: invalid parameters, exiting\n");
291  exit(EXIT_FAILURE);
292  }
293 
294  if (!init(settings)) {
295  fprintf(stderr, "Error: initialization error, exiting\n");
296  exit(EXIT_FAILURE);
297  }
298 
299  if (!compile_cuda(settings)) {
300  fprintf(stderr, "Error: compilation error, exiting\n");
301  exit(EXIT_FAILURE);
302  }
303 
304  if (!settings.ptx_only) {
305  if (!link_ptxas(settings)) {
306  exit(EXIT_FAILURE);
307  }
308  }
309 
310  return 0;
311 }
in reality light always falls off quadratically Particle Retrieve the data of the particle that spawned the object for example to give variation to multiple instances of an object Point Retrieve information about points in a point cloud Retrieve the edges of an object as it appears to Cycles topology will always appear triangulated Convert a blackbody temperature to an RGB value Normal Generate a perturbed normal from an RGB normal map image Typically used for faking highly detailed surfaces Generate an OSL shader from a file or text data block Image Sample an image file as a texture Sky Generate a procedural sky texture Noise Generate fractal Perlin noise Wave Generate procedural bands or rings with noise Voronoi Generate Worley noise based on the distance to random points Typically used to generate textures such as or biological cells Brick Generate a procedural texture producing bricks Texture Retrieve multiple types of texture coordinates nTypically used as inputs for texture nodes Vector Convert a vector
vector< string > includes
vector< string > defines
int main(int argc, const char **argv)
static bool init(CompilationSettings &settings)
static bool link_ptxas(CompilationSettings &settings)
static bool parse_parameters(int argc, const char **argv, CompilationSettings &settings)
static bool compile_cuda(CompilationSettings &settings)
CCL_NAMESPACE_BEGIN struct Options options
SyclQueue void void size_t num_bytes void
ccl_device_inline float3 log(float3 v)
Definition: math_float3.h:397
#define T
bool remove(void *owner, const AttributeIDRef &attribute_id)
std::string to_string(const T &n)