Проблема синхронизации деформации CUDA
Обобщая ядро, которое сдвигает значения 2D-массива на один пробел вправо (обтекание границ строк), я столкнулся с проблемой синхронизации деформации. Полный код прилагается и включен ниже.
Код предназначен для работы с произвольной шириной массива, высотой массива, количеством блоков потоков и числом потоков в блоке. При выборе размера потока 33 (то есть на один поток больше, чем полный перекос), 33-й поток не синхронизируется с __syncthreads() называется. Это вызывает проблемы с выходными данными. Проблема возникает только тогда, когда существует более одного перекоса, а ширина массива превышает количество потоков (например, с шириной =35 и 34 потока).
Ниже приведен уменьшенный пример того, что происходит (в действительности массив должен иметь больше элементов, чтобы ядро могло выдать ошибку).
Начальный массив:
0 1 2 3 4
5 6 7 8 9
Ожидаемый результат:
4 0 1 2 3
9 5 6 7 8
Ядро Производит:
4 0 1 2 3
8 5 6 7 8
Первая строка выполняется правильно (для каждого блока, если их более одного), причем все последующие строки имеют повторное второе последнее значение. Я проверил эту две разные карты (8600GT и GTX280) и получаю одинаковые результаты. Я хотел бы знать, является ли это просто ошибкой в моем ядре, или проблема, которую нельзя исправить, корректируя мой код?
Полный исходный файл включен ниже.
Спасибо.
#include <cstdio>
#include <cstdlib>
// A method to ensure all reads use the same logical layout.
inline __device__ __host__ int loc(int x, int y, int width)
{
return y*width + x;
}
//kernel to shift all items in a 2D array one position to the right (wrapping around rows)
__global__ void shiftRight ( int* globalArray, int width, int height)
{
int temp1=0; //temporary swap variables
int temp2=0;
int blockRange=0; //the number of rows that a single block will shift
if (height%gridDim.x==0) //logic to account for awkward array sizes
blockRange = height/gridDim.x;
else
blockRange = (1+height/gridDim.x);
int yStart = blockIdx.x*blockRange;
int yEnd = yStart+blockRange; //the end condition for the y-loop
yEnd = min(height,yEnd); //make sure that the array doesn't go out of bounds
for (int y = yStart; y < yEnd ; ++y)
{
//do the first read so the swap variables are loaded for the x-loop
temp1 = globalArray[loc(threadIdx.x,y,width)];
//Each block shifts an entire row by itself, even if there are more columns than threads
for (int threadXOffset = threadIdx.x ; threadXOffset < width ; threadXOffset+=blockDim.x)
{
//blockDim.x is added so that we store the next round of values
//this has to be done now, because the next operation will
//overwrite one of these values
temp2 = globalArray[loc((threadXOffset + blockDim.x)%width,y,width)];
__syncthreads(); //sync before the write to ensure all the values have been read
globalArray[loc((threadXOffset +1)%width,y,width)] = temp1;
__syncthreads(); //sync after the write so ensure all the values have been written
temp1 = temp2; //swap the storage variables.
}
if (threadIdx.x == 0 && y == 0)
globalArray[loc(12,2,width)]=globalArray[67];
}
}
int main (int argc, char* argv[])
{
//set the parameters to be used
int width = 34;
int height = 3;
int threadsPerBlock=33;
int numBlocks = 1;
int memSizeInBytes = width*height*sizeof(int);
//create the host data and assign each element of the array to equal its index
int* hostData = (int*) malloc (memSizeInBytes);
for (int y = 0 ; y < height ; ++y)
for (int x = 0 ; x < width ; ++x)
hostData [loc(x,y,width)] = loc(x,y,width);
//create an allocate the device pointers
int* deviceData;
cudaMalloc ( &deviceData ,memSizeInBytes);
cudaMemset ( deviceData,0,memSizeInBytes);
cudaMemcpy ( deviceData, hostData, memSizeInBytes, cudaMemcpyHostToDevice);
cudaThreadSynchronize();
//launch the kernel
shiftRight<<<numBlocks,threadsPerBlock>>> (deviceData, width, height);
cudaThreadSynchronize();
//copy the device data to a host array
int* hostDeviceOutput = (int*) malloc (memSizeInBytes);
cudaMemcpy (hostDeviceOutput, deviceData, memSizeInBytes, cudaMemcpyDeviceToHost);
cudaFree (deviceData);
//Print out the expected/desired device output
printf("---- Expected Device Output ----\n");
printf(" | ");
for (int x = 0 ; x < width ; ++x)
printf("%4d ",x);
printf("\n---|-");
for (int x = 0 ; x < width ; ++x)
printf("-----");
for (int y = 0 ; y < height ; ++y)
{
printf("\n%2d | ",y);
for (int x = 0 ; x < width ; ++x)
printf("%4d ",hostData[loc((x-1+width)%width,y,width)]);
}
printf("\n\n");
printf("---- Actual Device Output ----\n");
printf(" | ");
for (int x = 0 ; x < width ; ++x)
printf("%4d ",x);
printf("\n---|-");
for (int x = 0 ; x < width ; ++x)
printf("-----");
for (int y = 0 ; y < height ; ++y)
{
printf("\n%2d | ",y);
for (int x = 0 ; x < width ; ++x)
printf("%4d ",hostDeviceOutput[loc(x,y,width)]);
}
printf("\n\n");
}
2 ответа
Поскольку не все потоки выполняют одинаковое количество итераций цикла, синхронизация является проблемой! Все потоки должны постоянно попадать в одни и те же __syncthreads().
Я бы предложил преобразовать ваш внутренний цикл for в что-то вроде этого:
for(int blockXOffset=0; blockXOffset < width; blockXOffset+=blockDim.x) {
int threadXOffset=blockXOffset+threadIdx.x;
bool isActive=(threadXOffset < width);
if (isActive) temp2 = globalArray[loc((threadXOffset + blockDim.x)%width,y,width)];
__syncthreads();
if (isActive) globalArray[loc((threadXOffset +1)%width,y,width)] = temp1;
__syncthreads();
temp1 = temp2;
}
Из руководства по программированию:
__syncthreads()допускается в условном коде, но только в том случае, если условное вычисление выполняется одинаково по всему блоку потока, в противном случае выполнение кода может зависнуть или привести к непреднамеренным побочным эффектам.
В моем примере не все потоки выполняют одинаковое количество итераций цикла, поэтому синхронизация не происходит.