Skip to content

Instantly share code, notes, and snippets.

@fjarri

fjarri/test.c Secret

Created Feb 10, 2017
Embed
What would you like to do?
#include <OpenCL/opencl.h>
#include <stdio.h>
#include <string.h>
#include <stdlib.h>
cl_int err;
const char* errStr(cl_int error)
{
static const char* errorString[] = {
"CL_SUCCESS",
"CL_DEVICE_NOT_FOUND",
"CL_DEVICE_NOT_AVAILABLE",
"CL_COMPILER_NOT_AVAILABLE",
"CL_MEM_OBJECT_ALLOCATION_FAILURE",
"CL_OUT_OF_RESOURCES",
"CL_OUT_OF_HOST_MEMORY",
"CL_PROFILING_INFO_NOT_AVAILABLE",
"CL_MEM_COPY_OVERLAP",
"CL_IMAGE_FORMAT_MISMATCH",
"CL_IMAGE_FORMAT_NOT_SUPPORTED",
"CL_BUILD_PROGRAM_FAILURE",
"CL_MAP_FAILURE",
"",
"",
"",
"",
"",
"",
"",
"",
"",
"",
"",
"",
"",
"",
"",
"",
"",
"CL_INVALID_VALUE",
"CL_INVALID_DEVICE_TYPE",
"CL_INVALID_PLATFORM",
"CL_INVALID_DEVICE",
"CL_INVALID_CONTEXT",
"CL_INVALID_QUEUE_PROPERTIES",
"CL_INVALID_COMMAND_QUEUE",
"CL_INVALID_HOST_PTR",
"CL_INVALID_MEM_OBJECT",
"CL_INVALID_IMAGE_FORMAT_DESCRIPTOR",
"CL_INVALID_IMAGE_SIZE",
"CL_INVALID_SAMPLER",
"CL_INVALID_BINARY",
"CL_INVALID_BUILD_OPTIONS",
"CL_INVALID_PROGRAM",
"CL_INVALID_PROGRAM_EXECUTABLE",
"CL_INVALID_KERNEL_NAME",
"CL_INVALID_KERNEL_DEFINITION",
"CL_INVALID_KERNEL",
"CL_INVALID_ARG_INDEX",
"CL_INVALID_ARG_VALUE",
"CL_INVALID_ARG_SIZE",
"CL_INVALID_KERNEL_ARGS",
"CL_INVALID_WORK_DIMENSION",
"CL_INVALID_WORK_GROUP_SIZE",
"CL_INVALID_WORK_ITEM_SIZE",
"CL_INVALID_GLOBAL_OFFSET",
"CL_INVALID_EVENT_WAIT_LIST",
"CL_INVALID_EVENT",
"CL_INVALID_OPERATION",
"CL_INVALID_GL_OBJECT",
"CL_INVALID_BUFFER_SIZE",
"CL_INVALID_MIP_LEVEL",
"CL_INVALID_GLOBAL_WORK_SIZE",
};
const int errorCount = sizeof(errorString) / sizeof(errorString[0]);
const int index = -error;
return (index >= 0 && index < errorCount) ? errorString[index] : "Unspecified Error";
}
void checkErr(int res, int ref, cl_int code, const char *msg)
{
if(res != ref)
{
printf("%s: %s\n", msg, errStr(code));
exit(1);
}
}
void test(cl_context context, cl_device_id device)
{
cl_command_queue queue = clCreateCommandQueue(context, device, 0, &err);
checkErr(queue != NULL, true, err, "Failed to create queue");
FILE *fp = fopen("test.cl", "r");
fseek(fp, 0, SEEK_END);
size_t f_size = ftell(fp);
fseek(fp, 0, SEEK_SET);
char *src = (char *)malloc(f_size + 1);
fread(src, 1, f_size, fp);
fclose(fp);
src[f_size] = '\0';
cl_program program = clCreateProgramWithSource(context, 1, (const char **)&src, NULL, NULL);
// build the compute program executable
cl_int err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
if (err != CL_SUCCESS)
{
char *build_log;
size_t ret_val_size;
clGetProgramBuildInfo(
program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
build_log = malloc((ret_val_size+1) * sizeof(char));
clGetProgramBuildInfo(
program, device, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
build_log[ret_val_size] = '\0';
printf("Build log:\n%s\n", build_log);
}
else
{
printf("Build succeeded\n");
}
checkErr(err, CL_SUCCESS, err, "Failed to build the program");
}
int main()
{
cl_int res;
cl_platform_id platforms[3];
cl_uint num_platforms;
res = clGetPlatformIDs(3, platforms, &num_platforms);
printf("Found %d platforms\n", num_platforms);
int plnum;
for(plnum = 0; plnum < num_platforms; plnum++)
{
char *name;
size_t param_value_size_ret;
res = clGetPlatformInfo(platforms[plnum], CL_PLATFORM_NAME,
0, NULL, &param_value_size_ret);
name = (char *)malloc(param_value_size_ret);
res = clGetPlatformInfo(platforms[plnum], CL_PLATFORM_NAME,
param_value_size_ret, name, &param_value_size_ret);
printf("Platform: %s\n", name);
free(name);
cl_uint num_devices;
res = clGetDeviceIDs(platforms[plnum], CL_DEVICE_TYPE_ALL, 0, NULL, &num_devices);
printf("Found %d devices\n", num_devices);
cl_device_id *devices = malloc(sizeof(cl_device_id) * num_devices);
res = clGetDeviceIDs(
platforms[plnum], CL_DEVICE_TYPE_ALL, num_devices, devices, &num_devices);
int devnum;
for(devnum = 0; devnum < num_devices; devnum++)
{
res = clGetDeviceInfo(devices[devnum], CL_DEVICE_NAME,
0, NULL, &param_value_size_ret);
name = (char *)malloc(param_value_size_ret);
res = clGetDeviceInfo(devices[devnum], CL_DEVICE_NAME,
param_value_size_ret, name, &param_value_size_ret);
printf("Device: %s\n", name);
free(name);
cl_context context = clCreateContext(NULL, 1, devices + devnum, NULL, NULL, &err);
checkErr(context != NULL, true, err, "Failed to create context");
test(context, devices[devnum]);
}
}
return 0;
}
// taken from pyopencl._cluda
#define LOCAL_BARRIER barrier(CLK_LOCAL_MEM_FENCE)
// 'static' helps to avoid the "no previous prototype for function" warning
#if __OPENCL_VERSION__ >= 120
#define WITHIN_KERNEL static
#else
#define WITHIN_KERNEL
#endif
#define KERNEL __kernel
#define GLOBAL_MEM __global
#define LOCAL_MEM __local
#define LOCAL_MEM_DYNAMIC __local
#define LOCAL_MEM_ARG __local
// INLINE is already defined in Beignet driver
#ifndef INLINE
#define INLINE inline
#endif
#define SIZE_T size_t
#define VSIZE_T size_t
// used to align fields in structures
#define ALIGN(bytes) __attribute__ ((aligned(bytes)))
#if defined(cl_khr_fp64)
#pragma OPENCL EXTENSION cl_khr_fp enable
#elif defined(cl_amd_fp64)
#pragma OPENCL EXTENSION cl_amd_fp enable
#endif
#define COMPLEX_CTR(T) (T)
WITHIN_KERNEL VSIZE_T virtual_local_id(unsigned int dim)
{
if (dim == 0)
{
SIZE_T flat_id =
get_local_id(0) * 1 +
0;
return (flat_id / 1);
}
return 0;
}
WITHIN_KERNEL VSIZE_T virtual_local_size(unsigned int dim)
{
if (dim == 0)
{
return 512;
}
return 1;
}
WITHIN_KERNEL VSIZE_T virtual_group_id(unsigned int dim)
{
if (dim == 0)
{
SIZE_T flat_id =
get_group_id(0) * 1 +
0;
return (flat_id / 1);
}
return 0;
}
WITHIN_KERNEL VSIZE_T virtual_num_groups(unsigned int dim)
{
if (dim == 0)
{
return 2;
}
return 1;
}
WITHIN_KERNEL VSIZE_T virtual_global_id(unsigned int dim)
{
return virtual_local_id(dim) + virtual_group_id(dim) * virtual_local_size(dim);
}
WITHIN_KERNEL VSIZE_T virtual_global_size(unsigned int dim)
{
if(dim == 0)
{
return 1000;
}
return 1;
}
WITHIN_KERNEL VSIZE_T virtual_global_flat_id()
{
return
virtual_global_id(0) * 1 +
0;
}
WITHIN_KERNEL VSIZE_T virtual_global_flat_size()
{
return
virtual_global_size(0) *
1;
}
WITHIN_KERNEL bool virtual_skip_local_threads()
{
return false;
}
WITHIN_KERNEL bool virtual_skip_groups()
{
return false;
}
WITHIN_KERNEL bool virtual_skip_global_threads()
{
if (virtual_global_id(0) >= 1000)
return true;
return false;
}
#define VIRTUAL_SKIP_THREADS if(virtual_skip_local_threads() || virtual_skip_groups() || virtual_skip_global_threads()) return
typedef struct _module0__ {
int val1;
int val2;
char pad;
} ALIGN(4) _module0_;
// leaf output macro for "output"
#define _module1_(_idx0, _val) _leaf_output[(_idx0) * 1] = (_val)
// input transformation node for "input"
INLINE WITHIN_KERNEL _module0_ _module2_func(
VSIZE_T _idx0)
{
_module0_ _val;
const _module0_ val = {83, 31, 96};
_val =(val);
return _val;
}
#define _module2_(_idx0) _module2_func( _idx0)
KERNEL void _kernel_func(GLOBAL_MEM _module0_ *_leaf_output)
{
VIRTUAL_SKIP_THREADS;
VSIZE_T _idx0 = virtual_global_id(0);
_module1_(_idx0, _module2_(_idx0));
}
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment