#define _CRT_SECURE_NO_WARNINGS

#include "private.h"

#include "private_opencl.h"

#include "gpuarray/buffer.h"

#include "gpuarray/util.h"

#include "gpuarray/error.h"

#include "gpuarray/buffer_blas.h"

#include <assert.h>

#include <stdlib.h>

#include <string.h>

#include <limits.h>

#include "loaders/libclblas.h"

#include "loaders/libclblast.h"

#include "cluda_opencl.h.c"

#define _unused(x) ((void)x)

#define SSIZE_MIN (-(SSIZE_MAX-1))

extern gpuarray_blas_ops clblas_ops;

extern gpuarray_blas_ops clblast_ops;

const gpuarray_buffer_ops opencl_ops;

static int cl_property(gpucontext *c, gpudata *b, gpukernel *k, int p, void *r);

static gpudata *cl_alloc(gpucontext *c, size_t size, void *data, int flags);

static void cl_release(gpudata *b);

static void cl_free_ctx(cl_ctx *ctx);

static int cl_newkernel(gpukernel **k, gpucontext *ctx, unsigned int count,

const char **strings, const size_t *lengths,

const char *fname, unsigned int argcount,

const int *types, int flags, char **err_str);

static const char CL_CONTEXT_PREAMBLE[] =

"-D __GA_WARP_SIZE=%lu"; // to be filled by cl_make_ctx()

static int setup_done = 0;

static int setup_lib(error *e) {

if (setup_done)

return GA_NO_ERROR;

GA_CHECK(load_libopencl(e));

setup_done = 1;

return GA_NO_ERROR;

}

static int cl_get_platform_count(unsigned int* platcount) {

cl_uint nump;

GA_CHECK(setup_lib(global_err));

CL_CHECK(global_err, clGetPlatformIDs(0, NULL, &nump));

*platcount = (unsigned int)nump;

return GA_NO_ERROR;

}

static int cl_get_device_count(unsigned int platform, unsigned int* devcount) {

cl_platform_id *ps;

cl_platform_id p;

cl_uint numd;

cl_int err;

unsigned int platcount = 0;

/* This will load the library if needed */

GA_CHECK(cl_get_platform_count(&platcount));

ps = calloc(sizeof(*ps), platcount);

if (ps == NULL)

return error_sys(global_err, "calloc");

err = clGetPlatformIDs(platcount, ps, NULL);

if (err != CL_SUCCESS) {

free(ps);

return error_cl(global_err, "clGetPlatformIDs", err);

}

p = ps[platform];

err = clGetDeviceIDs(p, CL_DEVICE_TYPE_ALL, 0, NULL, &numd);

free(ps);

if (err != CL_SUCCESS)

return error_cl(global_err, "clGetDeviceIds", err);

*devcount = (unsigned int)numd;

return GA_NO_ERROR;

}

static cl_device_id get_dev(cl_context ctx, error *e) {

size_t sz;

cl_device_id res;

cl_device_id *ids;

cl_int err;

CL_CHECKN(e, clGetContextInfo(ctx, CL_CONTEXT_DEVICES, 0, NULL, &sz));

ids = malloc(sz);

if (ids == NULL) {

error_sys(e, "malloc");

return NULL;

}

err = clGetContextInfo(ctx, CL_CONTEXT_DEVICES, sz, ids, NULL);

if (err != CL_SUCCESS) {

free(ids);

error_cl(e, "clContextGetInfo", err);

return NULL;

}

res = ids[0];

free(ids);

return res;

}

