cl_int WINAPI wine_clRetainProgram(cl_program program)
{
cl_int ret;
TRACE("\n");
ret = clRetainProgram(program);
return ret;
}
cl_kernel
__clCloneKernel (cl_kernel kernel)
{
_cl_kernel *original_kernel_data = (_cl_kernel *) kernel;
_cl_kernel *kernel_data = NULL;
cl_uint i;
if (!original_kernel_data)
{
#ifdef OCL_DEBUG_MESSAGES
__clDebugPrint (CL_LOG_ERROR, "__clCloneKernel: invalid parameter\n");
#endif // #ifdef OCL_DEBUG_MESSAGES
return 0;
}
kernel_data = (_cl_kernel *) __clAlloc (sizeof (_cl_kernel));
if (!kernel_data)
{
#ifdef OCL_DEBUG_MESSAGES
__clDebugPrint (CL_LOG_ERROR, "__clCloneKernel: out of memory\n");
#endif // #ifdef OCL_DEBUG_MESSAGES
return (cl_kernel) 0;
}
// copy all the rest from original kernel
memcpy (kernel_data, original_kernel_data, sizeof (_cl_kernel));
kernel_data->header.ref_count = 1;
if (kernel_data->program_data)
{
// clone will reference same program as original, so increase refcount
clRetainProgram ((cl_program) kernel_data->program_data);
}
for (i = 0; i < MAX_KERNEL_ARGS; i++)
{
kernel_data->arg[i].buffer = original_kernel_data->arg[i].buffer;
if (kernel_data->arg[i].buffer)
{
// clone will reference same argument buffers, so increase refcount
clRetainMemObject (kernel_data->arg[i].buffer);
}
if (original_kernel_data->arg[i].data)
{
// make a copy of a argument data
kernel_data->arg[i].data =
__clAlloc (original_kernel_data->arg[i].datalen);
memcpy (kernel_data->arg[i].data, original_kernel_data->arg[i].data,
original_kernel_data->arg[i].datalen);
kernel_data->arg[i].datalen = original_kernel_data->arg[i].datalen;
}
}
__clListAddNode (&g_kernel_data_root, kernel_data);
return (cl_kernel) kernel_data;
}
//A callback function which displays the log of the program build
void programCallback(cl_program program, void* user_data)
{
clRetainProgram(program);
cl::Program p(program);
cl::Device dev_id = p.getInfo<CL_PROGRAM_DEVICES>()[0];
std::string device_name = dev_id.getInfo<CL_DEVICE_NAME>();
std::string log = p.getBuildInfo<CL_PROGRAM_BUILD_LOG>(dev_id);
std::cout << device_name << " Log: " << log << std::endl;
}
void cb(cl_program p,void* data)
{
clRetainProgram(p);
cl_device_id devid[1];
clGetProgramInfo(p,CL_PROGRAM_DEVICES,sizeof(cl_device_id),(void*)devid,NULL);
char bug[65536];
clGetProgramBuildInfo(p,devid[0],CL_PROGRAM_BUILD_LOG,65536*sizeof(char),bug,NULL);
clReleaseProgram(p);
LOGE("Build log \n %s\n",bug);
}
Program& Program::operator=(const Program& other)
{
_ctx = other._ctx;
_options = other._options;
_built = other._built;
if(other._id)
clRetainProgram(other._id);
if(_id)
clReleaseProgram(_id);
_id = other._id;
return *this;
}
Program::Program(const Program& other)
: _ctx(other._ctx), _id(other._id), _options(other._options) ,_built(other._built)
{
if(_id)
clRetainProgram(_id);
}
cl_int
clEnqueueFillBuffer(cl_command_queue command_queue,
cl_mem buffer,
const void *pattern,
size_t pattern_size,
size_t offset,
size_t size,
cl_uint num_events_in_wait_list,
const cl_event *event_wait_list,
cl_event *event)
{
static pthread_mutex_t lock = PTHREAD_MUTEX_INITIALIZER;
static cl_program last_program = NULL;
static cl_context last_context = NULL;
cl_kernel kernel;
cl_program program;
cl_context context;
cl_device_id device;
size_t gworksz;
size_t lworksz;
char kernel_name[80];
cl_int rc;
union {
cl_char v_char;
cl_short v_short;
cl_int v_int;
cl_long v_long;
} pattern_value;
cl_uint pattern_nums;
switch (pattern_size)
{
case sizeof(cl_char):
pattern_value.v_char = *((cl_char *)pattern);
break;
case sizeof(cl_short):
pattern_value.v_short = *((cl_short *)pattern);
break;
case sizeof(cl_int):
pattern_value.v_int = *((cl_int *)pattern);
break;
case sizeof(cl_long):
pattern_value.v_long = *((cl_long *)pattern);
break;
default:
/*
* pattern_size was not support one, even though OpenCL 1.2
* spec says 16, 32, 64 or 128 bytes patterns are supported.
*/
return CL_INVALID_VALUE;
}
/* ensure alignment */
if (offset % pattern_size != 0)
return CL_INVALID_VALUE;
if (size % pattern_size != 0)
return CL_INVALID_VALUE;
/* fetch context and device_id associated with this command queue */
rc = clGetCommandQueueInfo(command_queue,
CL_QUEUE_CONTEXT,
sizeof(cl_context),
&context,
NULL);
if (rc != CL_SUCCESS)
return rc;
pthread_mutex_lock(&lock);
if (last_program && last_context == context)
{
rc = clRetainProgram(last_program);
if (rc != CL_SUCCESS)
goto out_unlock;
program = last_program;
}
else
{
char source[10240];
const char *prog_source[1];
size_t prog_length[1];
cl_uint num_devices;
cl_device_id *device_ids;
static struct {
const char *type_name;
size_t type_size;
} pattern_types[] = {
{ "char", sizeof(cl_char) },
{ "short", sizeof(cl_short) },
{ "int", sizeof(cl_int) },
{ "long", sizeof(cl_long) },
};
size_t i, ofs;
/* fetch properties of cl_context */
rc = clGetContextInfo(context,
CL_CONTEXT_NUM_DEVICES,
sizeof(cl_uint),
&num_devices,
NULL);
if (rc != CL_SUCCESS)
goto out_unlock;
Assert(num_devices > 0);
device_ids = calloc(num_devices, sizeof(cl_device_id));
if (!device_ids)
{
rc = CL_OUT_OF_HOST_MEMORY;
goto out_unlock;
}
rc = clGetContextInfo(context,
CL_CONTEXT_DEVICES,
sizeof(cl_device_id) * num_devices,
device_ids,
NULL);
if (rc != CL_SUCCESS)
{
free(device_ids);
goto out_unlock;
}
/* release the previous program */
if (last_program)
{
rc = clReleaseProgram(last_program);
Assert(rc == CL_SUCCESS);
last_program = NULL;
last_context = NULL;
}
/* create a program object */
for (i=0, ofs=0; i < lengthof(pattern_types); i++)
{
ofs += snprintf(
source + ofs, sizeof(source) - ofs,
"__kernel void\n"
"pgstromEnqueueFillBuffer_%zu(__global %s *buffer,\n"
" %s value, uint nums)\n"
"{\n"
" if (get_global_id(0) >= nums)\n"
" return;\n"
" buffer[get_global_id(0)] = value;\n"
"}\n",
pattern_types[i].type_size,
pattern_types[i].type_name,
pattern_types[i].type_name);
}
prog_source[0] = source;
prog_length[0] = ofs;
program = clCreateProgramWithSource(context,
1,
prog_source,
prog_length,
&rc);
if (rc != CL_SUCCESS)
{
free(device_ids);
goto out_unlock;
}
/* build this program object */
rc = clBuildProgram(program,
num_devices,
device_ids,
NULL,
NULL,
NULL);
free(device_ids);
if (rc != CL_SUCCESS)
{
clReleaseProgram(program);
goto out_unlock;
}
/* acquire the program object */
rc = clRetainProgram(program);
if (rc != CL_SUCCESS)
{
clReleaseProgram(program);
goto out_unlock;
}
last_program = program;
last_context = context;
}
pthread_mutex_unlock(&lock);
Assert(program != NULL);
/* fetch a device id of this command queue */
rc = clGetCommandQueueInfo(command_queue,
CL_QUEUE_DEVICE,
sizeof(cl_device_id),
&device,
NULL);
if (rc != CL_SUCCESS)
goto out_release_program;
/* fetch a kernel object to be called */
snprintf(kernel_name, sizeof(kernel_name),
"pgstromEnqueueFillBuffer_%zu", pattern_size);
kernel = clCreateKernel(program,
kernel_name,
&rc);
if (rc != CL_SUCCESS)
goto out_release_program;
/* 1st arg: __global <typename> *buffer */
rc = clSetKernelArg(kernel, 0, sizeof(cl_mem), &buffer);
if (rc != CL_SUCCESS)
goto out_release_kernel;
/* 2nd arg: <typename> value */
rc = clSetKernelArg(kernel, 1, pattern_size, &pattern_value);
if (rc != CL_SUCCESS)
goto out_release_kernel;
/* 3rd arg: size_t nums */
pattern_nums = (offset + size) / pattern_size;
rc = clSetKernelArg(kernel, 2, sizeof(cl_uint), &pattern_nums);
if (rc != CL_SUCCESS)
goto out_release_kernel;
/* calculate optimal workgroup size */
rc = clGetKernelWorkGroupInfo(kernel,
device,
CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE,
sizeof(size_t),
&lworksz,
NULL);
Assert((lworksz & (lworksz - 1)) == 0);
gworksz = ((size / pattern_size + lworksz - 1) / lworksz) * lworksz;
/* enqueue a kernel, instead of clEnqueueFillBuffer */
offset /= pattern_size;
rc = clEnqueueNDRangeKernel(command_queue,
kernel,
1,
&offset,
&gworksz,
&lworksz,
num_events_in_wait_list,
event_wait_list,
event);
if (rc != CL_SUCCESS)
goto out_release_kernel;
rc = clSetEventCallback(*event,
CL_COMPLETE,
pgstromEnqueueFillBufferCleanup,
kernel);
if (rc != CL_SUCCESS)
{
clWaitForEvents(1, event);
goto out_release_kernel;
}
return CL_SUCCESS;
out_unlock:
pthread_mutex_unlock(&lock);
return rc;
out_release_kernel:
clReleaseKernel(kernel);
out_release_program:
clReleaseProgram(program);
return rc;
}
static void inc(cl_program & something)
{
cl_int err = clRetainProgram(something);
VIENNACL_ERR_CHECK(err);
}
cl_kernel
clCreateKernel (cl_program program,
const char *kernel_name, cl_int * errcode_ret)
{
_cl_program *program_data = (_cl_program *) (program);
_cl_kernel *kernel_data = NULL;
int i;
cl_kernel_function_info *info = NULL;
_cl_program_binary *program_binary_data = NULL;
char *error = NULL;
if (!program_data)
{
#ifdef OCL_DEBUG_MESSAGES
__clDebugPrint (CL_LOG_ERROR,
"clCreateKernel: Invalid program parameter (program=0x%08x)\n",
(int) program);
#endif // #ifdef OCL_DEBUG_MESSAGES
if (errcode_ret)
{
*errcode_ret = CL_INVALID_PROGRAM;
}
return 0;
}
if (program_data->num_binaries == 0)
{
// program is not build
#ifdef OCL_DEBUG_MESSAGES
__clDebugPrint (CL_LOG_ERROR,
"clCreateKernel: Program is not build (program=0x%08x)\n",
(int) program);
#endif // #ifdef OCL_DEBUG_MESSAGES
if (errcode_ret)
{
*errcode_ret = CL_INVALID_PROGRAM;
}
return 0;
}
kernel_data = (_cl_kernel *) __clAlloc (sizeof (_cl_kernel));
if (!kernel_data)
{
#ifdef OCL_DEBUG_MESSAGES
__clDebugPrint (CL_LOG_ERROR, "clCreateKernel: out of memory\n");
#endif // #ifdef OCL_DEBUG_MESSAGES
if (errcode_ret)
{
*errcode_ret = CL_OUT_OF_HOST_MEMORY;
}
return (cl_kernel) 0;
}
memset (kernel_data, 0, sizeof (_cl_kernel));
kernel_data->header.ref_count = 1;
kernel_data->header.context = program_data->header.context;
for (i = 0; i < program_data->num_kernel_functions; i++)
{
info = &program_data->kernel_function_info[i];
if (strcmp (kernel_name, info->function_name) == 0)
{
// found the matching kernel function name
kernel_data->kernel_function_info = info;
#ifdef OCL_DEBUG_MESSAGES
__clDebugPrint (CL_LOG_DEBUG,
"clCreateKernel: Kernel '%s' found. Has %d parameters\n",
kernel_name,
kernel_data->kernel_function_info->num_parameters);
#endif // #ifdef OCL_DEBUG_MESSAGES
break;
}
}
if (!kernel_data->kernel_function_info)
{
#ifdef OCL_DEBUG_MESSAGES
__clDebugPrint (CL_LOG_ERROR,
"clCreateKernel: kernel function '%s' not found in program (0x%08x)\n",
kernel_name, (int) program_data);
#endif // #ifdef OCL_DEBUG_MESSAGES
if (errcode_ret)
{
*errcode_ret = CL_INVALID_VALUE;
}
return (cl_kernel) 0;
}
program_binary_data = &program_data->binaries[0];
if (program_binary_data->lib)
{
// find the function from dynamic library
kernel_data->entry_fn = dlsym (program_binary_data->lib, kernel_name);
if ((error = dlerror ()) != NULL)
{
#ifdef OCL_DEBUG_MESSAGES
__clDebugPrint (CL_LOG_ERROR,
"clCreateKernel: Kernel entry point function not found (dlerror=%s)\n",
error);
#endif // #ifdef OCL_DEBUG_MESSAGES
if (errcode_ret)
{
*errcode_ret = CL_INVALID_KERNEL_NAME;
}
clReleaseKernel ((cl_kernel) kernel_data);
return (cl_kernel) 0;
}
}
else
{
kernel_data->entry_fn = NULL;
}
memset (kernel_data->arg, 0, sizeof (CL_KERNEL_ARG_DATA) * MAX_KERNEL_ARGS);
kernel_data->program_data = program_data;
clRetainProgram (program);
__clListAddNode (&g_kernel_data_root, kernel_data);
if (errcode_ret)
{
*errcode_ret = CL_SUCCESS;
}
return (cl_kernel) kernel_data;
}
_clState *initCl(unsigned int gpu, char *name, size_t nameSize)
{
_clState *clState = calloc(1, sizeof(_clState));
bool patchbfi = false, prog_built = false;
cl_platform_id platform = NULL;
cl_platform_id* platforms;
cl_device_id *devices;
cl_uint numPlatforms;
cl_uint numDevices;
char pbuff[256];
cl_int status;
status = clGetPlatformIDs(0, NULL, &numPlatforms);
if (status != CL_SUCCESS) {
applog(LOG_ERR, "Error: Getting Platforms. (clGetPlatformsIDs)");
return NULL;
}
platforms = (cl_platform_id *)alloca(numPlatforms*sizeof(cl_platform_id));
status = clGetPlatformIDs(numPlatforms, platforms, NULL);
if (status != CL_SUCCESS) {
applog(LOG_ERR, "Error: Getting Platform Ids. (clGetPlatformsIDs)");
return NULL;
}
if (opt_platform_id >= numPlatforms) {
applog(LOG_ERR, "Specified platform that does not exist");
return NULL;
}
status = clGetPlatformInfo(platforms[opt_platform_id], CL_PLATFORM_VENDOR, sizeof(pbuff), pbuff, NULL);
if (status != CL_SUCCESS) {
applog(LOG_ERR, "Error: Getting Platform Info. (clGetPlatformInfo)");
return NULL;
}
platform = platforms[opt_platform_id];
if (platform == NULL) {
perror("NULL platform found!\n");
return NULL;
}
applog(LOG_INFO, "CL Platform vendor: %s", pbuff);
status = clGetPlatformInfo(platform, CL_PLATFORM_NAME, sizeof(pbuff), pbuff, NULL);
if (status == CL_SUCCESS)
applog(LOG_INFO, "CL Platform name: %s", pbuff);
status = clGetPlatformInfo(platform, CL_PLATFORM_VERSION, sizeof(pbuff), pbuff, NULL);
if (status == CL_SUCCESS)
applog(LOG_INFO, "CL Platform version: %s", pbuff);
status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 0, NULL, &numDevices);
if (status != CL_SUCCESS) {
applog(LOG_ERR, "Error: Getting Device IDs (num)");
return NULL;
}
if (numDevices > 0 ) {
devices = (cl_device_id *)malloc(numDevices*sizeof(cl_device_id));
/* Now, get the device list data */
status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, numDevices, devices, NULL);
if (status != CL_SUCCESS) {
applog(LOG_ERR, "Error: Getting Device IDs (list)");
return NULL;
}
applog(LOG_INFO, "List of devices:");
unsigned int i;
for (i = 0; i < numDevices; i++) {
status = clGetDeviceInfo(devices[i], CL_DEVICE_NAME, sizeof(pbuff), pbuff, NULL);
if (status != CL_SUCCESS) {
applog(LOG_ERR, "Error: Getting Device Info");
return NULL;
}
applog(LOG_INFO, "\t%i\t%s", i, pbuff);
}
if (gpu < numDevices) {
status = clGetDeviceInfo(devices[gpu], CL_DEVICE_NAME, sizeof(pbuff), pbuff, NULL);
if (status != CL_SUCCESS) {
applog(LOG_ERR, "Error: Getting Device Info");
return NULL;
}
applog(LOG_INFO, "Selected %i: %s", gpu, pbuff);
strncpy(name, pbuff, nameSize);
} else {
applog(LOG_ERR, "Invalid GPU %i", gpu);
return NULL;
}
} else return NULL;
cl_context_properties cps[3] = { CL_CONTEXT_PLATFORM, (cl_context_properties)platform, 0 };
clState->context = clCreateContextFromType(cps, CL_DEVICE_TYPE_GPU, NULL, NULL, &status);
if (status != CL_SUCCESS) {
applog(LOG_ERR, "Error: Creating Context. (clCreateContextFromType)");
return NULL;
}
/* Check for BFI INT support. Hopefully people don't mix devices with
* and without it! */
char * extensions = malloc(1024);
const char * camo = "cl_amd_media_ops";
char *find;
status = clGetDeviceInfo(devices[gpu], CL_DEVICE_EXTENSIONS, 1024, (void *)extensions, NULL);
if (status != CL_SUCCESS) {
applog(LOG_ERR, "Error: Failed to clGetDeviceInfo when trying to get CL_DEVICE_EXTENSIONS");
return NULL;
}
find = strstr(extensions, camo);
if (find)
clState->hasBitAlign = true;
status = clGetDeviceInfo(devices[gpu], CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT, sizeof(cl_uint), (void *)&clState->preferred_vwidth, NULL);
if (status != CL_SUCCESS) {
applog(LOG_ERR, "Error: Failed to clGetDeviceInfo when trying to get CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT");
return NULL;
}
if (opt_debug)
applog(LOG_DEBUG, "Preferred vector width reported %d", clState->preferred_vwidth);
status = clGetDeviceInfo(devices[gpu], CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(size_t), (void *)&clState->max_work_size, NULL);
if (status != CL_SUCCESS) {
applog(LOG_ERR, "Error: Failed to clGetDeviceInfo when trying to get CL_DEVICE_MAX_WORK_GROUP_SIZE");
return NULL;
}
if (opt_debug)
applog(LOG_DEBUG, "Max work group size reported %d", clState->max_work_size);
/* For some reason 2 vectors is still better even if the card says
* otherwise, and many cards lie about their max so use 256 as max
* unless explicitly set on the command line. 79x0 cards perform
* better without vectors */
if (clState->preferred_vwidth > 1) {
if (strstr(name, "Tahiti"))
clState->preferred_vwidth = 1;
else
clState->preferred_vwidth = 2;
}
if (opt_vectors)
clState->preferred_vwidth = opt_vectors;
if (opt_worksize && opt_worksize <= clState->max_work_size)
clState->work_size = opt_worksize;
else
clState->work_size = (clState->max_work_size <= 256 ? clState->max_work_size : 256) /
clState->preferred_vwidth;
/* Create binary filename based on parameters passed to opencl
* compiler to ensure we only load a binary that matches what would
* have otherwise created. The filename is:
* name + kernelname +/i bitalign + v + vectors + w + work_size + sizeof(long) + .bin
*/
char binaryfilename[255];
char numbuf[10];
char filename[16];
if (chosen_kernel == KL_NONE) {
if (!clState->hasBitAlign || strstr(name, "Tahiti"))
chosen_kernel = KL_POCLBM;
else
chosen_kernel = KL_PHATK;
}
switch (chosen_kernel) {
case KL_POCLBM:
strcpy(filename, "poclbm120203.cl");
strcpy(binaryfilename, "poclbm120203");
break;
case KL_NONE: /* Shouldn't happen */
case KL_PHATK:
strcpy(filename, "phatk120203.cl");
strcpy(binaryfilename, "phatk120203");
break;
}
FILE *binaryfile;
size_t *binary_sizes;
char **binaries;
int pl;
char *source = file_contents(filename, &pl);
size_t sourceSize[] = {(size_t)pl};
if (!source)
return NULL;
binary_sizes = (size_t *)malloc(sizeof(size_t)*numDevices);
if (unlikely(!binary_sizes)) {
applog(LOG_ERR, "Unable to malloc binary_sizes");
return NULL;
}
binaries = (char **)malloc(sizeof(char *)*numDevices);
if (unlikely(!binaries)) {
applog(LOG_ERR, "Unable to malloc binaries");
return NULL;
}
strcat(binaryfilename, name);
if (clState->hasBitAlign)
strcat(binaryfilename, "bitalign");
strcat(binaryfilename, "v");
sprintf(numbuf, "%d", clState->preferred_vwidth);
strcat(binaryfilename, numbuf);
strcat(binaryfilename, "w");
sprintf(numbuf, "%d", (int)clState->work_size);
strcat(binaryfilename, numbuf);
strcat(binaryfilename, "long");
sprintf(numbuf, "%d", (int)sizeof(long));
strcat(binaryfilename, numbuf);
strcat(binaryfilename, ".bin");
binaryfile = fopen(binaryfilename, "rb");
if (!binaryfile) {
if (opt_debug)
applog(LOG_DEBUG, "No binary found, generating from source");
} else {
struct stat binary_stat;
if (unlikely(stat(binaryfilename, &binary_stat))) {
if (opt_debug)
applog(LOG_DEBUG, "Unable to stat binary, generating from source");
fclose(binaryfile);
goto build;
}
if (!binary_stat.st_size)
goto build;
binary_sizes[gpu] = binary_stat.st_size;
binaries[gpu] = (char *)malloc(binary_sizes[gpu]);
if (unlikely(!binaries[gpu])) {
applog(LOG_ERR, "Unable to malloc binaries");
fclose(binaryfile);
return NULL;
}
if (fread(binaries[gpu], 1, binary_sizes[gpu], binaryfile) != binary_sizes[gpu]) {
applog(LOG_ERR, "Unable to fread binaries[gpu]");
fclose(binaryfile);
free(binaries[gpu]);
goto build;
}
clState->program = clCreateProgramWithBinary(clState->context, 1, &devices[gpu], &binary_sizes[gpu], (const unsigned char **)&binaries[gpu], &status, NULL);
if (status != CL_SUCCESS) {
applog(LOG_ERR, "Error: Loading Binary into cl_program (clCreateProgramWithBinary)");
fclose(binaryfile);
free(binaries[gpu]);
goto build;
}
fclose(binaryfile);
if (opt_debug)
applog(LOG_DEBUG, "Loaded binary image %s", binaryfilename);
/* We don't need to patch this already loaded image, but need to
* set the flag for status later */
if (clState->hasBitAlign)
patchbfi = true;
free(binaries[gpu]);
goto built;
}
/////////////////////////////////////////////////////////////////
// Load CL file, build CL program object, create CL kernel object
/////////////////////////////////////////////////////////////////
build:
clState->program = clCreateProgramWithSource(clState->context, 1, (const char **)&source, sourceSize, &status);
if (status != CL_SUCCESS) {
applog(LOG_ERR, "Error: Loading Binary into cl_program (clCreateProgramWithSource)");
return NULL;
}
clRetainProgram(clState->program);
if (status != CL_SUCCESS) {
applog(LOG_ERR, "Error: Retaining Program (clRetainProgram)");
return NULL;
}
/* create a cl program executable for all the devices specified */
char *CompilerOptions = calloc(1, 256);
sprintf(CompilerOptions, "-D WORKSIZE=%d -D VECTORS%d",
(int)clState->work_size, clState->preferred_vwidth);
if (opt_debug)
applog(LOG_DEBUG, "Setting worksize to %d", clState->work_size);
if (clState->preferred_vwidth > 1 && opt_debug)
applog(LOG_DEBUG, "Patched source to suit %d vectors", clState->preferred_vwidth);
if (clState->hasBitAlign) {
strcat(CompilerOptions, " -D BITALIGN");
if (opt_debug)
applog(LOG_DEBUG, "cl_amd_media_ops found, setting BITALIGN");
if (strstr(name, "Cedar") ||
strstr(name, "Redwood") ||
strstr(name, "Juniper") ||
strstr(name, "Cypress" ) ||
strstr(name, "Hemlock" ) ||
strstr(name, "Caicos" ) ||
strstr(name, "Turks" ) ||
strstr(name, "Barts" ) ||
strstr(name, "Cayman" ) ||
strstr(name, "Antilles" ) ||
strstr(name, "Wrestler" ) ||
strstr(name, "Zacate" ) ||
strstr(name, "WinterPark" ) ||
strstr(name, "BeaverCreek" ))
patchbfi = true;
} else if (opt_debug)
applog(LOG_DEBUG, "cl_amd_media_ops not found, will not set BITALIGN");
if (patchbfi) {
strcat(CompilerOptions, " -D BFI_INT");
if (opt_debug)
applog(LOG_DEBUG, "BFI_INT patch requiring device found, patched source with BFI_INT");
} else if (opt_debug)
applog(LOG_DEBUG, "BFI_INT patch requiring device not found, will not BFI_INT patch");
if (opt_debug)
applog(LOG_DEBUG, "CompilerOptions: %s", CompilerOptions);
status = clBuildProgram(clState->program, 1, &devices[gpu], CompilerOptions , NULL, NULL);
free(CompilerOptions);
if (status != CL_SUCCESS) {
applog(LOG_ERR, "Error: Building Program (clBuildProgram)");
size_t logSize;
status = clGetProgramBuildInfo(clState->program, devices[gpu], CL_PROGRAM_BUILD_LOG, 0, NULL, &logSize);
char *log = malloc(logSize);
status = clGetProgramBuildInfo(clState->program, devices[gpu], CL_PROGRAM_BUILD_LOG, logSize, log, NULL);
applog(LOG_INFO, "%s", log);
return NULL;
}
prog_built = true;
status = clGetProgramInfo( clState->program, CL_PROGRAM_BINARY_SIZES, sizeof(size_t)*numDevices, binary_sizes, NULL );
if (unlikely(status != CL_SUCCESS)) {
applog(LOG_ERR, "Error: Getting program info CL_PROGRAM_BINARY_SIZES. (clGetPlatformInfo)");
return NULL;
}
/* copy over all of the generated binaries. */
if (opt_debug)
applog(LOG_DEBUG, "binary size %d : %d", gpu, binary_sizes[gpu]);
if (!binary_sizes[gpu]) {
applog(LOG_ERR, "OpenCL compiler generated a zero sized binary, may need to reboot!");
return NULL;
}
binaries[gpu] = (char *)malloc( sizeof(char)*binary_sizes[gpu]);
status = clGetProgramInfo( clState->program, CL_PROGRAM_BINARIES, sizeof(char *)*numDevices, binaries, NULL );
if (unlikely(status != CL_SUCCESS)) {
applog(LOG_ERR, "Error: Getting program info. (clGetPlatformInfo)");
return NULL;
}
/* Patch the kernel if the hardware supports BFI_INT but it needs to
* be hacked in */
if (patchbfi) {
unsigned remaining = binary_sizes[gpu];
char *w = binaries[gpu];
unsigned int start, length;
/* Find 2nd incidence of .text, and copy the program's
* position and length at a fixed offset from that. Then go
* back and find the 2nd incidence of \x7ELF (rewind by one
* from ELF) and then patch the opcocdes */
if (!advance(&w, &remaining, ".text"))
{patchbfi = 0; goto build;}
w++; remaining--;
if (!advance(&w, &remaining, ".text")) {
/* 32 bit builds only one ELF */
w--; remaining++;
}
memcpy(&start, w + 285, 4);
memcpy(&length, w + 289, 4);
w = binaries[gpu]; remaining = binary_sizes[gpu];
if (!advance(&w, &remaining, "ELF"))
{patchbfi = 0; goto build;}
w++; remaining--;
if (!advance(&w, &remaining, "ELF")) {
/* 32 bit builds only one ELF */
w--; remaining++;
}
w--; remaining++;
w += start; remaining -= start;
if (opt_debug)
applog(LOG_DEBUG, "At %p (%u rem. bytes), to begin patching",
w, remaining);
patch_opcodes(w, length);
status = clReleaseProgram(clState->program);
if (status != CL_SUCCESS) {
applog(LOG_ERR, "Error: Releasing program. (clReleaseProgram)");
return NULL;
}
clState->program = clCreateProgramWithBinary(clState->context, 1, &devices[gpu], &binary_sizes[gpu], (const unsigned char **)&binaries[gpu], &status, NULL);
if (status != CL_SUCCESS) {
applog(LOG_ERR, "Error: Loading Binary into cl_program (clCreateProgramWithBinary)");
return NULL;
}
clRetainProgram(clState->program);
if (status != CL_SUCCESS) {
applog(LOG_ERR, "Error: Retaining Program (clRetainProgram)");
return NULL;
}
/* Program needs to be rebuilt */
prog_built = false;
}
free(source);
/* Save the binary to be loaded next time */
binaryfile = fopen(binaryfilename, "wb");
if (!binaryfile) {
/* Not a fatal problem, just means we build it again next time */
if (opt_debug)
applog(LOG_DEBUG, "Unable to create file %s", binaryfilename);
} else {
if (unlikely(fwrite(binaries[gpu], 1, binary_sizes[gpu], binaryfile) != binary_sizes[gpu])) {
applog(LOG_ERR, "Unable to fwrite to binaryfile");
return NULL;
}
fclose(binaryfile);
}
if (binaries[gpu])
free(binaries[gpu]);
built:
free(binaries);
free(binary_sizes);
applog(LOG_INFO, "Initialising kernel %s with%s BFI_INT, %d vectors and worksize %d",
filename, patchbfi ? "" : "out", clState->preferred_vwidth, clState->work_size);
if (!prog_built) {
/* create a cl program executable for all the devices specified */
status = clBuildProgram(clState->program, 1, &devices[gpu], NULL, NULL, NULL);
if (status != CL_SUCCESS) {
applog(LOG_ERR, "Error: Building Program (clBuildProgram)");
size_t logSize;
status = clGetProgramBuildInfo(clState->program, devices[gpu], CL_PROGRAM_BUILD_LOG, 0, NULL, &logSize);
char *log = malloc(logSize);
status = clGetProgramBuildInfo(clState->program, devices[gpu], CL_PROGRAM_BUILD_LOG, logSize, log, NULL);
applog(LOG_INFO, "%s", log);
return NULL;
}
clRetainProgram(clState->program);
if (status != CL_SUCCESS) {
applog(LOG_ERR, "Error: Retaining Program (clRetainProgram)");
return NULL;
}
}
/* get a kernel object handle for a kernel with the given name */
clState->kernel = clCreateKernel(clState->program, "search", &status);
if (status != CL_SUCCESS) {
applog(LOG_ERR, "Error: Creating Kernel from program. (clCreateKernel)");
return NULL;
}
/////////////////////////////////////////////////////////////////
// Create an OpenCL command queue
/////////////////////////////////////////////////////////////////
clState->commandQueue = clCreateCommandQueue(clState->context, devices[gpu],
CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &status);
if (status != CL_SUCCESS) /* Try again without OOE enable */
clState->commandQueue = clCreateCommandQueue(clState->context, devices[gpu], 0 , &status);
if (status != CL_SUCCESS) {
applog(LOG_ERR, "Creating Command Queue. (clCreateCommandQueue)");
return NULL;
}
clState->outputBuffer = clCreateBuffer(clState->context, CL_MEM_READ_WRITE, BUFFERSIZE, NULL, &status);
if (status != CL_SUCCESS) {
applog(LOG_ERR, "Error: clCreateBuffer (outputBuffer)");
return NULL;
}
return clState;
}
program( const program& rhs )
{
prog = rhs.prog;
clRetainProgram(prog);
}