#define _CRT_SECURE_NO_WARNINGS

#include "private.h"

#include "private_cuda.h"

#include <sys/types.h>

#include <assert.h>

#include <stdlib.h>

#include <cache.h>

#include "util/strb.h"

#include "util/xxhash.h"

#include "gpuarray/buffer.h"

#include "gpuarray/util.h"

#include "gpuarray/error.h"

#include "gpuarray/extension.h"

#include "gpuarray/buffer_blas.h"

#define BLOCK_SIZE (4 * 1024 * 1024)

#define FRAG_SIZE (64)

static CUresult err;

static int init_done = 0;

GPUARRAY_LOCAL const gpuarray_buffer_ops cuda_ops;

static void cuda_freekernel(gpukernel *);

static int cuda_property(gpucontext *, gpudata *, gpukernel *, int, void *);

static int cuda_waits(gpudata *, int, CUstream);

static int cuda_records(gpudata *, int, CUstream);

static int detect_arch(const char *prefix, char *ret, CUresult *err);

static gpudata *new_gpudata(cuda_context *ctx, CUdeviceptr ptr, size_t size);

static int strb_eq(void *_k1, void *_k2) {

strb *k1 = (strb *)_k1;

strb *k2 = (strb *)_k2;

return (k1->l == k2->l &&

memcmp(k1->s, k2->s, k1->l) == 0);

}

static uint32_t strb_hash(void *_k) {

strb *k = (strb *)_k;

return XXH32(k->s, k->l, 42);

}

static int cuda_get_platform_count(unsigned int* platcount) {

*platcount = 1; // CUDA works on NVIDIA's GPUs

return GA_NO_ERROR;

}

static int cuda_get_device_count(unsigned int platform,

unsigned int* devcount) {

int dv;

// platform number gets ignored in CUDA implementation

if (!init_done) {

err = cuInit(0);

if (err != CUDA_SUCCESS)

return GA_IMPL_ERROR;

init_done = 1;

}

err = cuDeviceGetCount(&dv);

if (err != CUDA_SUCCESS)

return GA_IMPL_ERROR;

*devcount = (unsigned int)dv;

return GA_NO_ERROR;

}

cuda_context *cuda_make_ctx(CUcontext ctx, int flags) {

cuda_context *res;

void *p;

res = calloc(1, sizeof(*res));

if (res == NULL)

return NULL;

res->ctx = ctx;

res->ops = &cuda_ops;

res->err = CUDA_SUCCESS;

res->refcnt = 1;

res->flags = flags;

res->enter = 0;

res->freeblocks = NULL;

if (detect_arch(ARCH_PREFIX, res->bin_id, &err)) {

goto fail_stream;

}

/* Don't add the nonblocking flags to help usage with other

err = cuStreamCreate(&res->s, 0);

if (err != CUDA_SUCCESS) {

goto fail_stream;

}

if (ISSET(res->flags, GA_CTX_SINGLE_STREAM)) {

res->mem_s = res->s;

} else {

/* Don't add the nonblocking flags to help usage with other

err = cuStreamCreate(&res->mem_s, 0);

if (err != CUDA_SUCCESS) {

goto fail_mem_stream;

}

}

res->kernel_cache = cache_twoq(64, 128, 64, 8, strb_eq, strb_hash,

(cache_freek_fn)strb_free,

(cache_freev_fn)cuda_freekernel);

if (res->kernel_cache == NULL)

goto fail_cache;

err = cuMemAllocHost(&p, 16);

if (err != CUDA_SUCCESS) {

goto fail_errbuf;

}

memset(p, 0, 16);

/* Need to tag for new_gpudata */

TAG_CTX(res);

res->errbuf = new_gpudata(res, (CUdeviceptr)p, 16);

if (res->errbuf == NULL) {

err = res->err;

goto fail_end;

}

res->errbuf->flags |= CUDA_MAPPED_PTR;

return res;

fail_end:

cuMemFreeHost(p);

fail_errbuf:

cache_destroy(res->kernel_cache);

fail_cache:

if (ISCLR(res->flags, GA_CTX_SINGLE_STREAM))

cuStreamDestroy(res->mem_s);

fail_mem_stream:

cuStreamDestroy(res->s);

fail_stream:

free(res);

return NULL;

}

