Хочу реализовать межблочный барьер на CUDA, но столкнулся с серьезной проблемой.
Я не могу понять, почему это не работает.
#include <iostream>
#include <cstdlib>
#include <ctime>
#define SIZE 10000000
#define BLOCKS 100
using namespace std;
struct Barrier {
int *count;
__device__ void wait() {
atomicSub(count, 1);
while(*count)
;
}
Barrier() {
int blocks = BLOCKS;
cudaMalloc((void**) &count, sizeof(int));
cudaMemcpy(count, &blocks, sizeof(int), cudaMemcpyHostToDevice);
}
~Barrier() {
cudaFree(count);
}
};
__global__ void sum(int* vec, int* cache, int *sum, Barrier barrier)
{
int tid = blockIdx.x;
int temp = 0;
while(tid < SIZE) {
temp += vec[tid];
tid += gridDim.x;
}
cache[blockIdx.x] = temp;
barrier.wait();
if(blockIdx.x == 0) {
for(int i = 0 ; i < BLOCKS; ++i)
*sum += cache[i];
}
}
int main()
{
int* vec_host = (int *) malloc(SIZE * sizeof(int));
for(int i = 0; i < SIZE; ++i)
vec_host[i] = 1;
int *vec_dev;
int *sum_dev;
int *cache;
int sum_gpu = 0;
cudaMalloc((void**) &vec_dev, SIZE * sizeof(int));
cudaMemcpy(vec_dev, vec_host, SIZE * sizeof(int), cudaMemcpyHostToDevice);
cudaMalloc((void**) &sum_dev, sizeof(int));
cudaMemcpy(sum_dev, &sum_gpu, sizeof(int), cudaMemcpyHostToDevice);
cudaMalloc((void**) &cache, BLOCKS * sizeof(int));
cudaMemset(cache, 0, BLOCKS * sizeof(int));
Barrier barrier;
sum<<<BLOCKS, 1>>>(vec_dev, cache, sum_dev, barrier);
cudaMemcpy(&sum_gpu, sum_dev, sizeof(int), cudaMemcpyDeviceToHost);
cudaFree(vec_dev);
cudaFree(sum_dev);
cudaFree(cache);
free(vec_host);
return 0;
}
На самом деле, даже если я перепишу wait() следующим образом
__device__ void wait() {
while(*count != 234124)
;
}
Программа завершается нормально. Но я ожидаю получить бесконечный цикл в этом случае.
thrust::reduce. 10.10.2011