/*
								+--------------------------------+
								|                                |
								|   ***  OpenCL utilities  ***   |
								|                                |
								|  Copyright © -tHE SWINe- 2010  |
								|                                |
								|          ClUtils.cpp           |
								|                                |
								+--------------------------------+
*/

/**
 *	@file gpgpu/ClUtils.cpp
 *	@date 2010
 *	@author -tHE SWINe-
 *	@brief OpenCL utilities
 */

#include "../NewFix.h"
#include "../CallStack.h"
#include <stdio.h>
#include <algorithm>
#if defined(_WIN32) || defined(_WIN64)
#include <io.h> //_access
#define access(f,m) _access(f,m)
#else // _WIN32 || _WIN64
#include <unistd.h>
#include <errno.h> // program_invocation_name
#endif // _WIN32 || _WIN64
#include "../Integer.h"
#include "../MinMax.h"
#include "../Dir.h" // PRIsizeB
#include "../Hash.h"
#if !defined(_MSC_VER) || defined(__MWERKS__) || _MSC_VER > 1200
#include <stdint.h>
#else // !_MSC_VER || __MWERKS__ || _MSC_VER > 1200
typedef ptrdiff_t intptr_t;
#endif // !_MSC_VER || __MWERKS__ || _MSC_VER > 1200
#include <limits.h>
//#include <CL/opencl.h> // don't
#include "ClUtils.h"

#if defined(_MSC_VER) && !defined(__MWERKS__) && !defined(for) && _MSC_VER <= 1200
#define for if(0) {} else for
#endif // _MSC_VER && !__MWERKS__ && !for && _MSC_VER <= 1200
// msvc 'for' scoping hack

/*
 *								=== CCLUtils::TDevice_GFLOPS ===
 */

double CCLUtils::TDevice_GFLOPS::operator ()(cl_device_id h_device) const
{
	cl_uint n_speed, n_processors;
	if(clGetDeviceInfo(h_device, CL_DEVICE_MAX_CLOCK_FREQUENCY,
	   sizeof(n_speed), &n_speed, NULL) != CL_SUCCESS ||
	   clGetDeviceInfo(h_device, CL_DEVICE_MAX_COMPUTE_UNITS,
	   sizeof(n_processors), &n_processors, NULL) != CL_SUCCESS)
		return -1;
	return double(n_speed) * n_processors;
}

/*
 *								=== ~CCLUtils::TDevice_GFLOPS ===
 */

/*
 *								=== CCLUtils::TDevice_GlobalMemSize ===
 */

cl_ulong CCLUtils::TDevice_GlobalMemSize::operator ()(cl_device_id h_device) const
{
	cl_ulong n_memory;
	if(clGetDeviceInfo(h_device, CL_DEVICE_GLOBAL_MEM_SIZE,
	   sizeof(n_memory), &n_memory, NULL) != CL_SUCCESS)
		return 0; // -1 would be maximum
	return n_memory;
}

/*
 *								=== ~CCLUtils::TDevice_GlobalMemSize ===
 */

/*
 *								=== CCLUtils::TDevice_ComputeCapability ===
 */

int CCLUtils::TDevice_ComputeCapability::operator ()(cl_device_id h_device) const
{
	CCLDeviceParams params(h_device);
	if(!params.b_Status() || !params.b_Is_NVIDIA())
		return -1;
	return params.n_NV_ComputeCap_Major() * 100 + params.n_NV_ComputeCap_Minor();
}

/*
 *								=== ~CCLUtils::TDevice_ComputeCapability ===
 */

/*
 *								=== CCLUtils::TDevice_OpenCL_C_Version ===
 */

int CCLUtils::TDevice_OpenCL_C_Version::operator ()(cl_device_id h_device) const
{
	int n_hi, n_lo;
#ifdef CL_DEVICE_OPENCL_C_VERSION
	CCLDeviceParams params(h_device);
	std::string s_info;
	if(!params.b_Status() || params.n_GetDeviceInfoString(s_info,
	   CL_DEVICE_OPENCL_C_VERSION) != CL_SUCCESS ||
#if defined(_MSC_VER) && !defined(__MWERKS__) && _MSC_VER >= 1400
	   sscanf_s(s_info.c_str(), "OpenCL C %d.%d", &n_hi, &n_lo) != 2) // use the safe function to avoid warnings
#else // _MSC_VER && !__MWERKS__ && _MSC_VER >= 1400
	   sscanf(s_info.c_str(), "OpenCL C %d.%d", &n_hi, &n_lo) != 2)
#endif // _MSC_VER && !__MWERKS__ && _MSC_VER >= 1400
		return -1;
#else // CL_DEVICE_OPENCL_C_VERSION
	n_hi = 1;
	n_lo = 0;
	// OpenCL 1.0 does not support CL_DEVICE_OPENCL_C_VERSION
#endif // CL_DEVICE_OPENCL_C_VERSION
	return n_hi * 100 + n_lo;
}

/*
 *								=== ~CCLUtils::TDevice_OpenCL_C_Version ===
 */

/*
 *								=== CCLUtils ===
 */

CLresult CCLUtils::n_Get_PlatformList(std::vector<cl_platform_id> &r_platform_list)
{
	r_platform_list.clear();
	// !!

	cl_uint n_platform_num = 0;
	CLresult n_result;
	if((n_result = (CLresult)clGetPlatformIDs(0, NULL, &n_platform_num)) != CL_SUCCESS)
		return n_result; // error
	if(n_platform_num == 0)
		return cl_Platform_Not_Found_KHR; // no OpenCL capable platform available
	// get number of OpenCL platforms

	try {
		r_platform_list.resize(n_platform_num);
	} catch(std::bad_alloc&) {
		return cl_Out_Of_Host_Memory;
	}
	_ASSERTE(r_platform_list.size() <= /*CL*/UINT_MAX);
	if((n_result = (CLresult)clGetPlatformIDs(cl_uint(r_platform_list.size()),
	   &r_platform_list[0], &n_platform_num)) != CL_SUCCESS)
		return n_result;
	_ASSERTE(r_platform_list.size() == n_platform_num);
	// get list of platform id's

	return cl_Success;
}

CLresult CCLUtils::n_Get_FullProfile_Platform(cl_platform_id *p_platform,
	cl_device_type n_device_type /*= CL_DEVICE_TYPE_GPU*/)
{
	std::vector<cl_platform_id> platform_list;
	CLresult n_result;
	if((n_result = n_Get_PlatformList(platform_list)) != CL_SUCCESS)
		return n_result;

	size_t n_full_platform = 0; // if no full platform is found, just choose the first one
	for(size_t i = 0, n = platform_list.size(); i < n; ++ i) {
		char p_s_profile_name[20] = {0};
		clGetPlatformInfo(platform_list[i], CL_PLATFORM_PROFILE, sizeof(p_s_profile_name), p_s_profile_name, NULL);
		p_s_profile_name[sizeof(p_s_profile_name) / sizeof(p_s_profile_name[0]) - 1] = 0; // make sure it's terminated

		if(!strcmp(p_s_profile_name, "FULL_PROFILE")) {
			{
				cl_uint n_device_num;
				CLresult n_result = (CLresult)clGetDeviceIDs(platform_list[i], n_device_type, 0, 0, &n_device_num);
				if(n_result != CL_SUCCESS || !n_device_num)
					continue;
				// make sure that the platform contains any devices that we are interested in
			}
			n_full_platform = i;
			break;
		}
	}
	// find "FULL_PROFILE" platform

	*p_platform = platform_list[n_full_platform];

	return cl_Success;
}

CLresult CCLUtils::n_OpenCL_Init(cl_context *p_context, cl_platform_id h_platform,
	size_t n_device_num, const cl_device_id *p_device)
{
	cl_context_properties p_props[3] = {cl_context_properties(CL_CONTEXT_PLATFORM),
		cl_context_properties(h_platform), 0};
	cl_int n_result;
	*p_context = clCreateContext((h_platform)? p_props : 0, n_device_num, p_device, 0, 0, &n_result);
	return (CLresult)n_result;
}

CLresult CCLUtils::n_OpenCL_Init(cl_context *p_context, int n_device_type, bool b_implementation_profile_selection)
{
	cl_context_properties p_props[3] = {cl_context_properties(CL_CONTEXT_PLATFORM), 0, 0};
	if(!b_implementation_profile_selection) {
		cl_platform_id h_platform;
		CLresult n_result;
		if((n_result = (CLresult)n_Get_FullProfile_Platform(&h_platform, n_device_type)) != CL_SUCCESS)
			return n_result;
		p_props[1] = cl_context_properties(h_platform);
	}
	// get the first available "FULL_PROFILE" OpenCL platform

	cl_int n_result;
	cl_context h_context = clCreateContextFromType((b_implementation_profile_selection)?
		NULL : p_props, n_device_type, 0, 0, &n_result);
	if(n_result != CL_SUCCESS)
		return (CLresult)n_result;
	// create context

	*p_context = h_context;

	return cl_Success;
}

CLresult CCLUtils::n_GetDeviceList(cl_context h_context, std::vector<cl_device_id> &r_device_list)
{
	{
		size_t n_device_size;
		CLresult n_result;
		if((n_result = (CLresult)clGetContextInfo(h_context, CL_CONTEXT_DEVICES, 0, 0, &n_device_size)) != CL_SUCCESS)
			return n_result;
		size_t n_device_num = n_device_size / sizeof(cl_device_id); // it's in bytes
		if(!n_device_num)
			return cl_Device_Not_Found; // no devices
		try {
			r_device_list.resize(n_device_num);
		} catch(std::bad_alloc&) {
			return cl_Out_Of_Host_Memory;
		}
		if((n_result = (CLresult)clGetContextInfo(h_context, CL_CONTEXT_DEVICES, n_device_size, &r_device_list[0], 0)) != CL_SUCCESS)
			return n_result;
	}
	// get list of devices

	return cl_Success;
}

CLresult CCLUtils::n_GetDeviceList(cl_platform_id h_platform, std::vector<cl_device_id> &r_device_list,
	cl_device_type n_device_type /*= CL_DEVICE_TYPE_GPU*/)
{
	{
		cl_uint n_device_num;
		CLresult n_result;
		if((n_result = (CLresult)clGetDeviceIDs(h_platform, n_device_type, 0, 0, &n_device_num)) != CL_SUCCESS)
			return n_result;
		if(!n_device_num)
			return cl_Device_Not_Found; // no devices
		try {
			r_device_list.resize(n_device_num);
		} catch(std::bad_alloc&) {
			return cl_Out_Of_Host_Memory;
		}
		if((n_result = (CLresult)clGetDeviceIDs(h_platform, n_device_type, n_device_num, &r_device_list[0], 0)) != CL_SUCCESS)
			return n_result;
	}
	// get list of devices

	return cl_Success;
}

size_t CCLUtils::n_Get_MaxGFlops_DeviceId(cl_device_id *p_device_id, cl_context h_context)
{
	return CCLUtils::n_Get_Best_DeviceId(p_device_id, h_context, CCLUtils::TDevice_DefaultScoring());
}

size_t CCLUtils::n_Get_MaxGFlops_DeviceId(cl_context h_context)
{
	cl_device_id h_dummy;
	return n_Get_MaxGFlops_DeviceId(&h_dummy, h_context);
}

/*size_t CCLUtils::n_Get_MaxGFlops_DeviceId(cl_device_id *p_device_id, cl_context h_context)
{
	std::vector<cl_device_id> device_list;
	if(n_GetDeviceList(h_context, device_list) != CL_SUCCESS)
		return -1;
	// get all the devices

	size_t n_best_device = 0;
	{
		double f_best_gflops = -1;
		for(size_t i = 0, n = device_list.size(); i < n; ++ i) {
			cl_uint n_speed, n_processors;
			if(clGetDeviceInfo(device_list[i], CL_DEVICE_MAX_CLOCK_FREQUENCY,
			   sizeof(n_speed), &n_speed, NULL) != CL_SUCCESS ||
			   clGetDeviceInfo(device_list[i], CL_DEVICE_MAX_COMPUTE_UNITS,
			   sizeof(n_processors), &n_processors, NULL) != CL_SUCCESS)
				return -1;
			double f_gflops = double(n_speed) * n_processors;
			if(!i || f_best_gflops < f_gflops) {
				f_best_gflops = f_gflops;
				n_best_device = i;
			}
		}
	}
	// go trough devices, and estimate computing power (not real gflops)

	*p_device_id = device_list[n_best_device];
	// write handle to the best device

	return n_best_device;
}*/

CLresult CCLUtils::n_GetPlatformInfoString(std::string &r_s_str, cl_platform_id h_platform, int n_name)
{
	size_t n_name_size;
	CLresult n_result;
	if((n_result = (CLresult)clGetPlatformInfo(h_platform, n_name, 0, NULL, &n_name_size)) != CL_SUCCESS)
		return n_result;
	// get string length

	try {
		r_s_str.resize(n_name_size + 1);
		// allocate string

		if((n_result = (CLresult)clGetPlatformInfo(h_platform, n_name, n_name_size, &r_s_str[0], NULL)) != CL_SUCCESS)
			return n_result;
		r_s_str[n_name_size] = 0;
		r_s_str.resize(strlen(r_s_str.c_str()));
		// get string
	} catch(std::bad_alloc&) {
		return cl_Out_Of_Host_Memory;
	}

	return cl_Success;
}