static void deallocate(gpudata *);

static void cuda_free_ctx(cuda_context *ctx) {

gpuarray_blas_ops *blas_ops;

gpudata *next, *curr;

#if CUDA_VERSION >= 7000

CUdevice dev;

ASSERT_CTX(ctx);

ctx->refcnt--;

if (ctx->refcnt == 0) {

assert(ctx->enter == 0 && "Context was active when freed!");

if (ctx->blas_handle != NULL) {

ctx->err = cuda_property((gpucontext *)ctx, NULL, NULL, GA_CTX_PROP_BLAS_OPS,

&blas_ops);

blas_ops->teardown((gpucontext *)ctx);

}

cuMemFreeHost((void *)ctx->errbuf->ptr);

deallocate(ctx->errbuf);

if (ISCLR(ctx->flags, GA_CTX_SINGLE_STREAM))

cuStreamDestroy(ctx->mem_s);

cuStreamDestroy(ctx->s);

/* Clear out the freelist */

for (curr = ctx->freeblocks; curr != NULL; curr = next) {

next = curr->next;

cuMemFree(curr->ptr);

deallocate(curr);

}

cache_destroy(ctx->kernel_cache);

if (!(ctx->flags & DONTFREE)) {

#if CUDA_VERSION < 7000

cuCtxDestroy(ctx->ctx);

cuCtxPushCurrent(ctx->ctx);

cuCtxGetDevice(&dev);

cuCtxPopCurrent(NULL);

cuDevicePrimaryCtxRelease(dev);

}

CLEAR(ctx);

free(ctx);

}

}

CUstream cuda_get_stream(cuda_context *ctx) {

ASSERT_CTX(ctx);

return ctx->s;

}

void cuda_enter(cuda_context *ctx) {

ASSERT_CTX(ctx);

if (!ctx->enter)

cuCtxPushCurrent(ctx->ctx);

ctx->enter++;

}

void cuda_exit(cuda_context *ctx) {

ASSERT_CTX(ctx);

assert(ctx->enter > 0);

ctx->enter--;

if (!ctx->enter)

cuCtxPopCurrent(NULL);

}

static gpudata *new_gpudata(cuda_context *ctx, CUdeviceptr ptr, size_t size) {

gpudata *res;

int fl = CU_EVENT_DISABLE_TIMING;

res = malloc(sizeof(*res));

if (res == NULL) return NULL;

res->refcnt = 0;

res->sz = size;

res->flags = 0;

res->ls = NULL;

cuda_enter(ctx);

if (ctx->flags & GA_CTX_MULTI_THREAD)

fl |= CU_EVENT_BLOCKING_SYNC;

ctx->err = cuEventCreate(&res->rev, fl);

if (ctx->err != CUDA_SUCCESS) {

cuda_exit(ctx);

free(res);

return NULL;

}

ctx->err = cuEventCreate(&res->wev, fl);

if (ctx->err != CUDA_SUCCESS) {

cuEventDestroy(res->rev);

cuda_exit(ctx);

free(res);

return NULL;

}

cuda_exit(ctx);

res->ptr = ptr;

res->next = NULL;

res->ctx = ctx;

TAG_BUF(res);

return res;

}

gpudata *cuda_make_buf(cuda_context *ctx, CUdeviceptr p, size_t sz) {

gpudata *res = new_gpudata(ctx, p, sz);

if (res == NULL) return NULL;

res->refcnt = 1;

res->flags |= DONTFREE;

res->ctx->refcnt++;

return res;

}

size_t cuda_get_sz(gpudata *g) { ASSERT_BUF(g); return g->sz; }

#define FAIL(v, e) { if (ret) *ret = e; return v; }

#define CHKFAIL(v) if (err != CUDA_SUCCESS) FAIL(v, GA_IMPL_ERROR)

