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