/*
 *								=== ~CCLUtils ===
 */

/*
 *								=== CCLDeviceParams ===
 */

CLresult CCLDeviceParams::n_GetDeviceInfoString(std::string &r_s_str, int n_name)
{
	return n_GetDeviceInfoString(r_s_str, m_h_device, n_name);
}

CLresult CCLDeviceParams::n_GetDeviceInfoString(std::string &r_s_str, cl_device_id h_device, int n_name)
{
	size_t n_name_size;
	CLresult n_result;
	if((n_result = (CLresult)clGetDeviceInfo(h_device, n_name, 0, NULL, &n_name_size)) != CL_SUCCESS)
		return n_result;
	// get string length

	try {
		r_s_str.resize(n_name_size + 1);
		// allocate string

		if((n_result = (CLresult)clGetDeviceInfo(h_device, n_name, n_name_size, &r_s_str[0], NULL)) != CL_SUCCESS)
			return n_result;
		r_s_str[n_name_size] = 0;
		r_s_str.resize(strlen(r_s_str.c_str()));
		// get string
	} catch(std::bad_alloc&) {
		return cl_Out_Of_Host_Memory;
	}

	return cl_Success;
}

bool CCLDeviceParams::b_IsExtensionSupported(const char *p_s_extension_name) const
{
	return b_IsExtensionSupported(m_h_device, p_s_extension_name);
}

bool CCLDeviceParams::b_IsExtensionSupported(cl_device_id h_device, const char *p_s_extension_name)
{
	std::string s_exts;
	if(n_GetDeviceInfoString(s_exts, h_device, CL_DEVICE_EXTENSIONS) != CL_SUCCESS)
		return false;
	// get device extension string

	size_t n_ext_name_length = strlen(p_s_extension_name);
	size_t n_pos = 0, n_end_pos = s_exts.length() - n_ext_name_length;
	while((n_pos = s_exts.find(p_s_extension_name, n_pos)) != std::string::npos) {
		if((!n_pos || isspace(s_exts[n_pos - 1])) &&
		   (n_pos >= n_end_pos || isspace(s_exts[n_pos + n_ext_name_length])))
			return true;
		// there must be space before and after it
	}
	// attempt to find the extension

	return false;
	// extension not found
}

void CCLDeviceParams::Dump(FILE *p_fw)
{
	if(b_Is_NVIDIA()) {
		fprintf(p_fw, "device(\'%s\', " PRIsizeB
			"B RAM, %.2f MHz, SM: " PRIsize ", max-threads-block: %d, max-block-size: %dx%dx%d, cap %d.%d)\n",
			p_s_Name(),
			PRIsizeBparams(n_Memory_Size()),
			t_Properties().clockRate / 1e3f,
			n_Multiprocessor_Num(),
			t_Properties().maxThreadsPerBlock,
			t_Properties().maxThreadsDim[0],
			t_Properties().maxThreadsDim[1],
			t_Properties().maxThreadsDim[2],
			m_p_device_caps[0], m_p_device_caps[1]);
	} else {
		fprintf(p_fw, "device(\'%s\', " PRIsizeB
			"B RAM, %.2f MHz, multiproc: " PRIsize ", max-threads-block: %d, max-block-size: %dx%dx%d)\n",
			p_s_Name(),
			PRIsizeBparams(n_Memory_Size()),
			t_Properties().clockRate / 1e3f,
			n_Multiprocessor_Num(),
			t_Properties().maxThreadsPerBlock,
			t_Properties().maxThreadsDim[0],
			t_Properties().maxThreadsDim[1],
			t_Properties().maxThreadsDim[2]);
	}
	// show some device parameters
}

CCLDeviceParams::CCLDeviceParams(cl_device_id h_device)
	:m_h_device(h_device), m_n_memory_size(0)
{
	memset(m_p_device_caps, 0, 2 * sizeof(int));
	memset(&m_t_devprop, 0, sizeof(CLdevprop));

	if(!QueryDeviceParams()) {
		m_h_device = 0; // mark error
		return;
	}
	// get device params
}

CCLDeviceParams::CCLDeviceParams(cl_context h_context, size_t n_device_index)
	:m_h_device(0), // mark error
	m_n_memory_size(0)
{
	memset(m_p_device_caps, 0, 2 * sizeof(int));
	memset(&m_t_devprop, 0, sizeof(CLdevprop));

	{
		std::vector<cl_device_id> device_list;
		if(CCLUtils::n_GetDeviceList(h_context, device_list) != CL_SUCCESS)
			return;
		// get all the devices

		if(n_device_index < 0 || n_device_index >= device_list.size())
			return;
		// check index

		m_h_device = device_list[n_device_index];
	}
	// query all devices available to the context, get device handle

	if(!QueryDeviceParams()) {
		m_h_device = 0; // mark error
		return;
	}
	// get device params
}

bool CCLDeviceParams::QueryDeviceParams()
{
	if(n_GetDeviceInfoString(m_s_name, CL_DEVICE_NAME) != CL_SUCCESS)
		return false;
	// get device name

	bool b_nv_device_attribute_query = b_IsExtensionSupported("cl_nv_device_attribute_query");
	// determines wheter we have cl_nv_device_attribute_query

	{
		/*std::string s_version;
		if(n_GetDeviceInfoString(s_version, CL_DEVICE_VERSION) != CL_SUCCESS)
			return;
		m_p_device_caps[0] = atol(s_version.c_str());
		m_p_device_caps[1] = (strchr(s_version.c_str(), '.'))? atol(strchr(s_version.c_str(), '.') + 1) : 0;*/ // todo

		if(b_nv_device_attribute_query) {
			cl_uint n_major, n_minor;
			if(clGetDeviceInfo(m_h_device, CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV, sizeof(cl_uint), &n_major, NULL) != CL_SUCCESS ||
			   clGetDeviceInfo(m_h_device, CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV, sizeof(cl_uint), &n_minor, NULL) != CL_SUCCESS)
				return false;
			m_p_device_caps[0] = n_major;
			m_p_device_caps[1] = n_minor;
		} else {
			m_p_device_caps[0] = 1;
			m_p_device_caps[1] = 0;
		}
	}
	// get device OpenCL revision

	{
		if(b_nv_device_attribute_query) {
			cl_uint n_value;
			if(clGetDeviceInfo(m_h_device, CL_DEVICE_KERNEL_EXEC_TIMEOUT_NV, sizeof(n_value), &n_value, NULL) != CL_SUCCESS)
				return false;
			m_b_kernel_exec_timeout = (n_value != 0);
		} else
			m_b_kernel_exec_timeout = false;
	}
	{
		cl_uint n_value;
		if(clGetDeviceInfo(m_h_device, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(n_value), &n_value, NULL) != CL_SUCCESS)
			return false;
		m_n_multiprocessor_num = n_value;
	}
	// get some interesting device attributes

	{
		cl_ulong n_value; // this is actually a 64bit integer
		if(clGetDeviceInfo(m_h_device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(n_value), &n_value, NULL) != CL_SUCCESS)
			return false;
		_ASSERTE(n_value <= UINT64_MAX); // this should be enough for a while
		m_n_memory_size = uint64_t(n_value);
	}
	// get device memory size

	return n_Query_DeviceProperties(m_t_devprop, m_h_device) == CL_SUCCESS;
}

CLresult CCLDeviceParams::n_Query_DeviceProperties(CLdevprop &r_t_devprop, cl_device_id h_device)
{
	memset(&r_t_devprop, 0, sizeof(CLdevprop));
	// clear output

	bool b_nv_device_attribute_query = b_IsExtensionSupported(h_device, "cl_nv_device_attribute_query");
	// determines wheter we have cl_nv_device_attribute_query

	{
		cl_uint n_value;
		CLresult n_result;
		if((n_result = (CLresult)clGetDeviceInfo(h_device, CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof(n_value), &n_value, NULL)) != CL_SUCCESS)
			return n_result;
		r_t_devprop.clockRate = n_value * 1000;
	}
	{
		cl_uint n_value;
		CLresult n_result;
		if((n_result = (CLresult)clGetDeviceInfo(h_device, CL_DEVICE_ADDRESS_BITS, sizeof(n_value), &n_value, NULL)) != CL_SUCCESS)
			return n_result;
		n_value /= 2;
		int n_size = int(min(uint64_t(INT_MAX), (uint64_t(1) << (n_value - 1)) | ((uint64_t(1) << (n_value - 1)) - 1)));
		r_t_devprop.maxGridSize[0] = n_size;
		r_t_devprop.maxGridSize[1] = n_size;

		if(b_nv_device_attribute_query) {
			cl_uint n_major, n_minor;
			CLresult n_result;
			if((n_result = (CLresult)clGetDeviceInfo(h_device, CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV, sizeof(cl_uint), &n_major, NULL)) != CL_SUCCESS ||
			   (n_result = (CLresult)clGetDeviceInfo(h_device, CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV, sizeof(cl_uint), &n_minor, NULL)) != CL_SUCCESS)
				return n_result;
			if(n_major < 2) {
				_ASSERTE(n_size == 65535); // CUDA_C_Programming_Guide.pdf, table 12 (page 185) ...
				r_t_devprop.maxGridSize[2] = n_size;
			} else
				r_t_devprop.maxGridSize[2] = 65535; // CUDA_C_Programming_Guide.pdf, table 12 (page 185). note that this *will* change with future models
		} else {
			r_t_devprop.maxGridSize[2] = n_size;
			// assume it is the same accross all the dimensions (according to OpenCL spec, it should be;
			// at http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/clEnqueueNDRangeKernel.html
			// it says: global_work_size values must be in the range 1 .. 2^32 - 1)
		}
		// handle the z-dimension
	}
	{
		size_t p_value[3]; // 3 is safe value, no device shall have CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS less than 3
		CLresult n_result;
		if((n_result = (CLresult)clGetDeviceInfo(h_device, CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(p_value), p_value, NULL)) != CL_SUCCESS)
			return n_result;
		_ASSERTE(p_value[0] <= INT_MAX && p_value[1] <= INT_MAX && p_value[2] <= INT_MAX);
		r_t_devprop.maxThreadsDim[0] = int(p_value[0]);
		r_t_devprop.maxThreadsDim[1] = int(p_value[1]);
		r_t_devprop.maxThreadsDim[2] = int(p_value[2]); // block size
	}
	{
		size_t n_value;
		CLresult n_result;
		if((n_result = (CLresult)clGetDeviceInfo(h_device, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(n_value), &n_value, NULL)) != CL_SUCCESS)
			return n_result;
		_ASSERTE(n_value <= INT_MAX);
		r_t_devprop.maxThreadsPerBlock = int(n_value);
	}
	{
		cl_ulong n_value;
		CLresult n_result;
		if((n_result = (CLresult)clGetDeviceInfo(h_device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(n_value), &n_value, NULL)) != CL_SUCCESS)
			return n_result;
		_ASSERTE(n_value <= INT_MAX);
		r_t_devprop.memPitch = int(n_value);
	}
	{
		if(b_nv_device_attribute_query) {
			cl_uint n_value;
			CLresult n_result;
			if((n_result = (CLresult)clGetDeviceInfo(h_device, CL_DEVICE_REGISTERS_PER_BLOCK_NV, sizeof(n_value), &n_value, NULL)) != CL_SUCCESS)
				return n_result;
			r_t_devprop.regsPerBlock = n_value;
		} else
			r_t_devprop.regsPerBlock = -1; // unknown
	}
	{
		cl_ulong n_value;
		CLresult n_result;
		if((n_result = (CLresult)clGetDeviceInfo(h_device, CL_DEVICE_LOCAL_MEM_SIZE, sizeof(n_value), &n_value, NULL)) != CL_SUCCESS)
			return n_result;
		_ASSERTE(n_value <= INT_MAX);
		r_t_devprop.sharedMemPerBlock = int(n_value);
	}
	{
		if(b_nv_device_attribute_query) {
			cl_uint n_value;
			CLresult n_result;
			if((n_result = (CLresult)clGetDeviceInfo(h_device, CL_DEVICE_WARP_SIZE_NV, sizeof(n_value), &n_value, NULL)) != CL_SUCCESS)
				return n_result;
			r_t_devprop.SIMDWidth = n_value;
		} else {
			r_t_devprop.SIMDWidth = -1;
			// unknown

			do {
				cl_int n_result;
				cl_context h_context = clCreateContext(0, 1, &h_device, 0, 0, &n_result);
				if(n_result != CL_SUCCESS)
					break;

				cl_command_queue h_cmd_queue = clCreateCommandQueue(h_context,
					h_device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &n_result);
				if(n_result != CL_SUCCESS) {
					clReleaseContext(h_context);
					break;
				}
				// create command queue

				cl_program h_program;
				if((n_result = CCLProgramCompiler::n_CompileProgram(h_context, &h_program,
				   "__kernel void k_dummy(__global int *x) { x[get_global_id(0)] = 0; }",
				   1, &h_device)) != CL_SUCCESS) {
					clReleaseCommandQueue(h_cmd_queue);
					clReleaseContext(h_context);
					break;
				}
				// compile the program

				cl_kernel h_dummy_kernel = clCreateKernel(h_program, "k_dummy", &n_result);
				if(n_result != CL_SUCCESS) {
					clReleaseCommandQueue(h_cmd_queue);
					clReleaseContext(h_context);
					clReleaseProgram(h_program);
					break;
				}
				// get the kernel

				size_t n_wavefront_size = 0, n_size;
				clGetKernelWorkGroupInfo(h_dummy_kernel, h_device,
					CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE,
					sizeof(size_t), &n_wavefront_size, &n_size);
				// get work-group size multiple

				if(n_result == CL_SUCCESS)
					r_t_devprop.SIMDWidth = int(n_wavefront_size);

				clReleaseKernel(h_dummy_kernel);
				clReleaseProgram(h_program);
				clReleaseCommandQueue(h_cmd_queue);
				clReleaseContext(h_context);
				// cleanup
			} while(false);
			// get SIMD width on non-NVIDIA platforms
		}
	}
	{
		r_t_devprop.textureAlign = -1; // unknown; todo
	}
	{
		cl_ulong n_value;
		CLresult n_result;
		if((n_result = (CLresult)clGetDeviceInfo(h_device, CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE, sizeof(n_value), &n_value, NULL)) != CL_SUCCESS)
			return n_result;
		_ASSERTE(n_value <= INT_MAX);
		r_t_devprop.totalConstantMemory = int(n_value);
	}
	// get device properties

	return cl_Success;
}

