Ваша основная проблема в том, что ядро вашего ядра не имеет смысла.То, что у вас есть:
for(int i = 0; i<256; i++)
C[i] = A[ix+iy*16] + B[ix+iy*16] + C[i]; // << I wish to store all in C
Это позволит каждому потоку пройти и прочитать каждую запись в C, добавить к ней свои собственные части A и B и записать обратно.Поскольку каждый поток делает это одновременно, они собираются наступить друг на друга.Если вы действительно хотите, чтобы каждая запись в C была суммой всех записей в A и всех записей в B, вы хотите, чтобы каждый поток отвечал заопределенная запись в C:
for(int i = 0; i<256; i++)
C[ix+iy*16] += A[i] + B[i];
Если вместо этого вы хотите, чтобы каждая запись в C была суммой соответствующих записей в A и B, что представляется более вероятным, то вы получитеизбавьтесь от цикла, и ваше ядро будет выглядеть следующим образом:
__global__ void ADD(float* A, float* B, float* C)
{
const int ix = blockDim.x * blockIdx.x + threadIdx.x;
const int iy = blockDim.y * blockIdx.y + threadIdx.y;
if(ix < 16 && iy < 16)
{
C[ix+iy*16] = A[ix+iy*16] + B[ix+iy*16];
}
}
Каждый поток захватывает одну запись из A и одну из B и записывает одну запись в C.
Ваша дополнительная проблемачто вы неправильно запускаете ядро.Вы делаете:
ADDD<<<16,16>>>(dev_A,dev_B,dev_C);
Это запускает сетку размером 1x16 блоков по 1x16 потоков каждый (из ядра typo'd).Если вы хотите, чтобы ваши потоки располагались в двух измерениях (с использованием обоих индексов x и y), вам нужно использовать dim3 в качестве типа спецификатора размера.Что-то вроде:
// Use a grid of 4x4 blocks
dim3 gridSize;
gridSize.x = 4;
gridSize.y = 4;
// Use blocks of 4x4 threads.
dim3 blockSize;
blockSize.x = 4;
blockSize.y = 4;
// Run a 4x4 grid of blocks, each with 4x4 threads.
// So you end up with a 16x16 group of threads, matching your data layout.
ADD<<<gridSize,blockSize>>>(dev_A,dev_B,dev_C);