ClangFormat: apply to source, most of intern
[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/device.h"
24 #include "device/device_denoising.h"
25 #include "device/device_intern.h"
26 #include "device/device_split_kernel.h"
27
28 #include "render/buffers.h"
29
30 #include "kernel/filter/filter_defines.h"
31
32 #ifdef WITH_CUDA_DYNLOAD
33 #  include "cuew.h"
34 #else
35 #  include "util/util_opengl.h"
36 #  include <cuda.h>
37 #  include <cudaGL.h>
38 #endif
39 #include "util/util_debug.h"
40 #include "util/util_foreach.h"
41 #include "util/util_logging.h"
42 #include "util/util_map.h"
43 #include "util/util_md5.h"
44 #include "util/util_opengl.h"
45 #include "util/util_path.h"
46 #include "util/util_string.h"
47 #include "util/util_system.h"
48 #include "util/util_types.h"
49 #include "util/util_time.h"
50
51 #include "kernel/split/kernel_split_data_types.h"
52
53 CCL_NAMESPACE_BEGIN
54
55 #ifndef WITH_CUDA_DYNLOAD
56
57 /* Transparently implement some functions, so majority of the file does not need
58  * to worry about difference between dynamically loaded and linked CUDA at all.
59  */
60
61 namespace {
62
63 const char *cuewErrorString(CUresult result)
64 {
65   /* We can only give error code here without major code duplication, that
66    * should be enough since dynamic loading is only being disabled by folks
67    * who knows what they're doing anyway.
68    *
69    * NOTE: Avoid call from several threads.
70    */
71   static string error;
72   error = string_printf("%d", result);
73   return error.c_str();
74 }
75
76 const char *cuewCompilerPath()
77 {
78   return CYCLES_CUDA_NVCC_EXECUTABLE;
79 }
80
81 int cuewCompilerVersion()
82 {
83   return (CUDA_VERSION / 100) + (CUDA_VERSION % 100 / 10);
84 }
85
86 } /* namespace */
87 #endif /* WITH_CUDA_DYNLOAD */
88
89 class CUDADevice;
90
91 class CUDASplitKernel : public DeviceSplitKernel {
92   CUDADevice *device;
93
94  public:
95   explicit CUDASplitKernel(CUDADevice *device);
96
97   virtual uint64_t state_buffer_size(device_memory &kg, device_memory &data, size_t num_threads);
98
99   virtual bool enqueue_split_kernel_data_init(const KernelDimensions &dim,
100                                               RenderTile &rtile,
101                                               int num_global_elements,
102                                               device_memory &kernel_globals,
103                                               device_memory &kernel_data_,
104                                               device_memory &split_data,
105                                               device_memory &ray_state,
106                                               device_memory &queue_index,
107                                               device_memory &use_queues_flag,
108                                               device_memory &work_pool_wgs);
109
110   virtual SplitKernelFunction *get_split_kernel_function(const string &kernel_name,
111                                                          const DeviceRequestedFeatures &);
112   virtual int2 split_kernel_local_size();
113   virtual int2 split_kernel_global_size(device_memory &kg, device_memory &data, DeviceTask *task);
114 };
115
116 /* Utility to push/pop CUDA context. */
117 class CUDAContextScope {
118  public:
119   CUDAContextScope(CUDADevice *device);
120   ~CUDAContextScope();
121
122  private:
123   CUDADevice *device;
124 };
125
126 class CUDADevice : public Device {
127  public:
128   DedicatedTaskPool task_pool;
129   CUdevice cuDevice;
130   CUcontext cuContext;
131   CUmodule cuModule, cuFilterModule;
132   size_t device_texture_headroom;
133   size_t device_working_headroom;
134   bool move_texture_to_host;
135   size_t map_host_used;
136   size_t map_host_limit;
137   int can_map_host;
138   int cuDevId;
139   int cuDevArchitecture;
140   bool first_error;
141   CUDASplitKernel *split_kernel;
142
143   struct CUDAMem {
144     CUDAMem() : texobject(0), array(0), map_host_pointer(0), free_map_host(false)
145     {
146     }
147
148     CUtexObject texobject;
149     CUarray array;
150     void *map_host_pointer;
151     bool free_map_host;
152   };
153   typedef map<device_memory *, CUDAMem> CUDAMemMap;
154   CUDAMemMap cuda_mem_map;
155
156   struct PixelMem {
157     GLuint cuPBO;
158     CUgraphicsResource cuPBOresource;
159     GLuint cuTexId;
160     int w, h;
161   };
162   map<device_ptr, PixelMem> pixel_mem_map;
163
164   /* Bindless Textures */
165   device_vector<TextureInfo> texture_info;
166   bool need_texture_info;
167
168   CUdeviceptr cuda_device_ptr(device_ptr mem)
169   {
170     return (CUdeviceptr)mem;
171   }
172
173   static bool have_precompiled_kernels()
174   {
175     string cubins_path = path_get("lib");
176     return path_exists(cubins_path);
177   }
178
179   virtual bool show_samples() const
180   {
181     /* The CUDADevice only processes one tile at a time, so showing samples is fine. */
182     return true;
183   }
184
185   virtual BVHLayoutMask get_bvh_layout_mask() const
186   {
187     return BVH_LAYOUT_BVH2;
188   }
189
190   /*#ifdef NDEBUG
191 #define cuda_abort()
192 #else
193 #define cuda_abort() abort()
194 #endif*/
195   void cuda_error_documentation()
196   {
197     if (first_error) {
198       fprintf(stderr,
199               "\nRefer to the Cycles GPU rendering documentation for possible solutions:\n");
200       fprintf(stderr,
201               "https://docs.blender.org/manual/en/dev/render/cycles/gpu_rendering.html\n\n");
202       first_error = false;
203     }
204   }
205
206 #define cuda_assert(stmt) \
207   { \
208     CUresult result = stmt; \
209 \
210     if (result != CUDA_SUCCESS) { \
211       string message = string_printf( \
212           "CUDA error: %s in %s, line %d", cuewErrorString(result), #stmt, __LINE__); \
213       if (error_msg == "") \
214         error_msg = message; \
215       fprintf(stderr, "%s\n", message.c_str()); \
216       /*cuda_abort();*/ \
217       cuda_error_documentation(); \
218     } \
219   } \
220   (void)0
221
222   bool cuda_error_(CUresult result, const string &stmt)
223   {
224     if (result == CUDA_SUCCESS)
225       return false;
226
227     string message = string_printf("CUDA error at %s: %s", stmt.c_str(), cuewErrorString(result));
228     if (error_msg == "")
229       error_msg = message;
230     fprintf(stderr, "%s\n", message.c_str());
231     cuda_error_documentation();
232     return true;
233   }
234
235 #define cuda_error(stmt) cuda_error_(stmt, #stmt)
236
237   void cuda_error_message(const string &message)
238   {
239     if (error_msg == "")
240       error_msg = message;
241     fprintf(stderr, "%s\n", message.c_str());
242     cuda_error_documentation();
243   }
244
245   CUDADevice(DeviceInfo &info, Stats &stats, Profiler &profiler, bool background_)
246       : Device(info, stats, profiler, background_),
247         texture_info(this, "__texture_info", MEM_TEXTURE)
248   {
249     first_error = true;
250     background = background_;
251
252     cuDevId = info.num;
253     cuDevice = 0;
254     cuContext = 0;
255
256     cuModule = 0;
257     cuFilterModule = 0;
258
259     split_kernel = NULL;
260
261     need_texture_info = false;
262
263     device_texture_headroom = 0;
264     device_working_headroom = 0;
265     move_texture_to_host = false;
266     map_host_limit = 0;
267     map_host_used = 0;
268     can_map_host = 0;
269
270     /* Intialize CUDA. */
271     if (cuda_error(cuInit(0)))
272       return;
273
274     /* Setup device and context. */
275     if (cuda_error(cuDeviceGet(&cuDevice, cuDevId)))
276       return;
277
278     /* CU_CTX_MAP_HOST for mapping host memory when out of device memory.
279      * CU_CTX_LMEM_RESIZE_TO_MAX for reserving local memory ahead of render,
280      * so we can predict which memory to map to host. */
281     cuda_assert(
282         cuDeviceGetAttribute(&can_map_host, CU_DEVICE_ATTRIBUTE_CAN_MAP_HOST_MEMORY, cuDevice));
283
284     unsigned int ctx_flags = CU_CTX_LMEM_RESIZE_TO_MAX;
285     if (can_map_host) {
286       ctx_flags |= CU_CTX_MAP_HOST;
287       init_host_memory();
288     }
289
290     /* Create context. */
291     CUresult result;
292
293     if (background) {
294       result = cuCtxCreate(&cuContext, ctx_flags, cuDevice);
295     }
296     else {
297       result = cuGLCtxCreate(&cuContext, ctx_flags, cuDevice);
298
299       if (result != CUDA_SUCCESS) {
300         result = cuCtxCreate(&cuContext, ctx_flags, cuDevice);
301         background = true;
302       }
303     }
304
305     if (cuda_error_(result, "cuCtxCreate"))
306       return;
307
308     int major, minor;
309     cuDeviceGetAttribute(&major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, cuDevId);
310     cuDeviceGetAttribute(&minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, cuDevId);
311     cuDevArchitecture = major * 100 + minor * 10;
312
313     /* Pop context set by cuCtxCreate. */
314     cuCtxPopCurrent(NULL);
315   }
316
317   ~CUDADevice()
318   {
319     task_pool.stop();
320
321     delete split_kernel;
322
323     texture_info.free();
324
325     cuda_assert(cuCtxDestroy(cuContext));
326   }
327
328   bool support_device(const DeviceRequestedFeatures & /*requested_features*/)
329   {
330     int major, minor;
331     cuDeviceGetAttribute(&major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, cuDevId);
332     cuDeviceGetAttribute(&minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, cuDevId);
333
334     /* We only support sm_30 and above */
335     if (major < 3) {
336       cuda_error_message(string_printf(
337           "CUDA device supported only with compute capability 3.0 or up, found %d.%d.",
338           major,
339           minor));
340       return false;
341     }
342
343     return true;
344   }
345
346   bool use_adaptive_compilation()
347   {
348     return DebugFlags().cuda.adaptive_compile;
349   }
350
351   bool use_split_kernel()
352   {
353     return DebugFlags().cuda.split_kernel;
354   }
355
356   /* Common NVCC flags which stays the same regardless of shading model,
357    * kernel sources md5 and only depends on compiler or compilation settings.
358    */
359   string compile_kernel_get_common_cflags(const DeviceRequestedFeatures &requested_features,
360                                           bool filter = false,
361                                           bool split = false)
362   {
363     const int machine = system_cpu_bits();
364     const string source_path = path_get("source");
365     const string include_path = source_path;
366     string cflags = string_printf(
367         "-m%d "
368         "--ptxas-options=\"-v\" "
369         "--use_fast_math "
370         "-DNVCC "
371         "-I\"%s\"",
372         machine,
373         include_path.c_str());
374     if (!filter && use_adaptive_compilation()) {
375       cflags += " " + requested_features.get_build_options();
376     }
377     const char *extra_cflags = getenv("CYCLES_CUDA_EXTRA_CFLAGS");
378     if (extra_cflags) {
379       cflags += string(" ") + string(extra_cflags);
380     }
381 #ifdef WITH_CYCLES_DEBUG
382     cflags += " -D__KERNEL_DEBUG__";
383 #endif
384
385     if (split) {
386       cflags += " -D__SPLIT__";
387     }
388
389     return cflags;
390   }
391
392   bool compile_check_compiler()
393   {
394     const char *nvcc = cuewCompilerPath();
395     if (nvcc == NULL) {
396       cuda_error_message(
397           "CUDA nvcc compiler not found. "
398           "Install CUDA toolkit in default location.");
399       return false;
400     }
401     const int cuda_version = cuewCompilerVersion();
402     VLOG(1) << "Found nvcc " << nvcc << ", CUDA version " << cuda_version << ".";
403     const int major = cuda_version / 10, minor = cuda_version % 10;
404     if (cuda_version == 0) {
405       cuda_error_message("CUDA nvcc compiler version could not be parsed.");
406       return false;
407     }
408     if (cuda_version < 80) {
409       printf(
410           "Unsupported CUDA version %d.%d detected, "
411           "you need CUDA 8.0 or newer.\n",
412           major,
413           minor);
414       return false;
415     }
416     else if (cuda_version != 101) {
417       printf(
418           "CUDA version %d.%d detected, build may succeed but only "
419           "CUDA 10.1 is officially supported.\n",
420           major,
421           minor);
422     }
423     return true;
424   }
425
426   string compile_kernel(const DeviceRequestedFeatures &requested_features,
427                         bool filter = false,
428                         bool split = false)
429   {
430     const char *name, *source;
431     if (filter) {
432       name = "filter";
433       source = "filter.cu";
434     }
435     else if (split) {
436       name = "kernel_split";
437       source = "kernel_split.cu";
438     }
439     else {
440       name = "kernel";
441       source = "kernel.cu";
442     }
443     /* Compute cubin name. */
444     int major, minor;
445     cuDeviceGetAttribute(&major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, cuDevId);
446     cuDeviceGetAttribute(&minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, cuDevId);
447
448     /* Attempt to use kernel provided with Blender. */
449     if (!use_adaptive_compilation()) {
450       const string cubin = path_get(string_printf("lib/%s_sm_%d%d.cubin", name, major, minor));
451       VLOG(1) << "Testing for pre-compiled kernel " << cubin << ".";
452       if (path_exists(cubin)) {
453         VLOG(1) << "Using precompiled kernel.";
454         return cubin;
455       }
456     }
457
458     const string common_cflags = compile_kernel_get_common_cflags(
459         requested_features, filter, split);
460
461     /* Try to use locally compiled kernel. */
462     const string source_path = path_get("source");
463     const string kernel_md5 = path_files_md5_hash(source_path);
464
465     /* We include cflags into md5 so changing cuda toolkit or changing other
466      * compiler command line arguments makes sure cubin gets re-built.
467      */
468     const string cubin_md5 = util_md5_string(kernel_md5 + common_cflags);
469
470     const string cubin_file = string_printf(
471         "cycles_%s_sm%d%d_%s.cubin", name, major, minor, cubin_md5.c_str());
472     const string cubin = path_cache_get(path_join("kernels", cubin_file));
473     VLOG(1) << "Testing for locally compiled kernel " << cubin << ".";
474     if (path_exists(cubin)) {
475       VLOG(1) << "Using locally compiled kernel.";
476       return cubin;
477     }
478
479 #ifdef _WIN32
480     if (have_precompiled_kernels()) {
481       if (major < 3) {
482         cuda_error_message(
483             string_printf("CUDA device requires compute capability 3.0 or up, "
484                           "found %d.%d. Your GPU is not supported.",
485                           major,
486                           minor));
487       }
488       else {
489         cuda_error_message(
490             string_printf("CUDA binary kernel for this graphics card compute "
491                           "capability (%d.%d) not found.",
492                           major,
493                           minor));
494       }
495       return "";
496     }
497 #endif
498
499     /* Compile. */
500     if (!compile_check_compiler()) {
501       return "";
502     }
503     const char *nvcc = cuewCompilerPath();
504     const string kernel = path_join(path_join(source_path, "kernel"),
505                                     path_join("kernels", path_join("cuda", source)));
506     double starttime = time_dt();
507     printf("Compiling CUDA kernel ...\n");
508
509     path_create_directories(cubin);
510
511     string command = string_printf(
512         "\"%s\" "
513         "-arch=sm_%d%d "
514         "--cubin \"%s\" "
515         "-o \"%s\" "
516         "%s ",
517         nvcc,
518         major,
519         minor,
520         kernel.c_str(),
521         cubin.c_str(),
522         common_cflags.c_str());
523
524     printf("%s\n", command.c_str());
525
526     if (system(command.c_str()) == -1) {
527       cuda_error_message(
528           "Failed to execute compilation command, "
529           "see console for details.");
530       return "";
531     }
532
533     /* Verify if compilation succeeded */
534     if (!path_exists(cubin)) {
535       cuda_error_message(
536           "CUDA kernel compilation failed, "
537           "see console for details.");
538       return "";
539     }
540
541     printf("Kernel compilation finished in %.2lfs.\n", time_dt() - starttime);
542
543     return cubin;
544   }
545
546   bool load_kernels(const DeviceRequestedFeatures &requested_features)
547   {
548     /* TODO(sergey): Support kernels re-load for CUDA devices.
549      *
550      * Currently re-loading kernel will invalidate memory pointers,
551      * causing problems in cuCtxSynchronize.
552      */
553     if (cuFilterModule && cuModule) {
554       VLOG(1) << "Skipping kernel reload, not currently supported.";
555       return true;
556     }
557
558     /* check if cuda init succeeded */
559     if (cuContext == 0)
560       return false;
561
562     /* check if GPU is supported */
563     if (!support_device(requested_features))
564       return false;
565
566     /* get kernel */
567     string cubin = compile_kernel(requested_features, false, use_split_kernel());
568     if (cubin == "")
569       return false;
570
571     string filter_cubin = compile_kernel(requested_features, true, false);
572     if (filter_cubin == "")
573       return false;
574
575     /* open module */
576     CUDAContextScope scope(this);
577
578     string cubin_data;
579     CUresult result;
580
581     if (path_read_text(cubin, cubin_data))
582       result = cuModuleLoadData(&cuModule, cubin_data.c_str());
583     else
584       result = CUDA_ERROR_FILE_NOT_FOUND;
585
586     if (cuda_error_(result, "cuModuleLoad"))
587       cuda_error_message(string_printf("Failed loading CUDA kernel %s.", cubin.c_str()));
588
589     if (path_read_text(filter_cubin, cubin_data))
590       result = cuModuleLoadData(&cuFilterModule, cubin_data.c_str());
591     else
592       result = CUDA_ERROR_FILE_NOT_FOUND;
593
594     if (cuda_error_(result, "cuModuleLoad"))
595       cuda_error_message(string_printf("Failed loading CUDA kernel %s.", filter_cubin.c_str()));
596
597     if (result == CUDA_SUCCESS) {
598       reserve_local_memory(requested_features);
599     }
600
601     return (result == CUDA_SUCCESS);
602   }
603
604   void reserve_local_memory(const DeviceRequestedFeatures &requested_features)
605   {
606     if (use_split_kernel()) {
607       /* Split kernel mostly uses global memory and adaptive compilation,
608        * difficult to predict how much is needed currently. */
609       return;
610     }
611
612     /* Together with CU_CTX_LMEM_RESIZE_TO_MAX, this reserves local memory
613      * needed for kernel launches, so that we can reliably figure out when
614      * to allocate scene data in mapped host memory. */
615     CUDAContextScope scope(this);
616
617     size_t total = 0, free_before = 0, free_after = 0;
618     cuMemGetInfo(&free_before, &total);
619
620     /* Get kernel function. */
621     CUfunction cuPathTrace;
622
623     if (requested_features.use_integrator_branched) {
624       cuda_assert(cuModuleGetFunction(&cuPathTrace, cuModule, "kernel_cuda_branched_path_trace"));
625     }
626     else {
627       cuda_assert(cuModuleGetFunction(&cuPathTrace, cuModule, "kernel_cuda_path_trace"));
628     }
629
630     cuda_assert(cuFuncSetCacheConfig(cuPathTrace, CU_FUNC_CACHE_PREFER_L1));
631
632     int min_blocks, num_threads_per_block;
633     cuda_assert(cuOccupancyMaxPotentialBlockSize(
634         &min_blocks, &num_threads_per_block, cuPathTrace, NULL, 0, 0));
635
636     /* Launch kernel, using just 1 block appears sufficient to reserve
637      * memory for all multiprocessors. It would be good to do this in
638      * parallel for the multi GPU case still to make it faster. */
639     CUdeviceptr d_work_tiles = 0;
640     uint total_work_size = 0;
641
642     void *args[] = {&d_work_tiles, &total_work_size};
643
644     cuda_assert(cuLaunchKernel(cuPathTrace, 1, 1, 1, num_threads_per_block, 1, 1, 0, 0, args, 0));
645
646     cuda_assert(cuCtxSynchronize());
647
648     cuMemGetInfo(&free_after, &total);
649     VLOG(1) << "Local memory reserved " << string_human_readable_number(free_before - free_after)
650             << " bytes. (" << string_human_readable_size(free_before - free_after) << ")";
651
652 #if 0
653     /* For testing mapped host memory, fill up device memory. */
654     const size_t keep_mb = 1024;
655
656     while(free_after > keep_mb * 1024 * 1024LL) {
657       CUdeviceptr tmp;
658       cuda_assert(cuMemAlloc(&tmp, 10 * 1024 * 1024LL));
659       cuMemGetInfo(&free_after, &total);
660     }
661 #endif
662   }
663
664   void init_host_memory()
665   {
666     /* Limit amount of host mapped memory, because allocating too much can
667      * cause system instability. Leave at least half or 4 GB of system
668      * memory free, whichever is smaller. */
669     size_t default_limit = 4 * 1024 * 1024 * 1024LL;
670     size_t system_ram = system_physical_ram();
671
672     if (system_ram > 0) {
673       if (system_ram / 2 > default_limit) {
674         map_host_limit = system_ram - default_limit;
675       }
676       else {
677         map_host_limit = system_ram / 2;
678       }
679     }
680     else {
681       VLOG(1) << "Mapped host memory disabled, failed to get system RAM";
682       map_host_limit = 0;
683     }
684
685     /* Amount of device memory to keep is free after texture memory
686      * and working memory allocations respectively. We set the working
687      * memory limit headroom lower so that some space is left after all
688      * texture memory allocations. */
689     device_working_headroom = 32 * 1024 * 1024LL;   // 32MB
690     device_texture_headroom = 128 * 1024 * 1024LL;  // 128MB
691
692     VLOG(1) << "Mapped host memory limit set to " << string_human_readable_number(map_host_limit)
693             << " bytes. (" << string_human_readable_size(map_host_limit) << ")";
694   }
695
696   void load_texture_info()
697   {
698     if (need_texture_info) {
699       texture_info.copy_to_device();
700       need_texture_info = false;
701     }
702   }
703
704   void move_textures_to_host(size_t size, bool for_texture)
705   {
706     /* Signal to reallocate textures in host memory only. */
707     move_texture_to_host = true;
708
709     while (size > 0) {
710       /* Find suitable memory allocation to move. */
711       device_memory *max_mem = NULL;
712       size_t max_size = 0;
713       bool max_is_image = false;
714
715       foreach (CUDAMemMap::value_type &pair, cuda_mem_map) {
716         device_memory &mem = *pair.first;
717         CUDAMem *cmem = &pair.second;
718
719         bool is_texture = (mem.type == MEM_TEXTURE) && (&mem != &texture_info);
720         bool is_image = is_texture && (mem.data_height > 1);
721
722         /* Can't move this type of memory. */
723         if (!is_texture || cmem->array) {
724           continue;
725         }
726
727         /* Already in host memory. */
728         if (cmem->map_host_pointer) {
729           continue;
730         }
731
732         /* For other textures, only move image textures. */
733         if (for_texture && !is_image) {
734           continue;
735         }
736
737         /* Try to move largest allocation, prefer moving images. */
738         if (is_image > max_is_image || (is_image == max_is_image && mem.device_size > max_size)) {
739           max_is_image = is_image;
740           max_size = mem.device_size;
741           max_mem = &mem;
742         }
743       }
744
745       /* Move to host memory. This part is mutex protected since
746        * multiple CUDA devices could be moving the memory. The
747        * first one will do it, and the rest will adopt the pointer. */
748       if (max_mem) {
749         VLOG(1) << "Move memory from device to host: " << max_mem->name;
750
751         static thread_mutex move_mutex;
752         thread_scoped_lock lock(move_mutex);
753
754         /* Preserve the original device pointer, in case of multi device
755          * we can't change it because the pointer mapping would break. */
756         device_ptr prev_pointer = max_mem->device_pointer;
757         size_t prev_size = max_mem->device_size;
758
759         tex_free(*max_mem);
760         tex_alloc(*max_mem);
761         size = (max_size >= size) ? 0 : size - max_size;
762
763         max_mem->device_pointer = prev_pointer;
764         max_mem->device_size = prev_size;
765       }
766       else {
767         break;
768       }
769     }
770
771     /* Update texture info array with new pointers. */
772     load_texture_info();
773
774     move_texture_to_host = false;
775   }
776
777   CUDAMem *generic_alloc(device_memory &mem, size_t pitch_padding = 0)
778   {
779     CUDAContextScope scope(this);
780
781     CUdeviceptr device_pointer = 0;
782     size_t size = mem.memory_size() + pitch_padding;
783
784     CUresult mem_alloc_result = CUDA_ERROR_OUT_OF_MEMORY;
785     const char *status = "";
786
787     /* First try allocating in device memory, respecting headroom. We make
788      * an exception for texture info. It is small and frequently accessed,
789      * so treat it as working memory.
790      *
791      * If there is not enough room for working memory, we will try to move
792      * textures to host memory, assuming the performance impact would have
793      * been worse for working memory. */
794     bool is_texture = (mem.type == MEM_TEXTURE) && (&mem != &texture_info);
795     bool is_image = is_texture && (mem.data_height > 1);
796
797     size_t headroom = (is_texture) ? device_texture_headroom : device_working_headroom;
798
799     size_t total = 0, free = 0;
800     cuMemGetInfo(&free, &total);
801
802     /* Move textures to host memory if needed. */
803     if (!move_texture_to_host && !is_image && (size + headroom) >= free) {
804       move_textures_to_host(size + headroom - free, is_texture);
805       cuMemGetInfo(&free, &total);
806     }
807
808     /* Allocate in device memory. */
809     if (!move_texture_to_host && (size + headroom) < free) {
810       mem_alloc_result = cuMemAlloc(&device_pointer, size);
811       if (mem_alloc_result == CUDA_SUCCESS) {
812         status = " in device memory";
813       }
814     }
815
816     /* Fall back to mapped host memory if needed and possible. */
817     void *map_host_pointer = 0;
818     bool free_map_host = false;
819
820     if (mem_alloc_result != CUDA_SUCCESS && can_map_host &&
821         map_host_used + size < map_host_limit) {
822       if (mem.shared_pointer) {
823         /* Another device already allocated host memory. */
824         mem_alloc_result = CUDA_SUCCESS;
825         map_host_pointer = mem.shared_pointer;
826       }
827       else {
828         /* Allocate host memory ourselves. */
829         mem_alloc_result = cuMemHostAlloc(
830             &map_host_pointer, size, CU_MEMHOSTALLOC_DEVICEMAP | CU_MEMHOSTALLOC_WRITECOMBINED);
831         mem.shared_pointer = map_host_pointer;
832         free_map_host = true;
833       }
834
835       if (mem_alloc_result == CUDA_SUCCESS) {
836         cuda_assert(cuMemHostGetDevicePointer_v2(&device_pointer, mem.shared_pointer, 0));
837         map_host_used += size;
838         status = " in host memory";
839
840         /* Replace host pointer with our host allocation. Only works if
841          * CUDA memory layout is the same and has no pitch padding. Also
842          * does not work if we move textures to host during a render,
843          * since other devices might be using the memory. */
844         if (!move_texture_to_host && pitch_padding == 0 && mem.host_pointer &&
845             mem.host_pointer != mem.shared_pointer) {
846           memcpy(mem.shared_pointer, mem.host_pointer, size);
847           mem.host_free();
848           mem.host_pointer = mem.shared_pointer;
849         }
850       }
851       else {
852         status = " failed, out of host memory";
853       }
854     }
855     else if (mem_alloc_result != CUDA_SUCCESS) {
856       status = " failed, out of device and host memory";
857     }
858
859     if (mem_alloc_result != CUDA_SUCCESS) {
860       cuda_assert(mem_alloc_result);
861     }
862
863     if (mem.name) {
864       VLOG(1) << "Buffer allocate: " << mem.name << ", "
865               << string_human_readable_number(mem.memory_size()) << " bytes. ("
866               << string_human_readable_size(mem.memory_size()) << ")" << status;
867     }
868
869     mem.device_pointer = (device_ptr)device_pointer;
870     mem.device_size = size;
871     stats.mem_alloc(size);
872
873     if (!mem.device_pointer) {
874       return NULL;
875     }
876
877     /* Insert into map of allocations. */
878     CUDAMem *cmem = &cuda_mem_map[&mem];
879     cmem->map_host_pointer = map_host_pointer;
880     cmem->free_map_host = free_map_host;
881     return cmem;
882   }
883
884   void generic_copy_to(device_memory &mem)
885   {
886     if (mem.host_pointer && mem.device_pointer) {
887       CUDAContextScope scope(this);
888
889       if (mem.host_pointer != mem.shared_pointer) {
890         cuda_assert(cuMemcpyHtoD(
891             cuda_device_ptr(mem.device_pointer), mem.host_pointer, mem.memory_size()));
892       }
893     }
894   }
895
896   void generic_free(device_memory &mem)
897   {
898     if (mem.device_pointer) {
899       CUDAContextScope scope(this);
900       const CUDAMem &cmem = cuda_mem_map[&mem];
901
902       if (cmem.map_host_pointer) {
903         /* Free host memory. */
904         if (cmem.free_map_host) {
905           cuMemFreeHost(cmem.map_host_pointer);
906           if (mem.host_pointer == mem.shared_pointer) {
907             mem.host_pointer = 0;
908           }
909           mem.shared_pointer = 0;
910         }
911
912         map_host_used -= mem.device_size;
913       }
914       else {
915         /* Free device memory. */
916         cuMemFree(mem.device_pointer);
917       }
918
919       stats.mem_free(mem.device_size);
920       mem.device_pointer = 0;
921       mem.device_size = 0;
922
923       cuda_mem_map.erase(cuda_mem_map.find(&mem));
924     }
925   }
926
927   void mem_alloc(device_memory &mem)
928   {
929     if (mem.type == MEM_PIXELS && !background) {
930       pixels_alloc(mem);
931     }
932     else if (mem.type == MEM_TEXTURE) {
933       assert(!"mem_alloc not supported for textures.");
934     }
935     else {
936       generic_alloc(mem);
937     }
938   }
939
940   void mem_copy_to(device_memory &mem)
941   {
942     if (mem.type == MEM_PIXELS) {
943       assert(!"mem_copy_to not supported for pixels.");
944     }
945     else if (mem.type == MEM_TEXTURE) {
946       tex_free(mem);
947       tex_alloc(mem);
948     }
949     else {
950       if (!mem.device_pointer) {
951         generic_alloc(mem);
952       }
953
954       generic_copy_to(mem);
955     }
956   }
957
958   void mem_copy_from(device_memory &mem, int y, int w, int h, int elem)
959   {
960     if (mem.type == MEM_PIXELS && !background) {
961       pixels_copy_from(mem, y, w, h);
962     }
963     else if (mem.type == MEM_TEXTURE) {
964       assert(!"mem_copy_from not supported for textures.");
965     }
966     else {
967       CUDAContextScope scope(this);
968       size_t offset = elem * y * w;
969       size_t size = elem * w * h;
970
971       if (mem.host_pointer && mem.device_pointer) {
972         cuda_assert(cuMemcpyDtoH(
973             (uchar *)mem.host_pointer + offset, (CUdeviceptr)(mem.device_pointer + offset), size));
974       }
975       else if (mem.host_pointer) {
976         memset((char *)mem.host_pointer + offset, 0, size);
977       }
978     }
979   }
980
981   void mem_zero(device_memory &mem)
982   {
983     if (!mem.device_pointer) {
984       mem_alloc(mem);
985     }
986
987     if (mem.host_pointer) {
988       memset(mem.host_pointer, 0, mem.memory_size());
989     }
990
991     if (mem.device_pointer && (!mem.host_pointer || mem.host_pointer != mem.shared_pointer)) {
992       CUDAContextScope scope(this);
993       cuda_assert(cuMemsetD8(cuda_device_ptr(mem.device_pointer), 0, mem.memory_size()));
994     }
995   }
996
997   void mem_free(device_memory &mem)
998   {
999     if (mem.type == MEM_PIXELS && !background) {
1000       pixels_free(mem);
1001     }
1002     else if (mem.type == MEM_TEXTURE) {
1003       tex_free(mem);
1004     }
1005     else {
1006       generic_free(mem);
1007     }
1008   }
1009
1010   virtual device_ptr mem_alloc_sub_ptr(device_memory &mem, int offset, int /*size*/)
1011   {
1012     return (device_ptr)(((char *)mem.device_pointer) + mem.memory_elements_size(offset));
1013   }
1014
1015   void const_copy_to(const char *name, void *host, size_t size)
1016   {
1017     CUDAContextScope scope(this);
1018     CUdeviceptr mem;
1019     size_t bytes;
1020
1021     cuda_assert(cuModuleGetGlobal(&mem, &bytes, cuModule, name));
1022     //assert(bytes == size);
1023     cuda_assert(cuMemcpyHtoD(mem, host, size));
1024   }
1025
1026   void tex_alloc(device_memory &mem)
1027   {
1028     CUDAContextScope scope(this);
1029
1030     /* General variables for both architectures */
1031     string bind_name = mem.name;
1032     size_t dsize = datatype_size(mem.data_type);
1033     size_t size = mem.memory_size();
1034
1035     CUaddress_mode address_mode = CU_TR_ADDRESS_MODE_WRAP;
1036     switch (mem.extension) {
1037       case EXTENSION_REPEAT:
1038         address_mode = CU_TR_ADDRESS_MODE_WRAP;
1039         break;
1040       case EXTENSION_EXTEND:
1041         address_mode = CU_TR_ADDRESS_MODE_CLAMP;
1042         break;
1043       case EXTENSION_CLIP:
1044         address_mode = CU_TR_ADDRESS_MODE_BORDER;
1045         break;
1046       default:
1047         assert(0);
1048         break;
1049     }
1050
1051     CUfilter_mode filter_mode;
1052     if (mem.interpolation == INTERPOLATION_CLOSEST) {
1053       filter_mode = CU_TR_FILTER_MODE_POINT;
1054     }
1055     else {
1056       filter_mode = CU_TR_FILTER_MODE_LINEAR;
1057     }
1058
1059     /* Data Storage */
1060     if (mem.interpolation == INTERPOLATION_NONE) {
1061       generic_alloc(mem);
1062       generic_copy_to(mem);
1063
1064       CUdeviceptr cumem;
1065       size_t cubytes;
1066
1067       cuda_assert(cuModuleGetGlobal(&cumem, &cubytes, cuModule, bind_name.c_str()));
1068
1069       if (cubytes == 8) {
1070         /* 64 bit device pointer */
1071         uint64_t ptr = mem.device_pointer;
1072         cuda_assert(cuMemcpyHtoD(cumem, (void *)&ptr, cubytes));
1073       }
1074       else {
1075         /* 32 bit device pointer */
1076         uint32_t ptr = (uint32_t)mem.device_pointer;
1077         cuda_assert(cuMemcpyHtoD(cumem, (void *)&ptr, cubytes));
1078       }
1079       return;
1080     }
1081
1082     /* Image Texture Storage */
1083     CUarray_format_enum format;
1084     switch (mem.data_type) {
1085       case TYPE_UCHAR:
1086         format = CU_AD_FORMAT_UNSIGNED_INT8;
1087         break;
1088       case TYPE_UINT16:
1089         format = CU_AD_FORMAT_UNSIGNED_INT16;
1090         break;
1091       case TYPE_UINT:
1092         format = CU_AD_FORMAT_UNSIGNED_INT32;
1093         break;
1094       case TYPE_INT:
1095         format = CU_AD_FORMAT_SIGNED_INT32;
1096         break;
1097       case TYPE_FLOAT:
1098         format = CU_AD_FORMAT_FLOAT;
1099         break;
1100       case TYPE_HALF:
1101         format = CU_AD_FORMAT_HALF;
1102         break;
1103       default:
1104         assert(0);
1105         return;
1106     }
1107
1108     CUDAMem *cmem = NULL;
1109     CUarray array_3d = NULL;
1110     size_t src_pitch = mem.data_width * dsize * mem.data_elements;
1111     size_t dst_pitch = src_pitch;
1112
1113     if (mem.data_depth > 1) {
1114       /* 3D texture using array, there is no API for linear memory. */
1115       CUDA_ARRAY3D_DESCRIPTOR desc;
1116
1117       desc.Width = mem.data_width;
1118       desc.Height = mem.data_height;
1119       desc.Depth = mem.data_depth;
1120       desc.Format = format;
1121       desc.NumChannels = mem.data_elements;
1122       desc.Flags = 0;
1123
1124       VLOG(1) << "Array 3D allocate: " << mem.name << ", "
1125               << string_human_readable_number(mem.memory_size()) << " bytes. ("
1126               << string_human_readable_size(mem.memory_size()) << ")";
1127
1128       cuda_assert(cuArray3DCreate(&array_3d, &desc));
1129
1130       if (!array_3d) {
1131         return;
1132       }
1133
1134       CUDA_MEMCPY3D param;
1135       memset(&param, 0, sizeof(param));
1136       param.dstMemoryType = CU_MEMORYTYPE_ARRAY;
1137       param.dstArray = array_3d;
1138       param.srcMemoryType = CU_MEMORYTYPE_HOST;
1139       param.srcHost = mem.host_pointer;
1140       param.srcPitch = src_pitch;
1141       param.WidthInBytes = param.srcPitch;
1142       param.Height = mem.data_height;
1143       param.Depth = mem.data_depth;
1144
1145       cuda_assert(cuMemcpy3D(&param));
1146
1147       mem.device_pointer = (device_ptr)array_3d;
1148       mem.device_size = size;
1149       stats.mem_alloc(size);
1150
1151       cmem = &cuda_mem_map[&mem];
1152       cmem->texobject = 0;
1153       cmem->array = array_3d;
1154     }
1155     else if (mem.data_height > 0) {
1156       /* 2D texture, using pitch aligned linear memory. */
1157       int alignment = 0;
1158       cuda_assert(
1159           cuDeviceGetAttribute(&alignment, CU_DEVICE_ATTRIBUTE_TEXTURE_PITCH_ALIGNMENT, cuDevice));
1160       dst_pitch = align_up(src_pitch, alignment);
1161       size_t dst_size = dst_pitch * mem.data_height;
1162
1163       cmem = generic_alloc(mem, dst_size - mem.memory_size());
1164       if (!cmem) {
1165         return;
1166       }
1167
1168       CUDA_MEMCPY2D param;
1169       memset(&param, 0, sizeof(param));
1170       param.dstMemoryType = CU_MEMORYTYPE_DEVICE;
1171       param.dstDevice = mem.device_pointer;
1172       param.dstPitch = dst_pitch;
1173       param.srcMemoryType = CU_MEMORYTYPE_HOST;
1174       param.srcHost = mem.host_pointer;
1175       param.srcPitch = src_pitch;
1176       param.WidthInBytes = param.srcPitch;
1177       param.Height = mem.data_height;
1178
1179       cuda_assert(cuMemcpy2DUnaligned(&param));
1180     }
1181     else {
1182       /* 1D texture, using linear memory. */
1183       cmem = generic_alloc(mem);
1184       if (!cmem) {
1185         return;
1186       }
1187
1188       cuda_assert(cuMemcpyHtoD(mem.device_pointer, mem.host_pointer, size));
1189     }
1190
1191     /* Kepler+, bindless textures. */
1192     int flat_slot = 0;
1193     if (string_startswith(mem.name, "__tex_image")) {
1194       int pos = string(mem.name).rfind("_");
1195       flat_slot = atoi(mem.name + pos + 1);
1196     }
1197     else {
1198       assert(0);
1199     }
1200
1201     CUDA_RESOURCE_DESC resDesc;
1202     memset(&resDesc, 0, sizeof(resDesc));
1203
1204     if (array_3d) {
1205       resDesc.resType = CU_RESOURCE_TYPE_ARRAY;
1206       resDesc.res.array.hArray = array_3d;
1207       resDesc.flags = 0;
1208     }
1209     else if (mem.data_height > 0) {
1210       resDesc.resType = CU_RESOURCE_TYPE_PITCH2D;
1211       resDesc.res.pitch2D.devPtr = mem.device_pointer;
1212       resDesc.res.pitch2D.format = format;
1213       resDesc.res.pitch2D.numChannels = mem.data_elements;
1214       resDesc.res.pitch2D.height = mem.data_height;
1215       resDesc.res.pitch2D.width = mem.data_width;
1216       resDesc.res.pitch2D.pitchInBytes = dst_pitch;
1217     }
1218     else {
1219       resDesc.resType = CU_RESOURCE_TYPE_LINEAR;
1220       resDesc.res.linear.devPtr = mem.device_pointer;
1221       resDesc.res.linear.format = format;
1222       resDesc.res.linear.numChannels = mem.data_elements;
1223       resDesc.res.linear.sizeInBytes = mem.device_size;
1224     }
1225
1226     CUDA_TEXTURE_DESC texDesc;
1227     memset(&texDesc, 0, sizeof(texDesc));
1228     texDesc.addressMode[0] = address_mode;
1229     texDesc.addressMode[1] = address_mode;
1230     texDesc.addressMode[2] = address_mode;
1231     texDesc.filterMode = filter_mode;
1232     texDesc.flags = CU_TRSF_NORMALIZED_COORDINATES;
1233
1234     cuda_assert(cuTexObjectCreate(&cmem->texobject, &resDesc, &texDesc, NULL));
1235
1236     /* Resize once */
1237     if (flat_slot >= texture_info.size()) {
1238       /* Allocate some slots in advance, to reduce amount
1239        * of re-allocations. */
1240       texture_info.resize(flat_slot + 128);
1241     }
1242
1243     /* Set Mapping and tag that we need to (re-)upload to device */
1244     TextureInfo &info = texture_info[flat_slot];
1245     info.data = (uint64_t)cmem->texobject;
1246     info.cl_buffer = 0;
1247     info.interpolation = mem.interpolation;
1248     info.extension = mem.extension;
1249     info.width = mem.data_width;
1250     info.height = mem.data_height;
1251     info.depth = mem.data_depth;
1252     need_texture_info = true;
1253   }
1254
1255   void tex_free(device_memory &mem)
1256   {
1257     if (mem.device_pointer) {
1258       CUDAContextScope scope(this);
1259       const CUDAMem &cmem = cuda_mem_map[&mem];
1260
1261       if (cmem.texobject) {
1262         /* Free bindless texture. */
1263         cuTexObjectDestroy(cmem.texobject);
1264       }
1265
1266       if (cmem.array) {
1267         /* Free array. */
1268         cuArrayDestroy(cmem.array);
1269         stats.mem_free(mem.device_size);
1270         mem.device_pointer = 0;
1271         mem.device_size = 0;
1272
1273         cuda_mem_map.erase(cuda_mem_map.find(&mem));
1274       }
1275       else {
1276         generic_free(mem);
1277       }
1278     }
1279   }
1280
1281 #define CUDA_GET_BLOCKSIZE(func, w, h) \
1282   int threads_per_block; \
1283   cuda_assert( \
1284       cuFuncGetAttribute(&threads_per_block, CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, func)); \
1285   int threads = (int)sqrt((float)threads_per_block); \
1286   int xblocks = ((w) + threads - 1) / threads; \
1287   int yblocks = ((h) + threads - 1) / threads;
1288
1289 #define CUDA_LAUNCH_KERNEL(func, args) \
1290   cuda_assert(cuLaunchKernel(func, xblocks, yblocks, 1, threads, threads, 1, 0, 0, args, 0));
1291
1292 /* Similar as above, but for 1-dimensional blocks. */
1293 #define CUDA_GET_BLOCKSIZE_1D(func, w, h) \
1294   int threads_per_block; \
1295   cuda_assert( \
1296       cuFuncGetAttribute(&threads_per_block, CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, func)); \
1297   int xblocks = ((w) + threads_per_block - 1) / threads_per_block; \
1298   int yblocks = h;
1299
1300 #define CUDA_LAUNCH_KERNEL_1D(func, args) \
1301   cuda_assert(cuLaunchKernel(func, xblocks, yblocks, 1, threads_per_block, 1, 1, 0, 0, args, 0));
1302
1303   bool denoising_non_local_means(device_ptr image_ptr,
1304                                  device_ptr guide_ptr,
1305                                  device_ptr variance_ptr,
1306                                  device_ptr out_ptr,
1307                                  DenoisingTask *task)
1308   {
1309     if (have_error())
1310       return false;
1311
1312     CUDAContextScope scope(this);
1313
1314     int stride = task->buffer.stride;
1315     int w = task->buffer.width;
1316     int h = task->buffer.h;
1317     int r = task->nlm_state.r;
1318     int f = task->nlm_state.f;
1319     float a = task->nlm_state.a;
1320     float k_2 = task->nlm_state.k_2;
1321
1322     int pass_stride = task->buffer.pass_stride;
1323     int num_shifts = (2 * r + 1) * (2 * r + 1);
1324     int channel_offset = task->nlm_state.is_color ? task->buffer.pass_stride : 0;
1325     int frame_offset = 0;
1326
1327     if (have_error())
1328       return false;
1329
1330     CUdeviceptr difference = cuda_device_ptr(task->buffer.temporary_mem.device_pointer);
1331     CUdeviceptr blurDifference = difference + sizeof(float) * pass_stride * num_shifts;
1332     CUdeviceptr weightAccum = difference + 2 * sizeof(float) * pass_stride * num_shifts;
1333     CUdeviceptr scale_ptr = 0;
1334
1335     cuda_assert(cuMemsetD8(weightAccum, 0, sizeof(float) * pass_stride));
1336     cuda_assert(cuMemsetD8(out_ptr, 0, sizeof(float) * pass_stride));
1337
1338     {
1339       CUfunction cuNLMCalcDifference, cuNLMBlur, cuNLMCalcWeight, cuNLMUpdateOutput;
1340       cuda_assert(cuModuleGetFunction(
1341           &cuNLMCalcDifference, cuFilterModule, "kernel_cuda_filter_nlm_calc_difference"));
1342       cuda_assert(cuModuleGetFunction(&cuNLMBlur, cuFilterModule, "kernel_cuda_filter_nlm_blur"));
1343       cuda_assert(cuModuleGetFunction(
1344           &cuNLMCalcWeight, cuFilterModule, "kernel_cuda_filter_nlm_calc_weight"));
1345       cuda_assert(cuModuleGetFunction(
1346           &cuNLMUpdateOutput, cuFilterModule, "kernel_cuda_filter_nlm_update_output"));
1347
1348       cuda_assert(cuFuncSetCacheConfig(cuNLMCalcDifference, CU_FUNC_CACHE_PREFER_L1));
1349       cuda_assert(cuFuncSetCacheConfig(cuNLMBlur, CU_FUNC_CACHE_PREFER_L1));
1350       cuda_assert(cuFuncSetCacheConfig(cuNLMCalcWeight, CU_FUNC_CACHE_PREFER_L1));
1351       cuda_assert(cuFuncSetCacheConfig(cuNLMUpdateOutput, CU_FUNC_CACHE_PREFER_L1));
1352
1353       CUDA_GET_BLOCKSIZE_1D(cuNLMCalcDifference, w * h, num_shifts);
1354
1355       void *calc_difference_args[] = {&guide_ptr,
1356                                       &variance_ptr,
1357                                       &scale_ptr,
1358                                       &difference,
1359                                       &w,
1360                                       &h,
1361                                       &stride,
1362                                       &pass_stride,
1363                                       &r,
1364                                       &channel_offset,
1365                                       &frame_offset,
1366                                       &a,
1367                                       &k_2};
1368       void *blur_args[] = {&difference, &blurDifference, &w, &h, &stride, &pass_stride, &r, &f};
1369       void *calc_weight_args[] = {
1370           &blurDifference, &difference, &w, &h, &stride, &pass_stride, &r, &f};
1371       void *update_output_args[] = {&blurDifference,
1372                                     &image_ptr,
1373                                     &out_ptr,
1374                                     &weightAccum,
1375                                     &w,
1376                                     &h,
1377                                     &stride,
1378                                     &pass_stride,
1379                                     &channel_offset,
1380                                     &r,
1381                                     &f};
1382
1383       CUDA_LAUNCH_KERNEL_1D(cuNLMCalcDifference, calc_difference_args);
1384       CUDA_LAUNCH_KERNEL_1D(cuNLMBlur, blur_args);
1385       CUDA_LAUNCH_KERNEL_1D(cuNLMCalcWeight, calc_weight_args);
1386       CUDA_LAUNCH_KERNEL_1D(cuNLMBlur, blur_args);
1387       CUDA_LAUNCH_KERNEL_1D(cuNLMUpdateOutput, update_output_args);
1388     }
1389
1390     {
1391       CUfunction cuNLMNormalize;
1392       cuda_assert(cuModuleGetFunction(
1393           &cuNLMNormalize, cuFilterModule, "kernel_cuda_filter_nlm_normalize"));
1394       cuda_assert(cuFuncSetCacheConfig(cuNLMNormalize, CU_FUNC_CACHE_PREFER_L1));
1395       void *normalize_args[] = {&out_ptr, &weightAccum, &w, &h, &stride};
1396       CUDA_GET_BLOCKSIZE(cuNLMNormalize, w, h);
1397       CUDA_LAUNCH_KERNEL(cuNLMNormalize, normalize_args);
1398       cuda_assert(cuCtxSynchronize());
1399     }
1400
1401     return !have_error();
1402   }
1403
1404   bool denoising_construct_transform(DenoisingTask *task)
1405   {
1406     if (have_error())
1407       return false;
1408
1409     CUDAContextScope scope(this);
1410
1411     CUfunction cuFilterConstructTransform;
1412     cuda_assert(cuModuleGetFunction(
1413         &cuFilterConstructTransform, cuFilterModule, "kernel_cuda_filter_construct_transform"));
1414     cuda_assert(cuFuncSetCacheConfig(cuFilterConstructTransform, CU_FUNC_CACHE_PREFER_SHARED));
1415     CUDA_GET_BLOCKSIZE(cuFilterConstructTransform, task->storage.w, task->storage.h);
1416
1417     void *args[] = {&task->buffer.mem.device_pointer,
1418                     &task->tile_info_mem.device_pointer,
1419                     &task->storage.transform.device_pointer,
1420                     &task->storage.rank.device_pointer,
1421                     &task->filter_area,
1422                     &task->rect,
1423                     &task->radius,
1424                     &task->pca_threshold,
1425                     &task->buffer.pass_stride,
1426                     &task->buffer.frame_stride,
1427                     &task->buffer.use_time};
1428     CUDA_LAUNCH_KERNEL(cuFilterConstructTransform, args);
1429     cuda_assert(cuCtxSynchronize());
1430
1431     return !have_error();
1432   }
1433
1434   bool denoising_accumulate(device_ptr color_ptr,
1435                             device_ptr color_variance_ptr,
1436                             device_ptr scale_ptr,
1437                             int frame,
1438                             DenoisingTask *task)
1439   {
1440     if (have_error())
1441       return false;
1442
1443     CUDAContextScope scope(this);
1444
1445     int r = task->radius;
1446     int f = 4;
1447     float a = 1.0f;
1448     float k_2 = task->nlm_k_2;
1449
1450     int w = task->reconstruction_state.source_w;
1451     int h = task->reconstruction_state.source_h;
1452     int stride = task->buffer.stride;
1453     int frame_offset = frame * task->buffer.frame_stride;
1454     int t = task->tile_info->frames[frame];
1455
1456     int pass_stride = task->buffer.pass_stride;
1457     int num_shifts = (2 * r + 1) * (2 * r + 1);
1458
1459     if (have_error())
1460       return false;
1461
1462     CUdeviceptr difference = cuda_device_ptr(task->buffer.temporary_mem.device_pointer);
1463     CUdeviceptr blurDifference = difference + sizeof(float) * pass_stride * num_shifts;
1464
1465     CUfunction cuNLMCalcDifference, cuNLMBlur, cuNLMCalcWeight, cuNLMConstructGramian;
1466     cuda_assert(cuModuleGetFunction(
1467         &cuNLMCalcDifference, cuFilterModule, "kernel_cuda_filter_nlm_calc_difference"));
1468     cuda_assert(cuModuleGetFunction(&cuNLMBlur, cuFilterModule, "kernel_cuda_filter_nlm_blur"));
1469     cuda_assert(cuModuleGetFunction(
1470         &cuNLMCalcWeight, cuFilterModule, "kernel_cuda_filter_nlm_calc_weight"));
1471     cuda_assert(cuModuleGetFunction(
1472         &cuNLMConstructGramian, cuFilterModule, "kernel_cuda_filter_nlm_construct_gramian"));
1473
1474     cuda_assert(cuFuncSetCacheConfig(cuNLMCalcDifference, CU_FUNC_CACHE_PREFER_L1));
1475     cuda_assert(cuFuncSetCacheConfig(cuNLMBlur, CU_FUNC_CACHE_PREFER_L1));
1476     cuda_assert(cuFuncSetCacheConfig(cuNLMCalcWeight, CU_FUNC_CACHE_PREFER_L1));
1477     cuda_assert(cuFuncSetCacheConfig(cuNLMConstructGramian, CU_FUNC_CACHE_PREFER_SHARED));
1478
1479     CUDA_GET_BLOCKSIZE_1D(cuNLMCalcDifference,
1480                           task->reconstruction_state.source_w *
1481                               task->reconstruction_state.source_h,
1482                           num_shifts);
1483
1484     void *calc_difference_args[] = {&color_ptr,
1485                                     &color_variance_ptr,
1486                                     &scale_ptr,
1487                                     &difference,
1488                                     &w,
1489                                     &h,
1490                                     &stride,
1491                                     &pass_stride,
1492                                     &r,
1493                                     &pass_stride,
1494                                     &frame_offset,
1495                                     &a,
1496                                     &k_2};
1497     void *blur_args[] = {&difference, &blurDifference, &w, &h, &stride, &pass_stride, &r, &f};
1498     void *calc_weight_args[] = {
1499         &blurDifference, &difference, &w, &h, &stride, &pass_stride, &r, &f};
1500     void *construct_gramian_args[] = {&t,
1501                                       &blurDifference,
1502                                       &task->buffer.mem.device_pointer,
1503                                       &task->storage.transform.device_pointer,
1504                                       &task->storage.rank.device_pointer,
1505                                       &task->storage.XtWX.device_pointer,
1506                                       &task->storage.XtWY.device_pointer,
1507                                       &task->reconstruction_state.filter_window,
1508                                       &w,
1509                                       &h,
1510                                       &stride,
1511                                       &pass_stride,
1512                                       &r,
1513                                       &f,
1514                                       &frame_offset,
1515                                       &task->buffer.use_time};
1516
1517     CUDA_LAUNCH_KERNEL_1D(cuNLMCalcDifference, calc_difference_args);
1518     CUDA_LAUNCH_KERNEL_1D(cuNLMBlur, blur_args);
1519     CUDA_LAUNCH_KERNEL_1D(cuNLMCalcWeight, calc_weight_args);
1520     CUDA_LAUNCH_KERNEL_1D(cuNLMBlur, blur_args);
1521     CUDA_LAUNCH_KERNEL_1D(cuNLMConstructGramian, construct_gramian_args);
1522     cuda_assert(cuCtxSynchronize());
1523
1524     return !have_error();
1525   }
1526
1527   bool denoising_solve(device_ptr output_ptr, DenoisingTask *task)
1528   {
1529     CUfunction cuFinalize;
1530     cuda_assert(cuModuleGetFunction(&cuFinalize, cuFilterModule, "kernel_cuda_filter_finalize"));
1531     cuda_assert(cuFuncSetCacheConfig(cuFinalize, CU_FUNC_CACHE_PREFER_L1));
1532     void *finalize_args[] = {&output_ptr,
1533                              &task->storage.rank.device_pointer,
1534                              &task->storage.XtWX.device_pointer,
1535                              &task->storage.XtWY.device_pointer,
1536                              &task->filter_area,
1537                              &task->reconstruction_state.buffer_params.x,
1538                              &task->render_buffer.samples};
1539     CUDA_GET_BLOCKSIZE(
1540         cuFinalize, task->reconstruction_state.source_w, task->reconstruction_state.source_h);
1541     CUDA_LAUNCH_KERNEL(cuFinalize, finalize_args);
1542     cuda_assert(cuCtxSynchronize());
1543
1544     return !have_error();
1545   }
1546
1547   bool denoising_combine_halves(device_ptr a_ptr,
1548                                 device_ptr b_ptr,
1549                                 device_ptr mean_ptr,
1550                                 device_ptr variance_ptr,
1551                                 int r,
1552                                 int4 rect,
1553                                 DenoisingTask *task)
1554   {
1555     if (have_error())
1556       return false;
1557
1558     CUDAContextScope scope(this);
1559
1560     CUfunction cuFilterCombineHalves;
1561     cuda_assert(cuModuleGetFunction(
1562         &cuFilterCombineHalves, cuFilterModule, "kernel_cuda_filter_combine_halves"));
1563     cuda_assert(cuFuncSetCacheConfig(cuFilterCombineHalves, CU_FUNC_CACHE_PREFER_L1));
1564     CUDA_GET_BLOCKSIZE(
1565         cuFilterCombineHalves, task->rect.z - task->rect.x, task->rect.w - task->rect.y);
1566
1567     void *args[] = {&mean_ptr, &variance_ptr, &a_ptr, &b_ptr, &rect, &r};
1568     CUDA_LAUNCH_KERNEL(cuFilterCombineHalves, args);
1569     cuda_assert(cuCtxSynchronize());
1570
1571     return !have_error();
1572   }
1573
1574   bool denoising_divide_shadow(device_ptr a_ptr,
1575                                device_ptr b_ptr,
1576                                device_ptr sample_variance_ptr,
1577                                device_ptr sv_variance_ptr,
1578                                device_ptr buffer_variance_ptr,
1579                                DenoisingTask *task)
1580   {
1581     if (have_error())
1582       return false;
1583
1584     CUDAContextScope scope(this);
1585
1586     CUfunction cuFilterDivideShadow;
1587     cuda_assert(cuModuleGetFunction(
1588         &cuFilterDivideShadow, cuFilterModule, "kernel_cuda_filter_divide_shadow"));
1589     cuda_assert(cuFuncSetCacheConfig(cuFilterDivideShadow, CU_FUNC_CACHE_PREFER_L1));
1590     CUDA_GET_BLOCKSIZE(
1591         cuFilterDivideShadow, task->rect.z - task->rect.x, task->rect.w - task->rect.y);
1592
1593     void *args[] = {&task->render_buffer.samples,
1594                     &task->tile_info_mem.device_pointer,
1595                     &a_ptr,
1596                     &b_ptr,
1597                     &sample_variance_ptr,
1598                     &sv_variance_ptr,
1599                     &buffer_variance_ptr,
1600                     &task->rect,
1601                     &task->render_buffer.pass_stride,
1602                     &task->render_buffer.offset};
1603     CUDA_LAUNCH_KERNEL(cuFilterDivideShadow, args);
1604     cuda_assert(cuCtxSynchronize());
1605
1606     return !have_error();
1607   }
1608
1609   bool denoising_get_feature(int mean_offset,
1610                              int variance_offset,
1611                              device_ptr mean_ptr,
1612                              device_ptr variance_ptr,
1613                              float scale,
1614                              DenoisingTask *task)
1615   {
1616     if (have_error())
1617       return false;
1618
1619     CUDAContextScope scope(this);
1620
1621     CUfunction cuFilterGetFeature;
1622     cuda_assert(cuModuleGetFunction(
1623         &cuFilterGetFeature, cuFilterModule, "kernel_cuda_filter_get_feature"));
1624     cuda_assert(cuFuncSetCacheConfig(cuFilterGetFeature, CU_FUNC_CACHE_PREFER_L1));
1625     CUDA_GET_BLOCKSIZE(
1626         cuFilterGetFeature, task->rect.z - task->rect.x, task->rect.w - task->rect.y);
1627
1628     void *args[] = {&task->render_buffer.samples,
1629                     &task->tile_info_mem.device_pointer,
1630                     &mean_offset,
1631                     &variance_offset,
1632                     &mean_ptr,
1633                     &variance_ptr,
1634                     &scale,
1635                     &task->rect,
1636                     &task->render_buffer.pass_stride,
1637                     &task->render_buffer.offset};
1638     CUDA_LAUNCH_KERNEL(cuFilterGetFeature, args);
1639     cuda_assert(cuCtxSynchronize());
1640
1641     return !have_error();
1642   }
1643
1644   bool denoising_write_feature(int out_offset,
1645                                device_ptr from_ptr,
1646                                device_ptr buffer_ptr,
1647                                DenoisingTask *task)
1648   {
1649     if (have_error())
1650       return false;
1651
1652     CUDAContextScope scope(this);
1653
1654     CUfunction cuFilterWriteFeature;
1655     cuda_assert(cuModuleGetFunction(
1656         &cuFilterWriteFeature, cuFilterModule, "kernel_cuda_filter_write_feature"));
1657     cuda_assert(cuFuncSetCacheConfig(cuFilterWriteFeature, CU_FUNC_CACHE_PREFER_L1));
1658     CUDA_GET_BLOCKSIZE(cuFilterWriteFeature, task->filter_area.z, task->filter_area.w);
1659
1660     void *args[] = {&task->render_buffer.samples,
1661                     &task->reconstruction_state.buffer_params,
1662                     &task->filter_area,
1663                     &from_ptr,
1664                     &buffer_ptr,
1665                     &out_offset,
1666                     &task->rect};
1667     CUDA_LAUNCH_KERNEL(cuFilterWriteFeature, args);
1668     cuda_assert(cuCtxSynchronize());
1669
1670     return !have_error();
1671   }
1672
1673   bool denoising_detect_outliers(device_ptr image_ptr,
1674                                  device_ptr variance_ptr,
1675                                  device_ptr depth_ptr,
1676                                  device_ptr output_ptr,
1677                                  DenoisingTask *task)
1678   {
1679     if (have_error())
1680       return false;
1681
1682     CUDAContextScope scope(this);
1683
1684     CUfunction cuFilterDetectOutliers;
1685     cuda_assert(cuModuleGetFunction(
1686         &cuFilterDetectOutliers, cuFilterModule, "kernel_cuda_filter_detect_outliers"));
1687     cuda_assert(cuFuncSetCacheConfig(cuFilterDetectOutliers, CU_FUNC_CACHE_PREFER_L1));
1688     CUDA_GET_BLOCKSIZE(
1689         cuFilterDetectOutliers, task->rect.z - task->rect.x, task->rect.w - task->rect.y);
1690
1691     void *args[] = {&image_ptr,
1692                     &variance_ptr,
1693                     &depth_ptr,
1694                     &output_ptr,
1695                     &task->rect,
1696                     &task->buffer.pass_stride};
1697
1698     CUDA_LAUNCH_KERNEL(cuFilterDetectOutliers, args);
1699     cuda_assert(cuCtxSynchronize());
1700
1701     return !have_error();
1702   }
1703
1704   void denoise(RenderTile &rtile, DenoisingTask &denoising)
1705   {
1706     denoising.functions.construct_transform = function_bind(
1707         &CUDADevice::denoising_construct_transform, this, &denoising);
1708     denoising.functions.accumulate = function_bind(
1709         &CUDADevice::denoising_accumulate, this, _1, _2, _3, _4, &denoising);
1710     denoising.functions.solve = function_bind(&CUDADevice::denoising_solve, this, _1, &denoising);
1711     denoising.functions.divide_shadow = function_bind(
1712         &CUDADevice::denoising_divide_shadow, this, _1, _2, _3, _4, _5, &denoising);
1713     denoising.functions.non_local_means = function_bind(
1714         &CUDADevice::denoising_non_local_means, this, _1, _2, _3, _4, &denoising);
1715     denoising.functions.combine_halves = function_bind(
1716         &CUDADevice::denoising_combine_halves, this, _1, _2, _3, _4, _5, _6, &denoising);
1717     denoising.functions.get_feature = function_bind(
1718         &CUDADevice::denoising_get_feature, this, _1, _2, _3, _4, _5, &denoising);
1719     denoising.functions.write_feature = function_bind(
1720         &CUDADevice::denoising_write_feature, this, _1, _2, _3, &denoising);
1721     denoising.functions.detect_outliers = function_bind(
1722         &CUDADevice::denoising_detect_outliers, this, _1, _2, _3, _4, &denoising);
1723
1724     denoising.filter_area = make_int4(rtile.x, rtile.y, rtile.w, rtile.h);
1725     denoising.render_buffer.samples = rtile.sample;
1726     denoising.buffer.gpu_temporary_mem = true;
1727
1728     denoising.run_denoising(&rtile);
1729   }
1730
1731   void path_trace(DeviceTask &task, RenderTile &rtile, device_vector<WorkTile> &work_tiles)
1732   {
1733     scoped_timer timer(&rtile.buffers->render_time);
1734
1735     if (have_error())
1736       return;
1737
1738     CUDAContextScope scope(this);
1739     CUfunction cuPathTrace;
1740
1741     /* Get kernel function. */
1742     if (task.integrator_branched) {
1743       cuda_assert(cuModuleGetFunction(&cuPathTrace, cuModule, "kernel_cuda_branched_path_trace"));
1744     }
1745     else {
1746       cuda_assert(cuModuleGetFunction(&cuPathTrace, cuModule, "kernel_cuda_path_trace"));
1747     }
1748
1749     if (have_error()) {
1750       return;
1751     }
1752
1753     cuda_assert(cuFuncSetCacheConfig(cuPathTrace, CU_FUNC_CACHE_PREFER_L1));
1754
1755     /* Allocate work tile. */
1756     work_tiles.alloc(1);
1757
1758     WorkTile *wtile = work_tiles.data();
1759     wtile->x = rtile.x;
1760     wtile->y = rtile.y;
1761     wtile->w = rtile.w;
1762     wtile->h = rtile.h;
1763     wtile->offset = rtile.offset;
1764     wtile->stride = rtile.stride;
1765     wtile->buffer = (float *)cuda_device_ptr(rtile.buffer);
1766
1767     /* Prepare work size. More step samples render faster, but for now we
1768      * remain conservative for GPUs connected to a display to avoid driver
1769      * timeouts and display freezing. */
1770     int min_blocks, num_threads_per_block;
1771     cuda_assert(cuOccupancyMaxPotentialBlockSize(
1772         &min_blocks, &num_threads_per_block, cuPathTrace, NULL, 0, 0));
1773     if (!info.display_device) {
1774       min_blocks *= 8;
1775     }
1776
1777     uint step_samples = divide_up(min_blocks * num_threads_per_block, wtile->w * wtile->h);
1778
1779     /* Render all samples. */
1780     int start_sample = rtile.start_sample;
1781     int end_sample = rtile.start_sample + rtile.num_samples;
1782
1783     for (int sample = start_sample; sample < end_sample; sample += step_samples) {
1784       /* Setup and copy work tile to device. */
1785       wtile->start_sample = sample;
1786       wtile->num_samples = min(step_samples, end_sample - sample);
1787       work_tiles.copy_to_device();
1788
1789       CUdeviceptr d_work_tiles = cuda_device_ptr(work_tiles.device_pointer);
1790       uint total_work_size = wtile->w * wtile->h * wtile->num_samples;
1791       uint num_blocks = divide_up(total_work_size, num_threads_per_block);
1792
1793       /* Launch kernel. */
1794       void *args[] = {&d_work_tiles, &total_work_size};
1795
1796       cuda_assert(cuLaunchKernel(
1797           cuPathTrace, num_blocks, 1, 1, num_threads_per_block, 1, 1, 0, 0, args, 0));
1798
1799       cuda_assert(cuCtxSynchronize());
1800
1801       /* Update progress. */
1802       rtile.sample = sample + wtile->num_samples;
1803       task.update_progress(&rtile, rtile.w * rtile.h * wtile->num_samples);
1804
1805       if (task.get_cancel()) {
1806         if (task.need_finish_queue == false)
1807           break;
1808       }
1809     }
1810   }
1811
1812   void film_convert(DeviceTask &task,
1813                     device_ptr buffer,
1814                     device_ptr rgba_byte,
1815                     device_ptr rgba_half)
1816   {
1817     if (have_error())
1818       return;
1819
1820     CUDAContextScope scope(this);
1821
1822     CUfunction cuFilmConvert;
1823     CUdeviceptr d_rgba = map_pixels((rgba_byte) ? rgba_byte : rgba_half);
1824     CUdeviceptr d_buffer = cuda_device_ptr(buffer);
1825
1826     /* get kernel function */
1827     if (rgba_half) {
1828       cuda_assert(
1829           cuModuleGetFunction(&cuFilmConvert, cuModule, "kernel_cuda_convert_to_half_float"));
1830     }
1831     else {
1832       cuda_assert(cuModuleGetFunction(&cuFilmConvert, cuModule, "kernel_cuda_convert_to_byte"));
1833     }
1834
1835     float sample_scale = 1.0f / (task.sample + 1);
1836
1837     /* pass in parameters */
1838     void *args[] = {&d_rgba,
1839                     &d_buffer,
1840                     &sample_scale,
1841                     &task.x,
1842                     &task.y,
1843                     &task.w,
1844                     &task.h,
1845                     &task.offset,
1846                     &task.stride};
1847
1848     /* launch kernel */
1849     int threads_per_block;
1850     cuda_assert(cuFuncGetAttribute(
1851         &threads_per_block, CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, cuFilmConvert));
1852
1853     int xthreads = (int)sqrt(threads_per_block);
1854     int ythreads = (int)sqrt(threads_per_block);
1855     int xblocks = (task.w + xthreads - 1) / xthreads;
1856     int yblocks = (task.h + ythreads - 1) / ythreads;
1857
1858     cuda_assert(cuFuncSetCacheConfig(cuFilmConvert, CU_FUNC_CACHE_PREFER_L1));
1859
1860     cuda_assert(cuLaunchKernel(cuFilmConvert,
1861                                xblocks,
1862                                yblocks,
1863                                1, /* blocks */
1864                                xthreads,
1865                                ythreads,
1866                                1, /* threads */
1867                                0,
1868                                0,
1869                                args,
1870                                0));
1871
1872     unmap_pixels((rgba_byte) ? rgba_byte : rgba_half);
1873
1874     cuda_assert(cuCtxSynchronize());
1875   }
1876
1877   void shader(DeviceTask &task)
1878   {
1879     if (have_error())
1880       return;
1881
1882     CUDAContextScope scope(this);
1883
1884     CUfunction cuShader;
1885     CUdeviceptr d_input = cuda_device_ptr(task.shader_input);
1886     CUdeviceptr d_output = cuda_device_ptr(task.shader_output);
1887
1888     /* get kernel function */
1889     if (task.shader_eval_type >= SHADER_EVAL_BAKE) {
1890       cuda_assert(cuModuleGetFunction(&cuShader, cuModule, "kernel_cuda_bake"));
1891     }
1892     else if (task.shader_eval_type == SHADER_EVAL_DISPLACE) {
1893       cuda_assert(cuModuleGetFunction(&cuShader, cuModule, "kernel_cuda_displace"));
1894     }
1895     else {
1896       cuda_assert(cuModuleGetFunction(&cuShader, cuModule, "kernel_cuda_background"));
1897     }
1898
1899     /* do tasks in smaller chunks, so we can cancel it */
1900     const int shader_chunk_size = 65536;
1901     const int start = task.shader_x;
1902     const int end = task.shader_x + task.shader_w;
1903     int offset = task.offset;
1904
1905     bool canceled = false;
1906     for (int sample = 0; sample < task.num_samples && !canceled; sample++) {
1907       for (int shader_x = start; shader_x < end; shader_x += shader_chunk_size) {
1908         int shader_w = min(shader_chunk_size, end - shader_x);
1909
1910         /* pass in parameters */
1911         void *args[8];
1912         int arg = 0;
1913         args[arg++] = &d_input;
1914         args[arg++] = &d_output;
1915         args[arg++] = &task.shader_eval_type;
1916         if (task.shader_eval_type >= SHADER_EVAL_BAKE) {
1917           args[arg++] = &task.shader_filter;
1918         }
1919         args[arg++] = &shader_x;
1920         args[arg++] = &shader_w;
1921         args[arg++] = &offset;
1922         args[arg++] = &sample;
1923
1924         /* launch kernel */
1925         int threads_per_block;
1926         cuda_assert(cuFuncGetAttribute(
1927             &threads_per_block, CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, cuShader));
1928
1929         int xblocks = (shader_w + threads_per_block - 1) / threads_per_block;
1930
1931         cuda_assert(cuFuncSetCacheConfig(cuShader, CU_FUNC_CACHE_PREFER_L1));
1932         cuda_assert(cuLaunchKernel(cuShader,
1933                                    xblocks,
1934                                    1,
1935                                    1, /* blocks */
1936                                    threads_per_block,
1937                                    1,
1938                                    1, /* threads */
1939                                    0,
1940                                    0,
1941                                    args,
1942                                    0));
1943
1944         cuda_assert(cuCtxSynchronize());
1945
1946         if (task.get_cancel()) {
1947           canceled = true;
1948           break;
1949         }
1950       }
1951
1952       task.update_progress(NULL);
1953     }
1954   }
1955
1956   CUdeviceptr map_pixels(device_ptr mem)
1957   {
1958     if (!background) {
1959       PixelMem pmem = pixel_mem_map[mem];
1960       CUdeviceptr buffer;
1961
1962       size_t bytes;
1963       cuda_assert(cuGraphicsMapResources(1, &pmem.cuPBOresource, 0));
1964       cuda_assert(cuGraphicsResourceGetMappedPointer(&buffer, &bytes, pmem.cuPBOresource));
1965
1966       return buffer;
1967     }
1968
1969     return cuda_device_ptr(mem);
1970   }
1971
1972   void unmap_pixels(device_ptr mem)
1973   {
1974     if (!background) {
1975       PixelMem pmem = pixel_mem_map[mem];
1976
1977       cuda_assert(cuGraphicsUnmapResources(1, &pmem.cuPBOresource, 0));
1978     }
1979   }
1980
1981   void pixels_alloc(device_memory &mem)
1982   {
1983     PixelMem pmem;
1984
1985     pmem.w = mem.data_width;
1986     pmem.h = mem.data_height;
1987
1988     CUDAContextScope scope(this);
1989
1990     glGenBuffers(1, &pmem.cuPBO);
1991     glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pmem.cuPBO);
1992     if (mem.data_type == TYPE_HALF)
1993       glBufferData(
1994           GL_PIXEL_UNPACK_BUFFER, pmem.w * pmem.h * sizeof(GLhalf) * 4, NULL, GL_DYNAMIC_DRAW);
1995     else
1996       glBufferData(
1997           GL_PIXEL_UNPACK_BUFFER, pmem.w * pmem.h * sizeof(uint8_t) * 4, NULL, GL_DYNAMIC_DRAW);
1998
1999     glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);
2000
2001     glActiveTexture(GL_TEXTURE0);
2002     glGenTextures(1, &pmem.cuTexId);
2003     glBindTexture(GL_TEXTURE_2D, pmem.cuTexId);
2004     if (mem.data_type == TYPE_HALF)
2005       glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA16F, pmem.w, pmem.h, 0, GL_RGBA, GL_HALF_FLOAT, NULL);
2006     else
2007       glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA8, pmem.w, pmem.h, 0, GL_RGBA, GL_UNSIGNED_BYTE, NULL);
2008     glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
2009     glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST);
2010     glBindTexture(GL_TEXTURE_2D, 0);
2011
2012     CUresult result = cuGraphicsGLRegisterBuffer(
2013         &pmem.cuPBOresource, pmem.cuPBO, CU_GRAPHICS_MAP_RESOURCE_FLAGS_NONE);
2014
2015     if (result == CUDA_SUCCESS) {
2016       mem.device_pointer = pmem.cuTexId;
2017       pixel_mem_map[mem.device_pointer] = pmem;
2018
2019       mem.device_size = mem.memory_size();
2020       stats.mem_alloc(mem.device_size);
2021
2022       return;
2023     }
2024     else {
2025       /* failed to register buffer, fallback to no interop */
2026       glDeleteBuffers(1, &pmem.cuPBO);
2027       glDeleteTextures(1, &pmem.cuTexId);
2028
2029       background = true;
2030     }
2031   }
2032
2033   void pixels_copy_from(device_memory &mem, int y, int w, int h)
2034   {
2035     PixelMem pmem = pixel_mem_map[mem.device_pointer];
2036
2037     CUDAContextScope scope(this);
2038
2039     glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pmem.cuPBO);
2040     uchar *pixels = (uchar *)glMapBuffer(GL_PIXEL_UNPACK_BUFFER, GL_READ_ONLY);
2041     size_t offset = sizeof(uchar) * 4 * y * w;
2042     memcpy((uchar *)mem.host_pointer + offset, pixels + offset, sizeof(uchar) * 4 * w * h);
2043     glUnmapBuffer(GL_PIXEL_UNPACK_BUFFER);
2044     glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);
2045   }
2046
2047   void pixels_free(device_memory &mem)
2048   {
2049     if (mem.device_pointer) {
2050       PixelMem pmem = pixel_mem_map[mem.device_pointer];
2051
2052       CUDAContextScope scope(this);
2053
2054       cuda_assert(cuGraphicsUnregisterResource(pmem.cuPBOresource));
2055       glDeleteBuffers(1, &pmem.cuPBO);
2056       glDeleteTextures(1, &pmem.cuTexId);
2057
2058       pixel_mem_map.erase(pixel_mem_map.find(mem.device_pointer));
2059       mem.device_pointer = 0;
2060
2061       stats.mem_free(mem.device_size);
2062       mem.device_size = 0;
2063     }
2064   }
2065
2066   void draw_pixels(device_memory &mem,
2067                    int y,
2068                    int w,
2069                    int h,
2070                    int width,
2071                    int height,
2072                    int dx,
2073                    int dy,
2074                    int dw,
2075                    int dh,
2076                    bool transparent,
2077                    const DeviceDrawParams &draw_params)
2078   {
2079     assert(mem.type == MEM_PIXELS);
2080
2081     if (!background) {
2082       const bool use_fallback_shader = (draw_params.bind_display_space_shader_cb == NULL);
2083       PixelMem pmem = pixel_mem_map[mem.device_pointer];
2084       float *vpointer;
2085
2086       CUDAContextScope scope(this);
2087
2088       /* for multi devices, this assumes the inefficient method that we allocate
2089        * all pixels on the device even though we only render to a subset */
2090       size_t offset = 4 * y * w;
2091
2092       if (mem.data_type == TYPE_HALF)
2093         offset *= sizeof(GLhalf);
2094       else
2095         offset *= sizeof(uint8_t);
2096
2097       glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pmem.cuPBO);
2098       glActiveTexture(GL_TEXTURE0);
2099       glBindTexture(GL_TEXTURE_2D, pmem.cuTexId);
2100       if (mem.data_type == TYPE_HALF) {
2101         glTexSubImage2D(GL_TEXTURE_2D, 0, 0, 0, w, h, GL_RGBA, GL_HALF_FLOAT, (void *)offset);
2102       }
2103       else {
2104         glTexSubImage2D(GL_TEXTURE_2D, 0, 0, 0, w, h, GL_RGBA, GL_UNSIGNED_BYTE, (void *)offset);
2105       }
2106       glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);
2107
2108       if (transparent) {
2109         glEnable(GL_BLEND);
2110         glBlendFunc(GL_ONE, GL_ONE_MINUS_SRC_ALPHA);
2111       }
2112
2113       GLint shader_program;
2114       if (use_fallback_shader) {
2115         if (!bind_fallback_display_space_shader(dw, dh)) {
2116           return;
2117         }
2118         shader_program = fallback_shader_program;
2119       }
2120       else {
2121         draw_params.bind_display_space_shader_cb();
2122         glGetIntegerv(GL_CURRENT_PROGRAM, &shader_program);
2123       }
2124
2125       if (!vertex_buffer) {
2126         glGenBuffers(1, &vertex_buffer);
2127       }
2128
2129       glBindBuffer(GL_ARRAY_BUFFER, vertex_buffer);
2130       /* invalidate old contents - avoids stalling if buffer is still waiting in queue to be rendered */
2131       glBufferData(GL_ARRAY_BUFFER, 16 * sizeof(float), NULL, GL_STREAM_DRAW);
2132
2133       vpointer = (float *)glMapBuffer(GL_ARRAY_BUFFER, GL_WRITE_ONLY);
2134
2135       if (vpointer) {
2136         /* texture coordinate - vertex pair */
2137         vpointer[0] = 0.0f;
2138         vpointer[1] = 0.0f;
2139         vpointer[2] = dx;
2140         vpointer[3] = dy;
2141
2142         vpointer[4] = (float)w / (float)pmem.w;
2143         vpointer[5] = 0.0f;
2144         vpointer[6] = (float)width + dx;
2145         vpointer[7] = dy;
2146
2147         vpointer[8] = (float)w / (float)pmem.w;
2148         vpointer[9] = (float)h / (float)pmem.h;
2149         vpointer[10] = (float)width + dx;
2150         vpointer[11] = (float)height + dy;
2151
2152         vpointer[12] = 0.0f;
2153         vpointer[13] = (float)h / (float)pmem.h;
2154         vpointer[14] = dx;
2155         vpointer[15] = (float)height + dy;
2156
2157         glUnmapBuffer(GL_ARRAY_BUFFER);
2158       }
2159
2160       GLuint vertex_array_object;
2161       GLuint position_attribute, texcoord_attribute;
2162
2163       glGenVertexArrays(1, &vertex_array_object);
2164       glBindVertexArray(vertex_array_object);
2165
2166       texcoord_attribute = glGetAttribLocation(shader_program, "texCoord");
2167       position_attribute = glGetAttribLocation(shader_program, "pos");
2168
2169       glEnableVertexAttribArray(texcoord_attribute);
2170       glEnableVertexAttribArray(position_attribute);
2171
2172       glVertexAttribPointer(
2173           texcoord_attribute, 2, GL_FLOAT, GL_FALSE, 4 * sizeof(float), (const GLvoid *)0);
2174       glVertexAttribPointer(position_attribute,
2175                             2,
2176                             GL_FLOAT,
2177                             GL_FALSE,
2178                             4 * sizeof(float),
2179                             (const GLvoid *)(sizeof(float) * 2));
2180
2181       glDrawArrays(GL_TRIANGLE_FAN, 0, 4);
2182
2183       if (use_fallback_shader) {
2184         glUseProgram(0);
2185       }
2186       else {
2187         draw_params.unbind_display_space_shader_cb();
2188       }
2189
2190       if (transparent) {
2191         glDisable(GL_BLEND);
2192       }
2193
2194       glBindTexture(GL_TEXTURE_2D, 0);
2195
2196       return;
2197     }
2198
2199     Device::draw_pixels(mem, y, w, h, width, height, dx, dy, dw, dh, transparent, draw_params);
2200   }
2201
2202   void thread_run(DeviceTask *task)
2203   {
2204     CUDAContextScope scope(this);
2205
2206     if (task->type == DeviceTask::RENDER) {
2207       DeviceRequestedFeatures requested_features;
2208       if (use_split_kernel()) {
2209         if (split_kernel == NULL) {
2210           split_kernel = new CUDASplitKernel(this);
2211           split_kernel->load_kernels(requested_features);
2212         }
2213       }
2214
2215       device_vector<WorkTile> work_tiles(this, "work_tiles", MEM_READ_ONLY);
2216
2217       /* keep rendering tiles until done */
2218       RenderTile tile;
2219       DenoisingTask denoising(this, *task);
2220
2221       while (task->acquire_tile(this, tile)) {
2222         if (tile.task == RenderTile::PATH_TRACE) {
2223           if (use_split_kernel()) {
2224             device_only_memory<uchar> void_buffer(this, "void_buffer");
2225             split_kernel->path_trace(task, tile, void_buffer, void_buffer);
2226           }
2227           else {
2228             path_trace(*task, tile, work_tiles);
2229           }
2230         }
2231         else if (tile.task == RenderTile::DENOISE) {
2232           tile.sample = tile.start_sample + tile.num_samples;
2233
2234           denoise(tile, denoising);
2235
2236           task->update_progress(&tile, tile.w * tile.h);
2237         }
2238
2239         task->release_tile(tile);
2240
2241         if (task->get_cancel()) {
2242           if (task->need_finish_queue == false)
2243             break;
2244         }
2245       }
2246
2247       work_tiles.free();
2248     }
2249     else if (task->type == DeviceTask::SHADER) {
2250       shader(*task);
2251
2252       cuda_assert(cuCtxSynchronize());
2253     }
2254   }
2255
2256   class CUDADeviceTask : public DeviceTask {
2257    public:
2258     CUDADeviceTask(CUDADevice *device, DeviceTask &task) : DeviceTask(task)
2259     {
2260       run = function_bind(&CUDADevice::thread_run, device, this);
2261     }
2262   };
2263
2264   int get_split_task_count(DeviceTask & /*task*/)
2265   {
2266     return 1;
2267   }
2268
2269   void task_add(DeviceTask &task)
2270   {
2271     CUDAContextScope scope(this);
2272
2273     /* Load texture info. */
2274     load_texture_info();
2275
2276     /* Synchronize all memory copies before executing task. */
2277     cuda_assert(cuCtxSynchronize());
2278
2279     if (task.type == DeviceTask::FILM_CONVERT) {
2280       /* must be done in main thread due to opengl access */
2281       film_convert(task, task.buffer, task.rgba_byte, task.rgba_half);
2282     }
2283     else {
2284       task_pool.push(new CUDADeviceTask(this, task));
2285     }
2286   }
2287
2288   void task_wait()
2289   {
2290     task_pool.wait();
2291   }
2292
2293   void task_cancel()
2294   {
2295     task_pool.cancel();
2296   }
2297
2298   friend class CUDASplitKernelFunction;
2299   friend class CUDASplitKernel;
2300   friend class CUDAContextScope;
2301 };
2302
2303 /* redefine the cuda_assert macro so it can be used outside of the CUDADevice class
2304  * now that the definition of that class is complete
2305  */
2306 #undef cuda_assert
2307 #define cuda_assert(stmt) \
2308   { \
2309     CUresult result = stmt; \
2310 \
2311     if (result != CUDA_SUCCESS) { \
2312       string message = string_printf("CUDA error: %s in %s", cuewErrorString(result), #stmt); \
2313       if (device->error_msg == "") \
2314         device->error_msg = message; \
2315       fprintf(stderr, "%s\n", message.c_str()); \
2316       /*cuda_abort();*/ \
2317       device->cuda_error_documentation(); \
2318     } \
2319   } \
2320   (void)0
2321
2322 /* CUDA context scope. */
2323
2324 CUDAContextScope::CUDAContextScope(CUDADevice *device) : device(device)
2325 {
2326   cuda_assert(cuCtxPushCurrent(device->cuContext));
2327 }
2328
2329 CUDAContextScope::~CUDAContextScope()
2330 {
2331   cuda_assert(cuCtxPopCurrent(NULL));
2332 }
2333
2334 /* split kernel */
2335
2336 class CUDASplitKernelFunction : public SplitKernelFunction {
2337   CUDADevice *device;
2338   CUfunction func;
2339
2340  public:
2341   CUDASplitKernelFunction(CUDADevice *device, CUfunction func) : device(device), func(func)
2342   {
2343   }
2344
2345   /* enqueue the kernel, returns false if there is an error */
2346   bool enqueue(const KernelDimensions &dim, device_memory & /*kg*/, device_memory & /*data*/)
2347   {
2348     return enqueue(dim, NULL);
2349   }
2350
2351   /* enqueue the kernel, returns false if there is an error */
2352   bool enqueue(const KernelDimensions &dim, void *args[])
2353   {
2354     if (device->have_error())
2355       return false;
2356
2357     CUDAContextScope scope(device);
2358
2359     /* we ignore dim.local_size for now, as this is faster */
2360     int threads_per_block;
2361     cuda_assert(
2362         cuFuncGetAttribute(&threads_per_block, CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, func));
2363
2364     int xblocks = (dim.global_size[0] * dim.global_size[1] + threads_per_block - 1) /
2365                   threads_per_block;
2366
2367     cuda_assert(cuFuncSetCacheConfig(func, CU_FUNC_CACHE_PREFER_L1));
2368
2369     cuda_assert(cuLaunchKernel(func,
2370                                xblocks,
2371                                1,
2372                                1, /* blocks */
2373                                threads_per_block,
2374                                1,
2375                                1, /* threads */
2376                                0,
2377                                0,
2378                                args,
2379                                0));
2380
2381     return !device->have_error();
2382   }
2383 };
2384
2385 CUDASplitKernel::CUDASplitKernel(CUDADevice *device) : DeviceSplitKernel(device), device(device)
2386 {
2387 }
2388
2389 uint64_t CUDASplitKernel::state_buffer_size(device_memory & /*kg*/,
2390                                             device_memory & /*data*/,
2391                                             size_t num_threads)
2392 {
2393   CUDAContextScope scope(device);
2394
2395   device_vector<uint64_t> size_buffer(device, "size_buffer", MEM_READ_WRITE);
2396   size_buffer.alloc(1);
2397   size_buffer.zero_to_device();
2398
2399   uint threads = num_threads;
2400   CUdeviceptr d_size = device->cuda_device_ptr(size_buffer.device_pointer);
2401
2402   struct args_t {
2403     uint *num_threads;
2404     CUdeviceptr *size;
2405   };
2406
2407   args_t args = {&threads, &d_size};
2408
2409   CUfunction state_buffer_size;
2410   cuda_assert(
2411       cuModuleGetFunction(&state_buffer_size, device->cuModule, "kernel_cuda_state_buffer_size"));
2412
2413   cuda_assert(cuLaunchKernel(state_buffer_size, 1, 1, 1, 1, 1, 1, 0, 0, (void **)&args, 0));
2414
2415   size_buffer.copy_from_device(0, 1, 1);
2416   size_t size = size_buffer[0];
2417   size_buffer.free();
2418
2419   return size;
2420 }
2421
2422 bool CUDASplitKernel::enqueue_split_kernel_data_init(const KernelDimensions &dim,
2423                                                      RenderTile &rtile,
2424                                                      int num_global_elements,
2425                                                      device_memory & /*kernel_globals*/,
2426                                                      device_memory & /*kernel_data*/,
2427                                                      device_memory &split_data,
2428                                                      device_memory &ray_state,
2429                                                      device_memory &queue_index,
2430                                                      device_memory &use_queues_flag,
2431                                                      device_memory &work_pool_wgs)
2432 {
2433   CUDAContextScope scope(device);
2434
2435   CUdeviceptr d_split_data = device->cuda_device_ptr(split_data.device_pointer);
2436   CUdeviceptr d_ray_state = device->cuda_device_ptr(ray_state.device_pointer);
2437   CUdeviceptr d_queue_index = device->cuda_device_ptr(queue_index.device_pointer);
2438   CUdeviceptr d_use_queues_flag = device->cuda_device_ptr(use_queues_flag.device_pointer);
2439   CUdeviceptr d_work_pool_wgs = device->cuda_device_ptr(work_pool_wgs.device_pointer);
2440
2441   CUdeviceptr d_buffer = device->cuda_device_ptr(rtile.buffer);
2442
2443   int end_sample = rtile.start_sample + rtile.num_samples;
2444   int queue_size = dim.global_size[0] * dim.global_size[1];
2445
2446   struct args_t {
2447     CUdeviceptr *split_data_buffer;
2448     int *num_elements;
2449     CUdeviceptr *ray_state;
2450     int *start_sample;
2451     int *end_sample;
2452     int *sx;
2453     int *sy;
2454     int *sw;
2455     int *sh;
2456     int *offset;
2457     int *stride;
2458     CUdeviceptr *queue_index;
2459     int *queuesize;
2460     CUdeviceptr *use_queues_flag;
2461     CUdeviceptr *work_pool_wgs;
2462     int *num_samples;
2463     CUdeviceptr *buffer;
2464   };
2465
2466   args_t args = {&d_split_data,
2467                  &num_global_elements,
2468                  &d_ray_state,
2469                  &rtile.start_sample,
2470                  &end_sample,
2471                  &rtile.x,
2472                  &rtile.y,
2473                  &rtile.w,
2474                  &rtile.h,
2475                  &rtile.offset,
2476                  &rtile.stride,
2477                  &d_queue_index,
2478                  &queue_size,
2479                  &d_use_queues_flag,
2480                  &d_work_pool_wgs,
2481                  &rtile.num_samples,
2482                  &d_buffer};
2483
2484   CUfunction data_init;
2485   cuda_assert(
2486       cuModuleGetFunction(&data_init, device->cuModule, "kernel_cuda_path_trace_data_init"));
2487   if (device->have_error()) {
2488     return false;
2489   }
2490
2491   CUDASplitKernelFunction(device, data_init).enqueue(dim, (void **)&args);
2492
2493   return !device->have_error();
2494 }
2495
2496 SplitKernelFunction *CUDASplitKernel::get_split_kernel_function(const string &kernel_name,
2497                                                                 const DeviceRequestedFeatures &)
2498 {
2499   CUDAContextScope scope(device);
2500   CUfunction func;
2501
2502   cuda_assert(
2503       cuModuleGetFunction(&func, device->cuModule, (string("kernel_cuda_") + kernel_name).data()));
2504   if (device->have_error()) {
2505     device->cuda_error_message(
2506         string_printf("kernel \"kernel_cuda_%s\" not found in module", kernel_name.data()));
2507     return NULL;
2508   }
2509
2510   return new CUDASplitKernelFunction(device, func);
2511 }
2512
2513 int2 CUDASplitKernel::split_kernel_local_size()
2514 {
2515   return make_int2(32, 1);
2516 }
2517
2518 int2 CUDASplitKernel::split_kernel_global_size(device_memory &kg,
2519                                                device_memory &data,
2520                                                DeviceTask * /*task*/)
2521 {
2522   CUDAContextScope scope(device);
2523   size_t free;
2524   size_t total;
2525
2526   cuda_assert(cuMemGetInfo(&free, &total));
2527
2528   VLOG(1) << "Maximum device allocation size: " << string_human_readable_number(free)
2529           << " bytes. (" << string_human_readable_size(free) << ").";
2530
2531   size_t num_elements = max_elements_for_max_buffer_size(kg, data, free / 2);
2532   size_t side = round_down((int)sqrt(num_elements), 32);
2533   int2 global_size = make_int2(side, round_down(num_elements / side, 16));
2534   VLOG(1) << "Global size: " << global_size << ".";
2535   return global_size;
2536 }
2537
2538 bool device_cuda_init()
2539 {
2540 #ifdef WITH_CUDA_DYNLOAD
2541   static bool initialized = false;
2542   static bool result = false;
2543
2544   if (initialized)
2545     return result;
2546
2547   initialized = true;
2548   int cuew_result = cuewInit(CUEW_INIT_CUDA);
2549   if (cuew_result == CUEW_SUCCESS) {
2550     VLOG(1) << "CUEW initialization succeeded";
2551     if (CUDADevice::have_precompiled_kernels()) {
2552       VLOG(1) << "Found precompiled kernels";
2553       result = true;
2554     }
2555 #  ifndef _WIN32
2556     else if (cuewCompilerPath() != NULL) {
2557       VLOG(1) << "Found CUDA compiler " << cuewCompilerPath();
2558       result = true;
2559     }
2560     else {
2561       VLOG(1) << "Neither precompiled kernels nor CUDA compiler was found,"
2562               << " unable to use CUDA";
2563     }
2564 #  endif
2565   }
2566   else {
2567     VLOG(1) << "CUEW initialization failed: "
2568             << ((cuew_result == CUEW_ERROR_ATEXIT_FAILED) ? "Error setting up atexit() handler" :
2569                                                             "Error opening the library");
2570   }
2571
2572   return result;
2573 #else  /* WITH_CUDA_DYNLOAD */
2574   return true;
2575 #endif /* WITH_CUDA_DYNLOAD */
2576 }
2577
2578 Device *device_cuda_create(DeviceInfo &info, Stats &stats, Profiler &profiler, bool background)
2579 {
2580   return new CUDADevice(info, stats, profiler, background);
2581 }
2582
2583 static CUresult device_cuda_safe_init()
2584 {
2585 #ifdef _WIN32
2586   __try {
2587     return cuInit(0);
2588   }
2589   __except (EXCEPTION_EXECUTE_HANDLER) {
2590     /* Ignore crashes inside the CUDA driver and hope we can
2591      * survive even with corrupted CUDA installs. */
2592     fprintf(stderr, "Cycles CUDA: driver crashed, continuing without CUDA.\n");
2593   }
2594
2595   return CUDA_ERROR_NO_DEVICE;
2596 #else
2597   return cuInit(0);
2598 #endif
2599 }
2600
2601 void device_cuda_info(vector<DeviceInfo> &devices)
2602 {
2603   CUresult result = device_cuda_safe_init();
2604   if (result != CUDA_SUCCESS) {
2605     if (result != CUDA_ERROR_NO_DEVICE)
2606       fprintf(stderr, "CUDA cuInit: %s\n", cuewErrorString(result));
2607     return;
2608   }
2609
2610   int count = 0;
2611   result = cuDeviceGetCount(&count);
2612   if (result != CUDA_SUCCESS) {
2613     fprintf(stderr, "CUDA cuDeviceGetCount: %s\n", cuewErrorString(result));
2614     return;
2615   }
2616
2617   vector<DeviceInfo> display_devices;
2618
2619   for (int num = 0; num < count; num++) {
2620     char name[256];
2621
2622     result = cuDeviceGetName(name, 256, num);
2623     if (result != CUDA_SUCCESS) {
2624       fprintf(stderr, "CUDA cuDeviceGetName: %s\n", cuewErrorString(result));
2625       continue;
2626     }
2627
2628     int major;
2629     cuDeviceGetAttribute(&major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, num);
2630     if (major < 3) {
2631       VLOG(1) << "Ignoring device \"" << name << "\", this graphics card is no longer supported.";
2632       continue;
2633     }
2634
2635     DeviceInfo info;
2636
2637     info.type = DEVICE_CUDA;
2638     info.description = string(name);
2639     info.num = num;
2640
2641     info.has_half_images = (major >= 3);
2642     info.has_volume_decoupled = false;
2643
2644     int pci_location[3] = {0, 0, 0};
2645     cuDeviceGetAttribute(&pci_location[0], CU_DEVICE_ATTRIBUTE_PCI_DOMAIN_ID, num);
2646     cuDeviceGetAttribute(&pci_location[1], CU_DEVICE_ATTRIBUTE_PCI_BUS_ID, num);
2647     cuDeviceGetAttribute(&pci_location[2], CU_DEVICE_ATTRIBUTE_PCI_DEVICE_ID, num);
2648     info.id = string_printf("CUDA_%s_%04x:%02x:%02x",
2649                             name,
2650                             (unsigned int)pci_location[0],
2651                             (unsigned int)pci_location[1],
2652                             (unsigned int)pci_location[2]);
2653
2654     /* If device has a kernel timeout and no compute preemption, we assume
2655      * it is connected to a display and will freeze the display while doing
2656      * computations. */
2657     int timeout_attr = 0, preempt_attr = 0;
2658     cuDeviceGetAttribute(&timeout_attr, CU_DEVICE_ATTRIBUTE_KERNEL_EXEC_TIMEOUT, num);
2659     cuDeviceGetAttribute(&preempt_attr, CU_DEVICE_ATTRIBUTE_COMPUTE_PREEMPTION_SUPPORTED, num);
2660
2661     if (timeout_attr && !preempt_attr) {
2662       VLOG(1) << "Device is recognized as display.";
2663       info.description += " (Display)";
2664       info.display_device = true;
2665       display_devices.push_back(info);
2666     }
2667     else {
2668       devices.push_back(info);
2669     }
2670     VLOG(1) << "Added device \"" << name << "\" with id \"" << info.id << "\".";
2671   }
2672
2673   if (!display_devices.empty())
2674     devices.insert(devices.end(), display_devices.begin(), display_devices.end());
2675 }
2676
2677 string device_cuda_capabilities()
2678 {
2679   CUresult result = device_cuda_safe_init();
2680   if (result != CUDA_SUCCESS) {
2681     if (result != CUDA_ERROR_NO_DEVICE) {
2682       return string("Error initializing CUDA: ") + cuewErrorString(result);
2683     }
2684     return "No CUDA device found\n";
2685   }
2686
2687   int count;
2688   result = cuDeviceGetCount(&count);
2689   if (result != CUDA_SUCCESS) {
2690     return string("Error getting devices: ") + cuewErrorString(result);
2691   }
2692
2693   string capabilities = "";
2694   for (int num = 0; num < count; num++) {
2695     char name[256];
2696     if (cuDeviceGetName(name, 256, num) != CUDA_SUCCESS) {
2697       continue;
2698     }
2699     capabilities += string("\t") + name + "\n";
2700     int value;
2701 #define GET_ATTR(attr) \
2702   { \
2703     if (cuDeviceGetAttribute(&value, CU_DEVICE_ATTRIBUTE_##attr, num) == CUDA_SUCCESS) { \
2704       capabilities += string_printf("\t\tCU_DEVICE_ATTRIBUTE_" #attr "\t\t\t%d\n", value); \
2705     } \
2706   } \
2707   (void)0
2708     /* TODO(sergey): Strip all attributes which are not useful for us
2709      * or does not depend on the driver.
2710      */
2711     GET_ATTR(MAX_THREADS_PER_BLOCK);
2712     GET_ATTR(MAX_BLOCK_DIM_X);
2713     GET_ATTR(MAX_BLOCK_DIM_Y);
2714     GET_ATTR(MAX_BLOCK_DIM_Z);
2715     GET_ATTR(MAX_GRID_DIM_X);
2716     GET_ATTR(MAX_GRID_DIM_Y);
2717     GET_ATTR(MAX_GRID_DIM_Z);
2718     GET_ATTR(MAX_SHARED_MEMORY_PER_BLOCK);
2719     GET_ATTR(SHARED_MEMORY_PER_BLOCK);
2720     GET_ATTR(TOTAL_CONSTANT_MEMORY);
2721     GET_ATTR(WARP_SIZE);
2722     GET_ATTR(MAX_PITCH);
2723     GET_ATTR(MAX_REGISTERS_PER_BLOCK);
2724     GET_ATTR(REGISTERS_PER_BLOCK);
2725     GET_ATTR(CLOCK_RATE);
2726     GET_ATTR(TEXTURE_ALIGNMENT);
2727     GET_ATTR(GPU_OVERLAP);
2728     GET_ATTR(MULTIPROCESSOR_COUNT);
2729     GET_ATTR(KERNEL_EXEC_TIMEOUT);
2730     GET_ATTR(INTEGRATED);
2731     GET_ATTR(CAN_MAP_HOST_MEMORY);
2732     GET_ATTR(COMPUTE_MODE);
2733     GET_ATTR(MAXIMUM_TEXTURE1D_WIDTH);
2734     GET_ATTR(MAXIMUM_TEXTURE2D_WIDTH);
2735     GET_ATTR(MAXIMUM_TEXTURE2D_HEIGHT);
2736     GET_ATTR(MAXIMUM_TEXTURE3D_WIDTH);
2737     GET_ATTR(MAXIMUM_TEXTURE3D_HEIGHT);
2738     GET_ATTR(MAXIMUM_TEXTURE3D_DEPTH);
2739     GET_ATTR(MAXIMUM_TEXTURE2D_LAYERED_WIDTH);
2740     GET_ATTR(MAXIMUM_TEXTURE2D_LAYERED_HEIGHT);
2741     GET_ATTR(MAXIMUM_TEXTURE2D_LAYERED_LAYERS);
2742     GET_ATTR(MAXIMUM_TEXTURE2D_ARRAY_WIDTH);
2743     GET_ATTR(MAXIMUM_TEXTURE2D_ARRAY_HEIGHT);
2744     GET_ATTR(MAXIMUM_TEXTURE2D_ARRAY_NUMSLICES);
2745     GET_ATTR(SURFACE_ALIGNMENT);
2746     GET_ATTR(CONCURRENT_KERNELS);
2747     GET_ATTR(ECC_ENABLED);
2748     GET_ATTR(TCC_DRIVER);
2749     GET_ATTR(MEMORY_CLOCK_RATE);
2750     GET_ATTR(GLOBAL_MEMORY_BUS_WIDTH);
2751     GET_ATTR(L2_CACHE_SIZE);
2752     GET_ATTR(MAX_THREADS_PER_MULTIPROCESSOR);
2753     GET_ATTR(ASYNC_ENGINE_COUNT);
2754     GET_ATTR(UNIFIED_ADDRESSING);
2755     GET_ATTR(MAXIMUM_TEXTURE1D_LAYERED_WIDTH);
2756     GET_ATTR(MAXIMUM_TEXTURE1D_LAYERED_LAYERS);
2757     GET_ATTR(CAN_TEX2D_GATHER);
2758     GET_ATTR(MAXIMUM_TEXTURE2D_GATHER_WIDTH);
2759     GET_ATTR(MAXIMUM_TEXTURE2D_GATHER_HEIGHT);
2760     GET_ATTR(MAXIMUM_TEXTURE3D_WIDTH_ALTERNATE);
2761     GET_ATTR(MAXIMUM_TEXTURE3D_HEIGHT_ALTERNATE);
2762     GET_ATTR(MAXIMUM_TEXTURE3D_DEPTH_ALTERNATE);
2763     GET_ATTR(TEXTURE_PITCH_ALIGNMENT);
2764     GET_ATTR(MAXIMUM_TEXTURECUBEMAP_WIDTH);
2765     GET_ATTR(MAXIMUM_TEXTURECUBEMAP_LAYERED_WIDTH);
2766     GET_ATTR(MAXIMUM_TEXTURECUBEMAP_LAYERED_LAYERS);
2767     GET_ATTR(MAXIMUM_SURFACE1D_WIDTH);
2768     GET_ATTR(MAXIMUM_SURFACE2D_WIDTH);
2769     GET_ATTR(MAXIMUM_SURFACE2D_HEIGHT);
2770     GET_ATTR(MAXIMUM_SURFACE3D_WIDTH);
2771     GET_ATTR(MAXIMUM_SURFACE3D_HEIGHT);
2772     GET_ATTR(MAXIMUM_SURFACE3D_DEPTH);
2773     GET_ATTR(MAXIMUM_SURFACE1D_LAYERED_WIDTH);
2774     GET_ATTR(MAXIMUM_SURFACE1D_LAYERED_LAYERS);
2775     GET_ATTR(MAXIMUM_SURFACE2D_LAYERED_WIDTH);
2776     GET_ATTR(MAXIMUM_SURFACE2D_LAYERED_HEIGHT);
2777     GET_ATTR(MAXIMUM_SURFACE2D_LAYERED_LAYERS);
2778     GET_ATTR(MAXIMUM_SURFACECUBEMAP_WIDTH);
2779     GET_ATTR(MAXIMUM_SURFACECUBEMAP_LAYERED_WIDTH);
2780     GET_ATTR(MAXIMUM_SURFACECUBEMAP_LAYERED_LAYERS);
2781     GET_ATTR(MAXIMUM_TEXTURE1D_LINEAR_WIDTH);
2782     GET_ATTR(MAXIMUM_TEXTURE2D_LINEAR_WIDTH);
2783     GET_ATTR(MAXIMUM_TEXTURE2D_LINEAR_HEIGHT);
2784     GET_ATTR(MAXIMUM_TEXTURE2D_LINEAR_PITCH);
2785     GET_ATTR(MAXIMUM_TEXTURE2D_MIPMAPPED_WIDTH);
2786     GET_ATTR(MAXIMUM_TEXTURE2D_MIPMAPPED_HEIGHT);
2787     GET_ATTR(COMPUTE_CAPABILITY_MAJOR);
2788     GET_ATTR(COMPUTE_CAPABILITY_MINOR);
2789     GET_ATTR(MAXIMUM_TEXTURE1D_MIPMAPPED_WIDTH);
2790     GET_ATTR(STREAM_PRIORITIES_SUPPORTED);
2791     GET_ATTR(GLOBAL_L1_CACHE_SUPPORTED);
2792     GET_ATTR(LOCAL_L1_CACHE_SUPPORTED);
2793     GET_ATTR(MAX_SHARED_MEMORY_PER_MULTIPROCESSOR);
2794     GET_ATTR(MAX_REGISTERS_PER_MULTIPROCESSOR);
2795     GET_ATTR(MANAGED_MEMORY);
2796     GET_ATTR(MULTI_GPU_BOARD);
2797     GET_ATTR(MULTI_GPU_BOARD_GROUP_ID);
2798 #undef GET_ATTR
2799     capabilities += "\n";
2800   }
2801
2802   return capabilities;
2803 }
2804
2805 CCL_NAMESPACE_END