/*
 *								=== ~CCLDeviceParams ===
 */

/*
 *								=== CCLProgramStorage::TProgramBinary::TBinary ===
 */

inline CCLProgramStorage::TProgramBinary::TBinary::TBinary()
	:n_size(0), p_data(0)
{}

/*
 *								=== ~CCLProgramStorage::TProgramBinary::TBinary ===
 */

/*
 *								=== CCLProgramStorage::TProgramBinary ===
 */

inline void CCLProgramStorage::TProgramBinary::Swap(TProgramBinary &r_t_other)
{
	binary_size_list.swap(r_t_other.binary_size_list);
	std::swap(t_data_buffer, r_t_other.t_data_buffer);
}

/*
 *								=== ~CCLProgramStorage::TProgramBinary ===
 */

/*
 *								=== CCLProgramStorage::_TSHA1 ===
 */

inline bool CCLProgramStorage::_TSHA1::operator ==(const _TSHA1 &r_t_other) const
{
	return !memcmp(p_data, r_t_other.p_data, 5 * sizeof(uint32_t));
}

/**
 *	@brief inequality comparator
 *	@param[in] r_t_other is hash to compare against
 *	@return Returns false if this hash equals r_t_other, otherwise returns true.
 */
inline bool CCLProgramStorage::_TSHA1::operator !=(const _TSHA1 &r_t_other) const
{
	return !(*this == r_t_other);
}

/*
 *								=== ~CCLProgramStorage::_TSHA1 ===
 */

/*
 *								=== CCLProgramStorage::TProgramBinary ===
 */

bool CCLProgramStorage::TProgramBinary::Download(cl_program h_program, size_t n_device_num)
{
	_ASSERTE(!t_data_buffer.p_data && binary_size_list.empty()); // make sure this is empty

	cl_uint _n_device_num;
	if(clGetProgramInfo(h_program, CL_PROGRAM_NUM_DEVICES,
	   sizeof(size_t), &_n_device_num, NULL) != CL_SUCCESS || _n_device_num != n_device_num)
		return false;
	// get number of devices, make sure it matches description

	if(!n_device_num || n_device_num > UINT32_MAX)
		return false;
	// we can't query binaries without devices

	std::vector<const unsigned char*> binary_ptr_list;
	try {
		binary_ptr_list.resize(n_device_num);
#ifdef _DEBUG
		_ASSERTE(n_device_num < SIZE_MAX);
		binary_size_list.resize(n_device_num + 1);
		if(n_device_num)
			binary_size_list[n_device_num - 1] = 0xbaadf00d; // paranoid check
		binary_size_list[n_device_num] = 0xbaadf00d; // "magic" word
#else // _DEBUG
		binary_size_list.resize(n_device_num);
#endif // _DEBUG
	} catch(std::bad_alloc&) {
		return false;
	}
	if(clGetProgramInfo(h_program, CL_PROGRAM_BINARY_SIZES,
	   n_device_num * sizeof(size_t), &binary_size_list[0], NULL) != CL_SUCCESS)
		return false;
#ifdef _DEBUG
	_ASSERTE(binary_size_list[n_device_num] == 0xbaadf00d); // make sure it's returning the right amount of data
	_ASSERTE(!n_device_num || binary_size_list[n_device_num - 1] != 0xbaadf00d); // paranoid check
	binary_size_list.erase(binary_size_list.end() - 1); // remove it from the list, length must be equal to n_device_num
#endif // _DEBUG
	// get binary sizes

	uint64_t n_size_total = 0;
	for(size_t i = 0; i < n_device_num; ++ i) {
		if(n_size_total > UINT64_MAX - binary_size_list[i])
			return false;
		n_size_total += binary_size_list[i];
	}
	// sum binary sizes up

	t_data_buffer.n_size = n_size_total;
	if(n_size_total > SIZE_MAX || !(t_data_buffer.p_data = new(std::nothrow) uint8_t[size_t(n_size_total)]))
		return false;
	// alloc data buffer

	{
		const uint8_t *p_data_ptr = t_data_buffer.p_data;
		for(size_t i = 0; i < n_device_num; ++ i) {
			binary_ptr_list[i] = (const unsigned char*)p_data_ptr;
			p_data_ptr += binary_size_list[i];
		}
		_ASSERTE(p_data_ptr == t_data_buffer.p_data + t_data_buffer.n_size);
	}
	// get pointers to individual binaries

	if(clGetProgramInfo(h_program, CL_PROGRAM_BINARIES,
	   size_t(n_size_total), &binary_ptr_list[0], NULL) != CL_SUCCESS)
		return false;
	// download binaries

	return true;
}

int CCLProgramStorage::TProgramBinary::n_Upload(cl_context h_context, cl_program *p_program,
	CLresult &r_n_result, size_t n_device_num, const cl_device_id *p_device_list) const
{
	_ASSERTE(p_program); // must not be null
	_ASSERTE(!binary_size_list.empty()); // make sure this is not empty

	if(binary_size_list.size() != n_device_num)
		return build_InvalidParams;
	// device list must match

	std::vector<const unsigned char*> binary_ptr_list;
	try {
		binary_ptr_list.resize(n_device_num);
	} catch(std::bad_alloc&) {
		return build_LowMemory;
	}
	{
		const uint8_t *p_data_ptr = t_data_buffer.p_data;
		for(size_t i = 0; i < n_device_num; ++ i) {
			binary_ptr_list[i] = (const unsigned char*)p_data_ptr;
			p_data_ptr += binary_size_list[i];
		}
		_ASSERTE(p_data_ptr == t_data_buffer.p_data + t_data_buffer.n_size);
	}
	// get pointers to individual binaries

	cl_int n_result;
	_ASSERTE(n_device_num <= /*CL*/UINT_MAX);
	cl_program h_prog = clCreateProgramWithBinary(h_context, cl_uint(n_device_num), p_device_list,
		&binary_size_list[0], &binary_ptr_list[0], NULL, &n_result);
	if(n_result != CL_SUCCESS) {
		r_n_result = (CLresult)n_result;
		return build_CreateProgramFailure;
	}
	if((n_result = clBuildProgram(h_prog, 0, NULL, NULL, NULL, NULL)) != CL_SUCCESS) { // note no build options here
		r_n_result = (CLresult)n_result;
		clReleaseProgram(h_prog);
		return build_BuildProgramFailure;
	}
	// use loaded binaries to create the program ...

#ifdef CL_PROGRAM_COMPILER_DISPLAY_SUCCESSFUL_BUILD_WARNINGS
	/*if(n_device_num) { // compiling from binary doesn't show output :'(
		std::string s_build_log, s_device_name;
		for(size_t i = 0; i < n_device_num; ++ i) {
			cl_build_status n_build_status;

			if(CCLProgramCompiler::Get_BuildLog(s_build_log, n_build_status, h_prog, p_device_list[i]) != CL_SUCCESS)
				return n_result;
			// get buld log and build status

			if(CCLDeviceParams::n_GetDeviceInfoString(s_device_name, p_device_list[i], CL_DEVICE_NAME) != CL_SUCCESS)
				s_device_name.erase();
			const char *p_s_device_name = (s_device_name.empty())? "(null)" : s_device_name.c_str();
			// get device name

			if(n_build_status == CL_BUILD_NONE) {
				fprintf(stderr, "warning: program wasn't built for device %d (%s)\n", i, p_s_device_name);
				continue;
			} else if(n_build_status == CL_BUILD_ERROR) {
				fprintf(stderr, "warning: there were errors while building program for device %d (%s)\n",
					i, p_s_device_name);
			} else if(n_build_status != CL_BUILD_SUCCESS)
				fprintf(stderr, "warning: unknown program build status for device %d (%s)\n", i, p_s_device_name);
			// show build status

			if(n_build_status != CL_BUILD_SUCCESS || !s_build_log.empty()) {
				fprintf(stderr, "=== OpenCL build log for device %d (%s) ===\n%s\n",
					i, p_s_device_name, s_build_log.c_str());
			}
			// show build log
		}
	}*/
	// show warnings
#endif // CL_PROGRAM_COMPILER_DISPLAY_SUCCESSFUL_BUILD_WARNINGS

	r_n_result = cl_Success; // !!
	*p_program = h_prog;
	// write output

	return build_Success;
	// success
}

void CCLProgramStorage::TProgramBinary::Free()
{
	if(t_data_buffer.p_data) {
		delete[] t_data_buffer.p_data;
		t_data_buffer.p_data = 0;
	}
	binary_size_list.clear();
}

/*
 *								=== CCLProgramStorage::~TProgramBinary ===
 */

/*
 *								=== CCLProgramStorage::TProgramInstance ===
 */

bool CCLProgramStorage::TProgramInstance::Download(cl_program h_program,
	size_t n_device_num, _TSHA1 t_build_options_hash)
{
	t_build_opts_hash = t_build_options_hash;
	// copy hash of build options

	return TProgramBinary::Download(h_program, n_device_num);
}

bool CCLProgramStorage::TProgramInstance::Read(FILE *p_fr, size_t n_device_num)
{
	_ASSERTE(!t_data_buffer.p_data && binary_size_list.empty()); // make sure this is empty

	if(fread(&t_build_opts_hash, sizeof(_TSHA1), 1, p_fr) != 1)
		return false;
	// read build options hash

	std::vector<uint64_t> binary_size_list64;
	try {
		binary_size_list64.resize(n_device_num);
		binary_size_list.resize(n_device_num);
	} catch(std::bad_alloc) {
		return false;
	}
	// alloc arrays

	if(fread(&binary_size_list64[0], sizeof(uint64_t), n_device_num, p_fr) != n_device_num)
		return false;
	uint64_t n_size_total = 0;
	for(size_t i = 0; i < n_device_num; ++ i) {
		if(binary_size_list64[i] > SIZE_MAX || n_size_total > UINT64_MAX - binary_size_list64[i])
			return false;
		binary_size_list[i] = size_t(binary_size_list64[i]);
		n_size_total += binary_size_list64[i];
	}
	// read binary lengths, sum up

	t_data_buffer.n_size = n_size_total;
	if(n_size_total > SIZE_MAX || !(t_data_buffer.p_data = new(std::nothrow) uint8_t[size_t(n_size_total)]))
		return false;
	{
		uint64_t n_size_remaining = n_size_total;
		uint8_t *p_data_ptr = t_data_buffer.p_data;
		while(n_size_remaining > 0) {
			int n_read = int(min(uint64_t(INT_MAX), n_size_remaining));
			if(fread(p_data_ptr, 1, n_read, p_fr) != n_read)
				return false;
			n_size_remaining -= n_read;
			p_data_ptr += n_read;
		}
		_ASSERTE(p_data_ptr == t_data_buffer.p_data + n_size_total);
	}
	// alloc and read data buffer

	return true;
}

bool CCLProgramStorage::TProgramInstance::Write(FILE *p_fw, size_t n_device_num)
{
	if(binary_size_list.size() != n_device_num)
		return false;
	// make sure there is correct number of binaries

	if(fwrite(&t_build_opts_hash, sizeof(_TSHA1), 1, p_fw) != 1)
		return false;
	// write build options hash

	for(size_t i = 0; i < n_device_num; ++ i) {
		_ASSERTE(SIZE_MAX <= UINT64_MAX);
		uint64_t n_size = binary_size_list[i];
		if(fwrite(&n_size, sizeof(uint64_t), 1, p_fw) != 1)
			return false;
	}
	// write binary lengths

	{
		uint64_t n_size_remaining = t_data_buffer.n_size;
		uint8_t *p_data_ptr = t_data_buffer.p_data;
		while(n_size_remaining) {
			int n_write = int(min(uint64_t(INT_MAX), n_size_remaining));
			if(fwrite(p_data_ptr, 1, n_write, p_fw) != n_write)
				return false;
			n_size_remaining -= n_write;
			p_data_ptr += n_write;
		}
		_ASSERTE(p_data_ptr == t_data_buffer.p_data + t_data_buffer.n_size);
	}
	// write binaries data buffer

	return true;
}

