Cycles: Log which device kernels are being loaded for
[blender-staging.git] / intern / cycles / device / opencl / 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.h"
20
21 #include "util_map.h"
22 #include "util_param.h"
23 #include "util_string.h"
24
25 #include "clew.h"
26
27 CCL_NAMESPACE_BEGIN
28
29 /* Define CYCLES_DISABLE_DRIVER_WORKAROUNDS to disable workaounds for testing */
30 #ifndef CYCLES_DISABLE_DRIVER_WORKAROUNDS
31 /* Work around AMD driver hangs by ensuring each command is finished before doing anything else. */
32 #  undef clEnqueueNDRangeKernel
33 #  define clEnqueueNDRangeKernel(a, b, c, d, e, f, g, h, i) \
34         clFinish(a); \
35         CLEW_GET_FUN(__clewEnqueueNDRangeKernel)(a, b, c, d, e, f, g, h, i); \
36         clFinish(a);
37
38 #  undef clEnqueueWriteBuffer
39 #  define clEnqueueWriteBuffer(a, b, c, d, e, f, g, h, i) \
40         clFinish(a); \
41         CLEW_GET_FUN(__clewEnqueueWriteBuffer)(a, b, c, d, e, f, g, h, i); \
42         clFinish(a);
43
44 #  undef clEnqueueReadBuffer
45 #  define clEnqueueReadBuffer(a, b, c, d, e, f, g, h, i) \
46         clFinish(a); \
47         CLEW_GET_FUN(__clewEnqueueReadBuffer)(a, b, c, d, e, f, g, h, i); \
48         clFinish(a);
49 #endif  /* CYCLES_DISABLE_DRIVER_WORKAROUNDS */
50
51 #define CL_MEM_PTR(p) ((cl_mem)(uintptr_t)(p))
52
53 struct OpenCLPlatformDevice {
54         OpenCLPlatformDevice(cl_platform_id platform_id,
55                              const string& platform_name,
56                              cl_device_id device_id,
57                              cl_device_type device_type,
58                              const string& device_name,
59                              const string& hardware_id)
60           : platform_id(platform_id),
61             platform_name(platform_name),
62             device_id(device_id),
63             device_type(device_type),
64             device_name(device_name),
65             hardware_id(hardware_id) {}
66         cl_platform_id platform_id;
67         string platform_name;
68         cl_device_id device_id;
69         cl_device_type device_type;
70         string device_name;
71         string hardware_id;
72 };
73
74 /* Contains all static OpenCL helper functions. */
75 class OpenCLInfo
76 {
77 public:
78         static cl_device_type device_type();
79         static bool use_debug();
80         static bool kernel_use_advanced_shading(const string& platform_name);
81         static bool kernel_use_split(const string& platform_name,
82                                      const cl_device_type device_type);
83         static bool device_supported(const string& platform_name,
84                                      const cl_device_id device_id);
85         static bool platform_version_check(cl_platform_id platform,
86                                            string *error = NULL);
87         static bool device_version_check(cl_device_id device,
88                                          string *error = NULL);
89         static string get_hardware_id(string platform_name,
90                                       cl_device_id device_id);
91         static void get_usable_devices(vector<OpenCLPlatformDevice> *usable_devices,
92                                        bool force_all = false);
93 };
94
95 /* Thread safe cache for contexts and programs.
96  */
97 class OpenCLCache
98 {
99         struct Slot
100         {
101                 struct ProgramEntry
102                 {
103                         ProgramEntry();
104                         ProgramEntry(const ProgramEntry& rhs);
105                         ~ProgramEntry();
106                         cl_program program;
107                         thread_mutex *mutex;
108                 };
109
110                 Slot();
111                 Slot(const Slot& rhs);
112                 ~Slot();
113
114                 thread_mutex *context_mutex;
115                 cl_context context;
116                 typedef map<ustring, ProgramEntry> EntryMap;
117                 EntryMap programs;
118
119         };
120
121         /* key is combination of platform ID and device ID */
122         typedef pair<cl_platform_id, cl_device_id> PlatformDevicePair;
123
124         /* map of Slot objects */
125         typedef map<PlatformDevicePair, Slot> CacheMap;
126         CacheMap cache;
127
128         /* MD5 hash of the kernel source. */
129         string kernel_md5;
130
131         thread_mutex cache_lock;
132         thread_mutex kernel_md5_lock;
133
134         /* lazy instantiate */
135         static OpenCLCache& global_instance();
136
137 public:
138
139         enum ProgramName {
140                 OCL_DEV_BASE_PROGRAM,
141                 OCL_DEV_MEGAKERNEL_PROGRAM,
142         };
143
144         /* Lookup context in the cache. If this returns NULL, slot_locker
145          * will be holding a lock for the cache. slot_locker should refer to a
146          * default constructed thread_scoped_lock. */
147         static cl_context get_context(cl_platform_id platform,
148                                       cl_device_id device,
149                                       thread_scoped_lock& slot_locker);
150         /* Same as above. */
151         static cl_program get_program(cl_platform_id platform,
152                                       cl_device_id device,
153                                       ustring key,
154                                       thread_scoped_lock& slot_locker);
155
156         /* Store context in the cache. You MUST have tried to get the item before storing to it. */
157         static void store_context(cl_platform_id platform,
158                                   cl_device_id device,
159                                   cl_context context,
160                                   thread_scoped_lock& slot_locker);
161         /* Same as above. */
162         static void store_program(cl_platform_id platform,
163                                   cl_device_id device,
164                                   cl_program program,
165                                   ustring key,
166                                   thread_scoped_lock& slot_locker);
167
168         static string get_kernel_md5();
169 };
170
171 #define opencl_assert(stmt) \
172         { \
173                 cl_int err = stmt; \
174                 \
175                 if(err != CL_SUCCESS) { \
176                         string message = string_printf("OpenCL error: %s in %s", clewErrorString(err), #stmt); \
177                         if(error_msg == "") \
178                                 error_msg = message; \
179                         fprintf(stderr, "%s\n", message.c_str()); \
180                 } \
181         } (void)0
182
183 class OpenCLDeviceBase : public Device
184 {
185 public:
186         DedicatedTaskPool task_pool;
187         cl_context cxContext;
188         cl_command_queue cqCommandQueue;
189         cl_platform_id cpPlatform;
190         cl_device_id cdDevice;
191         cl_int ciErr;
192
193         class OpenCLProgram {
194         public:
195                 OpenCLProgram() : loaded(false), device(NULL) {}
196                 OpenCLProgram(OpenCLDeviceBase *device,
197                               string program_name,
198                               string kernel_name,
199                               string kernel_build_options,
200                               bool use_stdout = true);
201                 ~OpenCLProgram();
202
203                 void add_kernel(ustring name);
204                 void load();
205
206                 bool is_loaded()    { return loaded; }
207                 string get_log()    { return log; }
208                 void report_error();
209
210                 cl_kernel operator()();
211                 cl_kernel operator()(ustring name);
212
213                 void release();
214
215         private:
216                 bool build_kernel(const string *debug_src);
217                 bool compile_kernel(const string *debug_src);
218                 bool load_binary(const string& clbin, const string *debug_src = NULL);
219                 bool save_binary(const string& clbin);
220
221                 void add_log(string msg, bool is_debug);
222                 void add_error(string msg);
223
224                 bool loaded;
225                 cl_program program;
226                 OpenCLDeviceBase *device;
227
228                 /* Used for the OpenCLCache key. */
229                 string program_name;
230
231                 string kernel_file, kernel_build_options, device_md5;
232
233                 bool use_stdout;
234                 string log, error_msg;
235                 string compile_output;
236
237                 map<ustring, cl_kernel> kernels;
238         };
239
240         OpenCLProgram base_program;
241
242         typedef map<string, device_vector<uchar>*> ConstMemMap;
243         typedef map<string, device_ptr> MemMap;
244
245         ConstMemMap const_mem_map;
246         MemMap mem_map;
247         device_ptr null_mem;
248
249         bool device_initialized;
250         string platform_name;
251         string device_name;
252
253         bool opencl_error(cl_int err);
254         void opencl_error(const string& message);
255         void opencl_assert_err(cl_int err, const char* where);
256
257         OpenCLDeviceBase(DeviceInfo& info, Stats &stats, bool background_);
258         ~OpenCLDeviceBase();
259
260         static void CL_CALLBACK context_notify_callback(const char *err_info,
261                 const void * /*private_info*/, size_t /*cb*/, void *user_data);
262
263         bool opencl_version_check();
264
265         string device_md5_hash(string kernel_custom_build_options = "");
266         bool load_kernels(const DeviceRequestedFeatures& requested_features);
267
268         /* Has to be implemented by the real device classes.
269          * The base device will then load all these programs. */
270         virtual bool load_kernels(const DeviceRequestedFeatures& requested_features,
271                                   vector<OpenCLProgram*> &programs) = 0;
272
273         void mem_alloc(const char *name, device_memory& mem, MemoryType type);
274         void mem_copy_to(device_memory& mem);
275         void mem_copy_from(device_memory& mem, int y, int w, int h, int elem);
276         void mem_zero(device_memory& mem);
277         void mem_free(device_memory& mem);
278         void const_copy_to(const char *name, void *host, size_t size);
279         void tex_alloc(const char *name,
280                        device_memory& mem,
281                        InterpolationType /*interpolation*/,
282                        ExtensionType /*extension*/);
283         void tex_free(device_memory& mem);
284
285         size_t global_size_round_up(int group_size, int global_size);
286         void enqueue_kernel(cl_kernel kernel, size_t w, size_t h);
287         void set_kernel_arg_mem(cl_kernel kernel, cl_uint *narg, const char *name);
288
289         void film_convert(DeviceTask& task, device_ptr buffer, device_ptr rgba_byte, device_ptr rgba_half);
290         void shader(DeviceTask& task);
291
292         class OpenCLDeviceTask : public DeviceTask {
293         public:
294                 OpenCLDeviceTask(OpenCLDeviceBase *device, DeviceTask& task)
295                 : DeviceTask(task)
296                 {
297                         run = function_bind(&OpenCLDeviceBase::thread_run,
298                                             device,
299                                             this);
300                 }
301         };
302
303         int get_split_task_count(DeviceTask& /*task*/)
304         {
305                 return 1;
306         }
307
308         void task_add(DeviceTask& task)
309         {
310                 task_pool.push(new OpenCLDeviceTask(this, task));
311         }
312
313         void task_wait()
314         {
315                 task_pool.wait();
316         }
317
318         void task_cancel()
319         {
320                 task_pool.cancel();
321         }
322
323         virtual void thread_run(DeviceTask * /*task*/) = 0;
324
325 protected:
326         string kernel_build_options(const string *debug_src = NULL);
327
328         class ArgumentWrapper {
329         public:
330                 ArgumentWrapper() : size(0), pointer(NULL)
331                 {
332                 }
333
334                 ArgumentWrapper(device_memory& argument) : size(sizeof(void*)),
335                                                            pointer((void*)(&argument.device_pointer))
336                 {
337                 }
338
339                 template<typename T>
340                 ArgumentWrapper(device_vector<T>& argument) : size(sizeof(void*)),
341                                                               pointer((void*)(&argument.device_pointer))
342                 {
343                 }
344
345                 template<typename T>
346                 ArgumentWrapper(T& argument) : size(sizeof(argument)),
347                                                pointer(&argument)
348                 {
349                 }
350
351                 ArgumentWrapper(int argument) : size(sizeof(int)),
352                                                 int_value(argument),
353                                                 pointer(&int_value)
354                 {
355                 }
356
357                 ArgumentWrapper(float argument) : size(sizeof(float)),
358                                                   float_value(argument),
359                                                   pointer(&float_value)
360                 {
361                 }
362
363                 size_t size;
364                 int int_value;
365                 float float_value;
366                 void *pointer;
367         };
368
369         /* TODO(sergey): In the future we can use variadic templates, once
370          * C++0x is allowed. Should allow to clean this up a bit.
371          */
372         int kernel_set_args(cl_kernel kernel,
373                             int start_argument_index,
374                             const ArgumentWrapper& arg1 = ArgumentWrapper(),
375                             const ArgumentWrapper& arg2 = ArgumentWrapper(),
376                             const ArgumentWrapper& arg3 = ArgumentWrapper(),
377                             const ArgumentWrapper& arg4 = ArgumentWrapper(),
378                             const ArgumentWrapper& arg5 = ArgumentWrapper(),
379                             const ArgumentWrapper& arg6 = ArgumentWrapper(),
380                             const ArgumentWrapper& arg7 = ArgumentWrapper(),
381                             const ArgumentWrapper& arg8 = ArgumentWrapper(),
382                             const ArgumentWrapper& arg9 = ArgumentWrapper(),
383                             const ArgumentWrapper& arg10 = ArgumentWrapper(),
384                             const ArgumentWrapper& arg11 = ArgumentWrapper(),
385                             const ArgumentWrapper& arg12 = ArgumentWrapper(),
386                             const ArgumentWrapper& arg13 = ArgumentWrapper(),
387                             const ArgumentWrapper& arg14 = ArgumentWrapper(),
388                             const ArgumentWrapper& arg15 = ArgumentWrapper(),
389                             const ArgumentWrapper& arg16 = ArgumentWrapper(),
390                             const ArgumentWrapper& arg17 = ArgumentWrapper(),
391                             const ArgumentWrapper& arg18 = ArgumentWrapper(),
392                             const ArgumentWrapper& arg19 = ArgumentWrapper(),
393                             const ArgumentWrapper& arg20 = ArgumentWrapper(),
394                             const ArgumentWrapper& arg21 = ArgumentWrapper(),
395                             const ArgumentWrapper& arg22 = ArgumentWrapper(),
396                             const ArgumentWrapper& arg23 = ArgumentWrapper(),
397                             const ArgumentWrapper& arg24 = ArgumentWrapper(),
398                             const ArgumentWrapper& arg25 = ArgumentWrapper(),
399                             const ArgumentWrapper& arg26 = ArgumentWrapper(),
400                             const ArgumentWrapper& arg27 = ArgumentWrapper(),
401                             const ArgumentWrapper& arg28 = ArgumentWrapper(),
402                             const ArgumentWrapper& arg29 = ArgumentWrapper(),
403                             const ArgumentWrapper& arg30 = ArgumentWrapper(),
404                             const ArgumentWrapper& arg31 = ArgumentWrapper(),
405                             const ArgumentWrapper& arg32 = ArgumentWrapper(),
406                             const ArgumentWrapper& arg33 = ArgumentWrapper());
407
408         void release_kernel_safe(cl_kernel kernel);
409         void release_mem_object_safe(cl_mem mem);
410         void release_program_safe(cl_program program);
411
412         /* ** Those guys are for workign around some compiler-specific bugs ** */
413
414         virtual cl_program load_cached_kernel(
415                 ustring key,
416                 thread_scoped_lock& cache_locker);
417
418         virtual void store_cached_kernel(
419                 cl_program program,
420                 ustring key,
421                 thread_scoped_lock& cache_locker);
422
423         virtual string build_options_for_base_program(
424                 const DeviceRequestedFeatures& /*requested_features*/);
425 };
426
427 Device *opencl_create_mega_device(DeviceInfo& info, Stats& stats, bool background);
428 Device *opencl_create_split_device(DeviceInfo& info, Stats& stats, bool background);
429
430 CCL_NAMESPACE_END
431
432 #endif