__device__ float produce(int iteration, unsigned idx);
__device__ void consume(int iteration, unsigned idx, float value);
__global__ void kernel(int iterations)
{
/*
* Tuning parameters: Number of samples to buffer
* and number of producer vs. consumer threads
*/
constexpr unsigned bufsize = 128;
const unsigned producer_count = 32;
const unsigned consumer_count = blockDim.x - producer_count;
/*
* Simple double-buffer setup to make a producer-consumer pipeline
*/
__shared__ float buf[2][bufsize];
const bool is_producer = threadIdx.x < producer_count;
/*
* Fill first buffer. This can use all threads
*/
for(unsigned i = threadIdx.x; i < bufsize; i += blockDim.x)
buf[0][i] = produce(0, i);
for(int i = 1; i < iterations; ++i) {
__syncthreads();
int cur_buf = i & 1;
int last_buf = (i - 1) & 1;
if(is_producer) {
// fill next buffer
for(unsigned j = threadIdx.x; j < bufsize; j += producer_count)
buf[cur_buf][j] = produce(i, j);
}
else {
const unsigned consumer_idx = threadIdx.x - producer_count;
// consume last buffer
for(unsigned j = consumer_idx; j < bufsize; j += consumer_count)
consume(i - 1, j, buf[last_buf][j]);
}
}
__syncthreads();
/*
* Consume last buffer. Can use all threads
*/
int last_buf = (iterations - 1) & 1;
for(unsigned i = threadIdx.x; i < bufsize; i += blockDim.x)
consume(iterations - 1, i, buf[last_buf][i]);
}
Comments
0 B
|0 👍
/0 👎