static const char CUDA_PREAMBLE[] =

"#define local_barrier() __syncthreads()\n"

"#define WITHIN_KERNEL extern \"C\" __device__\n"

"#define KERNEL extern \"C\" __global__\n"

"#define GLOBAL_MEM /* empty */\n"

"#define LOCAL_MEM __shared__\n"

"#define LOCAL_MEM_ARG /* empty */\n"

"#define REQD_WG_SIZE(X,Y,Z) __launch_bounds__(X*Y, Z)\n"

"#ifdef NAN\n"

"#undef NAN\n"

"#endif\n"

"#define NAN __int_as_float(0x7fffffff)\n"

"#define LID_0 threadIdx.x\n"

"#define LID_1 threadIdx.y\n"

"#define LID_2 threadIdx.z\n"

"#define LDIM_0 blockDim.x\n"

"#define LDIM_1 blockDim.y\n"

"#define LDIM_2 blockDim.z\n"

"#define GID_0 blockIdx.x\n"

"#define GID_1 blockIdx.y\n"

"#define GID_2 blockIdx.z\n"

"#define GDIM_0 gridDim.x\n"

"#define GDIM_1 gridDim.y\n"

"#define GDIM_2 gridDim.z\n"

"#define ga_bool unsigned char\n"

"#define ga_byte signed char\n"

"#define ga_ubyte unsigned char\n"

"#define ga_short short\n"

"#define ga_ushort unsigned short\n"

"#define ga_int int\n"

"#define ga_uint unsigned int\n"

"#define ga_long long long\n"

"#define ga_ulong unsigned long long\n"

"#define ga_float float\n"

"#define ga_double double\n"

"#define ga_half ga_ushort\n"

"#define ga_size size_t\n"

"#define ga_ssize ptrdiff_t\n"

"#define load_half(p) __half2float(*(p))\n"

"#define store_half(p, v) (*(p) = __float2half_rn(v))\n"

"#define GA_DECL_SHARED_PARAM(type, name)\n"

"#define GA_DECL_SHARED_BODY(type, name) extern __shared__ type name[];\n"

"#define GA_WARP_SIZE warpSize\n";

static cuda_context *do_init(CUdevice dev, int flags, int *ret) {

cuda_context *res;

CUcontext ctx;

unsigned int fl = CU_CTX_SCHED_AUTO;

#if CUDA_VERSION >= 7000

unsigned int cur_fl;

int act;

int i;

CHKFAIL(NULL);

if (flags & GA_CTX_SINGLE_THREAD)

fl = CU_CTX_SCHED_SPIN;

if (flags & GA_CTX_MULTI_THREAD)

fl = CU_CTX_SCHED_YIELD;

err = cuDeviceGetAttribute(&i, CU_DEVICE_ATTRIBUTE_UNIFIED_ADDRESSING, dev);

CHKFAIL(NULL);

if (i != 1)

FAIL(NULL, GA_UNSUPPORTED_ERROR);

#if CUDA_VERSION < 7000

err = cuCtxCreate(&ctx, fl, dev);

CHKFAIL(NULL);

err = cuDevicePrimaryCtxGetState(dev, &cur_fl, &act);

CHKFAIL(NULL);

if (act == 1) {

if ((cur_fl & fl) != fl)

FAIL(NULL, GA_INVALID_ERROR);

} else {

err = cuDevicePrimaryCtxSetFlags(dev, fl);

CHKFAIL(NULL);

}

err = cuDevicePrimaryCtxRetain(&ctx, dev);

CHKFAIL(NULL);

err = cuCtxPushCurrent(ctx);

CHKFAIL(NULL);

res = cuda_make_ctx(ctx, flags);

if (res == NULL) {

#if CUDA_VERSION < 7000

cuCtxDestroy(ctx);

cuDevicePrimaryCtxRelease(dev);

FAIL(NULL, GA_IMPL_ERROR);

}

/* Don't leave the context on the thread stack */

cuCtxPopCurrent(NULL);

return res;

}

