test.c

#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;
}

test.cl

    // 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));

            }