0

CUDA プログラミングを始めたばかりで、行列乗算のカーネル設計について何かを学んでいます。オンラインで見つかったメイン コードをコピーし、A(MxM) と B(MxN) の行列乗算を実装する部分を追加しました。

#include <iostream>
#include <iomanip>
#include <fstream>
#include <vector>
#include <cuda_runtime.h>
#include <cuComplex.h>
#include <cusp/complex.h>

using namespace std;

const int M=55, N=73;

typedef cusp::complex<double> Complex;

__global__ void kernelFunc(Complex* ad, Complex* bd, Complex* cd, int n) 
{
  int x = (blockIdx.x * blockDim.x) + threadIdx.x;
  int y = (blockIdx.y * blockDim.y) + threadIdx.y;

  if(x < n && y < n) 
  {   
    Complex v = Complex(0.0, 0.0);
    for(int i=0; i<n; i++) v += ad[y * n + i] * bd[i * n + x];
    cd[y * n + x] = v;
  }
}

int main(int argc, char *argv[])
{
  std::vector< Complex > _A(M*M);
  std::vector< Complex > _B(M*N);
  Complex *A, *B, *C;

  cudaMalloc((void**)&A, M*M*sizeof(Complex));
  cudaMalloc((void**)&B, M*N*sizeof(Complex));
  cudaMalloc((void**)&C, M*N*sizeof(Complex));

  for (int i=0; i<M*M; i++) _A[i] = Complex((double)i, (double)i);
  for (int i=0; i<M*N; i++) _B[i] = Complex(1.0, 0.0);

  cudaMemcpy( A, &_A[0], (M*M)*sizeof(Complex), cudaMemcpyHostToDevice );
  cudaMemcpy( B, &_B[0], (M*N)*sizeof(Complex), cudaMemcpyHostToDevice );

  dim3 block(32, 32);           
  dim3 grid((M+31)/32, (M+31)/32);

  kernelFunc<<<grid, block>>>(A, B, C, M);
  cudaMemcpy(&_B[0], &C[0], (M*N)*sizeof(Complex), cudaMemcpyDeviceToHost);

  cudaFree(A);
  cudaFree(B);
  cudaFree(C);

  return 0;
}

しかし、関連する行列は正方行列でなければならないとオンラインで述べているので、このコードは任意の次元の行列では使用できないということですか? 問題に合わせてブロック数とグリッド数を定義する方法がわかりません。私の問題の行列の次元は MxM で、M は奇数です。上記のコードを小さなマトリックスで試してみましたが、うまくいくようですが、それをかなり大きなマトリックスに適用する必要があります。それが大きなマトリックスでも機能するかどうかはわかりません。

4

1 に答える 1

2

指定されたコードは、正方行列のみに適用されます。一般化するには少し修正が必要です。カーネルは次のように変更できます。

__global__ void kernelFunc(Complex* ad, Complex* bd, Complex* cd, int m1, int n1, int n2) 
{
  int x = (blockIdx.x * blockDim.x) + threadIdx.x;
  int y = (blockIdx.y * blockDim.y) + threadIdx.y;

  if(x < n2 && y < m1) 
  {   
    Complex v = Complex(0.0, 0.0);
    for(int i=0; i<n1; i++) v += ad[y * n1 + i] * bd[i * n2 + x];
    cd[y * n2 + x] = v;
  }
}

は最初の行列m1の行、 は最初の行列のn1列n2、 は 2 番目の行列の列です。グリッド サイズは、次のように変更されます。

dim3 grid((N+31)/32, (M+31)/32);

最後に、カーネル呼び出し:

kernelFunc<<<grid, block>>>(A, B, C, M, M, N);
于 2013-09-11T18:49:05.680 に答える