inline void CCLProgramStorage::TProgramInstance::Swap(TProgramInstance &r_t_other)
{
	std::swap(t_build_opts_hash, r_t_other.t_build_opts_hash);
	TProgramBinary::Swap(r_t_other);
}

/*
 *								=== ~CCLProgramStorage::TProgramInstance ===
 */

/*
 *								=== CCLProgramStorage ===
 */

CCLProgramStorage::CCLProgramStorage()
	:m_b_dirty(false)
{
	memset(&m_t_header, 0, sizeof(TFileHeader));

	m_t_header.p_magic[0] = 'L';
	m_t_header.p_magic[1] = 'C';
	m_t_header.p_magic[2] = 'L';
	m_t_header.p_magic[3] = 's';
	// set magic word ("Lame CL storage") - so b_Status() doesn't fail
}

CCLProgramStorage::CCLProgramStorage(const char *p_s_filename)
	:m_b_dirty(false)
{
	if(!Load(p_s_filename))
		memset(&m_t_header, 0, sizeof(TFileHeader)); // to mark error
}

CCLProgramStorage::~CCLProgramStorage()
{
	std::for_each(instance_list.begin(), instance_list.end(), DeleteInstance);
}

bool CCLProgramStorage::b_Status() const
{
	return m_t_header.p_magic[0] == 'L' && m_t_header.p_magic[1] == 'C' &&
		m_t_header.p_magic[2] == 'L' && m_t_header.p_magic[3] != 's';
}

bool CCLProgramStorage::b_Dirty() const
{
	/*_TSHA1 t_ctrl_hash;
	if(!Calc_ControlHash(t_ctrl_hash))
		fprintf(stderr, "hash check failed ...\n"); // debug
	if(m_b_dirty != (t_ctrl_hash != m_t_header.t_control_hash)) {
		fprintf(stderr, "dirty discrepancy (dirty %d, hash-check %d) ...\n",
			m_b_dirty, t_ctrl_hash != m_t_header.t_control_hash); // debug
	}*/ // debug

	return m_b_dirty;
}

bool CCLProgramStorage::b_Check_DeviceIds(size_t n_device_num, const cl_device_id *p_device_list) const
{
	if(device_id_list.size() != n_device_num)
		return false;
	// counts must match

	std::string s_tmp0, s_tmp1;
	for(size_t i = 0, n = device_id_list.size(); i < n; ++ i) {
		if(CCLDeviceParams::n_GetDeviceInfoString(s_tmp0,
		   p_device_list[i], CL_DEVICE_NAME) != CL_SUCCESS ||
		   CCLDeviceParams::n_GetDeviceInfoString(s_tmp1,
		   p_device_list[i], CL_DRIVER_VERSION) != CL_SUCCESS ||
		   !stl_ut::AppendCStr(s_tmp0, "|") ||
		   !stl_ut::Append(s_tmp0, s_tmp1))
			return false; // this is, in fact, error, not npos
		if(s_tmp0 != device_id_list[i])
			return false;
	}
	// check device identification

	return true;
}

size_t CCLProgramStorage::n_Find_ProgramBinaries(const char *p_s_source_code, size_t n_device_num,
	const cl_device_id *p_device_list, const char *p_s_build_options) const
{
	if(p_s_source_code && strstr(p_s_source_code, "#include"))
		return npos;
	// in case there are #includes, need to compile every time (the included files may change)
	// todo - hash all the includes as well

	if(device_id_list.size() != n_device_num ||
	   m_t_header.t_source_hash != t_Hash_String(p_s_source_code) ||
	   !b_Check_DeviceIds(n_device_num, p_device_list))
		return npos;
	// make sure there are the same devices (that's kind of tricky though, this is limited
	// to number of devices as there's no device id that could be written to the file (except
	// maybe hash of name and drivers)) and the same source code // @t_odo - implement better device identification

	_TSHA1 t_build_opts_hash = t_Hash_String(p_s_build_options);
	// hash build options

	for(size_t i = 0, n = instance_list.size(); i < n; ++ i) {
		if(instance_list[i].t_build_opts_hash == t_build_opts_hash) {
			if(i) {
				/*fprintf(stderr, "promoting binaries %d\n", i);*/ // debug
				instance_list[i].Swap(instance_list[i - 1]);
				-- i;
				// promote this instance, effectively invalidating any indices obtained by previous calls to this functions

				m_b_dirty = true;
				// we just changed order of binaries
			}
			/*fprintf(stderr, "requested binaries found: %d (\'%s\')\n", i, p_s_build_options);*/ // debug
			return i;
		}
	}
	// try to find the program binary

	/*fprintf(stderr, "requested binaries not found (\'%s\')\n", p_s_build_options);*/ // debug

	return npos;
	// not found
}

int CCLProgramStorage::n_Get_ProgramBinaries(cl_context h_context, cl_program *p_program, CLresult &r_n_result,
	size_t n_binaries_index, size_t n_device_num, const cl_device_id *p_device_list) const
{
	if(n_binaries_index == npos || n_binaries_index >= instance_list.size() || device_id_list.size() != n_device_num)
		return build_InvalidParams;
	// check binaries index and devices

	_ASSERTE(b_Check_DeviceIds(n_device_num, p_device_list));
	// device identification isn't checked in release, it's checked in n_Find_ProgramBinaries(),
	// this function is supposed to get the same parameters

	const TProgramInstance &r_t_inst = instance_list[n_binaries_index];
	// get program instance

	return r_t_inst.n_Upload(h_context, p_program, r_n_result, n_device_num, p_device_list);
}

bool CCLProgramStorage::Put_ProgramBinaries(cl_program h_program, const char *p_s_source_code,
	size_t n_device_num, const cl_device_id *p_device_list, const char *p_s_build_options)
{
	// puts compiled program binaries to the storage
	// number of instances is limited (LRU-type policy), but not here
	// note this might erase all other instances in case source code changed

	for(size_t i = 0; i < n_device_num; ++ i) {
		cl_build_status n_status;
		if(clGetProgramBuildInfo(h_program, p_device_list[i], CL_PROGRAM_BUILD_STATUS,
		   sizeof(cl_build_status), &n_status, NULL) != CL_SUCCESS || n_status != CL_BUILD_SUCCESS)
			return false;
	}
	// the program must be completely and successfuly built

	if(n_Find_ProgramBinaries(p_s_source_code, n_device_num, p_device_list, p_s_build_options) != npos)
		return true;
	// it's already there

	m_b_dirty = true;
	// we're going to modify the binaries

	_TSHA1 t_source_hash = t_Hash_String(p_s_source_code);
	// calculate hash of source code

	if(device_id_list.size() != n_device_num || m_t_header.t_source_hash != t_source_hash ||
	   !b_Check_DeviceIds(n_device_num, p_device_list)) { // this is costy, make it at the end
		std::for_each(instance_list.begin(), instance_list.end(), DeleteInstance);
		instance_list.clear();
		device_id_list.clear();

		m_t_header.t_source_hash = t_source_hash;

		try {
			device_id_list.resize(n_device_num);
		} catch(std::bad_alloc&) {
			return false;
		}
		for(size_t i = 0; i < n_device_num; ++ i) {
			std::string s_dev_driver;
			if(CCLDeviceParams::n_GetDeviceInfoString(device_id_list[i],
			   p_device_list[i], CL_DEVICE_NAME) != CL_SUCCESS ||
			   CCLDeviceParams::n_GetDeviceInfoString(s_dev_driver,
			   p_device_list[i], CL_DRIVER_VERSION) != CL_SUCCESS ||
			   !stl_ut::AppendCStr(device_id_list[i], "|") ||
			   !stl_ut::Append(device_id_list[i], s_dev_driver)) {
				device_id_list.clear();
				return false;
			}
		}
		// get list of device id's (supposedly unique enough id's)
	}
	// in case it is completely different program, it invalidates all the instances

	try {
		instance_list.insert(instance_list.begin(), 1, TProgramInstance()); // add new instances at the beginning of the list
	} catch(std::bad_alloc&) {
		return false;
	}
	_TSHA1 t_build_opts_hash = t_Hash_String(p_s_build_options);
	if(!instance_list.front().Download(h_program, n_device_num, t_build_opts_hash)) {
		instance_list.front().Free();
		instance_list.erase(instance_list.begin()); // try to clean up so this would still be useful
		return false;
	}
	// add and fill a new program instance

	return true;
}

bool CCLProgramStorage::Save(const char *p_s_filename, size_t n_max_program_instance_num)
{
	if(!Calc_ControlHash(m_t_header.t_control_hash, n_max_program_instance_num))
		return false;
	// update hash

	FILE *p_fw;
#if defined(_MSC_VER) && !defined(__MWERKS__) && _MSC_VER >= 1400
	if(fopen_s(&p_fw, p_s_filename, "wb"))
#else //_MSC_VER && !__MWERKS__ && _MSC_VER >= 1400
	if(!(p_fw = fopen(p_s_filename, "wb")))
#endif //_MSC_VER && !__MWERKS__ && _MSC_VER >= 1400
		return false;
	CFILE_PtrGuard guard(p_fw); // closes the file when leaving the scope
	// open file

	m_t_header.p_magic[0] = 'L';
	m_t_header.p_magic[1] = 'C';
	m_t_header.p_magic[2] = 'L';
	m_t_header.p_magic[3] = 's';
	// set magic word ("Lame CL storage")

	_ASSERTE(device_id_list.size() < UINT32_MAX);
	uint32_t n_device_num = uint32_t(device_id_list.size());
	// set device count

	if(min(n_max_program_instance_num, instance_list.size()) > UINT32_MAX)
		return false;
	uint32_t n_instance_num = uint32_t(min(n_max_program_instance_num, instance_list.size()));
	if(fwrite(&m_t_header, sizeof(TFileHeader), 1, p_fw) != 1 ||
	   fwrite(&n_device_num, sizeof(uint32_t), 1, p_fw) != 1 ||
	   fwrite(&n_instance_num, sizeof(uint32_t), 1, p_fw) != 1)
		return false;
	// write header

	for(size_t i = 0, n = device_id_list.size(); i < n; ++ i) {
		_ASSERTE(device_id_list[i].length() < SIZE_MAX);
		_ASSERTE(sizeof(char) == 1); // this better be consistent
		size_t n_length = device_id_list[i].length();
		if(fwrite(device_id_list[i].c_str(), sizeof(char), n_length + 1, p_fw) != n_length + 1)
			return false;
	}
	// write device names (null-terminated)

	for(size_t i = 0; i < n_instance_num; ++ i) {
		if(!instance_list[i].Write(p_fw, device_id_list.size()))
			return false;
	}
	// write instances

	if(ferror(p_fw))
		return false;
	// close file

	m_b_dirty = (n_max_program_instance_num < instance_list.size());
	// the file now contains actual data (in case we didn't trim it)

	/*fprintf(stderr, "saving binaries ...\n");*/ // debug

	return true;
}

bool CCLProgramStorage::Load(const char *p_s_filename)
{
	/*fprintf(stderr, "loading binaries ...\n");*/ // debug

	std::for_each(instance_list.begin(), instance_list.end(), DeleteInstance);
	instance_list.clear();
	device_id_list.clear();
	// clear first

	{
		FILE *p_fr;
#if defined(_MSC_VER) && !defined(__MWERKS__) && _MSC_VER >= 1400
		if(fopen_s(&p_fr, p_s_filename, "rb"))
#else //_MSC_VER && !__MWERKS__ && _MSC_VER >= 1400
		if(!(p_fr = fopen(p_s_filename, "rb")))
#endif //_MSC_VER && !__MWERKS__ && _MSC_VER >= 1400
			return false;
		CFILE_PtrGuard guard(p_fr); // closes the file when leaving the scope
		// open file

		uint32_t n_instance_num, n_device_num;
		if(fread(&m_t_header, sizeof(TFileHeader), 1, p_fr) != 1 ||
		   fread(&n_device_num, sizeof(uint32_t), 1, p_fr) != 1 ||
		   fread(&n_instance_num, sizeof(uint32_t), 1, p_fr) != 1 ||
		   n_device_num > SIZE_MAX || n_instance_num > SIZE_MAX ||
		   m_t_header.p_magic[0] != 'L' || m_t_header.p_magic[1] != 'C' ||
		   m_t_header.p_magic[2] != 'L' || m_t_header.p_magic[3] != 's')
			return false;
		// read header

		try {
			_ASSERTE(n_device_num <= SIZE_MAX);
			device_id_list.resize(size_t(n_device_num));
			_ASSERTE(n_instance_num <= SIZE_MAX);
			instance_list.resize(size_t(n_instance_num));
		} catch(std::bad_alloc&) {
			return false;
		}
		// alloc devices and instances

		{
			for(size_t i = 0, n = device_id_list.size(); i < n; ++ i) {
				std::string s_id;
				for(;;) {
					char c;
					if(fread(&c, sizeof(char), 1, p_fr) != 1) {
						device_id_list.clear();
						return false;
					}
					// read a single character

					if(!c)
						break;
					// detect terminating null

					try {
						s_id += c;
					} catch(std::bad_alloc&) {
						device_id_list.clear();
						return false;
					}
					// add character to the string
				}

				device_id_list[i].swap(s_id);
			}
		}
		// read device id's

		for(size_t i = 0, n = instance_list.size(); i < n; ++ i) {
			if(!instance_list[i].Read(p_fr, size_t(n_device_num))) {
				std::for_each(instance_list.begin(), instance_list.end(), DeleteInstance);
				instance_list.clear();
				device_id_list.clear();
				// !!

				return false;
			}
		}
		// read instances
	}

	_TSHA1 t_ctrl_hash;
	if(!Calc_ControlHash(t_ctrl_hash) || t_ctrl_hash != m_t_header.t_control_hash) {
		/*fprintf(stderr, "hash check failed ...\n");*/ // debug

		std::for_each(instance_list.begin(), instance_list.end(), DeleteInstance);
		instance_list.clear();
		device_id_list.clear();
		// !!

		return false;
	}
	// check hash

	m_b_dirty = false;
	// actual data are now loaded

	return true;
}