cl_ctx *cl_make_ctx(cl_context ctx, gpucontext_props *p) {

cl_ctx *res;

cl_device_id id;

cl_command_queue_properties qprop;

char vendor[32];

char driver_version[64];

char *device_version = NULL;

size_t device_version_size = 0;

cl_uint vendor_id;

cl_int err;

size_t len;

int64_t v = 0;

int e = 0;

size_t warp_size;

int ret;

const char dummy_kern[] = "__kernel void kdummy(__global float *f) { f[0] = 0; }\n";

strb context_preamble = STRB_STATIC_INIT;

const char *rlk[1];

gpukernel *m;

e = setup_lib(global_err);

if (e != GA_NO_ERROR)

return NULL;

id = get_dev(ctx, global_err);

if (id == NULL) return NULL;

/* Query device version string size */

CL_CHECKN(global_err, clGetDeviceInfo(id, CL_DEVICE_VERSION,

0, NULL, &device_version_size));

if (device_version_size > 1024) {

error_set(global_err, GA_UNSUPPORTED_ERROR,

"device version buffer too large");

return NULL;

}

device_version = alloca(device_version_size);

CL_CHECKN(global_err, clGetDeviceInfo(id, CL_DEVICE_VERSION,

device_version_size,

device_version, NULL));

if (device_version[7] == '1' && device_version[9] < '2') {

error_set(global_err, GA_UNSUPPORTED_ERROR,

"We only support OpenCL 1.2 and up");

return NULL;

}

CL_CHECKN(global_err, clGetDeviceInfo(id, CL_DEVICE_QUEUE_PROPERTIES,

sizeof(qprop), &qprop, NULL));

CL_CHECKN(global_err, clGetDeviceInfo(id, CL_DEVICE_VENDOR, sizeof(vendor),

vendor, NULL));

CL_CHECKN(global_err, clGetDeviceInfo(id, CL_DEVICE_VENDOR_ID,

sizeof(vendor_id), &vendor_id, NULL));

CL_CHECKN(global_err, clGetDeviceInfo(id, CL_DRIVER_VERSION,

sizeof(driver_version),

driver_version, NULL));

res = malloc(sizeof(*res));

if (res == NULL) {

error_sys(global_err, "malloc");

return NULL;

}

res->ctx = ctx;

res->ops = &opencl_ops;

if (error_alloc(&res->err)) {

error_set(global_err, GA_SYS_ERROR, "Could not create error context");

free(res);

return NULL;

}

res->refcnt = 1;

res->exts = NULL;

res->blas_handle = NULL;

res->options = NULL;

res->q = clCreateCommandQueue(

ctx, id,

ISSET(p->flags, GA_CTX_SINGLE_STREAM) ? 0 : qprop&CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE,

&err);

if (res->q == NULL) {

error_cl(global_err, "clCreateCommandQueue", err);

error_free(res->err);

free(res);

return NULL;

}

/* Can't overflow (source is 32 + 16 + 12 and buffer is 64) */

len = strlcpy(res->bin_id, vendor, sizeof(res->bin_id));

snprintf(res->bin_id + len, sizeof(res->bin_id) - len, " %#x ", vendor_id);

strlcat(res->bin_id, driver_version, sizeof(res->bin_id));

clRetainContext(res->ctx);

TAG_CTX(res);

res->errbuf = cl_alloc((gpucontext *)res, 8, &v, GA_BUFFER_INIT);

if (res->errbuf == NULL)

goto fail;

res->refcnt--; /* Prevent ref loop */

/* Create per-context OpenCL preamble */

// Create a dummy kernel and check GA_KERNEL_PROP_PREFLSIZE

rlk[0] = dummy_kern;

len = sizeof(dummy_kern);

// this dummy kernel does not require a CLUDA preamble

if (cl_newkernel(&m, (gpucontext *)res, 1, rlk, &len, "kdummy", 0, NULL, 0, NULL) != GA_NO_ERROR)

goto fail;

ret = cl_property((gpucontext *)res, NULL, m, GA_KERNEL_PROP_PREFLSIZE, &warp_size);

if (ret != GA_NO_ERROR)

goto fail;

// Write the preferred workgroup multiple as GA_WARP_SIZE in preamble

strb_appendf(&context_preamble, CL_CONTEXT_PREAMBLE, (unsigned long)warp_size);

res->options = strb_cstr(&context_preamble);

if (res->options == NULL)

goto fail;

res->blas_handle = NULL;

if (load_libclblas(res->err) == GA_NO_ERROR) {

res->blas_ops = &clblas_ops;

} else if (load_libclblast(res->err) == GA_NO_ERROR) {

res->blas_ops = &clblast_ops;

} else {

res->blas_ops = NULL;

}

res->comm_ops = NULL;

return res;

fail:

error_set(global_err, res->err->code, res->err->msg);

cl_free_ctx(res);

return NULL;

}

cl_command_queue cl_get_stream(gpucontext *ctx) {

ASSERT_CTX((cl_ctx *)ctx);

return ((cl_ctx *)ctx)->q;

}

static void cl_free_ctx(cl_ctx *ctx) {

ASSERT_CTX(ctx);

assert(ctx->refcnt != 0);

ctx->refcnt--;

if (ctx->refcnt == 0) {

if (ctx->errbuf != NULL) {

ctx->refcnt = 2; /* Avoid recursive release */

cl_release(ctx->errbuf);

}

clReleaseCommandQueue(ctx->q);

clReleaseContext(ctx->ctx);

if (ctx->options != NULL)

free(ctx->options);

error_free(ctx->err);

CLEAR(ctx);

free(ctx);

}

}

