mirror-github-bspeice/cuburn
1
0
Fork 0
mirror of https://github.com/stevenrobertson/cuburn.git synced 2025-02-05 11:40:04 -05:00
Code Issues Actions Packages Projects Releases Wiki Activity

Fix rb_incr() when blockDim.y == 1.

Browse Source
This commit is contained in:
Steven Robertson 2012-02-15 10:06:35 -05:00
parent 8210c8fc73
commit 88abefa4f4
2 changed files with 5 additions and 6 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
cuburn/code/interp.py
View File

@ -325,7 +325,6 @@ __global__ void interp_palette_flat(
float tstart, float tstep) float tstart, float tstep)
{ {
mwc_st rctx = rctxs[rb_incr(rb->head, threadIdx.x)]; mwc_st rctx = rctxs[rb_incr(rb->head, threadIdx.x)];
int gid = blockIdx.x * blockDim.x + threadIdx.x;
float time = tstart + blockIdx.x * tstep; float time = tstart + blockIdx.x * tstep;
float4 yuva = interp_color(times, sources, time); float4 yuva = interp_color(times, sources, time);

10
cuburn/code/util.py
View File

@ -280,13 +280,13 @@ def mkringbuflib(rb_size):
""" """
return devlib(headers="#define RB_SIZE_MASK %d" % (rb_size - 1), decls=''' return devlib(headers="#define RB_SIZE_MASK %d" % (rb_size - 1), decls='''
typedef struct { typedef struct {
int head; uint32_t head;
int tail; uint32_t tail;
} ringbuf; } ringbuf;
''', defs=r''' ''', defs=r'''
__shared__ int rb_idx; __shared__ uint32_t rb_idx;
__device__ int rb_incr(int &rb_base, int tidx) { __device__ uint32_t rb_incr(uint32_t &rb_base, int tidx) {
if (threadIdx.y == 1 && threadIdx.x == 1) if (threadIdx.y == 0 && threadIdx.x == 0)
rb_idx = 256 * (atomicAdd(&rb_base, 1) & RB_SIZE_MASK); rb_idx = 256 * (atomicAdd(&rb_base, 1) & RB_SIZE_MASK);
__syncthreads(); __syncthreads();
return rb_idx + tidx; return rb_idx + tidx;