CCLProgramStorage::_TSHA1 CCLProgramStorage::t_Hash_String(const char *p_s_str)
{
	CStreamHash<TSHA1> hash;
	hash.Process_Data(p_s_str, strlen(p_s_str) * sizeof(char));
	// calculate

	TSHA1 t_hash = hash.t_Result();
	_TSHA1 t_result;
	memcpy(&t_result, &t_hash[0], sizeof(_TSHA1)); // TSHA1 contains byte counter which is not needed
	// convert result

	return t_result;
}

bool CCLProgramStorage::Calc_ControlHash(_TSHA1 &r_t_result, size_t n_instance_limit) const
{
	CStreamHash<TSHA1> hash;
	hash.Process_Data(&m_t_header.t_source_hash, sizeof(_TSHA1));
	_ASSERTE(device_id_list.size() <= UINT32_MAX);
	uint32_t n_device_num = uint32_t(device_id_list.size());
	hash.Process_Data(&n_device_num, sizeof(uint32_t));
	_ASSERTE(instance_list.size() <= UINT32_MAX);
	uint32_t n_instance_num = uint32_t(instance_list.size());
	hash.Process_Data(&n_instance_num, sizeof(uint32_t));
	// hash header ...

	for(size_t i = 0, n = device_id_list.size(); i < n; ++ i) {
		_ASSERTE(device_id_list[i].length() < SIZE_MAX);
		hash.Process_Data(device_id_list[i].c_str(), (device_id_list[i].length() + 1) * sizeof(char));
	}
	// hash device id's (including the terminating null - to mimic layout in the file)

	for(size_t i = 0, n = min(n_instance_limit, instance_list.size()); i < n; ++ i) {
		const TProgramInstance &r_t_inst = instance_list[i];
		// get program instance ...

		hash.Process_Data(&r_t_inst.t_build_opts_hash, sizeof(_TSHA1));
		if(r_t_inst.binary_size_list.size() != device_id_list.size())
			return false;
		for(size_t j = 0, m = r_t_inst.binary_size_list.size(); j < m; ++ j) {
			_ASSERTE(SIZE_MAX <= UINT64_MAX);
			uint64_t n_size = r_t_inst.binary_size_list[j];
			hash.Process_Data(&n_size, sizeof(uint64_t));
		}
		// hash header

		_ASSERTE(r_t_inst.t_data_buffer.n_size <= SIZE_MAX);
		hash.Process_Data(r_t_inst.t_data_buffer.p_data, size_t(r_t_inst.t_data_buffer.n_size)); // note this might overflow
		// hash data
	}
	// hash instances

	TSHA1 t_hash = hash.t_Result();
	memcpy(&r_t_result, &t_hash[0], sizeof(_TSHA1)); // TSHA1 contains byte counter which is not needed
	// write result

	return true;
}

inline void CCLProgramStorage::DeleteInstance(TProgramInstance &r_t_inst)
{
	r_t_inst.Free();
}

/*
 *								=== ~CCLProgramStorage ===
 */

/*
 *								=== CCLProgramCompiler ===
 */

#define CL_COMPILER_UTIL_BOGUS_DATA_DETECTION

void CCLProgramCompiler::Dump_StatusWord(int n_status_word, FILE *p_fw)
{
	if(n_status_word & cache_ReadAttempted) {
		if(n_status_word & cache_ReadSucceeded)
			fprintf(p_fw, "program loaded from cache\n");
		else {
			if(n_status_word & cache_ReadFailed_FileNotFound)
				fprintf(p_fw, "program not cached\n");
			else if(n_status_word & cache_ReadFailed_IO)
				fprintf(p_fw, "i/o error while loading program from cache\n");
			else if(n_status_word & cache_ReadFailed_OutOfMemory)
				fprintf(p_fw, "not enough memory while loading program from cache\n");
			else if(n_status_word & cache_ReadFailed_SourceChecksum)
				fprintf(p_fw, "program cache miss\n");
			else if(n_status_word & cache_ReadFailed_BinaryChecksum)
				fprintf(p_fw, "program cache corrupt\n");
			else if(n_status_word & cache_ReadFailed_CreateProgram)
				fprintf(p_fw, "program cached for different device\n");
			else
				fprintf(p_fw, "unknown error while loading program from cache\n");
		}
	}
	if(n_status_word & prog_CompiledFromSource) {
		if(!(n_status_word & prog_CreateSucceeded))
			fprintf(p_fw, "clCreateProgramWithSource() failed while compiling program from source\n");
		else if(!(n_status_word & prog_BuildSucceeded))
			fprintf(p_fw, "clBuildProgram() failed while compiling program from source\n");
		else
			fprintf(p_fw, "program was successfully compiled from source\n");
	}
	if(n_status_word & cache_WriteAttempted) {
		if(n_status_word & cache_WriteSucceeded) {
			if(n_status_word & cache_WriteTrimmed) {
				fprintf(p_fw, "program binaries were successfully written to cache, binary versions dropped\n");
			} else
				fprintf(p_fw, "program binaries were successfully written to cache\n");
		} else
			fprintf(p_fw, "failed to cache program binaries\n");
	}
}