gpudata *cl_make_buf(gpucontext *c, cl_mem buf) {

cl_ctx *ctx = (cl_ctx *)c;

gpudata *res;

cl_context buf_ctx;

cl_int err;

ASSERT_CTX(ctx);

CL_CHECKN(ctx->err, clGetMemObjectInfo(buf, CL_MEM_CONTEXT, sizeof(buf_ctx),

&buf_ctx, NULL));

if (buf_ctx != ctx->ctx) {

error_set(ctx->err, GA_VALUE_ERROR, "Requested context doesn't match object context");

return NULL;

}

res = malloc(sizeof(*res));

if (res == NULL) {

error_sys(ctx->err, "malloc");

return NULL;

}

res->buf = buf;

res->ev = NULL;

res->refcnt = 1;

err = clRetainMemObject(buf);

if (err != CL_SUCCESS) {

free(res);

error_cl(ctx->err, "clRetainMemObject", err);

return NULL;

}

res->ctx = ctx;

res->ctx->refcnt++;

TAG_BUF(res);

return res;

}

cl_mem cl_get_buf(gpudata *g) { ASSERT_BUF(g); return g->buf; }

#define PRAGMA "#pragma OPENCL EXTENSION "

#define ENABLE " : enable\n"

#define CL_SMALL "cl_khr_byte_addressable_store"

#define CL_DOUBLE "cl_khr_fp64"

#define CL_HALF "cl_khr_fp16"

static void cl_releasekernel(gpukernel *k);

static int cl_callkernel(gpukernel *k, unsigned int n,

const size_t *gs, const size_t *ls,

size_t shared, void **args);

const char *cl_error_string(cl_int err) {

switch (err) {

case CL_SUCCESS: return "Success!";

case CL_DEVICE_NOT_FOUND: return "Device not found.";

case CL_DEVICE_NOT_AVAILABLE: return "Device not available";

case CL_COMPILER_NOT_AVAILABLE: return "Compiler not available";

case CL_MEM_OBJECT_ALLOCATION_FAILURE: return "Memory object allocation failure";

case CL_OUT_OF_RESOURCES: return "Out of resources";

case CL_OUT_OF_HOST_MEMORY: return "Out of host memory";

case CL_PROFILING_INFO_NOT_AVAILABLE: return "Profiling information not available";

case CL_MEM_COPY_OVERLAP: return "Memory copy overlap";

case CL_IMAGE_FORMAT_MISMATCH: return "Image format mismatch";

case CL_IMAGE_FORMAT_NOT_SUPPORTED: return "Image format not supported";

case CL_BUILD_PROGRAM_FAILURE: return "Program build failure";

case CL_MAP_FAILURE: return "Map failure";

case CL_MISALIGNED_SUB_BUFFER_OFFSET: return "Buffer offset improperly aligned";

case CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST: return "Event in wait list has an error status";

case CL_INVALID_VALUE: return "Invalid value";

case CL_INVALID_DEVICE_TYPE: return "Invalid device type";

case CL_INVALID_PLATFORM: return "Invalid platform";

case CL_INVALID_DEVICE: return "Invalid device";

case CL_INVALID_CONTEXT: return "Invalid context";

case CL_INVALID_QUEUE_PROPERTIES: return "Invalid queue properties";

case CL_INVALID_COMMAND_QUEUE: return "Invalid command queue";

case CL_INVALID_HOST_PTR: return "Invalid host pointer";

case CL_INVALID_MEM_OBJECT: return "Invalid memory object";

case CL_INVALID_IMAGE_FORMAT_DESCRIPTOR:return "Invalid image format descriptor";

case CL_INVALID_IMAGE_SIZE: return "Invalid image size";

case CL_INVALID_SAMPLER: return "Invalid sampler";

case CL_INVALID_BINARY: return "Invalid binary";

case CL_INVALID_BUILD_OPTIONS: return "Invalid build options";

case CL_INVALID_PROGRAM: return "Invalid program";

case CL_INVALID_PROGRAM_EXECUTABLE: return "Invalid program executable";

case CL_INVALID_KERNEL_NAME: return "Invalid kernel name";

case CL_INVALID_KERNEL_DEFINITION: return "Invalid kernel definition";

case CL_INVALID_KERNEL: return "Invalid kernel";

case CL_INVALID_ARG_INDEX: return "Invalid argument index";

case CL_INVALID_ARG_VALUE: return "Invalid argument value";

case CL_INVALID_ARG_SIZE: return "Invalid argument size";

case CL_INVALID_KERNEL_ARGS: return "Invalid kernel arguments";

case CL_INVALID_WORK_DIMENSION: return "Invalid work dimension";

case CL_INVALID_WORK_GROUP_SIZE: return "Invalid work group size";

case CL_INVALID_WORK_ITEM_SIZE: return "Invalid work item size";

case CL_INVALID_GLOBAL_OFFSET: return "Invalid global offset";

case CL_INVALID_EVENT_WAIT_LIST: return "Invalid event wait list";

case CL_INVALID_EVENT: return "Invalid event";

case CL_INVALID_OPERATION: return "Invalid operation";

case CL_INVALID_GL_OBJECT: return "Invalid OpenGL object";

case CL_INVALID_BUFFER_SIZE: return "Invalid buffer size";

case CL_INVALID_MIP_LEVEL: return "Invalid mip-map level";

case CL_INVALID_GLOBAL_WORK_SIZE: return "Invalid global work size";

case CL_INVALID_PROPERTY: return "Invalid property";

default: return "Unknown error";

}

}

