Cycles: Calculate size of split state buffer kernel side
[blender.git] / intern / cycles / device / device_cuda.cpp
1 /*
2  * Copyright 2011-2013 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 <climits>
18 #include <limits.h>
19 #include <stdio.h>
20 #include <stdlib.h>
21 #include <string.h>
22
23 #include "device.h"
24 #include "device_intern.h"
25 #include "device_split_kernel.h"
26
27 #include "buffers.h"
28
29 #ifdef WITH_CUDA_DYNLOAD
30 #  include "cuew.h"
31 #else
32 #  include "util_opengl.h"
33 #  include <cuda.h>
34 #  include <cudaGL.h>
35 #endif
36 #include "util_debug.h"
37 #include "util_logging.h"
38 #include "util_map.h"
39 #include "util_md5.h"
40 #include "util_opengl.h"
41 #include "util_path.h"
42 #include "util_string.h"
43 #include "util_system.h"
44 #include "util_types.h"
45 #include "util_time.h"
46
47 #include "split/kernel_split_data.h"
48
49 CCL_NAMESPACE_BEGIN
50
51 #ifndef WITH_CUDA_DYNLOAD
52
53 /* Transparently implement some functions, so majority of the file does not need
54  * to worry about difference between dynamically loaded and linked CUDA at all.
55  */
56
57 namespace {
58
59 const char *cuewErrorString(CUresult result)
60 {
61         /* We can only give error code here without major code duplication, that
62          * should be enough since dynamic loading is only being disabled by folks
63          * who knows what they're doing anyway.
64          *
65          * NOTE: Avoid call from several threads.
66          */
67         static string error;
68         error = string_printf("%d", result);
69         return error.c_str();
70 }
71
72 const char *cuewCompilerPath(void)
73 {
74         return CYCLES_CUDA_NVCC_EXECUTABLE;
75 }
76
77 int cuewCompilerVersion(void)
78 {
79         return (CUDA_VERSION / 100) + (CUDA_VERSION % 100 / 10);
80 }
81
82 }  /* namespace */
83 #endif  /* WITH_CUDA_DYNLOAD */
84
85 class CUDADevice;
86
87 class CUDASplitKernel : public DeviceSplitKernel {
88         CUDADevice *device;
89 public:
90         explicit CUDASplitKernel(CUDADevice *device);
91
92         virtual size_t state_buffer_size(device_memory& kg, device_memory& data, size_t num_threads);
93
94         virtual bool enqueue_split_kernel_data_init(const KernelDimensions& dim,
95                                                     RenderTile& rtile,
96                                                     int num_global_elements,
97                                                     device_memory& kernel_globals,
98                                                     device_memory& kernel_data_,
99                                                     device_memory& split_data,
100                                                     device_memory& ray_state,
101                                                     device_memory& queue_index,
102                                                     device_memory& use_queues_flag,
103                                                     device_memory& work_pool_wgs);
104
105         virtual SplitKernelFunction* get_split_kernel_function(string kernel_name, const DeviceRequestedFeatures&);
106         virtual int2 split_kernel_local_size();
107         virtual int2 split_kernel_global_size(device_memory& kg, device_memory& data, DeviceTask *task);
108 };
109
110 class CUDADevice : public Device
111 {
112 public:
113         DedicatedTaskPool task_pool;
114         CUdevice cuDevice;
115         CUcontext cuContext;
116         CUmodule cuModule;
117         map<device_ptr, bool> tex_interp_map;
118         map<device_ptr, uint> tex_bindless_map;
119         int cuDevId;
120         int cuDevArchitecture;
121         bool first_error;
122
123         struct PixelMem {
124                 GLuint cuPBO;
125                 CUgraphicsResource cuPBOresource;
126                 GLuint cuTexId;
127                 int w, h;
128         };
129
130         map<device_ptr, PixelMem> pixel_mem_map;
131
132         /* Bindless Textures */
133         device_vector<uint> bindless_mapping;
134         bool need_bindless_mapping;
135
136         CUdeviceptr cuda_device_ptr(device_ptr mem)
137         {
138                 return (CUdeviceptr)mem;
139         }
140
141         static bool have_precompiled_kernels()
142         {
143                 string cubins_path = path_get("lib");
144                 return path_exists(cubins_path);
145         }
146
147         virtual bool show_samples() const
148         {
149                 /* The CUDADevice only processes one tile at a time, so showing samples is fine. */
150                 return true;
151         }
152
153 /*#ifdef NDEBUG
154 #define cuda_abort()
155 #else
156 #define cuda_abort() abort()
157 #endif*/
158         void cuda_error_documentation()
159         {
160                 if(first_error) {
161                         fprintf(stderr, "\nRefer to the Cycles GPU rendering documentation for possible solutions:\n");
162                         fprintf(stderr, "https://docs.blender.org/manual/en/dev/render/cycles/gpu_rendering.html\n\n");
163                         first_error = false;
164                 }
165         }
166
167 #define cuda_assert(stmt) \
168         { \
169                 CUresult result = stmt; \
170                 \
171                 if(result != CUDA_SUCCESS) { \
172                         string message = string_printf("CUDA error: %s in %s", cuewErrorString(result), #stmt); \
173                         if(error_msg == "") \
174                                 error_msg = message; \
175                         fprintf(stderr, "%s\n", message.c_str()); \
176                         /*cuda_abort();*/ \
177                         cuda_error_documentation(); \
178                 } \
179         } (void)0
180
181         bool cuda_error_(CUresult result, const string& stmt)
182         {
183                 if(result == CUDA_SUCCESS)
184                         return false;
185
186                 string message = string_printf("CUDA error at %s: %s", stmt.c_str(), cuewErrorString(result));
187                 if(error_msg == "")
188                         error_msg = message;
189                 fprintf(stderr, "%s\n", message.c_str());
190                 cuda_error_documentation();
191                 return true;
192         }
193
194 #define cuda_error(stmt) cuda_error_(stmt, #stmt)
195
196         void cuda_error_message(const string& message)
197         {
198                 if(error_msg == "")
199                         error_msg = message;
200                 fprintf(stderr, "%s\n", message.c_str());
201                 cuda_error_documentation();
202         }
203
204         void cuda_push_context()
205         {
206                 cuda_assert(cuCtxSetCurrent(cuContext));
207         }
208
209         void cuda_pop_context()
210         {
211                 cuda_assert(cuCtxSetCurrent(NULL));
212         }
213
214         CUDADevice(DeviceInfo& info, Stats &stats, bool background_)
215         : Device(info, stats, background_)
216         {
217                 first_error = true;
218                 background = background_;
219
220                 cuDevId = info.num;
221                 cuDevice = 0;
222                 cuContext = 0;
223
224                 need_bindless_mapping = false;
225
226                 /* intialize */
227                 if(cuda_error(cuInit(0)))
228                         return;
229
230                 /* setup device and context */
231                 if(cuda_error(cuDeviceGet(&cuDevice, cuDevId)))
232                         return;
233
234                 CUresult result;
235
236                 if(background) {
237                         result = cuCtxCreate(&cuContext, 0, cuDevice);
238                 }
239                 else {
240                         result = cuGLCtxCreate(&cuContext, 0, cuDevice);
241
242                         if(result != CUDA_SUCCESS) {
243                                 result = cuCtxCreate(&cuContext, 0, cuDevice);
244                                 background = true;
245                         }
246                 }
247
248                 if(cuda_error_(result, "cuCtxCreate"))
249                         return;
250
251                 int major, minor;
252                 cuDeviceGetAttribute(&major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, cuDevId);
253                 cuDeviceGetAttribute(&minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, cuDevId);
254                 cuDevArchitecture = major*100 + minor*10;
255
256                 cuda_pop_context();
257         }
258
259         ~CUDADevice()
260         {
261                 task_pool.stop();
262
263                 if(info.has_bindless_textures) {
264                         tex_free(bindless_mapping);
265                 }
266
267                 cuda_assert(cuCtxDestroy(cuContext));
268         }
269
270         bool support_device(const DeviceRequestedFeatures& /*requested_features*/)
271         {
272                 int major, minor;
273                 cuDeviceGetAttribute(&major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, cuDevId);
274                 cuDeviceGetAttribute(&minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, cuDevId);
275
276                 /* We only support sm_20 and above */
277                 if(major < 2) {
278                         cuda_error_message(string_printf("CUDA device supported only with compute capability 2.0 or up, found %d.%d.", major, minor));
279                         return false;
280                 }
281
282                 return true;
283         }
284
285         bool use_adaptive_compilation()
286         {
287                 return DebugFlags().cuda.adaptive_compile;
288         }
289
290         bool use_split_kernel()
291         {
292                 return DebugFlags().cuda.split_kernel;
293         }
294
295         /* Common NVCC flags which stays the same regardless of shading model,
296          * kernel sources md5 and only depends on compiler or compilation settings.
297          */
298         string compile_kernel_get_common_cflags(
299                 const DeviceRequestedFeatures& requested_features, bool split=false)
300         {
301                 const int cuda_version = cuewCompilerVersion();
302                 const int machine = system_cpu_bits();
303                 const string kernel_path = path_get("kernel");
304                 const string include = kernel_path;
305                 string cflags = string_printf("-m%d "
306                                               "--ptxas-options=\"-v\" "
307                                               "--use_fast_math "
308                                               "-DNVCC "
309                                               "-D__KERNEL_CUDA_VERSION__=%d "
310                                                "-I\"%s\"",
311                                               machine,
312                                               cuda_version,
313                                               include.c_str());
314                 if(use_adaptive_compilation()) {
315                         cflags += " " + requested_features.get_build_options();
316                 }
317                 const char *extra_cflags = getenv("CYCLES_CUDA_EXTRA_CFLAGS");
318                 if(extra_cflags) {
319                         cflags += string(" ") + string(extra_cflags);
320                 }
321 #ifdef WITH_CYCLES_DEBUG
322                 cflags += " -D__KERNEL_DEBUG__";
323 #endif
324
325                 if(split) {
326                         cflags += " -D__SPLIT__";
327                 }
328
329                 return cflags;
330         }
331
332         bool compile_check_compiler() {
333                 const char *nvcc = cuewCompilerPath();
334                 if(nvcc == NULL) {
335                         cuda_error_message("CUDA nvcc compiler not found. "
336                                            "Install CUDA toolkit in default location.");
337                         return false;
338                 }
339                 const int cuda_version = cuewCompilerVersion();
340                 VLOG(1) << "Found nvcc " << nvcc
341                         << ", CUDA version " << cuda_version
342                         << ".";
343                 const int major = cuda_version / 10, minor = cuda_version & 10;
344                 if(cuda_version == 0) {
345                         cuda_error_message("CUDA nvcc compiler version could not be parsed.");
346                         return false;
347                 }
348                 if(cuda_version < 75) {
349                         printf("Unsupported CUDA version %d.%d detected, "
350                                "you need CUDA 7.5 or newer.\n",
351                                major, minor);
352                         return false;
353                 }
354                 else if(cuda_version != 75 && cuda_version != 80) {
355                         printf("CUDA version %d.%d detected, build may succeed but only "
356                                "CUDA 7.5 and 8.0 are officially supported.\n",
357                                major, minor);
358                 }
359                 return true;
360         }
361
362         string compile_kernel(const DeviceRequestedFeatures& requested_features, bool split=false)
363         {
364                 /* Compute cubin name. */
365                 int major, minor;
366                 cuDeviceGetAttribute(&major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, cuDevId);
367                 cuDeviceGetAttribute(&minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, cuDevId);
368
369                 /* Attempt to use kernel provided with Blender. */
370                 if(!use_adaptive_compilation()) {
371                         const string cubin = path_get(string_printf(split ? "lib/kernel_split_sm_%d%d.cubin"
372                                                                           : "lib/kernel_sm_%d%d.cubin",
373                                                                     major, minor));
374                         VLOG(1) << "Testing for pre-compiled kernel " << cubin << ".";
375                         if(path_exists(cubin)) {
376                                 VLOG(1) << "Using precompiled kernel.";
377                                 return cubin;
378                         }
379                 }
380
381                 const string common_cflags =
382                         compile_kernel_get_common_cflags(requested_features, split);
383
384                 /* Try to use locally compiled kernel. */
385                 const string kernel_path = path_get("kernel");
386                 const string kernel_md5 = path_files_md5_hash(kernel_path);
387
388                 /* We include cflags into md5 so changing cuda toolkit or changing other
389                  * compiler command line arguments makes sure cubin gets re-built.
390                  */
391                 const string cubin_md5 = util_md5_string(kernel_md5 + common_cflags);
392
393                 const string cubin_file = string_printf(split ? "cycles_kernel_split_sm%d%d_%s.cubin"
394                                                               : "cycles_kernel_sm%d%d_%s.cubin",
395                                                         major, minor,
396                                                         cubin_md5.c_str());
397                 const string cubin = path_cache_get(path_join("kernels", cubin_file));
398                 VLOG(1) << "Testing for locally compiled kernel " << cubin << ".";
399                 if(path_exists(cubin)) {
400                         VLOG(1) << "Using locally compiled kernel.";
401                         return cubin;
402                 }
403
404 #ifdef _WIN32
405                 if(have_precompiled_kernels()) {
406                         if(major < 2) {
407                                 cuda_error_message(string_printf(
408                                         "CUDA device requires compute capability 2.0 or up, "
409                                         "found %d.%d. Your GPU is not supported.",
410                                         major, minor));
411                         }
412                         else {
413                                 cuda_error_message(string_printf(
414                                         "CUDA binary kernel for this graphics card compute "
415                                         "capability (%d.%d) not found.",
416                                         major, minor));
417                         }
418                         return "";
419                 }
420 #endif
421
422                 /* Compile. */
423                 if(!compile_check_compiler()) {
424                         return "";
425                 }
426                 const char *nvcc = cuewCompilerPath();
427                 const string kernel = path_join(kernel_path,
428                                           path_join("kernels",
429                                                     path_join("cuda", split ? "kernel_split.cu" : "kernel.cu")));
430                 double starttime = time_dt();
431                 printf("Compiling CUDA kernel ...\n");
432
433                 path_create_directories(cubin);
434
435                 string command = string_printf("\"%s\" "
436                                                "-arch=sm_%d%d "
437                                                "--cubin \"%s\" "
438                                                "-o \"%s\" "
439                                                "%s ",
440                                                nvcc,
441                                                major, minor,
442                                                kernel.c_str(),
443                                                cubin.c_str(),
444                                                common_cflags.c_str());
445
446                 printf("%s\n", command.c_str());
447
448                 if(system(command.c_str()) == -1) {
449                         cuda_error_message("Failed to execute compilation command, "
450                                            "see console for details.");
451                         return "";
452                 }
453
454                 /* Verify if compilation succeeded */
455                 if(!path_exists(cubin)) {
456                         cuda_error_message("CUDA kernel compilation failed, "
457                                            "see console for details.");
458                         return "";
459                 }
460
461                 printf("Kernel compilation finished in %.2lfs.\n", time_dt() - starttime);
462
463                 return cubin;
464         }
465
466         bool load_kernels(const DeviceRequestedFeatures& requested_features)
467         {
468                 /* check if cuda init succeeded */
469                 if(cuContext == 0)
470                         return false;
471
472                 /* check if GPU is supported */
473                 if(!support_device(requested_features))
474                         return false;
475
476                 /* get kernel */
477                 string cubin = compile_kernel(requested_features, use_split_kernel());
478
479                 if(cubin == "")
480                         return false;
481
482                 /* open module */
483                 cuda_push_context();
484
485                 string cubin_data;
486                 CUresult result;
487
488                 if(path_read_text(cubin, cubin_data))
489                         result = cuModuleLoadData(&cuModule, cubin_data.c_str());
490                 else
491                         result = CUDA_ERROR_FILE_NOT_FOUND;
492
493                 if(cuda_error_(result, "cuModuleLoad"))
494                         cuda_error_message(string_printf("Failed loading CUDA kernel %s.", cubin.c_str()));
495
496                 cuda_pop_context();
497
498                 return (result == CUDA_SUCCESS);
499         }
500
501         void load_bindless_mapping()
502         {
503                 if(info.has_bindless_textures && need_bindless_mapping) {
504                         tex_free(bindless_mapping);
505                         tex_alloc("__bindless_mapping", bindless_mapping, INTERPOLATION_NONE, EXTENSION_REPEAT);
506                         need_bindless_mapping = false;
507                 }
508         }
509
510         void mem_alloc(const char *name, device_memory& mem, MemoryType /*type*/)
511         {
512                 if(name) {
513                         VLOG(1) << "Buffer allocate: " << name << ", "
514                                     << string_human_readable_number(mem.memory_size()) << " bytes. ("
515                                     << string_human_readable_size(mem.memory_size()) << ")";
516                 }
517
518                 cuda_push_context();
519                 CUdeviceptr device_pointer;
520                 size_t size = mem.memory_size();
521                 cuda_assert(cuMemAlloc(&device_pointer, size));
522                 mem.device_pointer = (device_ptr)device_pointer;
523                 mem.device_size = size;
524                 stats.mem_alloc(size);
525                 cuda_pop_context();
526         }
527
528         void mem_copy_to(device_memory& mem)
529         {
530                 cuda_push_context();
531                 if(mem.device_pointer)
532                         cuda_assert(cuMemcpyHtoD(cuda_device_ptr(mem.device_pointer), (void*)mem.data_pointer, mem.memory_size()));
533                 cuda_pop_context();
534         }
535
536         void mem_copy_from(device_memory& mem, int y, int w, int h, int elem)
537         {
538                 size_t offset = elem*y*w;
539                 size_t size = elem*w*h;
540
541                 cuda_push_context();
542                 if(mem.device_pointer) {
543                         cuda_assert(cuMemcpyDtoH((uchar*)mem.data_pointer + offset,
544                                                  (CUdeviceptr)(mem.device_pointer + offset), size));
545                 }
546                 else {
547                         memset((char*)mem.data_pointer + offset, 0, size);
548                 }
549                 cuda_pop_context();
550         }
551
552         void mem_zero(device_memory& mem)
553         {
554                 if(mem.data_pointer) {
555                         memset((void*)mem.data_pointer, 0, mem.memory_size());
556                 }
557
558                 cuda_push_context();
559                 if(mem.device_pointer)
560                         cuda_assert(cuMemsetD8(cuda_device_ptr(mem.device_pointer), 0, mem.memory_size()));
561                 cuda_pop_context();
562         }
563
564         void mem_free(device_memory& mem)
565         {
566                 if(mem.device_pointer) {
567                         cuda_push_context();
568                         cuda_assert(cuMemFree(cuda_device_ptr(mem.device_pointer)));
569                         cuda_pop_context();
570
571                         mem.device_pointer = 0;
572
573                         stats.mem_free(mem.device_size);
574                         mem.device_size = 0;
575                 }
576         }
577
578         void const_copy_to(const char *name, void *host, size_t size)
579         {
580                 CUdeviceptr mem;
581                 size_t bytes;
582
583                 cuda_push_context();
584                 cuda_assert(cuModuleGetGlobal(&mem, &bytes, cuModule, name));
585                 //assert(bytes == size);
586                 cuda_assert(cuMemcpyHtoD(mem, host, size));
587                 cuda_pop_context();
588         }
589
590         void tex_alloc(const char *name,
591                        device_memory& mem,
592                        InterpolationType interpolation,
593                        ExtensionType extension)
594         {
595                 VLOG(1) << "Texture allocate: " << name << ", "
596                         << string_human_readable_number(mem.memory_size()) << " bytes. ("
597                         << string_human_readable_size(mem.memory_size()) << ")";
598
599                 /* Check if we are on sm_30 or above.
600                  * We use arrays and bindles textures for storage there */
601                 bool has_bindless_textures = info.has_bindless_textures;
602
603                 /* General variables for both architectures */
604                 string bind_name = name;
605                 size_t dsize = datatype_size(mem.data_type);
606                 size_t size = mem.memory_size();
607
608                 CUaddress_mode address_mode = CU_TR_ADDRESS_MODE_WRAP;
609                 switch(extension) {
610                         case EXTENSION_REPEAT:
611                                 address_mode = CU_TR_ADDRESS_MODE_WRAP;
612                                 break;
613                         case EXTENSION_EXTEND:
614                                 address_mode = CU_TR_ADDRESS_MODE_CLAMP;
615                                 break;
616                         case EXTENSION_CLIP:
617                                 address_mode = CU_TR_ADDRESS_MODE_BORDER;
618                                 break;
619                         default:
620                                 assert(0);
621                                 break;
622                 }
623
624                 CUfilter_mode filter_mode;
625                 if(interpolation == INTERPOLATION_CLOSEST) {
626                         filter_mode = CU_TR_FILTER_MODE_POINT;
627                 }
628                 else {
629                         filter_mode = CU_TR_FILTER_MODE_LINEAR;
630                 }
631
632                 CUarray_format_enum format;
633                 switch(mem.data_type) {
634                         case TYPE_UCHAR: format = CU_AD_FORMAT_UNSIGNED_INT8; break;
635                         case TYPE_UINT: format = CU_AD_FORMAT_UNSIGNED_INT32; break;
636                         case TYPE_INT: format = CU_AD_FORMAT_SIGNED_INT32; break;
637                         case TYPE_FLOAT: format = CU_AD_FORMAT_FLOAT; break;
638                         case TYPE_HALF: format = CU_AD_FORMAT_HALF; break;
639                         default: assert(0); return;
640                 }
641
642                 /* General variables for Fermi */
643                 CUtexref texref = NULL;
644
645                 if(!has_bindless_textures) {
646                         if(mem.data_depth > 1) {
647                                 /* Kernel uses different bind names for 2d and 3d float textures,
648                                  * so we have to adjust couple of things here.
649                                  */
650                                 vector<string> tokens;
651                                 string_split(tokens, name, "_");
652                                 bind_name = string_printf("__tex_image_%s_3d_%s",
653                                                           tokens[2].c_str(),
654                                                           tokens[3].c_str());
655                         }
656
657                         cuda_push_context();
658                         cuda_assert(cuModuleGetTexRef(&texref, cuModule, bind_name.c_str()));
659                         cuda_pop_context();
660
661                         if(!texref) {
662                                 return;
663                         }
664                 }
665
666                 /* Data Storage */
667                 if(interpolation == INTERPOLATION_NONE) {
668                         if(has_bindless_textures) {
669                                 mem_alloc(NULL, mem, MEM_READ_ONLY);
670                                 mem_copy_to(mem);
671
672                                 cuda_push_context();
673
674                                 CUdeviceptr cumem;
675                                 size_t cubytes;
676
677                                 cuda_assert(cuModuleGetGlobal(&cumem, &cubytes, cuModule, bind_name.c_str()));
678
679                                 if(cubytes == 8) {
680                                         /* 64 bit device pointer */
681                                         uint64_t ptr = mem.device_pointer;
682                                         cuda_assert(cuMemcpyHtoD(cumem, (void*)&ptr, cubytes));
683                                 }
684                                 else {
685                                         /* 32 bit device pointer */
686                                         uint32_t ptr = (uint32_t)mem.device_pointer;
687                                         cuda_assert(cuMemcpyHtoD(cumem, (void*)&ptr, cubytes));
688                                 }
689
690                                 cuda_pop_context();
691                         }
692                         else {
693                                 mem_alloc(NULL, mem, MEM_READ_ONLY);
694                                 mem_copy_to(mem);
695
696                                 cuda_push_context();
697
698                                 cuda_assert(cuTexRefSetAddress(NULL, texref, cuda_device_ptr(mem.device_pointer), size));
699                                 cuda_assert(cuTexRefSetFilterMode(texref, CU_TR_FILTER_MODE_POINT));
700                                 cuda_assert(cuTexRefSetFlags(texref, CU_TRSF_READ_AS_INTEGER));
701
702                                 cuda_pop_context();
703                         }
704                 }
705                 /* Texture Storage */
706                 else {
707                         CUarray handle = NULL;
708
709                         cuda_push_context();
710
711                         if(mem.data_depth > 1) {
712                                 CUDA_ARRAY3D_DESCRIPTOR desc;
713
714                                 desc.Width = mem.data_width;
715                                 desc.Height = mem.data_height;
716                                 desc.Depth = mem.data_depth;
717                                 desc.Format = format;
718                                 desc.NumChannels = mem.data_elements;
719                                 desc.Flags = 0;
720
721                                 cuda_assert(cuArray3DCreate(&handle, &desc));
722                         }
723                         else {
724                                 CUDA_ARRAY_DESCRIPTOR desc;
725
726                                 desc.Width = mem.data_width;
727                                 desc.Height = mem.data_height;
728                                 desc.Format = format;
729                                 desc.NumChannels = mem.data_elements;
730
731                                 cuda_assert(cuArrayCreate(&handle, &desc));
732                         }
733
734                         if(!handle) {
735                                 cuda_pop_context();
736                                 return;
737                         }
738
739                         /* Allocate 3D, 2D or 1D memory */
740                         if(mem.data_depth > 1) {
741                                 CUDA_MEMCPY3D param;
742                                 memset(&param, 0, sizeof(param));
743                                 param.dstMemoryType = CU_MEMORYTYPE_ARRAY;
744                                 param.dstArray = handle;
745                                 param.srcMemoryType = CU_MEMORYTYPE_HOST;
746                                 param.srcHost = (void*)mem.data_pointer;
747                                 param.srcPitch = mem.data_width*dsize*mem.data_elements;
748                                 param.WidthInBytes = param.srcPitch;
749                                 param.Height = mem.data_height;
750                                 param.Depth = mem.data_depth;
751
752                                 cuda_assert(cuMemcpy3D(&param));
753                         }
754                         else if(mem.data_height > 1) {
755                                 CUDA_MEMCPY2D param;
756                                 memset(&param, 0, sizeof(param));
757                                 param.dstMemoryType = CU_MEMORYTYPE_ARRAY;
758                                 param.dstArray = handle;
759                                 param.srcMemoryType = CU_MEMORYTYPE_HOST;
760                                 param.srcHost = (void*)mem.data_pointer;
761                                 param.srcPitch = mem.data_width*dsize*mem.data_elements;
762                                 param.WidthInBytes = param.srcPitch;
763                                 param.Height = mem.data_height;
764
765                                 cuda_assert(cuMemcpy2D(&param));
766                         }
767                         else
768                                 cuda_assert(cuMemcpyHtoA(handle, 0, (void*)mem.data_pointer, size));
769
770                         /* Fermi and Kepler */
771                         mem.device_pointer = (device_ptr)handle;
772                         mem.device_size = size;
773
774                         stats.mem_alloc(size);
775
776                         /* Bindless Textures - Kepler */
777                         if(has_bindless_textures) {
778                                 int flat_slot = 0;
779                                 if(string_startswith(name, "__tex_image")) {
780                                         int pos =  string(name).rfind("_");
781                                         flat_slot = atoi(name + pos + 1);
782                                 }
783                                 else {
784                                         assert(0);
785                                 }
786
787                                 CUDA_RESOURCE_DESC resDesc;
788                                 memset(&resDesc, 0, sizeof(resDesc));
789                                 resDesc.resType = CU_RESOURCE_TYPE_ARRAY;
790                                 resDesc.res.array.hArray = handle;
791                                 resDesc.flags = 0;
792
793                                 CUDA_TEXTURE_DESC texDesc;
794                                 memset(&texDesc, 0, sizeof(texDesc));
795                                 texDesc.addressMode[0] = address_mode;
796                                 texDesc.addressMode[1] = address_mode;
797                                 texDesc.addressMode[2] = address_mode;
798                                 texDesc.filterMode = filter_mode;
799                                 texDesc.flags = CU_TRSF_NORMALIZED_COORDINATES;
800
801                                 CUtexObject tex = 0;
802                                 cuda_assert(cuTexObjectCreate(&tex, &resDesc, &texDesc, NULL));
803
804                                 /* Safety check */
805                                 if((uint)tex > UINT_MAX) {
806                                         assert(0);
807                                 }
808
809                                 /* Resize once */
810                                 if(flat_slot >= bindless_mapping.size()) {
811                                         /* Allocate some slots in advance, to reduce amount
812                                          * of re-allocations.
813                                          */
814                                         bindless_mapping.resize(flat_slot + 128);
815                                 }
816
817                                 /* Set Mapping and tag that we need to (re-)upload to device */
818                                 bindless_mapping.get_data()[flat_slot] = (uint)tex;
819                                 tex_bindless_map[mem.device_pointer] = (uint)tex;
820                                 need_bindless_mapping = true;
821                         }
822                         /* Regular Textures - Fermi */
823                         else {
824                                 cuda_assert(cuTexRefSetArray(texref, handle, CU_TRSA_OVERRIDE_FORMAT));
825                                 cuda_assert(cuTexRefSetFilterMode(texref, filter_mode));
826                                 cuda_assert(cuTexRefSetFlags(texref, CU_TRSF_NORMALIZED_COORDINATES));
827                         }
828
829                         cuda_pop_context();
830                 }
831
832                 /* Fermi, Data and Image Textures */
833                 if(!has_bindless_textures) {
834                         cuda_push_context();
835
836                         cuda_assert(cuTexRefSetAddressMode(texref, 0, address_mode));
837                         cuda_assert(cuTexRefSetAddressMode(texref, 1, address_mode));
838                         if(mem.data_depth > 1) {
839                                 cuda_assert(cuTexRefSetAddressMode(texref, 2, address_mode));
840                         }
841
842                         cuda_assert(cuTexRefSetFormat(texref, format, mem.data_elements));
843
844                         cuda_pop_context();
845                 }
846
847                 /* Fermi and Kepler */
848                 tex_interp_map[mem.device_pointer] = (interpolation != INTERPOLATION_NONE);
849         }
850
851         void tex_free(device_memory& mem)
852         {
853                 if(mem.device_pointer) {
854                         if(tex_interp_map[mem.device_pointer]) {
855                                 cuda_push_context();
856                                 cuArrayDestroy((CUarray)mem.device_pointer);
857                                 cuda_pop_context();
858
859                                 /* Free CUtexObject (Bindless Textures) */
860                                 if(info.has_bindless_textures && tex_bindless_map[mem.device_pointer]) {
861                                         uint flat_slot = tex_bindless_map[mem.device_pointer];
862                                         cuTexObjectDestroy(flat_slot);
863                                 }
864
865                                 tex_interp_map.erase(tex_interp_map.find(mem.device_pointer));
866                                 mem.device_pointer = 0;
867
868                                 stats.mem_free(mem.device_size);
869                                 mem.device_size = 0;
870                         }
871                         else {
872                                 tex_interp_map.erase(tex_interp_map.find(mem.device_pointer));
873                                 mem_free(mem);
874                         }
875                 }
876         }
877
878         void path_trace(RenderTile& rtile, int sample, bool branched)
879         {
880                 if(have_error())
881                         return;
882
883                 cuda_push_context();
884
885                 CUfunction cuPathTrace;
886                 CUdeviceptr d_buffer = cuda_device_ptr(rtile.buffer);
887                 CUdeviceptr d_rng_state = cuda_device_ptr(rtile.rng_state);
888
889                 /* get kernel function */
890                 if(branched) {
891                         cuda_assert(cuModuleGetFunction(&cuPathTrace, cuModule, "kernel_cuda_branched_path_trace"));
892                 }
893                 else {
894                         cuda_assert(cuModuleGetFunction(&cuPathTrace, cuModule, "kernel_cuda_path_trace"));
895                 }
896
897                 if(have_error())
898                         return;
899
900                 /* pass in parameters */
901                 void *args[] = {&d_buffer,
902                                 &d_rng_state,
903                                 &sample,
904                                 &rtile.x,
905                                 &rtile.y,
906                                 &rtile.w,
907                                 &rtile.h,
908                                 &rtile.offset,
909                                 &rtile.stride};
910
911                 /* launch kernel */
912                 int threads_per_block;
913                 cuda_assert(cuFuncGetAttribute(&threads_per_block, CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, cuPathTrace));
914
915                 /*int num_registers;
916                 cuda_assert(cuFuncGetAttribute(&num_registers, CU_FUNC_ATTRIBUTE_NUM_REGS, cuPathTrace));
917
918                 printf("threads_per_block %d\n", threads_per_block);
919                 printf("num_registers %d\n", num_registers);*/
920
921                 int xthreads = (int)sqrt(threads_per_block);
922                 int ythreads = (int)sqrt(threads_per_block);
923                 int xblocks = (rtile.w + xthreads - 1)/xthreads;
924                 int yblocks = (rtile.h + ythreads - 1)/ythreads;
925
926                 cuda_assert(cuFuncSetCacheConfig(cuPathTrace, CU_FUNC_CACHE_PREFER_L1));
927
928                 cuda_assert(cuLaunchKernel(cuPathTrace,
929                                            xblocks , yblocks, 1, /* blocks */
930                                            xthreads, ythreads, 1, /* threads */
931                                            0, 0, args, 0));
932
933                 cuda_assert(cuCtxSynchronize());
934
935                 cuda_pop_context();
936         }
937
938         void film_convert(DeviceTask& task, device_ptr buffer, device_ptr rgba_byte, device_ptr rgba_half)
939         {
940                 if(have_error())
941                         return;
942
943                 cuda_push_context();
944
945                 CUfunction cuFilmConvert;
946                 CUdeviceptr d_rgba = map_pixels((rgba_byte)? rgba_byte: rgba_half);
947                 CUdeviceptr d_buffer = cuda_device_ptr(buffer);
948
949                 /* get kernel function */
950                 if(rgba_half) {
951                         cuda_assert(cuModuleGetFunction(&cuFilmConvert, cuModule, "kernel_cuda_convert_to_half_float"));
952                 }
953                 else {
954                         cuda_assert(cuModuleGetFunction(&cuFilmConvert, cuModule, "kernel_cuda_convert_to_byte"));
955                 }
956
957
958                 float sample_scale = 1.0f/(task.sample + 1);
959
960                 /* pass in parameters */
961                 void *args[] = {&d_rgba,
962                                 &d_buffer,
963                                 &sample_scale,
964                                 &task.x,
965                                 &task.y,
966                                 &task.w,
967                                 &task.h,
968                                 &task.offset,
969                                 &task.stride};
970
971                 /* launch kernel */
972                 int threads_per_block;
973                 cuda_assert(cuFuncGetAttribute(&threads_per_block, CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, cuFilmConvert));
974
975                 int xthreads = (int)sqrt(threads_per_block);
976                 int ythreads = (int)sqrt(threads_per_block);
977                 int xblocks = (task.w + xthreads - 1)/xthreads;
978                 int yblocks = (task.h + ythreads - 1)/ythreads;
979
980                 cuda_assert(cuFuncSetCacheConfig(cuFilmConvert, CU_FUNC_CACHE_PREFER_L1));
981
982                 cuda_assert(cuLaunchKernel(cuFilmConvert,
983                                            xblocks , yblocks, 1, /* blocks */
984                                            xthreads, ythreads, 1, /* threads */
985                                            0, 0, args, 0));
986
987                 unmap_pixels((rgba_byte)? rgba_byte: rgba_half);
988
989                 cuda_pop_context();
990         }
991
992         void shader(DeviceTask& task)
993         {
994                 if(have_error())
995                         return;
996
997                 cuda_push_context();
998
999                 CUfunction cuShader;
1000                 CUdeviceptr d_input = cuda_device_ptr(task.shader_input);
1001                 CUdeviceptr d_output = cuda_device_ptr(task.shader_output);
1002                 CUdeviceptr d_output_luma = cuda_device_ptr(task.shader_output_luma);
1003
1004                 /* get kernel function */
1005                 if(task.shader_eval_type >= SHADER_EVAL_BAKE) {
1006                         cuda_assert(cuModuleGetFunction(&cuShader, cuModule, "kernel_cuda_bake"));
1007                 }
1008                 else {
1009                         cuda_assert(cuModuleGetFunction(&cuShader, cuModule, "kernel_cuda_shader"));
1010                 }
1011
1012                 /* do tasks in smaller chunks, so we can cancel it */
1013                 const int shader_chunk_size = 65536;
1014                 const int start = task.shader_x;
1015                 const int end = task.shader_x + task.shader_w;
1016                 int offset = task.offset;
1017
1018                 bool canceled = false;
1019                 for(int sample = 0; sample < task.num_samples && !canceled; sample++) {
1020                         for(int shader_x = start; shader_x < end; shader_x += shader_chunk_size) {
1021                                 int shader_w = min(shader_chunk_size, end - shader_x);
1022
1023                                 /* pass in parameters */
1024                                 void *args[8];
1025                                 int arg = 0;
1026                                 args[arg++] = &d_input;
1027                                 args[arg++] = &d_output;
1028                                 if(task.shader_eval_type < SHADER_EVAL_BAKE) {
1029                                         args[arg++] = &d_output_luma;
1030                                 }
1031                                 args[arg++] = &task.shader_eval_type;
1032                                 if(task.shader_eval_type >= SHADER_EVAL_BAKE) {
1033                                         args[arg++] = &task.shader_filter;
1034                                 }
1035                                 args[arg++] = &shader_x;
1036                                 args[arg++] = &shader_w;
1037                                 args[arg++] = &offset;
1038                                 args[arg++] = &sample;
1039
1040                                 /* launch kernel */
1041                                 int threads_per_block;
1042                                 cuda_assert(cuFuncGetAttribute(&threads_per_block, CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, cuShader));
1043
1044                                 int xblocks = (shader_w + threads_per_block - 1)/threads_per_block;
1045
1046                                 cuda_assert(cuFuncSetCacheConfig(cuShader, CU_FUNC_CACHE_PREFER_L1));
1047                                 cuda_assert(cuLaunchKernel(cuShader,
1048                                                            xblocks , 1, 1, /* blocks */
1049                                                            threads_per_block, 1, 1, /* threads */
1050                                                            0, 0, args, 0));
1051
1052                                 cuda_assert(cuCtxSynchronize());
1053
1054                                 if(task.get_cancel()) {
1055                                         canceled = true;
1056                                         break;
1057                                 }
1058                         }
1059
1060                         task.update_progress(NULL);
1061                 }
1062
1063                 cuda_pop_context();
1064         }
1065
1066         CUdeviceptr map_pixels(device_ptr mem)
1067         {
1068                 if(!background) {
1069                         PixelMem pmem = pixel_mem_map[mem];
1070                         CUdeviceptr buffer;
1071
1072                         size_t bytes;
1073                         cuda_assert(cuGraphicsMapResources(1, &pmem.cuPBOresource, 0));
1074                         cuda_assert(cuGraphicsResourceGetMappedPointer(&buffer, &bytes, pmem.cuPBOresource));
1075
1076                         return buffer;
1077                 }
1078
1079                 return cuda_device_ptr(mem);
1080         }
1081
1082         void unmap_pixels(device_ptr mem)
1083         {
1084                 if(!background) {
1085                         PixelMem pmem = pixel_mem_map[mem];
1086
1087                         cuda_assert(cuGraphicsUnmapResources(1, &pmem.cuPBOresource, 0));
1088                 }
1089         }
1090
1091         void pixels_alloc(device_memory& mem)
1092         {
1093                 if(!background) {
1094                         PixelMem pmem;
1095
1096                         pmem.w = mem.data_width;
1097                         pmem.h = mem.data_height;
1098
1099                         cuda_push_context();
1100
1101                         glGenBuffers(1, &pmem.cuPBO);
1102                         glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pmem.cuPBO);
1103                         if(mem.data_type == TYPE_HALF)
1104                                 glBufferData(GL_PIXEL_UNPACK_BUFFER, pmem.w*pmem.h*sizeof(GLhalf)*4, NULL, GL_DYNAMIC_DRAW);
1105                         else
1106                                 glBufferData(GL_PIXEL_UNPACK_BUFFER, pmem.w*pmem.h*sizeof(uint8_t)*4, NULL, GL_DYNAMIC_DRAW);
1107
1108                         glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);
1109
1110                         glGenTextures(1, &pmem.cuTexId);
1111                         glBindTexture(GL_TEXTURE_2D, pmem.cuTexId);
1112                         if(mem.data_type == TYPE_HALF)
1113                                 glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA16F_ARB, pmem.w, pmem.h, 0, GL_RGBA, GL_HALF_FLOAT, NULL);
1114                         else
1115                                 glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA8, pmem.w, pmem.h, 0, GL_RGBA, GL_UNSIGNED_BYTE, NULL);
1116                         glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
1117                         glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST);
1118                         glBindTexture(GL_TEXTURE_2D, 0);
1119
1120                         CUresult result = cuGraphicsGLRegisterBuffer(&pmem.cuPBOresource, pmem.cuPBO, CU_GRAPHICS_MAP_RESOURCE_FLAGS_NONE);
1121
1122                         if(result == CUDA_SUCCESS) {
1123                                 cuda_pop_context();
1124
1125                                 mem.device_pointer = pmem.cuTexId;
1126                                 pixel_mem_map[mem.device_pointer] = pmem;
1127
1128                                 mem.device_size = mem.memory_size();
1129                                 stats.mem_alloc(mem.device_size);
1130
1131                                 return;
1132                         }
1133                         else {
1134                                 /* failed to register buffer, fallback to no interop */
1135                                 glDeleteBuffers(1, &pmem.cuPBO);
1136                                 glDeleteTextures(1, &pmem.cuTexId);
1137
1138                                 cuda_pop_context();
1139
1140                                 background = true;
1141                         }
1142                 }
1143
1144                 Device::pixels_alloc(mem);
1145         }
1146
1147         void pixels_copy_from(device_memory& mem, int y, int w, int h)
1148         {
1149                 if(!background) {
1150                         PixelMem pmem = pixel_mem_map[mem.device_pointer];
1151
1152                         cuda_push_context();
1153
1154                         glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pmem.cuPBO);
1155                         uchar *pixels = (uchar*)glMapBuffer(GL_PIXEL_UNPACK_BUFFER, GL_READ_ONLY);
1156                         size_t offset = sizeof(uchar)*4*y*w;
1157                         memcpy((uchar*)mem.data_pointer + offset, pixels + offset, sizeof(uchar)*4*w*h);
1158                         glUnmapBuffer(GL_PIXEL_UNPACK_BUFFER);
1159                         glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);
1160
1161                         cuda_pop_context();
1162
1163                         return;
1164                 }
1165
1166                 Device::pixels_copy_from(mem, y, w, h);
1167         }
1168
1169         void pixels_free(device_memory& mem)
1170         {
1171                 if(mem.device_pointer) {
1172                         if(!background) {
1173                                 PixelMem pmem = pixel_mem_map[mem.device_pointer];
1174
1175                                 cuda_push_context();
1176
1177                                 cuda_assert(cuGraphicsUnregisterResource(pmem.cuPBOresource));
1178                                 glDeleteBuffers(1, &pmem.cuPBO);
1179                                 glDeleteTextures(1, &pmem.cuTexId);
1180
1181                                 cuda_pop_context();
1182
1183                                 pixel_mem_map.erase(pixel_mem_map.find(mem.device_pointer));
1184                                 mem.device_pointer = 0;
1185
1186                                 stats.mem_free(mem.device_size);
1187                                 mem.device_size = 0;
1188
1189                                 return;
1190                         }
1191
1192                         Device::pixels_free(mem);
1193                 }
1194         }
1195
1196         void draw_pixels(device_memory& mem, int y, int w, int h, int dx, int dy, int width, int height, bool transparent,
1197                 const DeviceDrawParams &draw_params)
1198         {
1199                 if(!background) {
1200                         PixelMem pmem = pixel_mem_map[mem.device_pointer];
1201                         float *vpointer;
1202
1203                         cuda_push_context();
1204
1205                         /* for multi devices, this assumes the inefficient method that we allocate
1206                          * all pixels on the device even though we only render to a subset */
1207                         size_t offset = 4*y*w;
1208
1209                         if(mem.data_type == TYPE_HALF)
1210                                 offset *= sizeof(GLhalf);
1211                         else
1212                                 offset *= sizeof(uint8_t);
1213
1214                         glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pmem.cuPBO);
1215                         glBindTexture(GL_TEXTURE_2D, pmem.cuTexId);
1216                         if(mem.data_type == TYPE_HALF)
1217                                 glTexSubImage2D(GL_TEXTURE_2D, 0, 0, 0, w, h, GL_RGBA, GL_HALF_FLOAT, (void*)offset);
1218                         else
1219                                 glTexSubImage2D(GL_TEXTURE_2D, 0, 0, 0, w, h, GL_RGBA, GL_UNSIGNED_BYTE, (void*)offset);
1220                         glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);
1221
1222                         glEnable(GL_TEXTURE_2D);
1223
1224                         if(transparent) {
1225                                 glEnable(GL_BLEND);
1226                                 glBlendFunc(GL_ONE, GL_ONE_MINUS_SRC_ALPHA);
1227                         }
1228
1229                         glColor3f(1.0f, 1.0f, 1.0f);
1230
1231                         if(draw_params.bind_display_space_shader_cb) {
1232                                 draw_params.bind_display_space_shader_cb();
1233                         }
1234
1235                         if(!vertex_buffer)
1236                                 glGenBuffers(1, &vertex_buffer);
1237
1238                         glBindBuffer(GL_ARRAY_BUFFER, vertex_buffer);
1239                         /* invalidate old contents - avoids stalling if buffer is still waiting in queue to be rendered */
1240                         glBufferData(GL_ARRAY_BUFFER, 16 * sizeof(float), NULL, GL_STREAM_DRAW);
1241
1242                         vpointer = (float *)glMapBuffer(GL_ARRAY_BUFFER, GL_WRITE_ONLY);
1243
1244                         if(vpointer) {
1245                                 /* texture coordinate - vertex pair */
1246                                 vpointer[0] = 0.0f;
1247                                 vpointer[1] = 0.0f;
1248                                 vpointer[2] = dx;
1249                                 vpointer[3] = dy;
1250
1251                                 vpointer[4] = (float)w/(float)pmem.w;
1252                                 vpointer[5] = 0.0f;
1253                                 vpointer[6] = (float)width + dx;
1254                                 vpointer[7] = dy;
1255
1256                                 vpointer[8] = (float)w/(float)pmem.w;
1257                                 vpointer[9] = (float)h/(float)pmem.h;
1258                                 vpointer[10] = (float)width + dx;
1259                                 vpointer[11] = (float)height + dy;
1260
1261                                 vpointer[12] = 0.0f;
1262                                 vpointer[13] = (float)h/(float)pmem.h;
1263                                 vpointer[14] = dx;
1264                                 vpointer[15] = (float)height + dy;
1265
1266                                 glUnmapBuffer(GL_ARRAY_BUFFER);
1267                         }
1268
1269                         glTexCoordPointer(2, GL_FLOAT, 4 * sizeof(float), 0);
1270                         glVertexPointer(2, GL_FLOAT, 4 * sizeof(float), (char *)NULL + 2 * sizeof(float));
1271
1272                         glEnableClientState(GL_VERTEX_ARRAY);
1273                         glEnableClientState(GL_TEXTURE_COORD_ARRAY);
1274
1275                         glDrawArrays(GL_TRIANGLE_FAN, 0, 4);
1276
1277                         glDisableClientState(GL_TEXTURE_COORD_ARRAY);
1278                         glDisableClientState(GL_VERTEX_ARRAY);
1279
1280                         glBindBuffer(GL_ARRAY_BUFFER, 0);
1281
1282                         if(draw_params.unbind_display_space_shader_cb) {
1283                                 draw_params.unbind_display_space_shader_cb();
1284                         }
1285
1286                         if(transparent)
1287                                 glDisable(GL_BLEND);
1288
1289                         glBindTexture(GL_TEXTURE_2D, 0);
1290                         glDisable(GL_TEXTURE_2D);
1291
1292                         cuda_pop_context();
1293
1294                         return;
1295                 }
1296
1297                 Device::draw_pixels(mem, y, w, h, dx, dy, width, height, transparent, draw_params);
1298         }
1299
1300         void thread_run(DeviceTask *task)
1301         {
1302                 if(task->type == DeviceTask::PATH_TRACE) {
1303                         RenderTile tile;
1304
1305                         bool branched = task->integrator_branched;
1306
1307                         /* Upload Bindless Mapping */
1308                         load_bindless_mapping();
1309
1310                         if(!use_split_kernel()) {
1311                                 /* keep rendering tiles until done */
1312                                 while(task->acquire_tile(this, tile)) {
1313                                         int start_sample = tile.start_sample;
1314                                         int end_sample = tile.start_sample + tile.num_samples;
1315
1316                                         for(int sample = start_sample; sample < end_sample; sample++) {
1317                                                 if(task->get_cancel()) {
1318                                                         if(task->need_finish_queue == false)
1319                                                                 break;
1320                                                 }
1321
1322                                                 path_trace(tile, sample, branched);
1323
1324                                                 tile.sample = sample + 1;
1325
1326                                                 task->update_progress(&tile, tile.w*tile.h);
1327                                         }
1328
1329                                         task->release_tile(tile);
1330                                 }
1331                         }
1332                         else {
1333                                 DeviceRequestedFeatures requested_features;
1334                                 if(!use_adaptive_compilation()) {
1335                                         requested_features.max_closure = 64;
1336                                 }
1337
1338                                 CUDASplitKernel split_kernel(this);
1339                                 split_kernel.load_kernels(requested_features);
1340
1341                                 while(task->acquire_tile(this, tile)) {
1342                                         device_memory void_buffer;
1343                                         split_kernel.path_trace(task, tile, void_buffer, void_buffer);
1344
1345                                         task->release_tile(tile);
1346
1347                                         if(task->get_cancel()) {
1348                                                 if(task->need_finish_queue == false)
1349                                                         break;
1350                                         }
1351                                 }
1352                         }
1353                 }
1354                 else if(task->type == DeviceTask::SHADER) {
1355                         /* Upload Bindless Mapping */
1356                         load_bindless_mapping();
1357
1358                         shader(*task);
1359
1360                         cuda_push_context();
1361                         cuda_assert(cuCtxSynchronize());
1362                         cuda_pop_context();
1363                 }
1364         }
1365
1366         class CUDADeviceTask : public DeviceTask {
1367         public:
1368                 CUDADeviceTask(CUDADevice *device, DeviceTask& task)
1369                 : DeviceTask(task)
1370                 {
1371                         run = function_bind(&CUDADevice::thread_run, device, this);
1372                 }
1373         };
1374
1375         int get_split_task_count(DeviceTask& /*task*/)
1376         {
1377                 return 1;
1378         }
1379
1380         void task_add(DeviceTask& task)
1381         {
1382                 if(task.type == DeviceTask::FILM_CONVERT) {
1383                         /* must be done in main thread due to opengl access */
1384                         film_convert(task, task.buffer, task.rgba_byte, task.rgba_half);
1385
1386                         cuda_push_context();
1387                         cuda_assert(cuCtxSynchronize());
1388                         cuda_pop_context();
1389                 }
1390                 else {
1391                         task_pool.push(new CUDADeviceTask(this, task));
1392                 }
1393         }
1394
1395         void task_wait()
1396         {
1397                 task_pool.wait();
1398         }
1399
1400         void task_cancel()
1401         {
1402                 task_pool.cancel();
1403         }
1404
1405         friend class CUDASplitKernelFunction;
1406         friend class CUDASplitKernel;
1407 };
1408
1409 /* redefine the cuda_assert macro so it can be used outside of the CUDADevice class
1410  * now that the definition of that class is complete
1411  */
1412 #undef cuda_assert
1413 #define cuda_assert(stmt) \
1414         { \
1415                 CUresult result = stmt; \
1416                 \
1417                 if(result != CUDA_SUCCESS) { \
1418                         string message = string_printf("CUDA error: %s in %s", cuewErrorString(result), #stmt); \
1419                         if(device->error_msg == "") \
1420                                 device->error_msg = message; \
1421                         fprintf(stderr, "%s\n", message.c_str()); \
1422                         /*cuda_abort();*/ \
1423                         device->cuda_error_documentation(); \
1424                 } \
1425         } (void)0
1426
1427 /* split kernel */
1428
1429 class CUDASplitKernelFunction : public SplitKernelFunction{
1430         CUDADevice* device;
1431         CUfunction func;
1432 public:
1433         CUDASplitKernelFunction(CUDADevice *device, CUfunction func) : device(device), func(func) {}
1434
1435         /* enqueue the kernel, returns false if there is an error */
1436         bool enqueue(const KernelDimensions &dim, device_memory &/*kg*/, device_memory &/*data*/)
1437         {
1438                 return enqueue(dim, NULL);
1439         }
1440
1441         /* enqueue the kernel, returns false if there is an error */
1442         bool enqueue(const KernelDimensions &dim, void *args[])
1443         {
1444                 device->cuda_push_context();
1445
1446                 if(device->have_error())
1447                         return false;
1448
1449                 /* we ignore dim.local_size for now, as this is faster */
1450                 int threads_per_block;
1451                 cuda_assert(cuFuncGetAttribute(&threads_per_block, CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, func));
1452
1453                 int xthreads = (int)sqrt(threads_per_block);
1454                 int ythreads = (int)sqrt(threads_per_block);
1455
1456                 int xblocks = (dim.global_size[0] + xthreads - 1)/xthreads;
1457                 int yblocks = (dim.global_size[1] + ythreads - 1)/ythreads;
1458
1459                 cuda_assert(cuFuncSetCacheConfig(func, CU_FUNC_CACHE_PREFER_L1));
1460
1461                 cuda_assert(cuLaunchKernel(func,
1462                                                xblocks , yblocks, 1, /* blocks */
1463                                                xthreads, ythreads, 1, /* threads */
1464                                                0, 0, args, 0));
1465
1466                 device->cuda_pop_context();
1467
1468                 return !device->have_error();
1469         }
1470 };
1471
1472 CUDASplitKernel::CUDASplitKernel(CUDADevice *device) : DeviceSplitKernel(device), device(device)
1473 {
1474 }
1475
1476 size_t CUDASplitKernel::state_buffer_size(device_memory& /*kg*/, device_memory& /*data*/, size_t num_threads)
1477 {
1478         device_vector<uint> size_buffer;
1479         size_buffer.resize(1);
1480         device->mem_alloc(NULL, size_buffer, MEM_READ_WRITE);
1481
1482         device->cuda_push_context();
1483
1484         uint threads = num_threads;
1485         CUdeviceptr d_size = device->cuda_device_ptr(size_buffer.device_pointer);
1486
1487         struct args_t {
1488                 uint* num_threads;
1489                 CUdeviceptr* size;
1490         };
1491
1492         args_t args = {
1493                 &threads,
1494                 &d_size
1495         };
1496
1497         CUfunction state_buffer_size;
1498         cuda_assert(cuModuleGetFunction(&state_buffer_size, device->cuModule, "kernel_cuda_state_buffer_size"));
1499
1500         cuda_assert(cuLaunchKernel(state_buffer_size,
1501                                        1, 1, 1,
1502                                        1, 1, 1,
1503                                        0, 0, &args, 0));
1504
1505         device->cuda_pop_context();
1506
1507         device->mem_copy_from(size_buffer, 0, 1, 1, sizeof(uint));
1508         device->mem_free(size_buffer);
1509
1510         return *size_buffer.get_data();
1511 }
1512
1513 bool CUDASplitKernel::enqueue_split_kernel_data_init(const KernelDimensions& dim,
1514                                     RenderTile& rtile,
1515                                     int num_global_elements,
1516                                     device_memory& /*kernel_globals*/,
1517                                     device_memory& /*kernel_data*/,
1518                                     device_memory& split_data,
1519                                     device_memory& ray_state,
1520                                     device_memory& queue_index,
1521                                     device_memory& use_queues_flag,
1522                                     device_memory& work_pool_wgs)
1523 {
1524         device->cuda_push_context();
1525
1526         CUdeviceptr d_split_data = device->cuda_device_ptr(split_data.device_pointer);
1527         CUdeviceptr d_ray_state = device->cuda_device_ptr(ray_state.device_pointer);
1528         CUdeviceptr d_queue_index = device->cuda_device_ptr(queue_index.device_pointer);
1529         CUdeviceptr d_use_queues_flag = device->cuda_device_ptr(use_queues_flag.device_pointer);
1530         CUdeviceptr d_work_pool_wgs = device->cuda_device_ptr(work_pool_wgs.device_pointer);
1531
1532         CUdeviceptr d_rng_state = device->cuda_device_ptr(rtile.rng_state);
1533         CUdeviceptr d_buffer = device->cuda_device_ptr(rtile.buffer);
1534
1535         int end_sample = rtile.start_sample + rtile.num_samples;
1536         int queue_size = dim.global_size[0] * dim.global_size[1];
1537
1538         struct args_t {
1539                 CUdeviceptr* split_data_buffer;
1540                 int* num_elements;
1541                 CUdeviceptr* ray_state;
1542                 CUdeviceptr* rng_state;
1543                 int* start_sample;
1544                 int* end_sample;
1545                 int* sx;
1546                 int* sy;
1547                 int* sw;
1548                 int* sh;
1549                 int* offset;
1550                 int* stride;
1551                 CUdeviceptr* queue_index;
1552                 int* queuesize;
1553                 CUdeviceptr* use_queues_flag;
1554                 CUdeviceptr* work_pool_wgs;
1555                 int* num_samples;
1556                 CUdeviceptr* buffer;
1557         };
1558
1559         args_t args = {
1560                 &d_split_data,
1561                 &num_global_elements,
1562                 &d_ray_state,
1563                 &d_rng_state,
1564                 &rtile.start_sample,
1565                 &end_sample,
1566                 &rtile.x,
1567                 &rtile.y,
1568                 &rtile.w,
1569                 &rtile.h,
1570                 &rtile.offset,
1571                 &rtile.stride,
1572                 &d_queue_index,
1573                 &queue_size,
1574                 &d_use_queues_flag,
1575                 &d_work_pool_wgs,
1576                 &rtile.num_samples,
1577                 &d_buffer
1578         };
1579
1580         CUfunction data_init;
1581         cuda_assert(cuModuleGetFunction(&data_init, device->cuModule, "kernel_cuda_path_trace_data_init"));
1582         if(device->have_error()) {
1583                 return false;
1584         }
1585
1586         CUDASplitKernelFunction(device, data_init).enqueue(dim, (void**)&args);
1587
1588         device->cuda_pop_context();
1589
1590         return !device->have_error();
1591 }
1592
1593 SplitKernelFunction* CUDASplitKernel::get_split_kernel_function(string kernel_name, const DeviceRequestedFeatures&)
1594 {
1595         CUfunction func;
1596
1597         device->cuda_push_context();
1598
1599         cuda_assert(cuModuleGetFunction(&func, device->cuModule, (string("kernel_cuda_") + kernel_name).data()));
1600         if(device->have_error()) {
1601                 device->cuda_error_message(string_printf("kernel \"kernel_cuda_%s\" not found in module", kernel_name.data()));
1602                 return NULL;
1603         }
1604
1605         device->cuda_pop_context();
1606
1607         return new CUDASplitKernelFunction(device, func);
1608 }
1609
1610 int2 CUDASplitKernel::split_kernel_local_size()
1611 {
1612         return make_int2(32, 1);
1613 }
1614
1615 int2 CUDASplitKernel::split_kernel_global_size(device_memory& /*kg*/, device_memory& /*data*/, DeviceTask */*task*/)
1616 {
1617         /* TODO(mai): implement something here to detect ideal work size */
1618         return make_int2(256, 256);
1619 }
1620
1621 bool device_cuda_init(void)
1622 {
1623 #ifdef WITH_CUDA_DYNLOAD
1624         static bool initialized = false;
1625         static bool result = false;
1626
1627         if(initialized)
1628                 return result;
1629
1630         initialized = true;
1631         int cuew_result = cuewInit();
1632         if(cuew_result == CUEW_SUCCESS) {
1633                 VLOG(1) << "CUEW initialization succeeded";
1634                 if(CUDADevice::have_precompiled_kernels()) {
1635                         VLOG(1) << "Found precompiled kernels";
1636                         result = true;
1637                 }
1638 #ifndef _WIN32
1639                 else if(cuewCompilerPath() != NULL) {
1640                         VLOG(1) << "Found CUDA compiler " << cuewCompilerPath();
1641                         result = true;
1642                 }
1643                 else {
1644                         VLOG(1) << "Neither precompiled kernels nor CUDA compiler wad found,"
1645                                 << " unable to use CUDA";
1646                 }
1647 #endif
1648         }
1649         else {
1650                 VLOG(1) << "CUEW initialization failed: "
1651                         << ((cuew_result == CUEW_ERROR_ATEXIT_FAILED)
1652                             ? "Error setting up atexit() handler"
1653                             : "Error opening the library");
1654         }
1655
1656         return result;
1657 #else  /* WITH_CUDA_DYNLOAD */
1658         return true;
1659 #endif /* WITH_CUDA_DYNLOAD */
1660 }
1661
1662 Device *device_cuda_create(DeviceInfo& info, Stats &stats, bool background)
1663 {
1664         return new CUDADevice(info, stats, background);
1665 }
1666
1667 void device_cuda_info(vector<DeviceInfo>& devices)
1668 {
1669         CUresult result;
1670         int count = 0;
1671
1672         result = cuInit(0);
1673         if(result != CUDA_SUCCESS) {
1674                 if(result != CUDA_ERROR_NO_DEVICE)
1675                         fprintf(stderr, "CUDA cuInit: %s\n", cuewErrorString(result));
1676                 return;
1677         }
1678
1679         result = cuDeviceGetCount(&count);
1680         if(result != CUDA_SUCCESS) {
1681                 fprintf(stderr, "CUDA cuDeviceGetCount: %s\n", cuewErrorString(result));
1682                 return;
1683         }
1684
1685         vector<DeviceInfo> display_devices;
1686
1687         for(int num = 0; num < count; num++) {
1688                 char name[256];
1689                 int attr;
1690
1691                 if(cuDeviceGetName(name, 256, num) != CUDA_SUCCESS)
1692                         continue;
1693
1694                 int major;
1695                 cuDeviceGetAttribute(&major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, num);
1696                 if(major < 2) {
1697                         continue;
1698                 }
1699
1700                 DeviceInfo info;
1701
1702                 info.type = DEVICE_CUDA;
1703                 info.description = string(name);
1704                 info.num = num;
1705
1706                 info.advanced_shading = (major >= 2);
1707                 info.has_bindless_textures = (major >= 3);
1708                 info.pack_images = false;
1709
1710                 int pci_location[3] = {0, 0, 0};
1711                 cuDeviceGetAttribute(&pci_location[0], CU_DEVICE_ATTRIBUTE_PCI_DOMAIN_ID, num);
1712                 cuDeviceGetAttribute(&pci_location[1], CU_DEVICE_ATTRIBUTE_PCI_BUS_ID, num);
1713                 cuDeviceGetAttribute(&pci_location[2], CU_DEVICE_ATTRIBUTE_PCI_DEVICE_ID, num);
1714                 info.id = string_printf("CUDA_%s_%04x:%02x:%02x",
1715                                         name,
1716                                         (unsigned int)pci_location[0],
1717                                         (unsigned int)pci_location[1],
1718                                         (unsigned int)pci_location[2]);
1719
1720                 /* if device has a kernel timeout, assume it is used for display */
1721                 if(cuDeviceGetAttribute(&attr, CU_DEVICE_ATTRIBUTE_KERNEL_EXEC_TIMEOUT, num) == CUDA_SUCCESS && attr == 1) {
1722                         info.description += " (Display)";
1723                         info.display_device = true;
1724                         display_devices.push_back(info);
1725                 }
1726                 else
1727                         devices.push_back(info);
1728         }
1729
1730         if(!display_devices.empty())
1731                 devices.insert(devices.end(), display_devices.begin(), display_devices.end());
1732 }
1733
1734 string device_cuda_capabilities(void)
1735 {
1736         CUresult result = cuInit(0);
1737         if(result != CUDA_SUCCESS) {
1738                 if(result != CUDA_ERROR_NO_DEVICE) {
1739                         return string("Error initializing CUDA: ") + cuewErrorString(result);
1740                 }
1741                 return "No CUDA device found\n";
1742         }
1743
1744         int count;
1745         result = cuDeviceGetCount(&count);
1746         if(result != CUDA_SUCCESS) {
1747                 return string("Error getting devices: ") + cuewErrorString(result);
1748         }
1749
1750         string capabilities = "";
1751         for(int num = 0; num < count; num++) {
1752                 char name[256];
1753                 if(cuDeviceGetName(name, 256, num) != CUDA_SUCCESS) {
1754                         continue;
1755                 }
1756                 capabilities += string("\t") + name + "\n";
1757                 int value;
1758 #define GET_ATTR(attr) \
1759                 { \
1760                         if(cuDeviceGetAttribute(&value, \
1761                                                 CU_DEVICE_ATTRIBUTE_##attr, \
1762                                                 num) == CUDA_SUCCESS) \
1763                         { \
1764                                 capabilities += string_printf("\t\tCU_DEVICE_ATTRIBUTE_" #attr "\t\t\t%d\n", \
1765                                                               value); \
1766                         } \
1767                 } (void)0
1768                 /* TODO(sergey): Strip all attributes which are not useful for us
1769                  * or does not depend on the driver.
1770                  */
1771                 GET_ATTR(MAX_THREADS_PER_BLOCK);
1772                 GET_ATTR(MAX_BLOCK_DIM_X);
1773                 GET_ATTR(MAX_BLOCK_DIM_Y);
1774                 GET_ATTR(MAX_BLOCK_DIM_Z);
1775                 GET_ATTR(MAX_GRID_DIM_X);
1776                 GET_ATTR(MAX_GRID_DIM_Y);
1777                 GET_ATTR(MAX_GRID_DIM_Z);
1778                 GET_ATTR(MAX_SHARED_MEMORY_PER_BLOCK);
1779                 GET_ATTR(SHARED_MEMORY_PER_BLOCK);
1780                 GET_ATTR(TOTAL_CONSTANT_MEMORY);
1781                 GET_ATTR(WARP_SIZE);
1782                 GET_ATTR(MAX_PITCH);
1783                 GET_ATTR(MAX_REGISTERS_PER_BLOCK);
1784                 GET_ATTR(REGISTERS_PER_BLOCK);
1785                 GET_ATTR(CLOCK_RATE);
1786                 GET_ATTR(TEXTURE_ALIGNMENT);
1787                 GET_ATTR(GPU_OVERLAP);
1788                 GET_ATTR(MULTIPROCESSOR_COUNT);
1789                 GET_ATTR(KERNEL_EXEC_TIMEOUT);
1790                 GET_ATTR(INTEGRATED);
1791                 GET_ATTR(CAN_MAP_HOST_MEMORY);
1792                 GET_ATTR(COMPUTE_MODE);
1793                 GET_ATTR(MAXIMUM_TEXTURE1D_WIDTH);
1794                 GET_ATTR(MAXIMUM_TEXTURE2D_WIDTH);
1795                 GET_ATTR(MAXIMUM_TEXTURE2D_HEIGHT);
1796                 GET_ATTR(MAXIMUM_TEXTURE3D_WIDTH);
1797                 GET_ATTR(MAXIMUM_TEXTURE3D_HEIGHT);
1798                 GET_ATTR(MAXIMUM_TEXTURE3D_DEPTH);
1799                 GET_ATTR(MAXIMUM_TEXTURE2D_LAYERED_WIDTH);
1800                 GET_ATTR(MAXIMUM_TEXTURE2D_LAYERED_HEIGHT);
1801                 GET_ATTR(MAXIMUM_TEXTURE2D_LAYERED_LAYERS);
1802                 GET_ATTR(MAXIMUM_TEXTURE2D_ARRAY_WIDTH);
1803                 GET_ATTR(MAXIMUM_TEXTURE2D_ARRAY_HEIGHT);
1804                 GET_ATTR(MAXIMUM_TEXTURE2D_ARRAY_NUMSLICES);
1805                 GET_ATTR(SURFACE_ALIGNMENT);
1806                 GET_ATTR(CONCURRENT_KERNELS);
1807                 GET_ATTR(ECC_ENABLED);
1808                 GET_ATTR(TCC_DRIVER);
1809                 GET_ATTR(MEMORY_CLOCK_RATE);
1810                 GET_ATTR(GLOBAL_MEMORY_BUS_WIDTH);
1811                 GET_ATTR(L2_CACHE_SIZE);
1812                 GET_ATTR(MAX_THREADS_PER_MULTIPROCESSOR);
1813                 GET_ATTR(ASYNC_ENGINE_COUNT);
1814                 GET_ATTR(UNIFIED_ADDRESSING);
1815                 GET_ATTR(MAXIMUM_TEXTURE1D_LAYERED_WIDTH);
1816                 GET_ATTR(MAXIMUM_TEXTURE1D_LAYERED_LAYERS);
1817                 GET_ATTR(CAN_TEX2D_GATHER);
1818                 GET_ATTR(MAXIMUM_TEXTURE2D_GATHER_WIDTH);
1819                 GET_ATTR(MAXIMUM_TEXTURE2D_GATHER_HEIGHT);
1820                 GET_ATTR(MAXIMUM_TEXTURE3D_WIDTH_ALTERNATE);
1821                 GET_ATTR(MAXIMUM_TEXTURE3D_HEIGHT_ALTERNATE);
1822                 GET_ATTR(MAXIMUM_TEXTURE3D_DEPTH_ALTERNATE);
1823                 GET_ATTR(TEXTURE_PITCH_ALIGNMENT);
1824                 GET_ATTR(MAXIMUM_TEXTURECUBEMAP_WIDTH);
1825                 GET_ATTR(MAXIMUM_TEXTURECUBEMAP_LAYERED_WIDTH);
1826                 GET_ATTR(MAXIMUM_TEXTURECUBEMAP_LAYERED_LAYERS);
1827                 GET_ATTR(MAXIMUM_SURFACE1D_WIDTH);
1828                 GET_ATTR(MAXIMUM_SURFACE2D_WIDTH);
1829                 GET_ATTR(MAXIMUM_SURFACE2D_HEIGHT);
1830                 GET_ATTR(MAXIMUM_SURFACE3D_WIDTH);
1831                 GET_ATTR(MAXIMUM_SURFACE3D_HEIGHT);
1832                 GET_ATTR(MAXIMUM_SURFACE3D_DEPTH);
1833                 GET_ATTR(MAXIMUM_SURFACE1D_LAYERED_WIDTH);
1834                 GET_ATTR(MAXIMUM_SURFACE1D_LAYERED_LAYERS);
1835                 GET_ATTR(MAXIMUM_SURFACE2D_LAYERED_WIDTH);
1836                 GET_ATTR(MAXIMUM_SURFACE2D_LAYERED_HEIGHT);
1837                 GET_ATTR(MAXIMUM_SURFACE2D_LAYERED_LAYERS);
1838                 GET_ATTR(MAXIMUM_SURFACECUBEMAP_WIDTH);
1839                 GET_ATTR(MAXIMUM_SURFACECUBEMAP_LAYERED_WIDTH);
1840                 GET_ATTR(MAXIMUM_SURFACECUBEMAP_LAYERED_LAYERS);
1841                 GET_ATTR(MAXIMUM_TEXTURE1D_LINEAR_WIDTH);
1842                 GET_ATTR(MAXIMUM_TEXTURE2D_LINEAR_WIDTH);
1843                 GET_ATTR(MAXIMUM_TEXTURE2D_LINEAR_HEIGHT);
1844                 GET_ATTR(MAXIMUM_TEXTURE2D_LINEAR_PITCH);
1845                 GET_ATTR(MAXIMUM_TEXTURE2D_MIPMAPPED_WIDTH);
1846                 GET_ATTR(MAXIMUM_TEXTURE2D_MIPMAPPED_HEIGHT);
1847                 GET_ATTR(COMPUTE_CAPABILITY_MAJOR);
1848                 GET_ATTR(COMPUTE_CAPABILITY_MINOR);
1849                 GET_ATTR(MAXIMUM_TEXTURE1D_MIPMAPPED_WIDTH);
1850                 GET_ATTR(STREAM_PRIORITIES_SUPPORTED);
1851                 GET_ATTR(GLOBAL_L1_CACHE_SUPPORTED);
1852                 GET_ATTR(LOCAL_L1_CACHE_SUPPORTED);
1853                 GET_ATTR(MAX_SHARED_MEMORY_PER_MULTIPROCESSOR);
1854                 GET_ATTR(MAX_REGISTERS_PER_MULTIPROCESSOR);
1855                 GET_ATTR(MANAGED_MEMORY);
1856                 GET_ATTR(MULTI_GPU_BOARD);
1857                 GET_ATTR(MULTI_GPU_BOARD_GROUP_ID);
1858 #undef GET_ATTR
1859                 capabilities += "\n";
1860         }
1861
1862         return capabilities;
1863 }
1864
1865 CCL_NAMESPACE_END