static gpucontext *cuda_init(int ord, int flags, int *ret) {

CUdevice dev;

cuda_context *res;

if (!init_done) {

err = cuInit(0);

CHKFAIL(NULL);

init_done = 1;

}

if (ord == -1) {

int i, c;

err = cuDeviceGetCount(&c);

CHKFAIL(NULL);

for (i = 0; i < c; i++) {

err = cuDeviceGet(&dev, i);

CHKFAIL(NULL);

res = do_init(dev, flags, NULL);

if (res != NULL)

return (gpucontext *)res;

}

FAIL(NULL, GA_NODEV_ERROR);

} else {

err = cuDeviceGet(&dev, ord);

CHKFAIL(NULL);

return (gpucontext *)do_init(dev, flags, ret);

}

}

static void cuda_deinit(gpucontext *c) {

cuda_free_ctx((cuda_context *)c);

}

static void find_best(cuda_context *ctx, gpudata **best, gpudata **prev,

size_t size) {

gpudata *temp, *tempPrev = NULL;

*best = NULL;

for (temp = ctx->freeblocks; temp; temp = temp->next) {

if (temp->sz >= size && (!*best || temp->sz < (*best)->sz)) {

*best = temp;

*prev = tempPrev;

}

tempPrev = temp;

}

}

static int allocate(cuda_context *ctx, gpudata **res, gpudata **prev,

size_t size) {

CUdeviceptr ptr;

gpudata *next;

*prev = NULL;

if (!(ctx->flags & GA_CTX_DISABLE_ALLOCATION_CACHE))

if (size < BLOCK_SIZE) size = BLOCK_SIZE;

cuda_enter(ctx);

ctx->err = cuMemAlloc(&ptr, size);

if (ctx->err != CUDA_SUCCESS) {

cuda_exit(ctx);

return GA_IMPL_ERROR;

}

*res = new_gpudata(ctx, ptr, size);

cuda_exit(ctx);

if (*res == NULL) {

cuMemFree(ptr);

return GA_MEMORY_ERROR;

}

(*res)->flags |= CUDA_HEAD_ALLOC;

/* Now that the block is allocated, enter it in the freelist */

next = ctx->freeblocks;

for (; next && next->ptr < (*res)->ptr; next = next->next) {

*prev = next;

}

(*res)->next = next;

if (*prev)

(*prev)->next = *res;

else

ctx->freeblocks = *res;

return GA_NO_ERROR;

}

static int extract(gpudata *curr, gpudata *prev, size_t size) {

gpudata *next, *split;

size_t remaining = curr->sz - size;

if (remaining < FRAG_SIZE) {

/* No need to split, the remaining block would be too small */

next = curr->next;

} else {

split = new_gpudata(curr->ctx, curr->ptr + size, remaining);

if (split == NULL)

return GA_MEMORY_ERROR;

/* Make sure the chain keeps going */

split->next = curr->next;

curr->next = NULL;

/* Make sure we don't start using the split buffer too soon */

cuda_records(split, CUDA_WAIT_ALL, curr->ls);

next = split;

curr->sz = size;

}

if (prev != NULL)

prev->next = next;

else

curr->ctx->freeblocks = next;

return GA_NO_ERROR;

}

static void cuda_free(gpudata *);

static int cuda_write(gpudata *dst, size_t dstoff, const void *src,

size_t sz);

static inline size_t roundup(size_t s, size_t m) {

return ((s + (m - 1)) / m) * m;

}