static int check_ext(cl_ctx *ctx, const char *name) {

cl_device_id dev;

if (ctx->exts == NULL) {

dev = get_dev(ctx->ctx, ctx->err);

if (dev == NULL) return ctx->err->code;

CL_GET_PROP(ctx->err, clGetDeviceInfo, dev, CL_DEVICE_EXTENSIONS, ctx->exts);

}

if (strstr(ctx->exts, name) == NULL)

return error_fmt(ctx->err, GA_DEVSUP_ERROR, "Unsupported extension %s", name);

else

return GA_NO_ERROR;

}

static void

#ifdef _MSC_VER

__stdcall

#endif

errcb(const char *errinfo, const void *pi, size_t cb, void *u) {

fprintf(stderr, "%s\n", errinfo);

}

static gpucontext *cl_init(gpucontext_props *pp) {

cl_device_id *ds;

cl_device_id d;

cl_platform_id *ps;

cl_platform_id p;

cl_uint nump, numd;

cl_context_properties props[3] = {

CL_CONTEXT_PLATFORM, 0,

0,

};

cl_context ctx;

cl_ctx *res;

cl_int err;

int platno;

int devno;

int e;

platno = pp->dev >> 16;

devno = pp->dev & 0xFFFF;

e = setup_lib(global_err);

if (e != GA_NO_ERROR)

return NULL;

CL_CHECKN(global_err, clGetPlatformIDs(0, NULL, &nump));

if ((unsigned int)platno >= nump || platno < 0) {

error_set(global_err, GA_VALUE_ERROR, "Platform ID out of range");

return NULL;

}

ps = calloc(sizeof(*ps), nump);

if (ps == NULL) {

error_sys(global_err, "calloc");

return NULL;

}

err = clGetPlatformIDs(nump, ps, NULL);

/* We may get garbage on failure here but it won't matter as we will

not use it */

p = ps[platno];

free(ps);

if (err != CL_SUCCESS) {

error_cl(global_err, "clGetPlatformIDs", err);

return NULL;

}

CL_CHECKN(global_err, clGetDeviceIDs(p, CL_DEVICE_TYPE_ALL, 0, NULL, &numd));

if ((unsigned int)devno >= numd || devno < 0) {

error_set(global_err, GA_VALUE_ERROR, "Device ID out of range");

return NULL;

}

ds = calloc(sizeof(*ds), numd);

if (ds == NULL) {

error_sys(global_err, "calloc");

return NULL;

}

err = clGetDeviceIDs(p, CL_DEVICE_TYPE_ALL, numd, ds, NULL);

d = ds[devno];

free(ds);

if (err != CL_SUCCESS) {

error_cl(global_err, "clGetDeviceIDs", err);

return NULL;

}

props[1] = (cl_context_properties)p;

ctx = clCreateContext(props, 1, &d, errcb, NULL, &err);

if (ctx == NULL) {

error_cl(global_err, "clCreateContext", err);

return NULL;

}

res = cl_make_ctx(ctx, pp);

clReleaseContext(ctx);

return (gpucontext *)res;

}

static void cl_deinit(gpucontext *c) {

ASSERT_CTX((cl_ctx *)c);

cl_free_ctx((cl_ctx *)c);

}

