CUDA: умножение плиточной матрицы - PullRequest
Новая
       26

CUDA: умножение плиточной матрицы

0 голосов
/ 22 апреля 2020

Я пытаюсь реализовать умножение матриц с общей памятью в CUDA C ++. Во-первых, у меня есть функция, которая принимает матрицы на CPU, чтобы выполнять работу на GPU: 

#define block_size 32 // 32*32 = 1024 = max nb of threads
#define tile_size 32


void matrix_mul_gpu(int *h_m1, int *h_m2, int *h_res, const int h1, const int w1, const int h2, const int w2) {
  // h_m1 and h_m2 are matrices of size h1*w1 and h2*w2, with w1=h2
  // GPU memory allocation
  int *d_m1, *d_m2, *d_res;
  cudaMalloc(&d_m1, h1*w1*sizeof(int));
  cudaMalloc(&d_m2, h2*w2*sizeof(int));
  cudaMalloc(&d_res, h1*w2*sizeof(int));

  cudaMemcpy(d_m1, h_m1, h1*w1*sizeof(int), cudaMemcpyHostToDevice);
  cudaMemcpy(d_m2, h_m2, h2*w2*sizeof(int), cudaMemcpyHostToDevice);


  // Matrix multiplication on GPU

  dim3 nb_threads(block_size,block_size,1);
  dim3 nb_blocs((int)ceil((float)w2/block_size),(int)ceil((float)h1/block_size),1);
  const int nb_loop = ceil((float)w1/tile_size);
  d_matrix_mul_tile<<<nb_blocs, nb_threads>>>(d_m1, d_m2, d_res, h1, w1, h2, w2, nb_loop);


  // Transfer GPU to CPU and free memory
  cudaMemcpy(h_res, d_res, h1*w2*sizeof(int), cudaMemcpyDeviceToHost);
  cudaFree(d_m1); cudaFree(d_m2); cudaFree(d_res);

}

Так что вся работа выполняется с помощью функции d_matrix_mul_tile . Эта версия работает: 

__global__
void d_matrix_mul_tile(int *d_m1, int *d_m2, int *d_res, const int h1, const int w1, const int h2, const int w2, const int nb_loop) {

  int tx = threadIdx.x;
  int ty = threadIdx.y;
  int col = blockIdx.x*blockDim.x + tx;
  int row = blockIdx.y*blockDim.y + ty;
  __shared__ int m1_tile[tile_size][tile_size];
  __shared__ int m2_tile[tile_size][tile_size];


  int trans=0; // value to feed in d_res[row,col]



  for (int i = 0 ; i < nb_loop ; i++) {
    // Step 1 : fill values in shared Memory

    if (i*tile_size+tx > w1) {
      m1_tile[ty][tx] = 0;
    }
    else {
      m1_tile[ty][tx] = d_m1[(row*w1)+ (i*tile_size + tx)];
    }


    if (i*tile_size+ty > h2) {
      m2_tile[ty][tx] = 0;
    }
    else {
      m2_tile[ty][tx] = d_m2[i*tile_size*w2 + ty*w2 + col];
    }

    __syncthreads();


    // Step 2 : Augment trans values
    for (int k = 0 ; k < tile_size ; k++) trans += m1_tile[ty][k]*m2_tile[k][tx];
    __syncthreads();
  }

  // if condition because of ceil in host function (last values)
  if ( (row < h1) && (col < w2) ) {d_res[row*w2+col] = trans;}


}

Но эта не работает, и я не могу понять, почему. 

__global__
void d_matrix_mul_tile(int *d_m1, int *d_m2, int *d_res, const int h1, const int w1, const int h2, const int w2, const int nb_loop) {

  int tx = threadIdx.x;
  int ty = threadIdx.y;
  int col = blockIdx.x*blockDim.x + tx;
  int row = blockIdx.y*blockDim.y + ty;
  __shared__ int m1_tile[tile_size][tile_size];
  __shared__ int m2_tile[tile_size][tile_size];


  int trans=0; // value to feed in d_res[row,col]

  // First if condition because of ceil in host function (last values)
  if ( (row < h1) && (col < w2) ) {

    for (int i = 0 ; i < nb_loop ; i++) {
      // Step 1 : fill values in shared Memory
      if (i*tile_size+tx > w1) {
        m1_tile[ty][tx] = 0;
      }
      else {
        m1_tile[ty][tx] = d_m1[(row*w1)+ (i*tile_size + tx)];
      }


      if (i*tile_size+ty > h2) {
        m2_tile[ty][tx] = 0;
      }
      else {
        m2_tile[ty][tx] = d_m2[i*tile_size*w2 + ty*w2 + col];
      }

      __syncthreads();


      // Step 2 : Augment trans values
      for (int k = 0 ; k < tile_size ; k++) trans += m1_tile[ty][k]*m2_tile[k][tx];
      __syncthreads();
    }

    d_res[row*w2+col] = trans;

  }


}

Например, для h1 = 3, w1 = 6, h2 = 6, w2 = 3 и случайных матриц я получил следующие результаты: 

m1 : 
5 4 4 2 2 2 
5 4 0 5 5 4 
4 0 4 3 0 0 

m2 : 
5 0 0 
3 4 5 
1 0 0 
2 1 5 
1 1 0 
1 0 4 

True res : 
49 20 38 
56 26 61 
30 3 15 

Res : 
41 16 20 
37 16 20 
24 0 0

Может кто-то увидеть скажите мне ошибку?

...