static gpudata *cuda_alloc(gpucontext *c, size_t size, void *data, int flags,

int *ret) {

gpudata *res = NULL, *prev = NULL;

cuda_context *ctx = (cuda_context *)c;

size_t asize;

int err;

if ((flags & GA_BUFFER_INIT) && data == NULL) FAIL(NULL, GA_VALUE_ERROR);

if ((flags & (GA_BUFFER_READ_ONLY|GA_BUFFER_WRITE_ONLY)) ==

(GA_BUFFER_READ_ONLY|GA_BUFFER_WRITE_ONLY)) FAIL(NULL, GA_VALUE_ERROR);

/* TODO: figure out how to make this work */

if (flags & GA_BUFFER_HOST) FAIL(NULL, GA_DEVSUP_ERROR);

/* We don't want to manage really small allocations so we round up

if (!(ctx->flags & GA_CTX_DISABLE_ALLOCATION_CACHE)) {

asize = roundup(size, FRAG_SIZE);

find_best(ctx, &res, &prev, asize);

} else {

asize = size;

}

if (res == NULL) {

err = allocate(ctx, &res, &prev, asize);

if (err != GA_NO_ERROR)

FAIL(NULL, err);

}

err = extract(res, prev, asize);

if (err != GA_NO_ERROR)

FAIL(NULL, err);

/* It's out of the freelist, so add a ref */

res->ctx->refcnt++;

/* We consider this buffer allocated and ready to go */

res->refcnt = 1;

if (flags & GA_BUFFER_INIT) {

err = cuda_write(res, 0, data, size);

if (err != GA_NO_ERROR) {

cuda_free(res);

FAIL(NULL, err);

}

}

return res;

}

static void cuda_retain(gpudata *d) {

ASSERT_BUF(d);

d->refcnt++;

}

static void deallocate(gpudata *d) {

cuda_enter(d->ctx);

cuEventDestroy(d->rev);

cuEventDestroy(d->wev);

cuda_exit(d->ctx);

CLEAR(d);

free(d);

}