static gpudata *cl_alloc(gpucontext *c, size_t size, void *data, int flags) {

cl_ctx *ctx = (cl_ctx *)c;

gpudata *res;

void *hostp = NULL;

cl_int err;

cl_mem_flags clflags = CL_MEM_READ_WRITE;

ASSERT_CTX(ctx);

if (flags & GA_BUFFER_INIT) {

if (data == NULL) {

error_set(ctx->err, GA_VALUE_ERROR, "Requested initialization, but no data provided");

return NULL;

}

hostp = data;

clflags |= CL_MEM_COPY_HOST_PTR;

}

if (flags & GA_BUFFER_HOST) {

clflags |= CL_MEM_ALLOC_HOST_PTR;

}

if (flags & GA_BUFFER_READ_ONLY) {

if (flags & GA_BUFFER_WRITE_ONLY) {

error_set(ctx->err, GA_VALUE_ERROR, "Invalid combinaison: READ_ONLY and WRITE_ONLY");

return NULL;

}

clflags &= ~CL_MEM_READ_WRITE;

clflags |= CL_MEM_READ_ONLY;

}

if (flags & GA_BUFFER_WRITE_ONLY) {

clflags &= ~CL_MEM_READ_WRITE;

clflags |= CL_MEM_WRITE_ONLY;

}

res = malloc(sizeof(*res));

if (res == NULL) {

error_sys(ctx->err, "malloc");

return NULL;

}

res->refcnt = 1;

if (size == 0) {

/* OpenCL doesn't like a zero-sized buffer */

size = 1;

}

res->buf = clCreateBuffer(ctx->ctx, clflags, size, hostp, &err);

res->ev = NULL;

if (err != CL_SUCCESS) {

free(res);

error_cl(ctx->err, "clCreateBuffer", err);

return NULL;

}

res->ctx = ctx;

ctx->refcnt++;

TAG_BUF(res);

return res;

}

static void cl_retain(gpudata *b) {

ASSERT_BUF(b);

b->refcnt++;

}

static void cl_release(gpudata *b) {

ASSERT_BUF(b);

b->refcnt--;

if (b->refcnt == 0) {

CLEAR(b);

clReleaseMemObject(b->buf);

if (b->ev != NULL)

clReleaseEvent(b->ev);

cl_free_ctx(b->ctx);

free(b);

}

}

static int cl_share(gpudata *a, gpudata *b) {

cl_ctx *ctx;

cl_mem aa, bb;

cl_int err;

ASSERT_BUF(a);

ASSERT_BUF(b);

if (a->buf == b->buf) return 1;

if (a->ctx != b->ctx) return 0;

ctx = a->ctx;

ASSERT_CTX(ctx);

err = clGetMemObjectInfo(a->buf, CL_MEM_ASSOCIATED_MEMOBJECT,

sizeof(aa), &aa, NULL);

if (err != CL_SUCCESS) {

error_cl(ctx->err, "clGetMemObjectInfo", err);

return -1;

}

err = clGetMemObjectInfo(b->buf, CL_MEM_ASSOCIATED_MEMOBJECT,

sizeof(bb), &bb, NULL);

if (err != CL_SUCCESS) {

error_cl(ctx->err, "clGetMemObjectInfo", err);

return -1;

}

if (aa == NULL) aa = a->buf;

if (bb == NULL) bb = b->buf;

if (aa == bb) return 1;

return 0;

}

static int cl_move(gpudata *dst, size_t dstoff, gpudata *src, size_t srcoff,

size_t sz) {

cl_ctx *ctx;

cl_event ev;

cl_event evw[2];

cl_event *evl = NULL;

cl_uint num_ev = 0;

ASSERT_BUF(dst);

ASSERT_BUF(src);

if (dst->ctx != src->ctx) {

error_set(src->ctx->err, GA_VALUE_ERROR, "Differing contexts for source and destination");

return error_set(dst->ctx->err, src->ctx->err->code, src->ctx->err->msg);

}

ctx = dst->ctx;

ASSERT_CTX(ctx);

if (sz == 0) return GA_NO_ERROR;

if (src->ev != NULL)

evw[num_ev++] = src->ev;

if (dst->ev != NULL && src != dst)

evw[num_ev++] = dst->ev;

if (num_ev > 0)

evl = evw;

CL_CHECK(ctx->err, clEnqueueCopyBuffer(ctx->q, src->buf, dst->buf, srcoff,

dstoff, sz, num_ev, evl, &ev));

if (src->ev != NULL)

clReleaseEvent(src->ev);

if (dst->ev != NULL && src != dst)

clReleaseEvent(dst->ev);

src->ev = ev;

dst->ev = ev;

clRetainEvent(ev);

return GA_NO_ERROR;

}