CLresult CCLProgramCompiler::n_CompileProgram(cl_context h_context, cl_program *p_program,
	const char *p_s_source, size_t n_device_num, const cl_device_id *p_device_list,
	const char *p_s_build_options, const char *p_s_cache_file, int *p_status_word,
	size_t n_max_program_instance_num)
{
	int n_dummy_status_word;
	if(!p_status_word)
		p_status_word = &n_dummy_status_word;
	// makes sure p_status_word is not null

	*p_status_word = 0;

	std::vector<cl_device_id> all_dev_list;
	if(!n_device_num) {
		CLresult n_result;
		if((n_result = CCLUtils::n_GetDeviceList(h_context, all_dev_list)) != CL_SUCCESS)
			return n_result;
		n_device_num = all_dev_list.size();
		p_device_list = (n_device_num)? &all_dev_list[0] : 0;
	} else
		_ASSERTE(p_device_list);
	// in case the devices are not specified, get the devices for this context;
	// otherwise will be unable to save program binaries

	std::string s_real_cache_file;
	if(p_s_cache_file && *p_s_cache_file == '%') {
		if(strstr(p_s_cache_file, "%temp_folder%") == p_s_cache_file) {
			p_s_cache_file += strlen("%temp_folder%");
			// skip

			_ASSERTE(strlen(p_s_cache_file) > 1 && (*p_s_cache_file == '/' ||
				*p_s_cache_file == CPath::path_Separator));
			// make sure that a slash and the actual name follows

			if(!CPath::Get_TempDirectory(s_real_cache_file) || // in Dir.h so far, don't have to include Dir.cpp to older builds
			   !stl_ut::AppendCStr(s_real_cache_file, p_s_cache_file))
				return cl_Out_Of_Host_Memory;
			p_s_cache_file = s_real_cache_file.c_str();
			// get temp, append the rest of the name
		} else if(strstr(p_s_cache_file, "%temp_default%") == p_s_cache_file) {
			p_s_cache_file += strlen("%temp_default%");
			// skip

			const char *p_s_filename;
#if defined(_WIN32) || defined(_WIN64)
			char p_s_filename_buffer[1024];
			if(!GetModuleFileNameA(GetModuleHandle(NULL), p_s_filename_buffer,
			   sizeof(p_s_filename_buffer) / sizeof(p_s_filename_buffer[0]) - 1))
				p_s_filename = 0; // error
			else
				p_s_filename = p_s_filename_buffer;
			// get name of the .exe
#else // _WIN32 || _WIN64
			p_s_filename = program_invocation_name;
#endif // _WIN32 || _WIN64
			if(p_s_filename && strchr(p_s_filename, CPath::path_Separator))
				p_s_filename = strrchr(p_s_filename, CPath::path_Separator) + 1;
			else if(!p_s_filename)
				p_s_filename = "null";
			// get filename

			char p_s_sep[2] = {CPath::path_Separator, 0}; // CPath::Get_TempDirectory() doesn't append slash!
			if(!CPath::Get_TempDirectory(s_real_cache_file) || // in Dir.h so far, don't have to include Dir.cpp to older builds
			   !stl_ut::AppendCStr(s_real_cache_file, p_s_sep) ||
			   !stl_ut::AppendCStr(s_real_cache_file, p_s_filename) ||
			   !stl_ut::AppendCStr(s_real_cache_file, "_") ||
			   !stl_ut::AppendCStr(s_real_cache_file, p_s_cache_file) ||
			   !stl_ut::AppendCStr(s_real_cache_file, ".clbin"))
				return cl_Out_Of_Host_Memory;
			p_s_cache_file = s_real_cache_file.c_str();
			// concat
		}

		//printf("debug: OpenCL compiler temp filename is \'%s\'\n", p_s_cache_file); // debug
	}
	// handle cache file in temp (so that the app doesn't need to have write privileges)

#ifdef CL_PROGRAM_COMPILER_USE_OLD_BINARIES_FORMAT
	TSHA1 t_hash = t_Hash_ProgramSource_BuildOptions(p_s_source, p_s_build_options);
	// calculate program and build options hash
#else // CL_PROGRAM_COMPILER_USE_OLD_BINARIES_FORMAT
	CCLProgramStorage program_storage;
	// prepare program storage
#endif // CL_PROGRAM_COMPILER_USE_OLD_BINARIES_FORMAT

	if(p_s_cache_file) {
		*p_status_word |= cache_ReadAttempted;

#ifdef CL_PROGRAM_COMPILER_USE_OLD_BINARIES_FORMAT
		int n_result;
		if((n_result = n_ReadProgramBinaries(p_s_cache_file, t_hash, h_context,
		   p_program, n_device_num, p_device_list)) == cache_ReadSucceeded) {
			*p_status_word |= cache_ReadSucceeded;
			return CL_SUCCESS;
		}
		// use the old function

		*p_status_word |= n_result;
#else // CL_PROGRAM_COMPILER_USE_OLD_BINARIES_FORMAT
		if(program_storage.Load(p_s_cache_file)) {
			size_t n_binaries_index;
			if((n_binaries_index = program_storage.n_Find_ProgramBinaries(p_s_source,
			   n_device_num, p_device_list, p_s_build_options)) != CCLProgramStorage::npos) {
				int n_result; CLresult n_cl_result;
				if((n_result = program_storage.n_Get_ProgramBinaries(h_context, p_program, n_cl_result, n_binaries_index,
				   n_device_num, p_device_list)) == CCLProgramStorage::build_Success) {
					*p_status_word |= cache_ReadSucceeded;

					if(program_storage.b_Dirty()) {
						*p_status_word |= cache_WriteAttempted;

#ifdef CL_PROGRAM_COMPILER_PRECAUTIOUS_BINARIES_SAVE
						std::string s_temp_name;
						char p_s_time[5];
						sprintf(p_s_time, "%04x", clock() & 0xffff); // "random" identifier
						try {
							s_temp_name = p_s_cache_file;
							s_temp_name += ".tmp";
							s_temp_name += p_s_time;
						} catch(std::bad_alloc&) {
							return CL_SUCCESS; // load succeeded, just the save failed. ignore it.
						}
						// create temp filename

						if(program_storage.Save(s_temp_name.c_str(), n_max_program_instance_num)) {
							remove(p_s_cache_file);
							rename(s_temp_name.c_str(), p_s_cache_file);

							*p_status_word |= cache_WriteSucceeded;
						}
						// save with new filename, if succeeds, rewrite the old file; this is important because of the lru policy
#else // CL_PROGRAM_COMPILER_PRECAUTIOUS_BINARIES_SAVE
						if(program_storage.Save(p_s_cache_file, n_max_program_instance_num)) {
							*p_status_word |= cache_WriteSucceeded;
							if(program_storage.b_Dirty())
								*p_status_word |= cache_WriteTrimmed; // it's dirty - file doesn't contain all the versions - n_max_program_instance_num limit
						} else
							remove(p_s_cache_file); // don't leave garbage
						// save with old filename, remove file on failure (no fallback)
#endif // CL_PROGRAM_COMPILER_PRECAUTIOUS_BINARIES_SAVE
					}

					return cl_Success;
				}
				switch(n_result) {
				case CCLProgramStorage::build_InvalidParams:
					*p_status_word |= cache_ReadFailed_OutOfMemory; // not actually the right error code (but it shouldn't happen in here, it's more of an internal program error)
					break;
				case CCLProgramStorage::build_LowMemory:
					*p_status_word |= cache_ReadFailed_OutOfMemory;
					break;
				case CCLProgramStorage::build_CreateProgramFailure:
					*p_status_word |= cache_ReadFailed_CreateProgram;
					break;
				case CCLProgramStorage::build_BuildProgramFailure:
					*p_status_word |= cache_ReadFailed_CreateProgram; // not actually the right error code
					break;
				}
			} else
				*p_status_word |= cache_ReadFailed_SourceChecksum;
		} else {
			if(access(p_s_cache_file, 00) != 0)
				*p_status_word |= cache_ReadFailed_FileNotFound;
			else
				*p_status_word |= cache_ReadFailed_IO;
		}
		// use the new program storage
#endif // CL_PROGRAM_COMPILER_USE_OLD_BINARIES_FORMAT
	}
	// attempt to load program binaries from file

	cl_program h_prog;
	{
		*p_status_word |= prog_CompiledFromSource;

		cl_int n_result;
		h_prog = clCreateProgramWithSource(h_context, 1, &p_s_source, NULL, &n_result);
		if(n_result != CL_SUCCESS)
			return (CLresult)n_result;

		*p_status_word |= prog_CreateSucceeded;

		_ASSERTE(n_device_num <= /*CL*/UINT_MAX);
		if((n_result = clBuildProgram(h_prog, cl_uint(n_device_num), p_device_list, p_s_build_options, NULL, NULL)) != CL_SUCCESS) {
#ifdef CL_PROGRAM_COMPILER_DISPLAY_BUILD_ERRORS
			if(n_device_num) {
				std::string s_build_log, s_device_name;
				for(size_t i = 0; i < n_device_num; ++ i) {
					cl_build_status n_build_status;

					CLresult n_cl_result;
					if((n_cl_result = n_Get_BuildLog(s_build_log, n_build_status, h_prog, p_device_list[i])) != CL_SUCCESS)
						return n_cl_result;
					// get buld log and build status

					if(CCLDeviceParams::n_GetDeviceInfoString(s_device_name, p_device_list[i], CL_DEVICE_NAME) != CL_SUCCESS)
						s_device_name.erase();
					const char *p_s_device_name = (s_device_name.empty())? "(null)" : s_device_name.c_str();
					// get device name

					if(n_build_status == CL_BUILD_NONE) {
						fprintf(stderr, "error: program wasn't built for device " PRIsize " (%s)\n", i, p_s_device_name);
						continue;
					} else if(n_build_status == CL_BUILD_ERROR) {
						fprintf(stderr, "error: there were errors while building program for device " PRIsize " (%s)\n",
							i, p_s_device_name);
					} else if(n_build_status != CL_BUILD_SUCCESS)
						fprintf(stderr, "error: unknown program build status for device " PRIsize " (%s)\n", i, p_s_device_name);
					// show build status

					if(n_build_status != CL_BUILD_SUCCESS || !s_build_log.empty()) {
						fprintf(stderr, "=== OpenCL build log for device " PRIsize " (%s) ===\n%s\n",
							i, p_s_device_name, s_build_log.c_str());
					}
					// show build log
				}
			}
			// show build error(s)
#endif // CL_PROGRAM_COMPILER_DISPLAY_BUILD_ERRORS

			clReleaseProgram(h_prog);
			return (CLresult)n_result;
		}
		// build program

#ifdef CL_PROGRAM_COMPILER_DISPLAY_SUCCESSFUL_BUILD_WARNINGS
		if(n_device_num) {
			std::string s_build_log, s_device_name;
			for(size_t i = 0; i < n_device_num; ++ i) {
				cl_build_status n_build_status;

				CLresult n_cl_result;
				if((n_cl_result = n_Get_BuildLog(s_build_log, n_build_status, h_prog, p_device_list[i])) != CL_SUCCESS)
					return n_cl_result;
				// get buld log and build status

				if(CCLDeviceParams::n_GetDeviceInfoString(s_device_name, p_device_list[i], CL_DEVICE_NAME) != CL_SUCCESS)
					s_device_name.erase();
				const char *p_s_device_name = (s_device_name.empty())? "(null)" : s_device_name.c_str();
				// get device name

				if(n_build_status == CL_BUILD_NONE) {
					fprintf(stderr, "warning: program wasn't built for device %d (%s)\n", i, p_s_device_name);
					continue;
				} else if(n_build_status == CL_BUILD_ERROR) {
					fprintf(stderr, "warning: there were errors while building program for device %d (%s)\n",
						i, p_s_device_name);
				} else if(n_build_status != CL_BUILD_SUCCESS)
					fprintf(stderr, "warning: unknown program build status for device %d (%s)\n", i, p_s_device_name);
				// show build status

				if(n_build_status != CL_BUILD_SUCCESS || !s_build_log.empty()) {
					fprintf(stderr, "=== OpenCL build log for device %d (%s) ===\n%s\n",
						i, p_s_device_name, s_build_log.c_str());
				}
				// show build log
			}
		}
		// show warnings
#endif // CL_PROGRAM_COMPILER_DISPLAY_SUCCESSFUL_BUILD_WARNINGS

		*p_status_word |= prog_BuildSucceeded;
	}
	// create program

	*p_program = h_prog;
	// output

	if(p_s_cache_file) {
		*p_status_word |= cache_WriteAttempted;

#ifdef CL_PROGRAM_COMPILER_USE_OLD_BINARIES_FORMAT
		if(WriteProgramBinaries(h_prog, t_hash, p_s_cache_file))
			*p_status_word |= cache_WriteSucceeded;
		else {
			remove(p_s_cache_file);
			// in case it didn't save correctly, do not leave it arround
		}
#else // CL_PROGRAM_COMPILER_USE_OLD_BINARIES_FORMAT
		if(program_storage.Put_ProgramBinaries(h_prog, p_s_source,
		   n_device_num, p_device_list, p_s_build_options)) {
			_ASSERTE(program_storage.b_Dirty());
			// should be "dirty" now

#ifdef CL_PROGRAM_COMPILER_PRECAUTIOUS_BINARIES_SAVE
			std::string s_temp_name;
			char p_s_time[5];
			sprintf(p_s_time, "%04x", clock() & 0xffff); // "random" identifier
			try {
				s_temp_name = p_s_cache_file;
				s_temp_name += ".tmp";
				s_temp_name += p_s_time;
			} catch(std::bad_alloc&) {
				return CL_SUCCESS; // load succeeded, just the save failed. ignore it.
			}
			// create temp filename

			if(program_storage.Save(s_temp_name.c_str(), n_max_program_instance_num)) {
				remove(p_s_cache_file);
				rename(s_temp_name.c_str(), p_s_cache_file);

				*p_status_word |= cache_WriteSucceeded;
			}
			// save with new filename, if succeeds, rewrite the old file; this is important because of the lru policy
#else // CL_PROGRAM_COMPILER_PRECAUTIOUS_BINARIES_SAVE
			if(program_storage.Save(p_s_cache_file, n_max_program_instance_num)) {
				*p_status_word |= cache_WriteSucceeded;
				if(program_storage.b_Dirty())
					*p_status_word |= cache_WriteTrimmed; // it's dirty - file doesn't contain all the versions - n_max_program_instance_num limit
			} else
				remove(p_s_cache_file); // don't leave garbage
			// save with old filename, remove file on failure (no fallback)
#endif // CL_PROGRAM_COMPILER_PRECAUTIOUS_BINARIES_SAVE
		}
#endif // CL_PROGRAM_COMPILER_USE_OLD_BINARIES_FORMAT
	}
	// write file with binaries

	return cl_Success;
}

CLresult CCLProgramCompiler::n_Get_BuildLog(std::string &r_s_build_log, cl_build_status &r_n_build_status,
	cl_program h_program, cl_device_id h_device)
{
	r_s_build_log.erase();
	r_n_build_status = CL_BUILD_NONE;
	// clear output

	for(;;) {
		CLresult n_result;
		if((n_result = (CLresult)clGetProgramBuildInfo(h_program, h_device, CL_PROGRAM_BUILD_STATUS,
		   sizeof(cl_build_status), &r_n_build_status, NULL)) != CL_SUCCESS)
			return n_result;
		if(r_n_build_status != CL_BUILD_IN_PROGRESS)
			break;
#if defined(_WIN32) || defined(_WIN64)
		Sleep(100);
#else // _WIN32 || _WIN64
		sleep(1);
#endif // _WIN32 || _WIN64
	}
	// wait while build is in progress

	if(r_n_build_status == CL_BUILD_NONE)
		return cl_Success;
	// check build status

	size_t n_build_log_size;
	{
		CLresult n_result;
		if((n_result = (CLresult)clGetProgramBuildInfo(h_program, h_device,
		   CL_PROGRAM_BUILD_LOG, 0, NULL, &n_build_log_size)) != CL_SUCCESS)
			return n_result;
	}
	if(!stl_ut::Resize_To_N(r_s_build_log, n_build_log_size + 1)) {
		fprintf(stderr, "error: not enough memory for build log\n");
		return cl_Out_Of_Host_Memory;
	}
	// get build log length, allocate string buffer

	{
		CLresult n_result;
		if((n_result = (CLresult)clGetProgramBuildInfo(h_program, h_device, CL_PROGRAM_BUILD_LOG,
		   n_build_log_size, &r_s_build_log[0], NULL)) != CL_SUCCESS)
			return n_result;
	}
	r_s_build_log.resize(strlen(r_s_build_log.c_str()));
	// get build log, erase terminating null from string

	return cl_Success;
}

CLresult CCLProgramCompiler::n_CompileProgramFile(cl_context h_context, cl_program *p_program,
	const char *p_s_source_file, size_t n_device_num, const cl_device_id *p_device_list,
	const char *p_s_build_options, const char *p_s_cache_file, int *p_status_word,
	size_t n_max_program_instance_num)
{
	int n_dummy_status_word;
	if(!p_status_word)
		p_status_word = &n_dummy_status_word;
	// makes sure p_status_word is not null

	*p_status_word = 0;
	// in case i/o errors occur below

	std::string s_program;
	{
		FILE *p_fr;
#if defined(_MSC_VER) && !defined(__MWERKS__) && _MSC_VER >= 1400
		if(fopen_s(&p_fr, p_s_source_file, "rb"))
#else //_MSC_VER && !__MWERKS__ && _MSC_VER >= 1400
		if(!(p_fr = fopen(p_s_source_file, "rb")))
#endif //_MSC_VER && !__MWERKS__ && _MSC_VER >= 1400
			return cl_Build_Program_Failure; // not really great, but there's no better error code
		CFILE_PtrGuard guard(p_fr); // closes the file when leaving the scope
		fseek(p_fr, 0, SEEK_END);
		size_t n_file_size = ftell(p_fr);
		try {
			s_program.resize(n_file_size);
		} catch(std::bad_alloc&) {
			return cl_Out_Of_Host_Memory;
		}
		fseek(p_fr, 0, SEEK_SET);
		if(fread(&s_program[0], 1, n_file_size, p_fr) != n_file_size)
			return cl_Build_Program_Failure; // not really great, but there's no better error code
	}
	// read program from a file

	return n_CompileProgram(h_context, p_program, s_program.c_str(), n_device_num,
		p_device_list, p_s_build_options, p_s_cache_file, p_status_word, n_max_program_instance_num);
	// use the other function
}

#ifdef CL_PROGRAM_COMPILER_USE_OLD_BINARIES_FORMAT

TSHA1 CCLProgramCompiler::t_Hash_ProgramSource_BuildOptions(const char *p_s_source, const char *p_s_build_options)
{
	CStreamHash<TSHA1> hash;
	hash.Process_Data(p_s_source, strlen(p_s_source) * sizeof(char));
	if(p_s_build_options && strcmp(p_s_build_options, ""))
		hash.Process_Data(p_s_build_options, strlen(p_s_build_options) * sizeof(char));
	return hash.t_Result();
}