static void cuda_free(gpudata *d) {

/* We ignore errors on free */

ASSERT_BUF(d);

d->refcnt--;

if (d->refcnt == 0) {

/* Keep a reference to the context since we deallocate the gpudata

cuda_context *ctx = d->ctx;

if (d->flags & DONTFREE) {

/* This is the path for "external" buffers */

deallocate(d);

} else if (ctx->flags & GA_CTX_DISABLE_ALLOCATION_CACHE) {

/* Just free the pointer */

cuMemFree(d->ptr);

deallocate(d);

} else {

/* Find the position in the freelist. Freelist is kept in order

gpudata *next = d->ctx->freeblocks, *prev = NULL;

for (; next && next->ptr < d->ptr; next = next->next) {

prev = next;

}

next = prev != NULL ? prev->next : d->ctx->freeblocks;

/* See if we can merge the block with the previous one */

if (!(d->flags & CUDA_HEAD_ALLOC) &&

prev != NULL && prev->ptr + prev->sz == d->ptr) {

prev->sz = prev->sz + d->sz;

cuda_waits(d, CUDA_WAIT_ALL, prev->ls);

cuda_records(prev, CUDA_WAIT_ALL, prev->ls);

deallocate(d);

d = prev;

} else if (prev != NULL) {

prev->next = d;

} else {

d->ctx->freeblocks = d;

}

/* See if we can merge with next */

if (next && !(next->flags & CUDA_HEAD_ALLOC) &&

d->ptr + d->sz == next->ptr) {

d->sz = d->sz + next->sz;

d->next = next->next;

cuda_wait(next, CUDA_WAIT_ALL);

cuda_record(d, CUDA_WAIT_ALL);

deallocate(next);

} else {

d->next = next;

}

}

/* We keep this at the end since the freed buffer could be the

cuda_free_ctx(ctx);

}

}

static int cuda_share(gpudata *a, gpudata *b, int *ret) {

ASSERT_BUF(a);

ASSERT_BUF(b);

return (a->ctx == b->ctx && a->sz != 0 && b->sz != 0 &&

((a->ptr <= b->ptr && a->ptr + a->sz > b->ptr) ||

(b->ptr <= a->ptr && b->ptr + b->sz > a->ptr)));

}

static int cuda_waits(gpudata *a, int flags, CUstream s) {

ASSERT_BUF(a);

/* Never skip the wait if CUDA_WAIT_FORCE */

if (ISCLR(flags, CUDA_WAIT_FORCE)) {

if (ISSET(a->ctx->flags, GA_CTX_SINGLE_STREAM))

return GA_NO_ERROR;

/* If the last stream to touch this buffer is the same, we don't

if (a->ls == s)

return GA_NO_ERROR;

}

cuda_enter(a->ctx);

/* We wait for writes that happened before since multiple reads at

if (ISSET(flags, CUDA_WAIT_READ) || ISSET(flags, CUDA_WAIT_WRITE))

CUDA_EXIT_ON_ERROR(a->ctx, cuStreamWaitEvent(s, a->wev, 0));

/* Make sure to not disturb previous reads */

if (ISSET(flags, CUDA_WAIT_WRITE))

CUDA_EXIT_ON_ERROR(a->ctx, cuStreamWaitEvent(s, a->rev, 0));

cuda_exit(a->ctx);

return GA_NO_ERROR;

}

int cuda_wait(gpudata *a, int flags) {

return cuda_waits(a, flags, a->ctx->s);

}

static int cuda_records(gpudata *a, int flags, CUstream s) {

ASSERT_BUF(a);

if (ISCLR(flags, CUDA_WAIT_FORCE) &&

ISSET(a->ctx->flags, GA_CTX_SINGLE_STREAM))

return GA_NO_ERROR;

cuda_enter(a->ctx);

if (ISSET(flags, CUDA_WAIT_READ))

CUDA_EXIT_ON_ERROR(a->ctx, cuEventRecord(a->rev, s));

if (ISSET(flags, CUDA_WAIT_WRITE))

CUDA_EXIT_ON_ERROR(a->ctx, cuEventRecord(a->wev, s));

cuda_exit(a->ctx);

a->ls = s;

return GA_NO_ERROR;

}

int cuda_record(gpudata *a, int flags) {

return cuda_records(a, flags, a->ctx->s);

}

static int cuda_move(gpudata *dst, size_t dstoff, gpudata *src,

size_t srcoff, size_t sz) {

cuda_context *ctx = dst->ctx;

int res = GA_NO_ERROR;

ASSERT_BUF(dst);

ASSERT_BUF(src);

if (src->ctx != dst->ctx) return GA_VALUE_ERROR;

if (sz == 0) return GA_NO_ERROR;

if ((dst->sz - dstoff) < sz || (src->sz - srcoff) < sz)

return GA_VALUE_ERROR;

cuda_enter(ctx);

GA_CUDA_EXIT_ON_ERROR(ctx,

cuda_wait(src, CUDA_WAIT_READ));

GA_CUDA_EXIT_ON_ERROR(ctx,

cuda_wait(dst, CUDA_WAIT_WRITE));

CUDA_EXIT_ON_ERROR(ctx,

cuMemcpyDtoDAsync(dst->ptr + dstoff, src->ptr + srcoff, sz, ctx->s));

GA_CUDA_EXIT_ON_ERROR(ctx,

cuda_record(src, CUDA_WAIT_READ));

GA_CUDA_EXIT_ON_ERROR(ctx,

cuda_record(dst, CUDA_WAIT_WRITE));

cuda_exit(ctx);

return res;

}

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

cuda_context *ctx = src->ctx;

ASSERT_BUF(src);

if (sz == 0) return GA_NO_ERROR;

if ((src->sz - srcoff) < sz)

return GA_VALUE_ERROR;

cuda_enter(ctx);

if (src->flags & CUDA_MAPPED_PTR) {

if (ISSET(ctx->flags, GA_CTX_SINGLE_STREAM))

ctx->err = cuStreamSynchronize(ctx->s);

else

ctx->err = cuEventSynchronize(src->wev);

if (ctx->err != CUDA_SUCCESS) {

cuda_exit(ctx);

return GA_IMPL_ERROR;

}

memcpy(dst, (void *)(src->ptr + srcoff), sz);

} else {

GA_CUDA_EXIT_ON_ERROR(ctx,

cuda_waits(src, CUDA_WAIT_READ, ctx->mem_s));

CUDA_EXIT_ON_ERROR(ctx,

cuMemcpyDtoHAsync(dst, src->ptr + srcoff, sz, ctx->mem_s));

GA_CUDA_EXIT_ON_ERROR(ctx,

cuda_records(src, CUDA_WAIT_READ, ctx->mem_s));

}

cuda_exit(ctx);

return GA_NO_ERROR;

}