static int cl_read(void *dst, gpudata *src, size_t srcoff, size_t sz) {

cl_ctx *ctx = src->ctx;

cl_event ev[1];

cl_event *evl = NULL;

cl_uint num_ev = 0;

ASSERT_BUF(src);

ASSERT_CTX(ctx);

if (sz == 0) return GA_NO_ERROR;

if (src->ev != NULL) {

ev[0] = src->ev;

evl = ev;

num_ev = 1;

}

CL_CHECK(ctx->err, clEnqueueReadBuffer(ctx->q, src->buf, CL_TRUE, srcoff, sz,

dst, num_ev, evl, NULL));

if (src->ev != NULL) clReleaseEvent(src->ev);

src->ev = NULL;

return GA_NO_ERROR;

}

static int cl_write(gpudata *dst, size_t dstoff, const void *src, size_t sz) {

cl_ctx *ctx = dst->ctx;

cl_event ev[1];

cl_event *evl = NULL;

cl_uint num_ev = 0;

ASSERT_BUF(dst);

ASSERT_CTX(ctx);

if (sz == 0) return GA_NO_ERROR;

if (dst->ev != NULL) {

ev[0] = dst->ev;

evl = ev;

num_ev = 1;

}

CL_CHECK(ctx->err, clEnqueueWriteBuffer(ctx->q, dst->buf, CL_TRUE, dstoff,

sz, src, num_ev, evl, NULL));

if (dst->ev != NULL) clReleaseEvent(dst->ev);

dst->ev = NULL;

return GA_NO_ERROR;

}

static int cl_memset(gpudata *dst, size_t offset, int data) {

char local_kern[256];

cl_ctx *ctx = dst->ctx;

const char *rlk[1];

void *args[1];

size_t sz, bytes, n, ls, gs;

gpukernel *m;

cl_mem_flags fl;

int type;

int r, res;

unsigned char val = (unsigned char)data;

cl_uint pattern = (cl_uint)val & (cl_uint)val >> 8 & \

(cl_uint)val >> 16 & (cl_uint)val >> 24;

ASSERT_BUF(dst);

ASSERT_CTX(ctx);

CL_CHECK(ctx->err, clGetMemObjectInfo(dst->buf, CL_MEM_FLAGS, sizeof(fl),

&fl, NULL));

if (fl & CL_MEM_READ_ONLY)

return error_set(ctx->err, GA_READONLY_ERROR, "destination is read only");

CL_CHECK(ctx->err, clGetMemObjectInfo(dst->buf, CL_MEM_SIZE, sizeof(bytes),

&bytes, NULL));

bytes -= offset;

if (bytes == 0) return GA_NO_ERROR;

if ((bytes % 16) == 0) {

n = bytes/16;

r = snprintf(local_kern, sizeof(local_kern),

"__kernel void kmemset(__global uint4 *mem) {"

"unsigned int i; __global char *tmp = (__global char *)mem;"

"tmp += %" SPREFIX "u; mem = (__global uint4 *)tmp;"

"for (i = get_global_id(0); i < %" SPREFIX "u; "

"i += get_global_size(0)) {mem[i] = (uint4)(%u,%u,%u,%u); }}",

offset, n, pattern, pattern, pattern, pattern);

} else if ((bytes % 8) == 0) {

n = bytes/8;

r = snprintf(local_kern, sizeof(local_kern),

"__kernel void kmemset(__global uint2 *mem) {"

"unsigned int i; __global char *tmp = (__global char *)mem;"

"tmp += %" SPREFIX "u; mem = (__global uint2 *)tmp;"

"for (i = get_global_id(0); i < %" SPREFIX "u;"

"i += get_global_size(0)) {mem[i] = (uint2)(%u,%u); }}",

offset, n, pattern, pattern);

} else if ((bytes % 4) == 0) {

n = bytes/4;

r = snprintf(local_kern, sizeof(local_kern),

"__kernel void kmemset(__global unsigned int *mem) {"

"unsigned int i; __global char *tmp = (__global char *)mem;"

"tmp += %" SPREFIX "u; mem = (__global unsigned int *)tmp;"

"for (i = get_global_id(0); i < %" SPREFIX "u;"

"i += get_global_size(0)) {mem[i] = %u; }}",

offset, n, pattern);

} else {

GA_CHECK(check_ext(ctx, CL_SMALL));

n = bytes;

r = snprintf(local_kern, sizeof(local_kern),

"__kernel void kmemset(__global unsigned char *mem) {"

"unsigned int i; mem += %" SPREFIX "u;"

"for (i = get_global_id(0); i < %" SPREFIX "u;"

"i += get_global_size(0)) {mem[i] = %u; }}",

offset, n, val);

}

/* If this assert fires, increase the size of local_kern above. */

assert(r <= (int)sizeof(local_kern));

_unused(r);

sz = strlen(local_kern);

rlk[0] = local_kern;

type = GA_BUFFER;

r = cl_newkernel(&m, (gpucontext *)ctx, 1, rlk, &sz, "kmemset", 1, &type, 0, NULL);

if (r != GA_NO_ERROR)

return r;

/* Cheap kernel scheduling */

res = cl_property(NULL, NULL, m, GA_KERNEL_PROP_MAXLSIZE, &ls);

if (res != GA_NO_ERROR) goto fail;

gs = ((n-1) / ls) + 1;

args[0] = dst;

res = cl_callkernel(m, 1, &gs, &ls, 0, args);

fail:

cl_releasekernel(m);

return res;

}

