Вычисление среднего 2000 2D-массивов с CUDA C

Question

У меня 2000 2D-массивов (каждый массив 1000x1000) Мне нужно вычислить среднее значение каждого из них и поместить результат в один вектор 2000.

Я попытался сделать это, вызвав ядро ​​для каждого 2D-массива, но это наивно. Я хочу сделать вычисление сразу.

Это то, что я сделал, является ядром для одного 2D-массива. Я хочу сделать ядро, чтобы сделать это для 2000 2D-массивов в одном ядре.

#include <stdio.h> 
#include <cuda.h> 
#include <time.h> 

void init_mat(float *a, const int N, const int M); 
void print_mat(float *a, const int N, const int M, char *d); 
void print_array(float *a, const int N, char *d); 









const int threadsPerBlock=256; 

__global__ 
void kernel(float *mat, float *out, const int N, const int M){ 

    __shared__ float cache[threadsPerBlock]; 

    int tid=threadIdx.x+blockIdx.x*blockDim.x; 
    int cacheIndex = threadIdx.x; 
    float sum=0; 



    if(tid<M){ 
     for(int i=0; i<N; i++) 
      sum += mat[(i*M)+tid]; 
      cache[cacheIndex] = sum; 
      out[tid] =cache[cacheIndex]; 
    } 

    __syncthreads(); 

    int i = blockDim.x/2; 
    while(i!=0){ 
     if(cacheIndex<i) 
      cache[cacheIndex]+= cache[cacheIndex +i]; 
     __syncthreads(); 
     i/=2; 
    } 

    if(cacheIndex==0) 
      out[blockIdx.x]=cache[0]; 


} 




int main (void) { 


    srand(time(NULL)); 

      float *a, *b, *c; 
       float *dev_a, *dev_b, *dev_c; 

      int N=1000; 
      int M=1000; 

      b=(float*)malloc(sizeof(float)*N*M); 
      c=(float*)malloc(sizeof(float)*M); 

       init_mat(b, N, M); 

      printf("<<<<<<<<<< initial data:\n"); 

       print_mat(b, N, M, "matrix"); 



       cudaMalloc((void**)&dev_b, sizeof(float)*N*M); 
       cudaMalloc((void**)&dev_c, sizeof(float)*M); 


       cudaMemcpy(dev_b, b, sizeof(float)*N*M, cudaMemcpyHostToDevice); 

      printf("\n\nRunning Kernel...\n\n"); 
       kernel<<<M/256+1, 256>>>(dev_b, dev_c, N, M); 


       cudaMemcpy(c, dev_c, sizeof(float)*M, cudaMemcpyDeviceToHost); 

       cudaFree(dev_a); 
       cudaFree(dev_b); 
       cudaFree(dev_c); 

      printf(">>>>>>>>>> final data:\n"); 
       print_array(c, M, "out-vector"); 
}; 




void init_mat(float *a, const int N, const int M) { 
     int i, j; 
     for(i=0; i<N; i++) 
      for(j=0; j<M; j++) 
        a[i*M+j] = rand() % 100 + 1; 
} 
void print_mat(float *a, const int N, const int M, char *d) { 
     int i, j; 
     for(i=0; i<N; i++){ 
     printf("\n%s[%d]:", d, i); 
     for (j=0; j<M; j++) 
        printf("\t%6.4f", a[i*M+j]); 
    } 
    printf("\n"); 
} 

void print_array(float *a, const int N, char *d) { 
     int i; 
     for(i=0; i<N; i++) 
       printf("\n%s[%d]: %f",d, i, a[i]); 
    printf("\n"); 
} 

Любая помощь Спасибо вам

Answer 1

Для достаточно большого количества массивов (например, 2000) и достаточно большого размера массивов (например, 2000), графический процессор может быть довольно эффективным, если мы относим блок для выполнения сокращение суммы (и средний расчет) для каждого массива. Это означает, что если у вас 2000 массивов, мы запустим 2000 блоков.