static int cuda_write(gpudata *dst, size_t dstoff, const void *src,

size_t sz) {

cuda_context *ctx = dst->ctx;

ASSERT_BUF(dst);

if (sz == 0) return GA_NO_ERROR;

if ((dst->sz - dstoff) < sz)

return GA_VALUE_ERROR;

cuda_enter(ctx);

if (dst->flags & CUDA_MAPPED_PTR) {

if (ISSET(ctx->flags, GA_CTX_SINGLE_STREAM))

ctx->err = cuStreamSynchronize(ctx->s);

else

ctx->err = cuEventSynchronize(dst->rev);

if (ctx->err != CUDA_SUCCESS) {

cuda_exit(ctx);

return GA_IMPL_ERROR;

}

memcpy((void *)(dst->ptr + dstoff), src, sz);

} else {

GA_CUDA_EXIT_ON_ERROR(ctx,

cuda_waits(dst, CUDA_WAIT_WRITE, ctx->mem_s));

CUDA_EXIT_ON_ERROR(ctx,

cuMemcpyHtoDAsync(dst->ptr + dstoff, src, sz, ctx->mem_s));

GA_CUDA_EXIT_ON_ERROR(ctx,

cuda_records(dst, CUDA_WAIT_WRITE, ctx->mem_s));

}

cuda_exit(ctx);

return GA_NO_ERROR;

}

static int cuda_memset(gpudata *dst, size_t dstoff, int data) {

cuda_context *ctx = dst->ctx;

ASSERT_BUF(dst);

if ((dst->sz - dstoff) == 0) return GA_NO_ERROR;

cuda_enter(ctx);

GA_CUDA_EXIT_ON_ERROR(ctx,

cuda_wait(dst, CUDA_WAIT_WRITE));

CUDA_EXIT_ON_ERROR(ctx,

cuMemsetD8Async(dst->ptr + dstoff, data, dst->sz - dstoff, ctx->s));

GA_CUDA_EXIT_ON_ERROR(ctx,

cuda_record(dst, CUDA_WAIT_WRITE));

cuda_exit(ctx);

return GA_NO_ERROR;

}

static CUresult get_cc(CUdevice dev, int *maj, int *min) {

#if CUDA_VERSION < 6500

return cuDeviceComputeCapability(maj, min, dev);

CUresult lerr;

lerr = cuDeviceGetAttribute(maj,

CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR,

dev);

if (lerr != CUDA_SUCCESS)

return lerr;

return cuDeviceGetAttribute(min,

CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR,

dev);

}

static int detect_arch(const char *prefix, char *ret, CUresult *err) {

CUdevice dev;

int major, minor;

int res;

size_t sz = strlen(prefix) + 3;

*err = cuCtxGetDevice(&dev);

if (*err != CUDA_SUCCESS) return GA_IMPL_ERROR;

*err = get_cc(dev, &major, &minor);

if (*err != CUDA_SUCCESS) return GA_IMPL_ERROR;

res = snprintf(ret, sz, "%s%d%d", prefix, major, minor);

if (res == -1 || res > sz) return GA_UNSUPPORTED_ERROR;

return GA_NO_ERROR;

}

#ifdef WITH_NVRTC

#include <nvrtc.h>

