сумма векторов значений с CUDA C ++

Question

Я пытаюсь суммировать значения многих векторов, используя CUDA c ++.Я нашел какое-то решение для двух векторов.Как видите, просто можно добавить два вектора, но я хочу динамически генерировать векторы одинаковой длины.

#include <stdio.h>
#include <stdlib.h>
#include <math.h>

// CUDA kernel. Each thread takes care of one element of c
 __global__ void vecAdd(double *a, double *b, double *c, int n)
{
// Get our global thread ID
int id = blockIdx.x*blockDim.x+threadIdx.x;
// Make sure we do not go out of bounds
if (id < n)
    c[id] = a[id] + b[id];
}

int main( int argc, char* argv[] )
{
// Size of vectors
int n = 100000;

// Host input vectors
double *h_a;
double *h_b;
//Host output vector
double *h_c;

// Device input vectors
double *d_a;
double *d_b;
//Device output vector
double *d_c;

// Size, in bytes, of each vector
size_t bytes = n*sizeof(double);

// Allocate memory for each vector on host
h_a = (double*)malloc(bytes);
h_b = (double*)malloc(bytes);
h_c = (double*)malloc(bytes);

// Allocate memory for each vector on GPU
cudaMalloc(&d_a, bytes);
cudaMalloc(&d_b, bytes);
cudaMalloc(&d_c, bytes);

int i;
// Initialize vectors on host
for( i = 0; i < n; i++ ) {
    h_a[i] = sin(i)*sin(i);
    h_b[i] = cos(i)*cos(i);
}

// Copy host vectors to device
cudaMemcpy( d_a, h_a, bytes, cudaMemcpyHostToDevice);
cudaMemcpy( d_b, h_b, bytes, cudaMemcpyHostToDevice);

int blockSize, gridSize;

// Number of threads in each thread block
blockSize = 1024;

// Number of thread blocks in grid
gridSize = (int)ceil((float)n/blockSize);

// Execute the kernel
vecAdd<<<gridSize, blockSize>>>(d_a, d_b, d_c, n);

// Copy array back to host
cudaMemcpy( h_c, d_c, bytes, cudaMemcpyDeviceToHost );

// Sum up vector c and the print result divided by n, this should equal 1 
within error
double sum = 0;
for(i=0; i<n; i++)
    sum += h_c[i];
printf("final result: %f\n", sum/n);

// Release device memory
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);

// Release host memory
free(h_a);
free(h_b);
free(h_c);

return 0;
}

Есть ли способ сделать это для многих векторов?Размер моего вектора:

#vector length
N = 1000 
#number of vectors
i = 300000
v[i] = [1,2,..., N]

В результате мне нужно получить:

out[i]= [sum(v[1]), sum(v[2]),..., sum(v[i])]

Спасибо за любой совет.

Answer 1

Суммирование нескольких векторов вместе способом, подобным показанному вами коду (т. Е. Генерация поэлементных сумм) эквивалентно суммированию столбцов матрицы.И эта идея представляет разумный способ реализовать решение.

Мы будем рассматривать ваши векторы как матрицу, где каждый вектор является строкой в ​​матрице.Ядро CUDA назначит по одному потоку для каждого столбца и суммирует элементы этого столбца, получая результат с одним числом.Этот результат с одним числом станет одним из элементов векторного результата всей задачи.

Вот полностью проработанный пример, демонстрирующий один из возможных подходов:

$ cat t2.cu
#include <iostream>

typedef double mt;
const int nTPB = 64;

template <typename T>
__global__ void column_sum(T *matrix, T *sums, unsigned n_vectors, unsigned vector_length){

  unsigned idx = threadIdx.x+blockDim.x*blockIdx.x;
  if (idx < vector_length){
    T temp = 0;
    for (unsigned i = 0; i < n_vectors; i++)
      temp += matrix[i*vector_length+idx];
    sums[idx] = temp;}
}

int main(){
  const unsigned vlen = 1000;
  const unsigned nvec = 300000;
  mt *h_matrix, *d_matrix, *h_sums, *d_sums;
  // create the desired number of vectors as a single matrix
  h_sums = new mt[vlen];
  h_matrix = new mt[vlen*nvec];
  cudaMalloc(&d_matrix, vlen*nvec*sizeof(mt));
  cudaMalloc(&d_sums, vlen*sizeof(mt));
  size_t count = 0;
  for (unsigned i = 0; i < nvec; i++)
    for (unsigned j = 0; j < vlen; j++)
      h_matrix[count++] = j;
  cudaMemcpy(d_matrix, h_matrix, vlen*nvec*sizeof(mt), cudaMemcpyHostToDevice);
  column_sum<<<(vlen+nTPB-1)/nTPB,nTPB>>>(d_matrix, d_sums, nvec, vlen);
  cudaMemcpy(h_sums, d_sums, vlen*sizeof(mt), cudaMemcpyDeviceToHost);
  for (unsigned i = 0; i < vlen; i++) if (h_sums[i] != ((mt)nvec)*i) {std::cout << " mismatch at " << i << " was: " << h_sums[i] << " should be: " << ((mt)nvec)*i << std::endl; return -1;}
  std::cout << cudaGetErrorString(cudaGetLastError()) << std::endl;
}
$ nvcc -o t2 t2.cu
$ cuda-memcheck ./t2
========= CUDA-MEMCHECK
no error
========= ERROR SUMMARY: 0 errors
$

Обратите внимание, что эта методология создает только какмного потоков на GPU, поскольку есть векторные элементы (1000 в приведенном выше примере).1000 потоков будет достаточно, чтобы заняты только самые маленькие графические процессоры.Однако этот алгоритм будет эффективен на большинстве графических процессоров, если длина вашего вектора составляет 10000 или больше.Если вы хотите исследовать создание более эффективных алгоритмов для задач небольшого размера, вы можете изучить идею классического параллельного сокращения .