Geometry Nodes: new Material input node
[blender.git] / intern / cycles / device / opencl / device_opencl.h
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 #ifdef WITH_OPENCL
18
19 #  include "device/device.h"
20 #  include "device/device_denoising.h"
21 #  include "device/device_split_kernel.h"
22
23 #  include "util/util_map.h"
24 #  include "util/util_param.h"
25 #  include "util/util_string.h"
26 #  include "util/util_task.h"
27
28 #  include "clew.h"
29
30 #  include "device/opencl/memory_manager.h"
31
32 CCL_NAMESPACE_BEGIN
33
34 /* Disable workarounds, seems to be working fine on latest drivers. */
35 #  define CYCLES_DISABLE_DRIVER_WORKAROUNDS
36
37 /* Define CYCLES_DISABLE_DRIVER_WORKAROUNDS to disable workarounds for testing. */
38 #  ifndef CYCLES_DISABLE_DRIVER_WORKAROUNDS
39 /* Work around AMD driver hangs by ensuring each command is finished before doing anything else. */
40 #    undef clEnqueueNDRangeKernel
41 #    define clEnqueueNDRangeKernel(a, b, c, d, e, f, g, h, i) \
42       CLEW_GET_FUN(__clewEnqueueNDRangeKernel)(a, b, c, d, e, f, g, h, i); \
43       clFinish(a);
44
45 #    undef clEnqueueWriteBuffer
46 #    define clEnqueueWriteBuffer(a, b, c, d, e, f, g, h, i) \
47       CLEW_GET_FUN(__clewEnqueueWriteBuffer)(a, b, c, d, e, f, g, h, i); \
48       clFinish(a);
49
50 #    undef clEnqueueReadBuffer
51 #    define clEnqueueReadBuffer(a, b, c, d, e, f, g, h, i) \
52       CLEW_GET_FUN(__clewEnqueueReadBuffer)(a, b, c, d, e, f, g, h, i); \
53       clFinish(a);
54 #  endif /* CYCLES_DISABLE_DRIVER_WORKAROUNDS */
55
56 #  define CL_MEM_PTR(p) ((cl_mem)(uintptr_t)(p))
57
58 struct OpenCLPlatformDevice {
59   OpenCLPlatformDevice(cl_platform_id platform_id,
60                        const string &platform_name,
61                        cl_device_id device_id,
62                        cl_device_type device_type,
63                        const string &device_name,
64                        const string &hardware_id,
65                        const string &device_extensions)
66       : platform_id(platform_id),
67         platform_name(platform_name),
68         device_id(device_id),
69         device_type(device_type),
70         device_name(device_name),
71         hardware_id(hardware_id),
72         device_extensions(device_extensions)
73   {
74   }
75   cl_platform_id platform_id;
76   string platform_name;
77   cl_device_id device_id;
78   cl_device_type device_type;
79   string device_name;
80   string hardware_id;
81   string device_extensions;
82 };
83
84 /* Contains all static OpenCL helper functions. */
85 class OpenCLInfo {
86  public:
87   static cl_device_type device_type();
88   static bool use_debug();
89   static bool device_supported(const string &platform_name, const cl_device_id device_id);
90   static bool platform_version_check(cl_platform_id platform, string *error = NULL);
91   static bool device_version_check(cl_device_id device, string *error = NULL);
92   static bool get_device_version(cl_device_id device,
93                                  int *r_major,
94                                  int *r_minor,
95                                  string *error = NULL);
96   static string get_hardware_id(const string &platform_name, cl_device_id device_id);
97   static void get_usable_devices(vector<OpenCLPlatformDevice> *usable_devices);
98
99   /* ** Some handy shortcuts to low level cl*GetInfo() functions. ** */
100
101   /* Platform information. */
102   static bool get_num_platforms(cl_uint *num_platforms, cl_int *error = NULL);
103   static cl_uint get_num_platforms();
104
105   static bool get_platforms(vector<cl_platform_id> *platform_ids, cl_int *error = NULL);
106   static vector<cl_platform_id> get_platforms();
107
108   static bool get_platform_name(cl_platform_id platform_id, string *platform_name);
109   static string get_platform_name(cl_platform_id platform_id);
110
111   static bool get_num_platform_devices(cl_platform_id platform_id,
112                                        cl_device_type device_type,
113                                        cl_uint *num_devices,
114                                        cl_int *error = NULL);
115   static cl_uint get_num_platform_devices(cl_platform_id platform_id, cl_device_type device_type);
116
117   static bool get_platform_devices(cl_platform_id platform_id,
118                                    cl_device_type device_type,
119                                    vector<cl_device_id> *device_ids,
120                                    cl_int *error = NULL);
121   static vector<cl_device_id> get_platform_devices(cl_platform_id platform_id,
122                                                    cl_device_type device_type);
123
124   /* Device information. */
125   static bool get_device_name(cl_device_id device_id, string *device_name, cl_int *error = NULL);
126
127   static string get_device_name(cl_device_id device_id);
128
129   static bool get_device_extensions(cl_device_id device_id,
130                                     string *device_extensions,
131                                     cl_int *error = NULL);
132
133   static string get_device_extensions(cl_device_id device_id);
134
135   static bool get_device_type(cl_device_id device_id,
136                               cl_device_type *device_type,
137                               cl_int *error = NULL);
138   static cl_device_type get_device_type(cl_device_id device_id);
139
140   static bool get_driver_version(cl_device_id device_id,
141                                  int *major,
142                                  int *minor,
143                                  cl_int *error = NULL);
144
145   static int mem_sub_ptr_alignment(cl_device_id device_id);
146
147   /* Get somewhat more readable device name.
148    * Main difference is AMD OpenCL here which only gives code name
149    * for the regular device name. This will give more sane device
150    * name using some extensions.
151    */
152   static string get_readable_device_name(cl_device_id device_id);
153 };
154
155 /* Thread safe cache for contexts and programs.
156  */
157 class OpenCLCache {
158   struct Slot {
159     struct ProgramEntry {
160       ProgramEntry();
161       ProgramEntry(const ProgramEntry &rhs);
162       ~ProgramEntry();
163       cl_program program;
164       thread_mutex *mutex;
165     };
166
167     Slot();
168     Slot(const Slot &rhs);
169     ~Slot();
170
171     thread_mutex *context_mutex;
172     cl_context context;
173     typedef map<ustring, ProgramEntry> EntryMap;
174     EntryMap programs;
175   };
176
177   /* key is combination of platform ID and device ID */
178   typedef pair<cl_platform_id, cl_device_id> PlatformDevicePair;
179
180   /* map of Slot objects */
181   typedef map<PlatformDevicePair, Slot> CacheMap;
182   CacheMap cache;
183
184   /* MD5 hash of the kernel source. */
185   string kernel_md5;
186
187   thread_mutex cache_lock;
188   thread_mutex kernel_md5_lock;
189
190   /* lazy instantiate */
191   static OpenCLCache &global_instance();
192
193  public:
194   enum ProgramName {
195     OCL_DEV_BASE_PROGRAM,
196     OCL_DEV_MEGAKERNEL_PROGRAM,
197   };
198
199   /* Lookup context in the cache. If this returns NULL, slot_locker
200    * will be holding a lock for the cache. slot_locker should refer to a
201    * default constructed thread_scoped_lock. */
202   static cl_context get_context(cl_platform_id platform,
203                                 cl_device_id device,
204                                 thread_scoped_lock &slot_locker);
205   /* Same as above. */
206   static cl_program get_program(cl_platform_id platform,
207                                 cl_device_id device,
208                                 ustring key,
209                                 thread_scoped_lock &slot_locker);
210
211   /* Store context in the cache. You MUST have tried to get the item before storing to it. */
212   static void store_context(cl_platform_id platform,
213                             cl_device_id device,
214                             cl_context context,
215                             thread_scoped_lock &slot_locker);
216   /* Same as above. */
217   static void store_program(cl_platform_id platform,
218                             cl_device_id device,
219                             cl_program program,
220                             ustring key,
221                             thread_scoped_lock &slot_locker);
222
223   static string get_kernel_md5();
224 };
225
226 #  define opencl_device_assert(device, stmt) \
227     { \
228       cl_int err = stmt; \
229 \
230       if (err != CL_SUCCESS) { \
231         string message = string_printf( \
232             "OpenCL error: %s in %s (%s:%d)", clewErrorString(err), #stmt, __FILE__, __LINE__); \
233         if ((device)->error_message() == "") { \
234           (device)->set_error(message); \
235         } \
236         fprintf(stderr, "%s\n", message.c_str()); \
237       } \
238     } \
239     (void)0
240
241 #  define opencl_assert(stmt) \
242     { \
243       cl_int err = stmt; \
244 \
245       if (err != CL_SUCCESS) { \
246         string message = string_printf( \
247             "OpenCL error: %s in %s (%s:%d)", clewErrorString(err), #stmt, __FILE__, __LINE__); \
248         if (error_msg == "") { \
249           error_msg = message; \
250         } \
251         fprintf(stderr, "%s\n", message.c_str()); \
252       } \
253     } \
254     (void)0
255
256 class OpenCLDevice : public Device {
257  public:
258   DedicatedTaskPool task_pool;
259
260   /* Task pool for required kernels (base, AO kernels during foreground rendering) */
261   TaskPool load_required_kernel_task_pool;
262   /* Task pool for optional kernels (feature kernels during foreground rendering) */
263   TaskPool load_kernel_task_pool;
264   std::atomic<int> load_kernel_num_compiling;
265
266   cl_context cxContext;
267   cl_command_queue cqCommandQueue;
268   cl_platform_id cpPlatform;
269   cl_device_id cdDevice;
270   cl_int ciErr;
271   int device_num;
272   bool use_preview_kernels;
273
274   class OpenCLProgram {
275    public:
276     OpenCLProgram() : loaded(false), needs_compiling(true), program(NULL), device(NULL)
277     {
278     }
279     OpenCLProgram(OpenCLDevice *device,
280                   const string &program_name,
281                   const string &kernel_name,
282                   const string &kernel_build_options,
283                   bool use_stdout = true);
284     ~OpenCLProgram();
285
286     void add_kernel(ustring name);
287
288     /* Try to load the program from device cache or disk */
289     bool load();
290     /* Compile the kernel (first separate, fail-back to local). */
291     void compile();
292     /* Create the OpenCL kernels after loading or compiling */
293     void create_kernels();
294
295     bool is_loaded() const
296     {
297       return loaded;
298     }
299     const string &get_log() const
300     {
301       return log;
302     }
303     void report_error();
304
305     /* Wait until this kernel is available to be used
306      * It will return true when the kernel is available.
307      * It will return false when the kernel is not available
308      * or could not be loaded. */
309     bool wait_for_availability();
310
311     cl_kernel operator()();
312     cl_kernel operator()(ustring name);
313
314     void release();
315
316    private:
317     bool build_kernel(const string *debug_src);
318     /* Build the program by calling the own process.
319      * This is required for multithreaded OpenCL compilation, since most Frameworks serialize
320      * build calls internally if they come from the same process.
321      * If that is not supported, this function just returns false.
322      */
323     bool compile_separate(const string &clbin);
324     /* Build the program by calling OpenCL directly. */
325     bool compile_kernel(const string *debug_src);
326     /* Loading and saving the program from/to disk. */
327     bool load_binary(const string &clbin, const string *debug_src = NULL);
328     bool save_binary(const string &clbin);
329
330     void add_log(const string &msg, bool is_debug);
331     void add_error(const string &msg);
332
333     bool loaded;
334     bool needs_compiling;
335
336     cl_program program;
337     OpenCLDevice *device;
338
339     /* Used for the OpenCLCache key. */
340     string program_name;
341
342     string kernel_file, kernel_build_options, device_md5;
343
344     bool use_stdout;
345     string log, error_msg;
346     string compile_output;
347
348     map<ustring, cl_kernel> kernels;
349   };
350
351   /* Container for all types of split programs. */
352   class OpenCLSplitPrograms {
353    public:
354     OpenCLDevice *device;
355     OpenCLProgram program_split;
356     OpenCLProgram program_lamp_emission;
357     OpenCLProgram program_do_volume;
358     OpenCLProgram program_indirect_background;
359     OpenCLProgram program_shader_eval;
360     OpenCLProgram program_holdout_emission_blurring_pathtermination_ao;
361     OpenCLProgram program_subsurface_scatter;
362     OpenCLProgram program_direct_lighting;
363     OpenCLProgram program_shadow_blocked_ao;
364     OpenCLProgram program_shadow_blocked_dl;
365
366     OpenCLSplitPrograms(OpenCLDevice *device);
367     ~OpenCLSplitPrograms();
368
369     /* Load the kernels and put the created kernels in the given
370      * `programs` parameter. */
371     void load_kernels(vector<OpenCLProgram *> &programs,
372                       const DeviceRequestedFeatures &requested_features,
373                       bool is_preview = false);
374   };
375
376   DeviceSplitKernel *split_kernel;
377
378   OpenCLProgram base_program;
379   OpenCLProgram bake_program;
380   OpenCLProgram displace_program;
381   OpenCLProgram background_program;
382   OpenCLProgram denoising_program;
383
384   OpenCLSplitPrograms kernel_programs;
385   OpenCLSplitPrograms preview_programs;
386
387   typedef map<string, device_vector<uchar> *> ConstMemMap;
388   typedef map<string, device_ptr> MemMap;
389
390   ConstMemMap const_mem_map;
391   MemMap mem_map;
392
393   bool device_initialized;
394   string platform_name;
395   string device_name;
396
397   bool opencl_error(cl_int err);
398   void opencl_error(const string &message);
399   void opencl_assert_err(cl_int err, const char *where);
400
401   OpenCLDevice(DeviceInfo &info, Stats &stats, Profiler &profiler, bool background);
402   ~OpenCLDevice();
403
404   static void CL_CALLBACK context_notify_callback(const char *err_info,
405                                                   const void * /*private_info*/,
406                                                   size_t /*cb*/,
407                                                   void *user_data);
408
409   bool opencl_version_check();
410   OpenCLSplitPrograms *get_split_programs();
411
412   string device_md5_hash(string kernel_custom_build_options = "");
413   bool load_kernels(const DeviceRequestedFeatures &requested_features);
414   void load_required_kernels(const DeviceRequestedFeatures &requested_features);
415   void load_preview_kernels();
416
417   bool wait_for_availability(const DeviceRequestedFeatures &requested_features);
418   DeviceKernelStatus get_active_kernel_switch_state();
419
420   /* Get the name of the opencl program for the given kernel */
421   const string get_opencl_program_name(const string &kernel_name);
422   /* Get the program file name to compile (*.cl) for the given kernel */
423   const string get_opencl_program_filename(const string &kernel_name);
424   string get_build_options(const DeviceRequestedFeatures &requested_features,
425                            const string &opencl_program_name,
426                            bool preview_kernel = false);
427   /* Enable the default features to reduce recompilation events */
428   void enable_default_features(DeviceRequestedFeatures &features);
429
430   void mem_alloc(device_memory &mem);
431   void mem_copy_to(device_memory &mem);
432   void mem_copy_from(device_memory &mem, int y, int w, int h, int elem);
433   void mem_zero(device_memory &mem);
434   void mem_free(device_memory &mem);
435
436   int mem_sub_ptr_alignment();
437
438   void const_copy_to(const char *name, void *host, size_t size);
439   void global_alloc(device_memory &mem);
440   void global_free(device_memory &mem);
441   void tex_alloc(device_texture &mem);
442   void tex_free(device_texture &mem);
443
444   size_t global_size_round_up(int group_size, int global_size);
445   void enqueue_kernel(cl_kernel kernel,
446                       size_t w,
447                       size_t h,
448                       bool x_workgroups = false,
449                       size_t max_workgroup_size = -1);
450   void set_kernel_arg_mem(cl_kernel kernel, cl_uint *narg, const char *name);
451   void set_kernel_arg_buffers(cl_kernel kernel, cl_uint *narg);
452
453   void film_convert(DeviceTask &task,
454                     device_ptr buffer,
455                     device_ptr rgba_byte,
456                     device_ptr rgba_half);
457   void shader(DeviceTask &task);
458   void update_adaptive(DeviceTask &task, RenderTile &tile, int sample);
459   void bake(DeviceTask &task, RenderTile &tile);
460
461   void denoise(RenderTile &tile, DenoisingTask &denoising);
462
463   int get_split_task_count(DeviceTask & /*task*/)
464   {
465     return 1;
466   }
467
468   void task_add(DeviceTask &task)
469   {
470     task_pool.push([=] {
471       DeviceTask task_copy = task;
472       thread_run(task_copy);
473     });
474   }
475
476   void task_wait()
477   {
478     task_pool.wait();
479   }
480
481   void task_cancel()
482   {
483     task_pool.cancel();
484   }
485
486   void thread_run(DeviceTask &task);
487
488   virtual BVHLayoutMask get_bvh_layout_mask() const
489   {
490     return BVH_LAYOUT_BVH2;
491   }
492
493   virtual bool show_samples() const
494   {
495     return true;
496   }
497
498  protected:
499   string kernel_build_options(const string *debug_src = NULL);
500
501   void mem_zero_kernel(device_ptr ptr, size_t size);
502
503   bool denoising_non_local_means(device_ptr image_ptr,
504                                  device_ptr guide_ptr,
505                                  device_ptr variance_ptr,
506                                  device_ptr out_ptr,
507                                  DenoisingTask *task);
508   bool denoising_construct_transform(DenoisingTask *task);
509   bool denoising_accumulate(device_ptr color_ptr,
510                             device_ptr color_variance_ptr,
511                             device_ptr scale_ptr,
512                             int frame,
513                             DenoisingTask *task);
514   bool denoising_solve(device_ptr output_ptr, DenoisingTask *task);
515   bool denoising_combine_halves(device_ptr a_ptr,
516                                 device_ptr b_ptr,
517                                 device_ptr mean_ptr,
518                                 device_ptr variance_ptr,
519                                 int r,
520                                 int4 rect,
521                                 DenoisingTask *task);
522   bool denoising_divide_shadow(device_ptr a_ptr,
523                                device_ptr b_ptr,
524                                device_ptr sample_variance_ptr,
525                                device_ptr sv_variance_ptr,
526                                device_ptr buffer_variance_ptr,
527                                DenoisingTask *task);
528   bool denoising_get_feature(int mean_offset,
529                              int variance_offset,
530                              device_ptr mean_ptr,
531                              device_ptr variance_ptr,
532                              float scale,
533                              DenoisingTask *task);
534   bool denoising_write_feature(int to_offset,
535                                device_ptr from_ptr,
536                                device_ptr buffer_ptr,
537                                DenoisingTask *task);
538   bool denoising_detect_outliers(device_ptr image_ptr,
539                                  device_ptr variance_ptr,
540                                  device_ptr depth_ptr,
541                                  device_ptr output_ptr,
542                                  DenoisingTask *task);
543
544   device_ptr mem_alloc_sub_ptr(device_memory &mem, int offset, int size);
545   void mem_free_sub_ptr(device_ptr ptr);
546
547   class ArgumentWrapper {
548    public:
549     ArgumentWrapper() : size(0), pointer(NULL)
550     {
551     }
552
553     ArgumentWrapper(device_memory &argument)
554         : size(sizeof(void *)), pointer((void *)(&argument.device_pointer))
555     {
556     }
557
558     template<typename T>
559     ArgumentWrapper(device_vector<T> &argument)
560         : size(sizeof(void *)), pointer((void *)(&argument.device_pointer))
561     {
562     }
563
564     template<typename T>
565     ArgumentWrapper(device_only_memory<T> &argument)
566         : size(sizeof(void *)), pointer((void *)(&argument.device_pointer))
567     {
568     }
569     template<typename T> ArgumentWrapper(T &argument) : size(sizeof(argument)), pointer(&argument)
570     {
571     }
572
573     ArgumentWrapper(int argument) : size(sizeof(int)), int_value(argument), pointer(&int_value)
574     {
575     }
576
577     ArgumentWrapper(float argument)
578         : size(sizeof(float)), float_value(argument), pointer(&float_value)
579     {
580     }
581
582     size_t size;
583     int int_value;
584     float float_value;
585     void *pointer;
586   };
587
588   /* TODO(sergey): In the future we can use variadic templates, once
589    * C++0x is allowed. Should allow to clean this up a bit.
590    */
591   int kernel_set_args(cl_kernel kernel,
592                       int start_argument_index,
593                       const ArgumentWrapper &arg1 = ArgumentWrapper(),
594                       const ArgumentWrapper &arg2 = ArgumentWrapper(),
595                       const ArgumentWrapper &arg3 = ArgumentWrapper(),
596                       const ArgumentWrapper &arg4 = ArgumentWrapper(),
597                       const ArgumentWrapper &arg5 = ArgumentWrapper(),
598                       const ArgumentWrapper &arg6 = ArgumentWrapper(),
599                       const ArgumentWrapper &arg7 = ArgumentWrapper(),
600                       const ArgumentWrapper &arg8 = ArgumentWrapper(),
601                       const ArgumentWrapper &arg9 = ArgumentWrapper(),
602                       const ArgumentWrapper &arg10 = ArgumentWrapper(),
603                       const ArgumentWrapper &arg11 = ArgumentWrapper(),
604                       const ArgumentWrapper &arg12 = ArgumentWrapper(),
605                       const ArgumentWrapper &arg13 = ArgumentWrapper(),
606                       const ArgumentWrapper &arg14 = ArgumentWrapper(),
607                       const ArgumentWrapper &arg15 = ArgumentWrapper(),
608                       const ArgumentWrapper &arg16 = ArgumentWrapper(),
609                       const ArgumentWrapper &arg17 = ArgumentWrapper(),
610                       const ArgumentWrapper &arg18 = ArgumentWrapper(),
611                       const ArgumentWrapper &arg19 = ArgumentWrapper(),
612                       const ArgumentWrapper &arg20 = ArgumentWrapper(),
613                       const ArgumentWrapper &arg21 = ArgumentWrapper(),
614                       const ArgumentWrapper &arg22 = ArgumentWrapper(),
615                       const ArgumentWrapper &arg23 = ArgumentWrapper(),
616                       const ArgumentWrapper &arg24 = ArgumentWrapper(),
617                       const ArgumentWrapper &arg25 = ArgumentWrapper(),
618                       const ArgumentWrapper &arg26 = ArgumentWrapper(),
619                       const ArgumentWrapper &arg27 = ArgumentWrapper(),
620                       const ArgumentWrapper &arg28 = ArgumentWrapper(),
621                       const ArgumentWrapper &arg29 = ArgumentWrapper(),
622                       const ArgumentWrapper &arg30 = ArgumentWrapper(),
623                       const ArgumentWrapper &arg31 = ArgumentWrapper(),
624                       const ArgumentWrapper &arg32 = ArgumentWrapper(),
625                       const ArgumentWrapper &arg33 = ArgumentWrapper());
626
627   void release_kernel_safe(cl_kernel kernel);
628   void release_mem_object_safe(cl_mem mem);
629   void release_program_safe(cl_program program);
630
631   /* ** Those guys are for working around some compiler-specific bugs ** */
632
633   cl_program load_cached_kernel(ustring key, thread_scoped_lock &cache_locker);
634
635   void store_cached_kernel(cl_program program, ustring key, thread_scoped_lock &cache_locker);
636
637  private:
638   MemoryManager memory_manager;
639   friend class MemoryManager;
640
641   static_assert_align(TextureInfo, 16);
642   device_vector<TextureInfo> texture_info;
643
644   typedef map<string, device_memory *> TexturesMap;
645   TexturesMap textures;
646
647   bool textures_need_update;
648
649  protected:
650   void flush_texture_buffers();
651
652   friend class OpenCLSplitKernel;
653   friend class OpenCLSplitKernelFunction;
654 };
655
656 Device *opencl_create_split_device(DeviceInfo &info,
657                                    Stats &stats,
658                                    Profiler &profiler,
659                                    bool background);
660
661 CCL_NAMESPACE_END
662
663 #endif