static void *call_compiler(const char *src, size_t len, const char *arch_arg,

size_t *bin_len, char **log, size_t *log_len,

int *ret) {

nvrtcProgram prog;

void *buf = NULL;

size_t buflen;

const char *opts[4] = {

"-arch", ""

, "-G", "-lineinfo"

};

nvrtcResult err, err2;

opts[1] = arch_arg;

err = nvrtcCreateProgram(&prog, src, NULL, 0, NULL, NULL);

if (err != NVRTC_SUCCESS) FAIL(NULL, GA_SYS_ERROR);

err = nvrtcCompileProgram(prog,

#ifdef DEBUG

4,

2,

opts);

if (log != NULL) {

err2 = nvrtcGetProgramLogSize(prog, &buflen);

if (err2 != NVRTC_SUCCESS) goto end2;

buf = malloc(buflen);

if (buf == NULL) goto end2;

err2 = nvrtcGetProgramLog(prog, (char *)buf);

if (err2 != NVRTC_SUCCESS) goto end2;

if (log_len != NULL) *log_len = buflen;

*log = (char *)buf;

buf = NULL;

}

end2:

if (err != NVRTC_SUCCESS) goto end;

err = nvrtcGetPTXSize(prog, &buflen);

if (err != NVRTC_SUCCESS) goto end;

buf = malloc(buflen);

if (buf == NULL) {

nvrtcDestroyProgram(&prog);

FAIL(NULL, GA_MEMORY_ERROR);

}

err = nvrtcGetPTX(prog, (char *)buf);

if (err != NVRTC_SUCCESS) goto end;

*bin_len = buflen;

end:

nvrtcDestroyProgram(&prog);

if (err != NVRTC_SUCCESS) {

free(buf);

FAIL(NULL, GA_SYS_ERROR);

}

return buf;

}

#else /* WITH_NVRTC */

#include <sys/stat.h>

#include <fcntl.h>

#include <limits.h>

#ifdef _WIN32

#include <process.h>

#define PATH_MAX 256

#include <sys/param.h>

#include <sys/wait.h>

#ifdef _MSC_VER

#include <io.h>

#define read _read

#define write _write

#define close _close

#define unlink _unlink

#define fstat _fstat

#define open _open

#include <unistd.h>

static const char *TMP_VAR_NAMES[] = {"GPUARRAY_TMPDIR", "TMPDIR", "TMP",

"TEMP", "USERPROFILE"};

static void *call_compiler(const char *src, size_t len, const char *arch_arg,

size_t *bin_len, char **log, size_t *log_len,

int *ret) {

char namebuf[PATH_MAX];

char outbuf[PATH_MAX];

char *tmpdir;

struct stat st;

ssize_t s;

#ifndef _WIN32

pid_t p;

unsigned int i;

int sys_err;

int fd;

char *buf;

for (i = 0; i < sizeof(TMP_VAR_NAMES)/sizeof(TMP_VAR_NAMES[0]); i++) {

tmpdir = getenv(TMP_VAR_NAMES[i]);

if (tmpdir != NULL) break;

}

if (tmpdir == NULL) {

#ifdef _WIN32

tmpdir = ".";

tmpdir = "/tmp";

}

strlcpy(namebuf, tmpdir, sizeof(namebuf));

strlcat(namebuf, "/gpuarray.cuda.XXXXXXXX", sizeof(namebuf));

fd = mkstemp(namebuf);

if (fd == -1) FAIL(NULL, GA_SYS_ERROR);

strlcpy(outbuf, namebuf, sizeof(outbuf));

strlcat(outbuf, ".cubin", sizeof(outbuf));

/* Don't want to write the final NUL */

s = write(fd, src, len-1);

close(fd);

/* fd is not non-blocking; should have complete write */

if (s == -1) {

unlink(namebuf);

FAIL(NULL, GA_SYS_ERROR);

}

/* This block executes nvcc on the written-out file */

#ifdef DEBUG

#define NVCC_ARGS NVCC_BIN, "-g", "-G", "-arch", arch_arg, "-x", "cu", \

"--cubin", namebuf, "-o", outbuf

#define NVCC_ARGS NVCC_BIN, "-arch", arch_arg, "-x", "cu", \

"--cubin", namebuf, "-o", outbuf

#ifdef _WIN32

sys_err = _spawnl(_P_WAIT, NVCC_BIN, NVCC_ARGS, NULL);

unlink(namebuf);

if (sys_err == -1) FAIL(NULL, GA_SYS_ERROR);

if (sys_err != 0) FAIL(NULL, GA_RUN_ERROR);