static int cl_check_extensions(const char **preamble, unsigned int *count,

int flags, cl_ctx *ctx) {

if (flags & GA_USE_SMALL) {

GA_CHECK(check_ext(ctx, CL_SMALL));

preamble[*count] = PRAGMA CL_SMALL ENABLE;

(*count)++;

}

if (flags & GA_USE_DOUBLE) {

if (check_ext(ctx, CL_DOUBLE) != GA_NO_ERROR) {

return error_set(ctx->err, GA_DEVSUP_ERROR, "This device does not support double precision (pygpu int/int, int32+float32, and floating point literals default to double precision)");

}

preamble[*count] = PRAGMA CL_DOUBLE ENABLE;

(*count)++;

}

if (flags & GA_USE_COMPLEX) {

return error_set(ctx->err, GA_UNSUPPORTED_ERROR, "Complex are not supported yet");

}

if (flags & GA_USE_CUDA) {

return error_set(ctx->err, GA_DEVSUP_ERROR, "Cuda kernels not supported on opencl devices");

}

return GA_NO_ERROR;

}

static int cl_newkernel(gpukernel **k, gpucontext *c, unsigned int count,

const char **strings, const size_t *lengths,

const char *fname, unsigned int argcount,

const int *types, int flags, char **err_str) {

cl_ctx *ctx = (cl_ctx *)c;

gpukernel *res;

cl_device_id dev;

cl_program p;

cl_program cluda;

cl_program tmp;

// Sync this table size with the number of flags that can add stuff

// at the beginning

const char *preamble[5];

const char *cluda_src[1];

const char *headers[1] = {"cluda.h"};

size_t *newl = NULL;

const char **news = NULL;

cl_int err;

unsigned int n = 0;

strb debug_msg = STRB_STATIC_INIT;

size_t log_size;

ASSERT_CTX(ctx);

if (count == 0)

return error_set(ctx->err, GA_VALUE_ERROR, "Empty kernel source list");

dev = get_dev(ctx->ctx, ctx->err);

if (dev == NULL) return ctx->err->code;

if (cl_check_extensions(preamble, &n, flags, ctx))

return ctx->err->code;

if (n != 0) {

news = calloc(count+n, sizeof(const char *));

if (news == NULL)

return error_sys(ctx->err, "calloc");

memcpy(news, preamble, n*sizeof(const char *));

memcpy(news+n, strings, count*sizeof(const char *));

if (lengths == NULL) {

newl = NULL;

} else {

newl = calloc(count+n, sizeof(size_t));

if (newl == NULL) {

free(news);

return error_sys(ctx->err, "calloc");

}

memcpy(newl+n, lengths, count*sizeof(size_t));

}

} else {

news = strings;

newl = (size_t *)lengths;

}

cluda_src[0] = cluda_opencl_h;

cluda = clCreateProgramWithSource(ctx->ctx, 1, cluda_src, NULL, &err);

if (err != CL_SUCCESS) {

if (n != 0) {

free(news);

free(newl);

}

return error_cl(ctx->err, "clCreateProgramWithSource (header)", err);

}

p = clCreateProgramWithSource(ctx->ctx, count+n, news, newl, &err);

if (err != CL_SUCCESS) {

if (n != 0) {

free(news);

free(newl);

clReleaseProgram(cluda);

}

return error_cl(ctx->err, "clCreateProgramWithSource (kernel)", err);

}

err = clCompileProgram(p, 0, NULL, ctx->options, 1, &cluda, headers, NULL, NULL);

if (err != CL_SUCCESS)

goto compile_error;

tmp = clLinkProgram(ctx->ctx, 0, NULL, NULL, 1, &p, NULL, NULL, &err);

if (tmp != NULL) {

clReleaseProgram(p);

p = tmp;

tmp = NULL;

}

compile_error:

if (err != CL_SUCCESS) {

if ((err == CL_COMPILE_PROGRAM_FAILURE || err == CL_LINK_PROGRAM_FAILURE)

&& err_str != NULL) {

*err_str = NULL; // Fallback, in case there's an error

// We're substituting debug_msg for a string with this first line:

strb_appends(&debug_msg, "Program build failure ::\n");

// Determine the size of the log

clGetProgramBuildInfo(p, dev, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);

if (strb_ensure(&debug_msg, log_size)!=-1 && log_size>=1) { // Checks strb has enough space

// Get the log directly into the debug_msg

clGetProgramBuildInfo(p, dev, CL_PROGRAM_BUILD_LOG, log_size, debug_msg.s+debug_msg.l, NULL);

debug_msg.l += (log_size-1); // Back off to before final '\0'

}

gpukernel_source_with_line_numbers(count+n, news, newl, &debug_msg);

strb_append0(&debug_msg); // Make sure a final '\0' is present

if (!strb_error(&debug_msg)) { // Make sure the strb is in a valid state

*err_str = memdup(debug_msg.s, debug_msg.l);

// If there's a memory alloc error, fall-through : announcing a compile error is more important

}

strb_clear(&debug_msg);

// *err_str will be free()d by the caller (see docs in kernel.h)

}

clReleaseProgram(p);

if (n != 0) {

free(news);

free(newl);

}

return error_cl(ctx->err, "clBuildProgram", err);

}

if (n != 0) {

free(news);

free(newl);

}

res = malloc(sizeof(*res));

if (res == NULL)

return error_sys(ctx->err, "malloc");

res->refcnt = 1;

res->ev = NULL;

res->argcount = argcount;

res->k = clCreateKernel(p, fname, &err);

res->types = NULL; /* This avoids a crash in cl_releasekernel */

res->evr = NULL; /* This avoids a crash in cl_releasekernel */

res->ctx = ctx;

ctx->refcnt++;

clReleaseProgram(p);

TAG_KER(res);

if (err != CL_SUCCESS) {

cl_releasekernel(res);

return error_cl(ctx->err, "clCreateKernel", err);

}

res->types = calloc(argcount, sizeof(int));

if (res->types == NULL) {

cl_releasekernel(res);

return error_sys(ctx->err, "calloc");

}

memcpy(res->types, types, argcount * sizeof(int));

res->evr = calloc(argcount, sizeof(cl_event *));

if (res->evr == NULL) {

cl_releasekernel(res);

return error_sys(ctx->err, "calloc");

}

*k = res;

return GA_NO_ERROR;

}