Для того, чтобы обрабатывать произвольные массивы размера с фиксированным числом витков на блок, мы будем использовать идею вроде grid-striding loop, но вместо этого мы будем вызывать каждый блок использовать блок-шагающие цикла, чтобы загрузить все данные, связанные с определенным массивом. Это означает, что потоки каждого блока будут «перемещаться» через назначенный массив, чтобы загрузить все элементы этого массива.

Кроме того, основная операция сокращения похожа на то, что вы написали, и вычисление среднего тривиально таким образом - мы можем вычислить среднее значение до того, как записать результат в глобальную память, как только мы рассчитаем сумму снижение.

Вот пример работы. При компиляции с -DMEAN код выводит среднее значение для каждого массива. Если вы опустите этот компилятор, код выведет сумму каждого массива. Пусть N - количество массивов, и пусть K будет размером каждого массива.

$ cat t1285.cu 
#include <stdio.h> 

const size_t N = 1000; // number of arrays 
const size_t K = 1000; // size of each array 
const int nTPB = 256; // number of threads per block, must be a power-of-2 
typedef float mytype; // type of data to be summed 

// produce the sum or mean of each array 
template <typename T> 
__global__ void breduce(const T * __restrict__ idata, T * __restrict__ odata, const int bsize){ 

    __shared__ T sdata[nTPB]; 

    T sum = 0; 
    //block-striding loop 
    size_t offset = blockIdx.x*bsize + threadIdx.x; 
    while (offset < (blockIdx.x+1)*bsize){ 
    sum += idata[offset]; 
    offset += blockDim.x;} 
    sdata[threadIdx.x] = sum; 
    __syncthreads(); 
    //shared memory reduction sweep 
    for (int i = nTPB>>1; i > 0; i>>=1){ 
    if (threadIdx.x < i) sdata[threadIdx.x] += sdata[threadIdx.x+i]; 
    __syncthreads();} 
    // write output sum for this block/array 
#ifndef MEAN 
    if (!threadIdx.x) odata[blockIdx.x] = sdata[0]; 
#else 
    if (!threadIdx.x) odata[blockIdx.x] = sdata[0]/bsize; 
#endif 
} 

int main(){ 

    mytype *h_idata, *h_odata, *d_idata, *d_odata; 
    h_idata=(mytype *)malloc(N*K*sizeof(mytype)); 
    h_odata=(mytype *)malloc(N*sizeof(mytype)); 
    cudaMalloc(&d_idata, N*K*sizeof(mytype)); 
    cudaMalloc(&d_odata, N*sizeof(mytype)); 
    for (size_t i = 0; i < N; i++) 
    for (size_t j = 0; j < K; j++) 
     h_idata[i*K+j] = 1 + (i&1); // fill alternating arrays with 1 and 2 
    memset(h_odata, 0, N*sizeof(mytype)); // zero out 
    cudaMemset(d_odata, 0, N*sizeof(mytype)); // zero out 
    cudaMemcpy(d_idata, h_idata, N*K*sizeof(mytype), cudaMemcpyHostToDevice); 
    breduce<<<N, nTPB>>>(d_idata, d_odata, K); 
    cudaMemcpy(h_odata, d_odata, N*sizeof(mytype), cudaMemcpyDeviceToHost); 
    // validate 
    for (size_t i = 0; i < N; i++) 
#ifndef MEAN 
    if (h_odata[i] != (K*(1 + (i&1)))) {printf("mismatch at %d, was: %f, should be: %f\n", i, (float)h_odata[i], (float)(K*(1 + (i&1)))); return 1;} 
#else 
    if (h_odata[i] != ((1 + (i&1)))) {printf("mismatch at %d, was: %f, should be: %f\n", i, (float)h_odata[i], (float)((1 + (i&1)))); return 1;} 
#endif 
    return 0; 
} 

$ nvcc -arch=sm_35 -o t1285 t1285.cu -DMEAN 
$ cuda-memcheck ./t1285 
========= CUDA-MEMCHECK 
========= ERROR SUMMARY: 0 errors 
$ nvcc -arch=sm_35 -o t1285 t1285.cu 
$ cuda-memcheck ./t1285 
========= CUDA-MEMCHECK 
========= ERROR SUMMARY: 0 errors 
$