int CCLProgramCompiler::n_ReadProgramBinaries(const char *p_s_filename, TSHA1 t_hash, cl_context h_context,
	cl_program *p_program, size_t n_device_num, const cl_device_id *p_device_list)
{
	FILE *p_fr;
	if(!(p_fr = fopen(p_s_filename, "rb")))
		return cache_ReadFailed_FileNotFound;
	// open the file

	TSHA1 t_ref_hash;
	if(fread(&t_ref_hash[0], sizeof(uint32_t), 5, p_fr) != 5) {
		fclose(p_fr);
		return cache_ReadFailed_IO;
	}
	// read reference hash

	if(memcmp(&t_hash[0], &t_ref_hash[0], 5 * sizeof(uint32_t))) {
		fclose(p_fr);
		return cache_ReadFailed_SourceChecksum;
	}
	// compare it to current hash

	size_t n_binary_num;
	std::vector<size_t> binary_size_list;
	{
		uint32_t n_binary_num32;
		if(fread(&n_binary_num32, sizeof(uint32_t), 1, p_fr) != 1 || n_binary_num32 > SIZE_MAX) {
			fclose(p_fr);
			return cache_ReadFailed_IO;
		}
		n_binary_num = size_t(n_binary_num32);
		try {
			binary_size_list.resize(n_binary_num);
		} catch(std::bad_alloc&) {
			fclose(p_fr);
			return cache_ReadFailed_OutOfMemory;
		}
		for(size_t i = 0; i < n_binary_num; ++ i) {
			uint32_t n_binary_size32;
			if(fread(&n_binary_size32, sizeof(uint32_t), 1, p_fr) != 1 || n_binary_size32 > SIZE_MAX) {
				fclose(p_fr);
				return cache_ReadFailed_IO;
			}
			binary_size_list[i] = size_t(n_binary_size32);
		}
	}
	// read binary sizes

	size_t n_binaries_size = 0;
	for(size_t i = 0; i < n_binary_num; ++ i) {
		_ASSERTE(n_binaries_size <= SIZE_MAX - binary_size_list[i]);
		n_binaries_size += binary_size_list[i];
	}
	// sum binary sizes up

	TSHA1 t_binaries_hash;
	std::vector<const unsigned char*> binary_ptr_list;
	unsigned char *p_binaries;
	{
		try {
			binary_ptr_list.resize(n_binary_num);
			p_binaries = new unsigned char[n_binaries_size];
		} catch(std::bad_alloc&) {
			fclose(p_fr);
			return cache_ReadFailed_OutOfMemory;
		}
		// alloc buffers

		CStreamHash<TSHA1> hash;

		unsigned char *p_binary_ptr = p_binaries;
		for(size_t i = 0; i < n_binary_num; ++ i) {
			binary_ptr_list[i] = p_binary_ptr;

			size_t n_binary_size = binary_size_list[i];

			if(fread(p_binary_ptr, n_binary_size, 1, p_fr) != 1) {
				fclose(p_fr);
				return cache_ReadFailed_IO;
			}

			hash.Process_Data(p_binary_ptr, n_binary_size);

			p_binary_ptr += n_binary_size;
		}
		// get pointers, read binaries (so read block size is 32-bit), calculate their hashes

		t_binaries_hash = hash.t_Result();
	}
	// read binaries

	TSHA1 t_binaries_hash_ref;
	if(fread(&t_binaries_hash_ref[0], sizeof(uint32_t), 5, p_fr) != 5) {
		fclose(p_fr);
		return cache_ReadFailed_IO;
	}
	// read checksum

	fclose(p_fr);
	// close file

	if(memcmp(&t_binaries_hash_ref[0], &t_binaries_hash_ref[0], 5 * sizeof(uint32_t)))
		return cache_ReadFailed_BinaryChecksum;
	// make sure checksum is correct

#if defined CL_COMPILER_UTIL_BOGUS_DATA_DETECTION && defined(__DIR_INCLUDED)
	TFileInfo t_file(p_s_filename);
	if(t_file.n_Size64() != 2 * 5 * sizeof(uint32_t) + sizeof(uint32_t) + // size of two hashes, number of binaries
	   n_binary_num * sizeof(uint32_t) + n_binaries_size) // lengths of binaries, binaries data
		return cache_ReadFailed_BinaryChecksum; // ...
	// make sure file doesn't contain bogus data (paranoid)
#endif // CL_COMPILER_UTIL_BOGUS_DATA_DETECTION && __DIR_INCLUDED

	cl_int n_result;
	cl_program h_prog = clCreateProgramWithBinary(h_context, n_device_num, p_device_list,
		&binary_size_list[0], &binary_ptr_list[0], NULL, &n_result);
	if(n_result != CL_SUCCESS)
		return cache_ReadFailed_CreateProgram;
	if((n_result = clBuildProgram(h_prog, 0, NULL, NULL, NULL, NULL)) != CL_SUCCESS) {
		clReleaseProgram(h_prog);
		return cache_ReadFailed_CreateProgram;
	}
	// use loaded binaries to create the program ...

	*p_program = h_prog;
	// write output

	return cache_ReadSucceeded;
}

bool CCLProgramCompiler::WriteProgramBinaries(cl_program h_program, TSHA1 t_hash, const char *p_s_filename)
{
	cl_uint n_device_num;
	if(clGetProgramInfo(h_program, CL_PROGRAM_NUM_DEVICES,
	   sizeof(size_t), &n_device_num, NULL) != CL_SUCCESS)
		return false;
	// get number of devices

	if(!n_device_num || n_device_num > UINT32_MAX)
		return false;
	// we can't query binaries without devices

	size_t n_binary_num = n_device_num;

	std::vector<size_t> binary_size_list;
	try {
#ifdef _DEBUG
		_ASSERTE(n_binary_num < SIZE_MAX);
		binary_size_list.resize(n_binary_num + 1);
		binary_size_list[n_binary_num] = 0xbaadf00d; // "magic" word
#else // _DEBUG
		binary_size_list.resize(n_binary_num);
#endif // _DEBUG
	} catch(std::bad_alloc&) {
		return false;
	}
	if(clGetProgramInfo(h_program, CL_PROGRAM_BINARY_SIZES,
	   n_binary_num * sizeof(size_t), &binary_size_list[0], NULL) != CL_SUCCESS)
		return false;
#ifdef _DEBUG
	_ASSERTE(binary_size_list[n_binary_num] == 0xbaadf00d); // make sure it's returning the right amount of data
#endif // _DEBUG
	// get binary sizes

	size_t n_binaries_size = 0;
	for(size_t i = 0; i < n_binary_num; ++ i) {
		_ASSERTE(n_binaries_size <= SIZE_MAX - binary_size_list[i]);
		n_binaries_size += binary_size_list[i];
	}
	// sum binary sizes up

	FILE *p_fw;
	if(!(p_fw = fopen(p_s_filename, "wb")))
		return false;
	// open file for writing ...

	if(fwrite(&t_hash[0], sizeof(uint32_t), 5, p_fw) != 5) {
		fclose(p_fw);
		return false;
	}
	// write hash

	uint32_t n_binary_num32 = uint32_t(n_binary_num);
	if(fwrite(&n_binary_num32, sizeof(uint32_t), 1, p_fw) != 1) {
		fclose(p_fw);
		return false;
	}
	// write number of binaries

	for(size_t i = 0; i < n_binary_num; ++ i) {
		_ASSERTE(binary_size_list[i] <= UINT32_MAX);
		uint32_t n_binary_size32 = binary_size_list[i];
		if(fwrite(&n_binary_size32, sizeof(uint32_t), 1, p_fw) != 1) {
			fclose(p_fw);
			return false;
		}
	}
	// write size of each binary

	unsigned char *p_binaries;
	std::vector<const unsigned char*> binary_ptr_list;
	{
		try {
			binary_ptr_list.resize(n_binary_num);
			p_binaries = new unsigned char[n_binaries_size];
		} catch(std::bad_alloc&) {
			fclose(p_fw);
			return false;
		}
		// alloc buffers

		unsigned char *p_binary_ptr = p_binaries;
		for(size_t i = 0; i < n_binary_num; ++ i) {
			binary_ptr_list[i] = p_binary_ptr;
			p_binary_ptr += binary_size_list[i];
		}
		// get pointers
	}
	if(clGetProgramInfo(h_program, CL_PROGRAM_BINARIES,
	   n_binaries_size, &binary_ptr_list[0], NULL) != CL_SUCCESS) {
		fclose(p_fw);
		delete[] p_binaries;
		return false;
	}
	// get binaries

	TSHA1 t_binaries_hash;
	{
		CStreamHash<TSHA1> hash;
		if(fwrite(p_binaries, n_binaries_size, 1, p_fw) != 1) {
			fclose(p_fw);
			delete[] p_binaries;
			return false;
		}
		hash.Process_Data(p_binaries, n_binaries_size);
		t_binaries_hash = hash.t_Result();
	}
	// write binaries, calculate hash

	delete[] p_binaries;
	// cleanup

	if(fwrite(&t_binaries_hash[0], sizeof(uint32_t), 5, p_fw) != 5) {
		fclose(p_fw);
		return false;
	}
	// write hash

	fclose(p_fw);
	// close output file

	return true;
}

#endif // CL_PROGRAM_COMPILER_USE_OLD_BINARIES_FORMAT

/*
 *								=== ~CCLProgramCompiler ===
 */

/*
 *								=== CCLUniqueProgram ===
 */

#if defined(_MSC_VER) && !defined(__MWERKS__) && _MSC_VER <= 1200
const CCLUniqueProgram::TBuildFromSource_Tag CCLUniqueProgram::from_source_code;
const CCLUniqueProgram::TBuildFromCompressedSource_Tag CCLUniqueProgram::from_compressed;
const CCLUniqueProgram::TBuildFromFile_Tag CCLUniqueProgram::from_file;
// MSVC 60 can't initialize empty structures
#else // _MSC_VER && !__MWERKS__ && _MSC_VER <= 1200
const CCLUniqueProgram::TBuildFromSource_Tag CCLUniqueProgram::from_source_code = {};
const CCLUniqueProgram::TBuildFromCompressedSource_Tag CCLUniqueProgram::from_compressed = {};
const CCLUniqueProgram::TBuildFromFile_Tag CCLUniqueProgram::from_file = {}; // g++ requires initialization
#endif // _MSC_VER && !__MWERKS__ && _MSC_VER <= 1200
// constructor tags

CCLUniqueProgram::CCLUniqueProgram(cl_context h_context, const char *p_s_source_code,
	TBuildFromSource_Tag UNUSED(t_tag), const char *p_s_compiler_options /*= ""*/,
	const char *p_s_cache_file /*= "%copykernelname%"*/)
	:m_h_program(0), m_n_compile_flags(0)
{
	m_n_last_result = n_Init(false, h_context, p_s_source_code, 0, 0,
		p_s_compiler_options, p_s_cache_file, 32);
}

CCLUniqueProgram::CCLUniqueProgram(CCLUniqueInstance &r_instance, const char *p_s_source_code,
	TBuildFromSource_Tag UNUSED(t_tag), const char *p_s_compiler_options /*= ""*/,
	const char *p_s_cache_file /*= "%copykernelname%"*/)
	:m_h_program(0), m_n_compile_flags(0)
{
	m_n_last_result = n_Init(false, r_instance.h_Context(), p_s_source_code,
		r_instance.n_Device_Num(), r_instance.p_Device(),
		p_s_compiler_options, p_s_cache_file, 32);
}

CCLUniqueProgram::CCLUniqueProgram(CCLUniqueInstance &r_instance, const char *p_s_filename,
	TBuildFromFile_Tag UNUSED(t_tag), const char *p_s_compiler_options /*= ""*/,
	const char *p_s_cache_file /*= "%copykernelname%"*/)
	:m_h_program(0), m_n_compile_flags(0)
{
	m_n_last_result = n_Init(true, r_instance.h_Context(), p_s_filename,
		r_instance.n_Device_Num(), r_instance.p_Device(),
		p_s_compiler_options, p_s_cache_file, 32);
}

CCLUniqueProgram::CCLUniqueProgram(cl_context h_context, const char *p_s_filename,
	TBuildFromFile_Tag UNUSED(t_tag), const char *p_s_compiler_options /*= ""*/,
	const char *p_s_cache_file /*= "%copykernelname%"*/)
	:m_h_program(0), m_n_compile_flags(0)
{
	m_n_last_result = n_Init(true, h_context, p_s_filename, 0, 0, // for all devices
		p_s_compiler_options, p_s_cache_file, 32);
}

CCLUniqueProgram::CCLUniqueProgram(cl_context h_context, const char *p_s_filename,
	TBuildFromFile_Tag UNUSED(t_tag), cl_device_id h_device,
	const char *p_s_compiler_options /*= ""*/, const char *p_s_cache_file /*= "%copykernelname%"*/)
	:m_h_program(0), m_n_compile_flags(0)
{
	m_n_last_result = n_Init(true, h_context, p_s_filename, 1, &h_device,
		p_s_compiler_options, p_s_cache_file, 32);
}

CCLUniqueProgram::CCLUniqueProgram(cl_context h_context, const char *p_s_filename,
	TBuildFromFile_Tag UNUSED(t_tag), size_t n_device_num,
	const cl_device_id *p_device, const char *p_s_compiler_options /*= ""*/,
	const char *p_s_cache_file /*= "%copykernelname%"*/, int n_max_cache_size /*= 32*/)
	:m_h_program(0), m_n_compile_flags(0)
{
	m_n_last_result = n_Init(true, h_context, p_s_filename, n_device_num, p_device,
		p_s_compiler_options, p_s_cache_file, n_max_cache_size);
}