static void cl_retainkernel(gpukernel *k) {

ASSERT_KER(k);

k->refcnt++;

}

static void cl_releasekernel(gpukernel *k) {

ASSERT_KER(k);

k->refcnt--;

if (k->refcnt == 0) {

CLEAR(k);

if (k->ev != NULL) clReleaseEvent(k->ev);

if (k->k) clReleaseKernel(k->k);

cl_free_ctx(k->ctx);

free(k->types);

free(k->evr);

free(k);

}

}

static int cl_setkernelarg(gpukernel *k, unsigned int i, void *a) {

cl_ctx *ctx = k->ctx;

gpudata *btmp;

cl_ulong temp;

cl_long stemp;

switch (k->types[i]) {

case GA_BUFFER:

btmp = (gpudata *)a;

CL_CHECK(ctx->err, clSetKernelArg(k->k, i, sizeof(cl_mem), &btmp->buf));

k->evr[i] = &btmp->ev;

break;

case GA_SIZE:

temp = *((size_t *)a);

CL_CHECK(ctx->err, clSetKernelArg(k->k, i, gpuarray_get_elsize(GA_ULONG), &temp));

k->evr[i] = NULL;

break;

case GA_SSIZE:

stemp = *((ssize_t *)a);

CL_CHECK(ctx->err, clSetKernelArg(k->k, i, gpuarray_get_elsize(GA_LONG), &stemp));

k->evr[i] = NULL;