Операция, которую необходимо использовать для выполнения глобального суммирования во всех потоках, называется «параллельным сокращением». Хотя вы могли бы использовать атомарные операции для этого, я бы не рекомендовал это. В CUDA SDK есть сокращенное ядро и очень хорошая статья, в которой обсуждается техника, которую стоит прочитать.
Если бы я писал код, чтобы делать то, что вы хотите, он, вероятно, выглядел бы так:
template <int blocksize>
__global__ void calcRatio(float *orig, float *modified, int size, float *result,
int *count, const float error)
{
__shared__ volatile float buff[blocksize];
int index = threadIdx.x + blockIdx.x * blockDim.x;
int stride = blockDim.x * gridDim.x;
int count = 0;
for(int i=index; i<n; i+=stride) {
val = orig[index] - modified[index];
count += (val < error);
result[index] = val;
}
buff[threadIdx.x] = count;
__syncthreads();
// Parallel reduction in shared memory using 1 warp
if (threadId.x < warpSize) {
for(int i=threadIdx.x + warpSize; i<blocksize; i+= warpSize) {
buff[threadIdx.x] += buff[i];
if (threadIdx.x < 16) buff[threadIdx.x] +=buff[threadIdx.x + 16];
if (threadIdx.x < 8) buff[threadIdx.x] +=buff[threadIdx.x + 8];
if (threadIdx.x < 4) buff[threadIdx.x] +=buff[threadIdx.x + 4];
if (threadIdx.x < 2) buff[threadIdx.x] +=buff[threadIdx.x + 2];
if (threadIdx.x == 0) count[blockIdx.x] = buff[0] + buff[1];
}
}
Первый раздел выполняет то, что делает ваш последовательный код - вычисляет разницу и локальный поток общее количество элементов, которые меньше ошибки. Примечание. Я написал эту версию так, что каждый поток предназначен для обработки более одной записи входных данных. Это было сделано, чтобы помочь компенсировать вычислительные затраты на параллельное сокращение, которое следует, и идея состоит в том, что вы будете использовать меньше блоков и потоков, чем было во входных наборах данных.
Второй раздел - это само сокращение, выполняемое в общей памяти. По сути, это «древовидная» операция, в которой размер набора локальных промежуточных итогов потока в пределах одного блока потоков сначала суммируется до 32 промежуточных итогов, затем промежуточные итоги объединяются до тех пор, пока не будет получен окончательный промежуточный итог для блока, и что затем сохраняется сумма для блока . Вы получите небольшой список промежуточных итогов, по одному для каждого блока, который вы запустили, который можно скопировать обратно на хост и рассчитать там конечный результат, который вам нужен.
Обратите внимание, что я закодировал это в браузере и не скомпилировал его, могут быть ошибки, но это должно дать представление о том, как будет работать "продвинутая" версия того, что вы пытаетесь сделать.