cb8b0783570651b1da2af9c0ef6c2bb31701aed7
[blender-staging.git] / intern / cycles / device / opencl / opencl_base.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 #ifdef WITH_OPENCL
18
19 #include "opencl.h"
20
21 #include "kernel_types.h"
22
23 #include "util_foreach.h"
24 #include "util_logging.h"
25 #include "util_md5.h"
26 #include "util_path.h"
27 #include "util_time.h"
28
29 CCL_NAMESPACE_BEGIN
30
31 bool OpenCLDeviceBase::opencl_error(cl_int err)
32 {
33         if(err != CL_SUCCESS) {
34                 string message = string_printf("OpenCL error (%d): %s", err, clewErrorString(err));
35                 if(error_msg == "")
36                         error_msg = message;
37                 fprintf(stderr, "%s\n", message.c_str());
38                 return true;
39         }
40
41         return false;
42 }
43
44 void OpenCLDeviceBase::opencl_error(const string& message)
45 {
46         if(error_msg == "")
47                 error_msg = message;
48         fprintf(stderr, "%s\n", message.c_str());
49 }
50
51 void OpenCLDeviceBase::opencl_assert_err(cl_int err, const char* where)
52 {
53         if(err != CL_SUCCESS) {
54                 string message = string_printf("OpenCL error (%d): %s in %s", err, clewErrorString(err), where);
55                 if(error_msg == "")
56                         error_msg = message;
57                 fprintf(stderr, "%s\n", message.c_str());
58 #ifndef NDEBUG
59                 abort();
60 #endif
61         }
62 }
63
64 OpenCLDeviceBase::OpenCLDeviceBase(DeviceInfo& info, Stats &stats, bool background_)
65 : Device(info, stats, background_)
66 {
67         cpPlatform = NULL;
68         cdDevice = NULL;
69         cxContext = NULL;
70         cqCommandQueue = NULL;
71         null_mem = 0;
72         device_initialized = false;
73
74         vector<OpenCLPlatformDevice> usable_devices;
75         OpenCLInfo::get_usable_devices(&usable_devices);
76         if(usable_devices.size() == 0) {
77                 opencl_error("OpenCL: no devices found.");
78                 return;
79         }
80         assert(info.num < usable_devices.size());
81         OpenCLPlatformDevice& platform_device = usable_devices[info.num];
82         cpPlatform = platform_device.platform_id;
83         cdDevice = platform_device.device_id;
84         platform_name = platform_device.platform_name;
85         VLOG(2) << "Creating new Cycles device for OpenCL platform "
86                 << platform_name << ", device "
87                 << platform_device.device_name << ".";
88
89         {
90                 /* try to use cached context */
91                 thread_scoped_lock cache_locker;
92                 cxContext = OpenCLCache::get_context(cpPlatform, cdDevice, cache_locker);
93
94                 if(cxContext == NULL) {
95                         /* create context properties array to specify platform */
96                         const cl_context_properties context_props[] = {
97                                 CL_CONTEXT_PLATFORM, (cl_context_properties)cpPlatform,
98                                 0, 0
99                         };
100
101                         /* create context */
102                         cxContext = clCreateContext(context_props, 1, &cdDevice,
103                                 context_notify_callback, cdDevice, &ciErr);
104
105                         if(opencl_error(ciErr)) {
106                                 opencl_error("OpenCL: clCreateContext failed");
107                                 return;
108                         }
109
110                         /* cache it */
111                         OpenCLCache::store_context(cpPlatform, cdDevice, cxContext, cache_locker);
112                 }
113         }
114
115         cqCommandQueue = clCreateCommandQueue(cxContext, cdDevice, 0, &ciErr);
116         if(opencl_error(ciErr))
117                 return;
118
119         null_mem = (device_ptr)clCreateBuffer(cxContext, CL_MEM_READ_ONLY, 1, NULL, &ciErr);
120         if(opencl_error(ciErr))
121                 return;
122
123         fprintf(stderr, "Device init success\n");
124         device_initialized = true;
125 }
126
127 OpenCLDeviceBase::~OpenCLDeviceBase()
128 {
129         task_pool.stop();
130
131         if(null_mem)
132                 clReleaseMemObject(CL_MEM_PTR(null_mem));
133
134         ConstMemMap::iterator mt;
135         for(mt = const_mem_map.begin(); mt != const_mem_map.end(); mt++) {
136                 mem_free(*(mt->second));
137                 delete mt->second;
138         }
139
140         base_program.release();
141         if(cqCommandQueue)
142                 clReleaseCommandQueue(cqCommandQueue);
143         if(cxContext)
144                 clReleaseContext(cxContext);
145 }
146
147 void CL_CALLBACK OpenCLDeviceBase::context_notify_callback(const char *err_info,
148         const void * /*private_info*/, size_t /*cb*/, void *user_data)
149 {
150         char name[256];
151         clGetDeviceInfo((cl_device_id)user_data, CL_DEVICE_NAME, sizeof(name), &name, NULL);
152
153         fprintf(stderr, "OpenCL error (%s): %s\n", name, err_info);
154 }
155
156 bool OpenCLDeviceBase::opencl_version_check()
157 {
158         string error;
159         if(!OpenCLInfo::platform_version_check(cpPlatform, &error)) {
160                 opencl_error(error);
161                 return false;
162         }
163         if(!OpenCLInfo::device_version_check(cdDevice, &error)) {
164                 opencl_error(error);
165                 return false;
166         }
167         return true;
168 }
169
170 string OpenCLDeviceBase::device_md5_hash(string kernel_custom_build_options)
171 {
172         MD5Hash md5;
173         char version[256], driver[256], name[256], vendor[256];
174
175         clGetPlatformInfo(cpPlatform, CL_PLATFORM_VENDOR, sizeof(vendor), &vendor, NULL);
176         clGetDeviceInfo(cdDevice, CL_DEVICE_VERSION, sizeof(version), &version, NULL);
177         clGetDeviceInfo(cdDevice, CL_DEVICE_NAME, sizeof(name), &name, NULL);
178         clGetDeviceInfo(cdDevice, CL_DRIVER_VERSION, sizeof(driver), &driver, NULL);
179
180         md5.append((uint8_t*)vendor, strlen(vendor));
181         md5.append((uint8_t*)version, strlen(version));
182         md5.append((uint8_t*)name, strlen(name));
183         md5.append((uint8_t*)driver, strlen(driver));
184
185         string options = kernel_build_options();
186         options += kernel_custom_build_options;
187         md5.append((uint8_t*)options.c_str(), options.size());
188
189         return md5.get_hex();
190 }
191
192 bool OpenCLDeviceBase::load_kernels(const DeviceRequestedFeatures& requested_features)
193 {
194         /* Verify if device was initialized. */
195         if(!device_initialized) {
196                 fprintf(stderr, "OpenCL: failed to initialize device.\n");
197                 return false;
198         }
199
200         /* Verify we have right opencl version. */
201         if(!opencl_version_check())
202                 return false;
203
204         base_program = OpenCLProgram(this, "base", "kernel.cl", build_options_for_base_program(requested_features));
205         base_program.add_kernel(ustring("convert_to_byte"));
206         base_program.add_kernel(ustring("convert_to_half_float"));
207         base_program.add_kernel(ustring("shader"));
208         base_program.add_kernel(ustring("bake"));
209         base_program.add_kernel(ustring("zero_buffer"));
210
211         vector<OpenCLProgram*> programs;
212         programs.push_back(&base_program);
213         /* Call actual class to fill the vector with its programs. */
214         if(!load_kernels(requested_features, programs)) {
215                 return false;
216         }
217
218         /* Parallel compilation is supported by Cycles, but currently all OpenCL frameworks
219          * serialize the calls internally, so it's not much use right now.
220          * Note: When enabling parallel compilation, use_stdout in the OpenCLProgram constructor
221          * should be set to false as well. */
222 #if 0
223         TaskPool task_pool;
224         foreach(OpenCLProgram *program, programs) {
225                 task_pool.push(function_bind(&OpenCLProgram::load, program));
226         }
227         task_pool.wait_work();
228
229         foreach(OpenCLProgram *program, programs) {
230                 VLOG(2) << program->get_log();
231                 if(!program->is_loaded()) {
232                         program->report_error();
233                         return false;
234                 }
235         }
236 #else
237         foreach(OpenCLProgram *program, programs) {
238                 program->load();
239                 if(!program->is_loaded()) {
240                         return false;
241                 }
242         }
243 #endif
244
245         return true;
246 }
247
248 void OpenCLDeviceBase::mem_alloc(const char *name, device_memory& mem, MemoryType type)
249 {
250         if(name) {
251                 VLOG(1) << "Buffer allocate: " << name << ", "
252                             << string_human_readable_number(mem.memory_size()) << " bytes. ("
253                             << string_human_readable_size(mem.memory_size()) << ")";
254         }
255
256         size_t size = mem.memory_size();
257
258         cl_mem_flags mem_flag;
259         void *mem_ptr = NULL;
260
261         if(type == MEM_READ_ONLY)
262                 mem_flag = CL_MEM_READ_ONLY;
263         else if(type == MEM_WRITE_ONLY)
264                 mem_flag = CL_MEM_WRITE_ONLY;
265         else
266                 mem_flag = CL_MEM_READ_WRITE;
267
268         /* Zero-size allocation might be invoked by render, but not really
269          * supported by OpenCL. Using NULL as device pointer also doesn't really
270          * work for some reason, so for the time being we'll use special case
271          * will null_mem buffer.
272          */
273         if(size != 0) {
274                 mem.device_pointer = (device_ptr)clCreateBuffer(cxContext,
275                                                                 mem_flag,
276                                                                 size,
277                                                                 mem_ptr,
278                                                                 &ciErr);
279                 opencl_assert_err(ciErr, "clCreateBuffer");
280         }
281         else {
282                 mem.device_pointer = null_mem;
283         }
284
285         stats.mem_alloc(size);
286         mem.device_size = size;
287 }
288
289 void OpenCLDeviceBase::mem_copy_to(device_memory& mem)
290 {
291         /* this is blocking */
292         size_t size = mem.memory_size();
293         if(size != 0) {
294                 opencl_assert(clEnqueueWriteBuffer(cqCommandQueue,
295                                                    CL_MEM_PTR(mem.device_pointer),
296                                                    CL_TRUE,
297                                                    0,
298                                                    size,
299                                                    (void*)mem.data_pointer,
300                                                    0,
301                                                    NULL, NULL));
302         }
303 }
304
305 void OpenCLDeviceBase::mem_copy_from(device_memory& mem, int y, int w, int h, int elem)
306 {
307         size_t offset = elem*y*w;
308         size_t size = elem*w*h;
309         assert(size != 0);
310         opencl_assert(clEnqueueReadBuffer(cqCommandQueue,
311                                           CL_MEM_PTR(mem.device_pointer),
312                                           CL_TRUE,
313                                           offset,
314                                           size,
315                                           (uchar*)mem.data_pointer + offset,
316                                           0,
317                                           NULL, NULL));
318 }
319
320 void OpenCLDeviceBase::mem_zero(device_memory& mem)
321 {
322         if(mem.device_pointer) {
323                 if(base_program.is_loaded()) {
324                         cl_kernel ckZeroBuffer = base_program(ustring("zero_buffer"));
325
326                         size_t global_size[] = {1024, 1024};
327                         size_t num_threads = global_size[0] * global_size[1];
328
329                         cl_mem d_buffer = CL_MEM_PTR(mem.device_pointer);
330                         unsigned long long d_offset = 0;
331                         unsigned long long d_size = 0;
332
333                         while(d_offset < mem.memory_size()) {
334                                 d_size = std::min<unsigned long long>(num_threads*sizeof(float4), mem.memory_size() - d_offset);
335
336                                 kernel_set_args(ckZeroBuffer, 0, d_buffer, d_size, d_offset);
337
338                                 ciErr = clEnqueueNDRangeKernel(cqCommandQueue,
339                                                                        ckZeroBuffer,
340                                                                        2,
341                                                                        NULL,
342                                                                        global_size,
343                                                                        NULL,
344                                                                        0,
345                                                                        NULL,
346                                                                        NULL);
347                                 opencl_assert_err(ciErr, "clEnqueueNDRangeKernel");
348
349                                 d_offset += d_size;
350                         }
351                 }
352
353                 if(mem.data_pointer) {
354                         memset((void*)mem.data_pointer, 0, mem.memory_size());
355                 }
356
357                 if(!base_program.is_loaded()) {
358                         void* zero = (void*)mem.data_pointer;
359
360                         if(!mem.data_pointer) {
361                                 zero = util_aligned_malloc(mem.memory_size(), 16);
362                                 memset(zero, 0, mem.memory_size());
363                         }
364
365                         opencl_assert(clEnqueueWriteBuffer(cqCommandQueue,
366                                    CL_MEM_PTR(mem.device_pointer),
367                                    CL_TRUE,
368                                    0,
369                                    mem.memory_size(),
370                                    zero,
371                                    0,
372                                    NULL, NULL));
373
374                         if(!mem.data_pointer) {
375                                 util_aligned_free(zero);
376                         }
377                 }
378         }
379 }
380
381 void OpenCLDeviceBase::mem_free(device_memory& mem)
382 {
383         if(mem.device_pointer) {
384                 if(mem.device_pointer != null_mem) {
385                         opencl_assert(clReleaseMemObject(CL_MEM_PTR(mem.device_pointer)));
386                 }
387                 mem.device_pointer = 0;
388
389                 stats.mem_free(mem.device_size);
390                 mem.device_size = 0;
391         }
392 }
393
394 void OpenCLDeviceBase::const_copy_to(const char *name, void *host, size_t size)
395 {
396         ConstMemMap::iterator i = const_mem_map.find(name);
397
398         if(i == const_mem_map.end()) {
399                 device_vector<uchar> *data = new device_vector<uchar>();
400                 data->copy((uchar*)host, size);
401
402                 mem_alloc(name, *data, MEM_READ_ONLY);
403                 i = const_mem_map.insert(ConstMemMap::value_type(name, data)).first;
404         }
405         else {
406                 device_vector<uchar> *data = i->second;
407                 data->copy((uchar*)host, size);
408         }
409
410         mem_copy_to(*i->second);
411 }
412
413 void OpenCLDeviceBase::tex_alloc(const char *name,
414                device_memory& mem,
415                InterpolationType /*interpolation*/,
416                ExtensionType /*extension*/)
417 {
418         VLOG(1) << "Texture allocate: " << name << ", "
419                 << string_human_readable_number(mem.memory_size()) << " bytes. ("
420                 << string_human_readable_size(mem.memory_size()) << ")";
421         mem_alloc(NULL, mem, MEM_READ_ONLY);
422         mem_copy_to(mem);
423         assert(mem_map.find(name) == mem_map.end());
424         mem_map.insert(MemMap::value_type(name, mem.device_pointer));
425 }
426
427 void OpenCLDeviceBase::tex_free(device_memory& mem)
428 {
429         if(mem.device_pointer) {
430                 foreach(const MemMap::value_type& value, mem_map) {
431                         if(value.second == mem.device_pointer) {
432                                 mem_map.erase(value.first);
433                                 break;
434                         }
435                 }
436
437                 mem_free(mem);
438         }
439 }
440
441 size_t OpenCLDeviceBase::global_size_round_up(int group_size, int global_size)
442 {
443         int r = global_size % group_size;
444         return global_size + ((r == 0)? 0: group_size - r);
445 }
446
447 void OpenCLDeviceBase::enqueue_kernel(cl_kernel kernel, size_t w, size_t h)
448 {
449         size_t workgroup_size, max_work_items[3];
450
451         clGetKernelWorkGroupInfo(kernel, cdDevice,
452                 CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &workgroup_size, NULL);
453         clGetDeviceInfo(cdDevice,
454                 CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(size_t)*3, max_work_items, NULL);
455
456         /* Try to divide evenly over 2 dimensions. */
457         size_t sqrt_workgroup_size = max((size_t)sqrt((double)workgroup_size), 1);
458         size_t local_size[2] = {sqrt_workgroup_size, sqrt_workgroup_size};
459
460         /* Some implementations have max size 1 on 2nd dimension. */
461         if(local_size[1] > max_work_items[1]) {
462                 local_size[0] = workgroup_size/max_work_items[1];
463                 local_size[1] = max_work_items[1];
464         }
465
466         size_t global_size[2] = {global_size_round_up(local_size[0], w),
467                                  global_size_round_up(local_size[1], h)};
468
469         /* Vertical size of 1 is coming from bake/shade kernels where we should
470          * not round anything up because otherwise we'll either be doing too
471          * much work per pixel (if we don't check global ID on Y axis) or will
472          * be checking for global ID to always have Y of 0.
473          */
474         if(h == 1) {
475                 global_size[h] = 1;
476         }
477
478         /* run kernel */
479         opencl_assert(clEnqueueNDRangeKernel(cqCommandQueue, kernel, 2, NULL, global_size, NULL, 0, NULL, NULL));
480         opencl_assert(clFlush(cqCommandQueue));
481 }
482
483 void OpenCLDeviceBase::set_kernel_arg_mem(cl_kernel kernel, cl_uint *narg, const char *name)
484 {
485         cl_mem ptr;
486
487         MemMap::iterator i = mem_map.find(name);
488         if(i != mem_map.end()) {
489                 ptr = CL_MEM_PTR(i->second);
490         }
491         else {
492                 /* work around NULL not working, even though the spec says otherwise */
493                 ptr = CL_MEM_PTR(null_mem);
494         }
495
496         opencl_assert(clSetKernelArg(kernel, (*narg)++, sizeof(ptr), (void*)&ptr));
497 }
498
499 void OpenCLDeviceBase::film_convert(DeviceTask& task, device_ptr buffer, device_ptr rgba_byte, device_ptr rgba_half)
500 {
501         /* cast arguments to cl types */
502         cl_mem d_data = CL_MEM_PTR(const_mem_map["__data"]->device_pointer);
503         cl_mem d_rgba = (rgba_byte)? CL_MEM_PTR(rgba_byte): CL_MEM_PTR(rgba_half);
504         cl_mem d_buffer = CL_MEM_PTR(buffer);
505         cl_int d_x = task.x;
506         cl_int d_y = task.y;
507         cl_int d_w = task.w;
508         cl_int d_h = task.h;
509         cl_float d_sample_scale = 1.0f/(task.sample + 1);
510         cl_int d_offset = task.offset;
511         cl_int d_stride = task.stride;
512
513
514         cl_kernel ckFilmConvertKernel = (rgba_byte)? base_program(ustring("convert_to_byte")): base_program(ustring("convert_to_half_float"));
515
516         cl_uint start_arg_index =
517                 kernel_set_args(ckFilmConvertKernel,
518                                 0,
519                                 d_data,
520                                 d_rgba,
521                                 d_buffer);
522
523 #define KERNEL_TEX(type, ttype, name) \
524 set_kernel_arg_mem(ckFilmConvertKernel, &start_arg_index, #name);
525 #include "kernel_textures.h"
526 #undef KERNEL_TEX
527
528         start_arg_index += kernel_set_args(ckFilmConvertKernel,
529                                            start_arg_index,
530                                            d_sample_scale,
531                                            d_x,
532                                            d_y,
533                                            d_w,
534                                            d_h,
535                                            d_offset,
536                                            d_stride);
537
538         enqueue_kernel(ckFilmConvertKernel, d_w, d_h);
539 }
540
541 void OpenCLDeviceBase::shader(DeviceTask& task)
542 {
543         /* cast arguments to cl types */
544         cl_mem d_data = CL_MEM_PTR(const_mem_map["__data"]->device_pointer);
545         cl_mem d_input = CL_MEM_PTR(task.shader_input);
546         cl_mem d_output = CL_MEM_PTR(task.shader_output);
547         cl_mem d_output_luma = CL_MEM_PTR(task.shader_output_luma);
548         cl_int d_shader_eval_type = task.shader_eval_type;
549         cl_int d_shader_filter = task.shader_filter;
550         cl_int d_shader_x = task.shader_x;
551         cl_int d_shader_w = task.shader_w;
552         cl_int d_offset = task.offset;
553
554         cl_kernel kernel;
555
556         if(task.shader_eval_type >= SHADER_EVAL_BAKE)
557                 kernel = base_program(ustring("bake"));
558         else
559                 kernel = base_program(ustring("shader"));
560
561         cl_uint start_arg_index =
562                 kernel_set_args(kernel,
563                                 0,
564                                 d_data,
565                                 d_input,
566                                 d_output);
567
568         if(task.shader_eval_type < SHADER_EVAL_BAKE) {
569                 start_arg_index += kernel_set_args(kernel,
570                                                    start_arg_index,
571                                                    d_output_luma);
572         }
573
574 #define KERNEL_TEX(type, ttype, name) \
575         set_kernel_arg_mem(kernel, &start_arg_index, #name);
576 #include "kernel_textures.h"
577 #undef KERNEL_TEX
578
579         start_arg_index += kernel_set_args(kernel,
580                                            start_arg_index,
581                                            d_shader_eval_type);
582         if(task.shader_eval_type >= SHADER_EVAL_BAKE) {
583                 start_arg_index += kernel_set_args(kernel,
584                                                    start_arg_index,
585                                                    d_shader_filter);
586         }
587         start_arg_index += kernel_set_args(kernel,
588                                            start_arg_index,
589                                            d_shader_x,
590                                            d_shader_w,
591                                            d_offset);
592
593         for(int sample = 0; sample < task.num_samples; sample++) {
594
595                 if(task.get_cancel())
596                         break;
597
598                 kernel_set_args(kernel, start_arg_index, sample);
599
600                 enqueue_kernel(kernel, task.shader_w, 1);
601
602                 clFinish(cqCommandQueue);
603
604                 task.update_progress(NULL);
605         }
606 }
607
608 string OpenCLDeviceBase::kernel_build_options(const string *debug_src)
609 {
610         string build_options = "-cl-fast-relaxed-math ";
611
612         if(platform_name == "NVIDIA CUDA") {
613                 build_options += "-D__KERNEL_OPENCL_NVIDIA__ "
614                                  "-cl-nv-maxrregcount=32 "
615                                  "-cl-nv-verbose ";
616
617                 uint compute_capability_major, compute_capability_minor;
618                 clGetDeviceInfo(cdDevice, CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV,
619                                 sizeof(cl_uint), &compute_capability_major, NULL);
620                 clGetDeviceInfo(cdDevice, CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV,
621                                 sizeof(cl_uint), &compute_capability_minor, NULL);
622
623                 build_options += string_printf("-D__COMPUTE_CAPABILITY__=%u ",
624                                                compute_capability_major * 100 +
625                                                compute_capability_minor * 10);
626         }
627
628         else if(platform_name == "Apple")
629                 build_options += "-D__KERNEL_OPENCL_APPLE__ ";
630
631         else if(platform_name == "AMD Accelerated Parallel Processing")
632                 build_options += "-D__KERNEL_OPENCL_AMD__ ";
633
634         else if(platform_name == "Intel(R) OpenCL") {
635                 build_options += "-D__KERNEL_OPENCL_INTEL_CPU__ ";
636
637                 /* Options for gdb source level kernel debugging.
638                  * this segfaults on linux currently.
639                  */
640                 if(OpenCLInfo::use_debug() && debug_src)
641                         build_options += "-g -s \"" + *debug_src + "\" ";
642         }
643
644         if(OpenCLInfo::use_debug())
645                 build_options += "-D__KERNEL_OPENCL_DEBUG__ ";
646
647 #ifdef WITH_CYCLES_DEBUG
648         build_options += "-D__KERNEL_DEBUG__ ";
649 #endif
650
651         return build_options;
652 }
653
654 /* TODO(sergey): In the future we can use variadic templates, once
655  * C++0x is allowed. Should allow to clean this up a bit.
656  */
657 int OpenCLDeviceBase::kernel_set_args(cl_kernel kernel,
658                     int start_argument_index,
659                     const ArgumentWrapper& arg1,
660                     const ArgumentWrapper& arg2,
661                     const ArgumentWrapper& arg3,
662                     const ArgumentWrapper& arg4,
663                     const ArgumentWrapper& arg5,
664                     const ArgumentWrapper& arg6,
665                     const ArgumentWrapper& arg7,
666                     const ArgumentWrapper& arg8,
667                     const ArgumentWrapper& arg9,
668                     const ArgumentWrapper& arg10,
669                     const ArgumentWrapper& arg11,
670                     const ArgumentWrapper& arg12,
671                     const ArgumentWrapper& arg13,
672                     const ArgumentWrapper& arg14,
673                     const ArgumentWrapper& arg15,
674                     const ArgumentWrapper& arg16,
675                     const ArgumentWrapper& arg17,
676                     const ArgumentWrapper& arg18,
677                     const ArgumentWrapper& arg19,
678                     const ArgumentWrapper& arg20,
679                     const ArgumentWrapper& arg21,
680                     const ArgumentWrapper& arg22,
681                     const ArgumentWrapper& arg23,
682                     const ArgumentWrapper& arg24,
683                     const ArgumentWrapper& arg25,
684                     const ArgumentWrapper& arg26,
685                     const ArgumentWrapper& arg27,
686                     const ArgumentWrapper& arg28,
687                     const ArgumentWrapper& arg29,
688                     const ArgumentWrapper& arg30,
689                     const ArgumentWrapper& arg31,
690                     const ArgumentWrapper& arg32,
691                     const ArgumentWrapper& arg33)
692 {
693         int current_arg_index = 0;
694 #define FAKE_VARARG_HANDLE_ARG(arg) \
695         do { \
696                 if(arg.pointer != NULL) { \
697                         opencl_assert(clSetKernelArg( \
698                                 kernel, \
699                                 start_argument_index + current_arg_index, \
700                                 arg.size, arg.pointer)); \
701                         ++current_arg_index; \
702                 } \
703                 else { \
704                         return current_arg_index; \
705                 } \
706         } while(false)
707         FAKE_VARARG_HANDLE_ARG(arg1);
708         FAKE_VARARG_HANDLE_ARG(arg2);
709         FAKE_VARARG_HANDLE_ARG(arg3);
710         FAKE_VARARG_HANDLE_ARG(arg4);
711         FAKE_VARARG_HANDLE_ARG(arg5);
712         FAKE_VARARG_HANDLE_ARG(arg6);
713         FAKE_VARARG_HANDLE_ARG(arg7);
714         FAKE_VARARG_HANDLE_ARG(arg8);
715         FAKE_VARARG_HANDLE_ARG(arg9);
716         FAKE_VARARG_HANDLE_ARG(arg10);
717         FAKE_VARARG_HANDLE_ARG(arg11);
718         FAKE_VARARG_HANDLE_ARG(arg12);
719         FAKE_VARARG_HANDLE_ARG(arg13);
720         FAKE_VARARG_HANDLE_ARG(arg14);
721         FAKE_VARARG_HANDLE_ARG(arg15);
722         FAKE_VARARG_HANDLE_ARG(arg16);
723         FAKE_VARARG_HANDLE_ARG(arg17);
724         FAKE_VARARG_HANDLE_ARG(arg18);
725         FAKE_VARARG_HANDLE_ARG(arg19);
726         FAKE_VARARG_HANDLE_ARG(arg20);
727         FAKE_VARARG_HANDLE_ARG(arg21);
728         FAKE_VARARG_HANDLE_ARG(arg22);
729         FAKE_VARARG_HANDLE_ARG(arg23);
730         FAKE_VARARG_HANDLE_ARG(arg24);
731         FAKE_VARARG_HANDLE_ARG(arg25);
732         FAKE_VARARG_HANDLE_ARG(arg26);
733         FAKE_VARARG_HANDLE_ARG(arg27);
734         FAKE_VARARG_HANDLE_ARG(arg28);
735         FAKE_VARARG_HANDLE_ARG(arg29);
736         FAKE_VARARG_HANDLE_ARG(arg30);
737         FAKE_VARARG_HANDLE_ARG(arg31);
738         FAKE_VARARG_HANDLE_ARG(arg32);
739         FAKE_VARARG_HANDLE_ARG(arg33);
740 #undef FAKE_VARARG_HANDLE_ARG
741         return current_arg_index;
742 }
743
744 void OpenCLDeviceBase::release_kernel_safe(cl_kernel kernel)
745 {
746         if(kernel) {
747                 clReleaseKernel(kernel);
748         }
749 }
750
751 void OpenCLDeviceBase::release_mem_object_safe(cl_mem mem)
752 {
753         if(mem != NULL) {
754                 clReleaseMemObject(mem);
755         }
756 }
757
758 void OpenCLDeviceBase::release_program_safe(cl_program program)
759 {
760         if(program) {
761                 clReleaseProgram(program);
762         }
763 }
764
765 /* ** Those guys are for workign around some compiler-specific bugs ** */
766
767 cl_program OpenCLDeviceBase::load_cached_kernel(
768         ustring key,
769         thread_scoped_lock& cache_locker)
770 {
771         return OpenCLCache::get_program(cpPlatform,
772                                         cdDevice,
773                                         key,
774                                         cache_locker);
775 }
776
777 void OpenCLDeviceBase::store_cached_kernel(
778         cl_program program,
779         ustring key,
780         thread_scoped_lock& cache_locker)
781 {
782         OpenCLCache::store_program(cpPlatform,
783                                    cdDevice,
784                                    program,
785                                    key,
786                                    cache_locker);
787 }
788
789 string OpenCLDeviceBase::build_options_for_base_program(
790         const DeviceRequestedFeatures& /*requested_features*/)
791 {
792         /* TODO(sergey): By default we compile all features, meaning
793          * mega kernel is not getting feature-based optimizations.
794          *
795          * Ideally we need always compile kernel with as less features
796          * enabled as possible to keep performance at it's max.
797          */
798         return "";
799 }
800
801 CCL_NAMESPACE_END
802
803 #endif