SHARE
TWEET

Untitled

a guest Feb 17th, 2019 57 Never
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
  1. #include <iostream>
  2. #include <chrono>
  3. #include <cuda_runtime.h>
  4.  
  5. // N x Nの行列を扱う
  6. constexpr std::size_t N = 1024;
  7.  
  8. // 1つのブロックでBLOCK_SIZE x BLOCK_SIZEのスレッドを管理する
  9. constexpr std::size_t BLOCK_SIZE = 16;
  10.  
  11. // ホスト(CPU)側の行列を定義
  12. static double hostMatrixA[N * N];
  13. static double hostMatrixB[N * N];
  14. static double hostMatrixC[N * N];
  15.  
  16. // デバイス(GPU)側の行列へのポインタ
  17. static double* deviceMatrixA;
  18. static double* deviceMatrixB;
  19. static double* deviceMatrixC;
  20.  
  21. // 行列の積を計算する関数
  22. __global__ void matrix_multiply(double*, double*, double*);
  23. __global__ void matrix_multiply_shared(double*, double*, double*);
  24.  
  25. int main(void)
  26. {
  27.     // デバイス側に行列用の記憶領域を確保する
  28.     cudaMalloc((void**)&deviceMatrixA, sizeof(hostMatrixA));
  29.     cudaMalloc((void**)&deviceMatrixB, sizeof(hostMatrixB));
  30.     cudaMalloc((void**)&deviceMatrixC, sizeof(hostMatrixC));
  31.  
  32.     // ホスト側の行列に値を設定する
  33.     for (std::size_t i = 0; i < N * N; i++) {
  34.         hostMatrixA[i] = static_cast<double>(i + 1);
  35.         hostMatrixB[i] = static_cast<double>(-i - 1);
  36.         hostMatrixC[i] = static_cast<double>(0);
  37.     }
  38.  
  39.     // タイマー開始
  40.     auto start = std::chrono::high_resolution_clock::now();
  41.  
  42.     // ホスト側の行列のデータをデバイス側の行列へ転送する
  43.     cudaMemcpy(deviceMatrixA, hostMatrixA, sizeof(hostMatrixA), cudaMemcpyHostToDevice);
  44.     cudaMemcpy(deviceMatrixB, hostMatrixB, sizeof(hostMatrixB), cudaMemcpyHostToDevice);
  45.  
  46.     // グリッドおよびブロックの定義
  47.     dim3 grid(N / BLOCK_SIZE, N / BLOCK_SIZE);
  48.     dim3 block(BLOCK_SIZE, BLOCK_SIZE);
  49.  
  50.     // GPU側の処理を起動させる
  51.     matrix_multiply <<< grid, block >>> (deviceMatrixA, deviceMatrixB, deviceMatrixC);
  52.  
  53.     // deviceMatrixCに格納されている計算結果をhostMatrixCへ転送する
  54.     cudaMemcpy(hostMatrixC, deviceMatrixC, sizeof(hostMatrixC), cudaMemcpyDeviceTohost);
  55.  
  56.     // タイマー終了
  57.     auto end = std::chrono::high_resolution_clock::now();
  58.  
  59.     // 計算結果発表
  60.     std::cout << "Calculator time is " << std::chrono::duration<double>(end - start).count() << " seconds." << std::endl;
  61.     std::cout << "Calculator result is " << std::setprecision(std::numeric_limits<double>::max_digits10) << hostMatrixC[N * N - 1] << "." << std::endl;
  62.  
  63.     // デバイス側の記憶領域を解放する
  64.     cudaFree(deviceMatrixC);
  65.     cudaFree(deviceMatrixB);
  66.     cudaFree(deviceMatrixA);
  67.  
  68.     return 0;
  69. }
  70.  
  71. // 行列の積を計算する関数
  72. __global__ void matrix_multiply(double* A, double* B, double* C)
  73. {
  74.     // 各スレッドが担当する行列の行yと列xを取得
  75.     std::size_t y = blockDim.y * blockIdx.y + threadIdx.y;
  76.     std::size_t x = blockDim.x * blockIdx.x + threadIdx.x;
  77.  
  78.     double c = 0;
  79.     for (std::size_t i = 0; i < N; i++) {
  80.         c += A[N * y + i] * B[N * i + x];
  81.     }
  82.     C[x + N * y] = c;
  83. }
  84.  
  85.  
  86. // シェアードメモリを利用した行列積計算関数
  87. __global__ void matrix_multiply_shared(double* A, double* B, double* C)
  88. {
  89.     // 各スレッドが担当する行列の行yと列xを取得
  90.     std::size_t y = blockDim.y * blockIdx.y + threadIdx.y;
  91.     std::size_t x = blockDim.x * blockIdx.x + threadIdx.x;
  92.  
  93.     __shared__ double s_A[BLOCK_SIZE][BLOCK_SIZE];
  94.     __shared__ double s_B[BLOCK_SIZE][BLOCK_SIZE];
  95.  
  96.     double c = 0;
  97.     for (std::size_t i = 0; i < N; i += BLOCK_SIZE) {
  98.  
  99.         // シェアードメモリに行列の一部をコピー
  100.         s_A[threadIdx.y][threadIdx.x] = A[N * y + i + threadIdx.x];
  101.         s_B[threadIdx.y][threadIdx.x] = B[N * (i + threadIdx.y) + x];
  102.         __syncthreads();
  103.  
  104.         // シェアードメモリで積を計算する
  105.         for (std::size_t j = 0; j < BLOCK_SIZE; j++) {
  106.             c += s_A[threadIdx.y][j] * s_B[j][threadIdx.x];
  107.         }
  108.         __syncthreads();
  109.     }
  110.     C[N * y + x] = c;
  111. }
RAW Paste Data
We use cookies for various purposes including analytics. By continuing to use Pastebin, you agree to our use of cookies as described in the Cookies Policy. OK, I Understand
 
Top