CCLUniqueProgram::~CCLUniqueProgram()
{
	if(m_h_program)
		clReleaseProgram(m_h_program);
}

/*int CCLUniqueProgram::n_Get_Kernel(cl_kernel &r_h_kernel,
	const char *p_s_kernel_name)
{
	cl_int n_result;
	r_h_kernel = clCreateKernel(m_h_program, p_s_kernel_name, &n_result);
	return n_result;
}

cl_kernel CCLUniqueProgram::h_Get_Kernel(const char *p_s_kernel_name)
{
	cl_kernel h_kernel;
	int n_result = n_Get_Kernel(h_kernel, p_s_kernel_name);
	_ASSERTE(n_result == CL_SUCCESS); // this would be programmer's error mostly, otherwise it seldom fails
	return (n_result == CL_SUCCESS)? h_kernel : 0;
}*/

cl_kernel CCLUniqueProgram::h_Get_Kernel(const char *p_s_kernel_name, CLresult &r_n_result)
{
	cl_int n_cl_result;
	cl_kernel h_kernel = clCreateKernel(m_h_program, p_s_kernel_name, &n_cl_result);
	r_n_result = (CLresult)n_cl_result;
	return h_kernel;
}

cl_kernel CCLUniqueProgram::h_Get_Kernel(const char *p_s_kernel_name)
{
	cl_int n_result;
	cl_kernel h_kernel = clCreateKernel(m_h_program, p_s_kernel_name, &n_result);
	_ASSERTE(n_result == CL_SUCCESS); // this would be programmer's error mostly, otherwise it seldom fails
	return (n_result == CL_SUCCESS)? h_kernel : 0;
}

CLresult CCLUniqueProgram::n_Init(bool b_is_filename, cl_context h_context, const char *p_s_filename,
	size_t n_device_num, const cl_device_id *p_device, const char *p_s_compiler_options /*= ""*/,
	const char *p_s_cache_file /*= "%copykernelname%"*/, int n_max_cache_size /*= 32*/)
{
	std::string s_cache_file, s_source_hash;
	if(!strcmp(p_s_cache_file, "%copykernelname%")) {
		TMD5 t_hash;
		if(!b_is_filename) {
			const char *p_s_source_code = p_s_filename;
			CStreamHash<TMD5> hash;
			hash.Process_Data(p_s_source_code, strlen(p_s_source_code) * sizeof(char));
			t_hash = hash.t_Result();
		}
		if(!stl_ut::AssignCStr(s_cache_file, "%temp_default%") ||
		   !stl_ut::AppendCStr(s_cache_file, (b_is_filename)? p_s_filename : t_hash.p_s_ToString()))
			return cl_Out_Of_Host_Memory;
		std::replace(s_cache_file.begin(), s_cache_file.end(), '/', '_');
		std::replace(s_cache_file.begin(), s_cache_file.end(), '\\', '_');
		std::replace(s_cache_file.begin(), s_cache_file.end(), ':', '_');
		p_s_cache_file = s_cache_file.c_str();
	}
	// mangle input filename to cache filename

	CLresult n_result;
	if(b_is_filename) {
		n_result = CCLProgramCompiler::n_CompileProgramFile(h_context, &m_h_program, p_s_filename,
			n_device_num, p_device, p_s_compiler_options, p_s_cache_file, &m_n_compile_flags,
			n_max_cache_size);
	} else {
		const char *p_s_source_code = p_s_filename;
		n_result = CCLProgramCompiler::n_CompileProgram(h_context, &m_h_program, p_s_source_code,
			n_device_num, p_device, p_s_compiler_options, p_s_cache_file, &m_n_compile_flags,
			n_max_cache_size);
	}
	// compile the program

	/*if(b_verbose) 
		CCLProgramCompiler::Dump_StatusWord(m_n_compile_flags);
	// verbose (see what steps did it take)

	if(n_result != CL_SUCCESS && b_stderr_output)
		fprintf(stderr, "error: failed to load OpenCL program (%d)\n", n_result);*/
	// write out

	return n_result;
}

/*
 *								=== ~CCLArgLoaderHelper ===
 */

/*
 *								=== CCLUniqueInstance ===
 */

CCLUniqueInstance::CCLUniqueInstance(int n_device_type /*= CL_DEVICE_TYPE_GPU*/,
	bool b_implementation_profile_selection /*= false*/, bool b_stderr_output /*= true*/)
{
	m_p_device[0] = 0;
	m_n_last_error = n_Init(CCLUtils::TDevice_DefaultScoring(), n_device_type,
		b_implementation_profile_selection, b_stderr_output);
}

CCLUniqueInstance::~CCLUniqueInstance()
{
	// note that both command queue and context now release itself
}

/*
 *								=== ~CCLUniqueInstance ===
 */

#if 0

static void PrintMacros()
{
	for(int i = 1; i < 17; ++ i) {
		printf(
			"/**\n"
			" *	@def clCall3D%d\n"
			" *	@brief sets arguments of a 3D kernel with %d arguments and calls it\n"
			" *\n"
			" *	@param[in] h_cmd_queue is handle to the command queue\n"
			" *	@param[in] h_kernel is handle to the kernel to call\n"
			" *	@param[in] n_work_size_x is global work size in the x dimension\n"
			" *	@param[in] n_work_size_y is global work size in the x dimension\n"
			" *	@param[in] n_work_size_z is global work size in the x dimension\n"
			" *	@param[in] n_block_size_x is thread block size in the x dimension\n"
			" *	@param[in] n_block_size_y is thread block size in the x dimension\n"
			" *	@param[in] n_block_size_z is thread block size in the x dimension\n"
			" *	@param[in] ... are kernel function argument (int, float or cl_mem)\n", i, i);
		printf(
			" */\n"
			"#define clCall3D%d(h_cmd_queue,h_kernel,n_work_size_x,n_block_size_x,...", i);
		printf(
			") \\\n"
			"	CCLKernelCall(clSetKernelArgs%d((h_kernel), __VA_ARGS__", i);
		printf(
			"), \\\n"
			"	(h_cmd_queue), (h_kernel), (n_work_size_x), (n_work_size_y), (n_work_size_z), (n_block_size_x), (n_block_size_y), (n_block_size_z)).n_Result()\n\n");
	}

	for(int i = 1; i < 17; ++ i) {
		printf(
			"/**\n"
			" *	@def clCall2D%d\n"
			" *	@brief sets arguments of a 2D kernel with %d arguments and calls it\n"
			" *\n"
			" *	@param[in] h_cmd_queue is handle to the command queue\n"
			" *	@param[in] h_kernel is handle to the kernel to call\n"
			" *	@param[in] n_work_size_x is global work size in the x dimension\n"
			" *	@param[in] n_work_size_y is global work size in the x dimension\n"
			" *	@param[in] n_block_size_x is thread block size in the x dimension\n"
			" *	@param[in] n_block_size_y is thread block size in the x dimension\n"
			" *	@param[in] ... are kernel function argument (int, float or cl_mem)\n", i, i);
		printf(
			" */\n"
			"#define clCall2D%d(h_cmd_queue,h_kernel,n_work_size_x,n_block_size_x,...", i);
		printf(
			") \\\n"
			"	CCLKernelCall(clSetKernelArgs%d((h_kernel), __VA_ARGS__", i);
		printf(
			"), \\\n"
			"	(h_cmd_queue), (h_kernel), (n_work_size_x), (n_work_size_y), (n_block_size_x), (n_block_size_y)).n_Result()\n\n");
	}

	for(int i = 1; i < 17; ++ i) {
		printf(
			"/**\n"
			" *	@def clCall1D%d\n"
			" *	@brief sets arguments of a 1D kernel with %d arguments and calls it\n"
			" *\n"
			" *	@param[in] h_cmd_queue is handle to the command queue\n"
			" *	@param[in] h_kernel is handle to the kernel to call\n"
			" *	@param[in] n_work_size_x is global work size in the x dimension\n"
			" *	@param[in] n_block_size_x is thread block size in the x dimension\n"
			" *	@param[in] ... are kernel function argument (int, float or cl_mem)\n", i, i);
		printf(
			" */\n"
			"#define clCall1D%d(h_cmd_queue,h_kernel,n_work_size_x,n_block_size_x,...", i);
		printf(
			") \\\n"
			"	CCLKernelCall(clSetKernelArgs%d((h_kernel), __VA_ARGS__", i);
		printf(
			"), \\\n"
			"	(h_cmd_queue), (h_kernel), (n_work_size_x), (n_block_size_x)).n_Result()\n\n");
	}

	for(int i = 1; i < 17; ++ i) {
		printf(
			"/**\n"
			" *	@def clCall3D%d\n"
			" *	@brief sets arguments of a 2D kernel with %d arguments and calls it\n"
			" *\n"
			" *	@param[in] h_cmd_queue is handle to the command queue\n"
			" *	@param[in] h_kernel is handle to the kernel to call\n"
			" *	@param[in] n_work_size_x is global work size in the x dimension\n"
			" *	@param[in] n_work_size_y is global work size in the y dimension\n"
			" *	@param[in] n_work_size_z is global work size in the z dimension\n"
			" *	@param[in] n_block_size_x is thread block size in the x dimension\n"
			" *	@param[in] n_block_size_y is thread block size in the y dimension\n"
			" *	@param[in] n_block_size_z is thread block size in the z dimension\n", i, i);
		for(int j = 0; j < i; ++ j)
			printf(" *	@param[in] %c is kernel function argument (int, float or cl_mem)\n", 'a' + j);
		printf(
			" */\n"
			"#define clCall3D%d(h_cmd_queue,h_kernel,n_work_size_x,n_work_size_y,n_work_size_z,n_block_size_x,n_block_size_y,n_block_size_z", i);
		for(int j = 0; j < i; ++ j)
			printf(",%c", 'a' + j);
		printf(
			") \\\n"
			"	CCLKernelCall(clSetKernelArgs%d((h_kernel)", i);
		for(int j = 0; j < i; ++ j)
			printf(", (%c)", 'a' + j);
		printf(
			"), \\\n"
			"	(h_cmd_queue), (h_kernel), (n_work_size_x), (n_work_size_y), (n_work_size_z), \\\n"
			"	(n_block_size_x), (n_block_size_y), (n_block_size_z)).n_Result()\n\n");
	}

	for(int i = 1; i < 17; ++ i) {
		printf(
			"/**\n"
			" *	@def clCall2D%d\n"
			" *	@brief sets arguments of a 2D kernel with %d arguments and calls it\n"
			" *\n"
			" *	@param[in] h_cmd_queue is handle to the command queue\n"
			" *	@param[in] h_kernel is handle to the kernel to call\n"
			" *	@param[in] n_work_size_x is global work size in the x dimension\n"
			" *	@param[in] n_work_size_y is global work size in the y dimension\n"
			" *	@param[in] n_block_size_x is thread block size in the x dimension\n"
			" *	@param[in] n_block_size_y is thread block size in the y dimension\n", i, i);
		for(int j = 0; j < i; ++ j)
			printf(" *	@param[in] %c is kernel function argument (int, float or cl_mem)\n", 'a' + j);
		printf(
			" */\n"
			"#define clCall2D%d(h_cmd_queue,h_kernel,n_work_size_x,n_work_size_y,n_block_size_x,n_block_size_y", i);
		for(int j = 0; j < i; ++ j)
			printf(",%c", 'a' + j);
		printf(
			") \\\n"
			"	CCLKernelCall(clSetKernelArgs%d((h_kernel)", i);
		for(int j = 0; j < i; ++ j)
			printf(", (%c)", 'a' + j);
		printf(
			"), \\\n"
			"	(h_cmd_queue), (h_kernel), (n_work_size_x), (n_work_size_y), (n_block_size_x), (n_block_size_y)).n_Result()\n\n");
	}

	for(int i = 1; i < 17; ++ i) {
		printf(
			"/**\n"
			" *	@def clCall1D%d\n"
			" *	@brief sets arguments of a 1D kernel with %d arguments and calls it\n"
			" *\n"
			" *	@param[in] h_cmd_queue is handle to the command queue\n"
			" *	@param[in] h_kernel is handle to the kernel to call\n"
			" *	@param[in] n_work_size_x is global work size in the x dimension\n"
			" *	@param[in] n_block_size_x is thread block size in the x dimension\n", i, i);
		for(int j = 0; j < i; ++ j)
			printf(" *	@param[in] %c is kernel function argument (int, float or cl_mem)\n", 'a' + j);
		printf(
			" */\n"
			"#define clCall1D%d(h_cmd_queue,h_kernel,n_work_size_x,n_block_size_x", i);
		for(int j = 0; j < i; ++ j)
			printf(",%c", 'a' + j);
		printf(
			") \\\n"
			"	CCLKernelCall(clSetKernelArgs%d((h_kernel)", i);
		for(int j = 0; j < i; ++ j)
			printf(", (%c)", 'a' + j);
		printf(
			"), \\\n"
			"	(h_cmd_queue), (h_kernel), (n_work_size_x), (n_block_size_x)).n_Result()\n\n");
	}
}

#endif // 0
