今回はこのサイトを参考にしながら、CUDA CとPyCUDAについてもう少し深く学習していきたい。この程度のCUDA Cソースコードを自力で書けるようになるにはまだまだ時間がかかりそうだ。

CUDA C source codeをload¶

cd git/gpu/Code/CUDA_C/vectorsums/

/home/workspace/1/git/gpu/Code/CUDA_C/vectorsums

ls

Makefile  vectorsums.cu*

# %load vectorsums.cu
#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>
#include <cuda_runtime.h> 

#define CUDA_CALL(x) {if((x) != cudaSuccess){ \
  printf("CUDA error at %s:%d\n",__FILE__,__LINE__); \
  printf("  %s\n", cudaGetErrorString(cudaGetLastError())); \
  exit(EXIT_FAILURE);}} 

#define N 10

__global__ void add(int *a, int *b, int *c){
  int bid = blockIdx.x;
  if(bid < N)
    c[bid] = a[bid] + b[bid];
}

int main(void) {
  int i, a[N], b[N], c[N];
  int *dev_a, *dev_b, *dev_c;

  CUDA_CALL(cudaMalloc((void**) &dev_a, N*sizeof(int)));
  CUDA_CALL(cudaMalloc((void**) &dev_b, N*sizeof(int)));
  CUDA_CALL(cudaMalloc((void**) &dev_c, N*sizeof(int)));

  for(i=0; i<N; i++){
    a[i] = -i;
    b[i] = i*i;
  }

  CUDA_CALL(cudaMemcpy(dev_a, a, N*sizeof(int), cudaMemcpyHostToDevice));
  CUDA_CALL(cudaMemcpy(dev_b, b, N*sizeof(int), cudaMemcpyHostToDevice));

  add<<<N,1>>>(dev_a, dev_b, dev_c);

  CUDA_CALL(cudaMemcpy(c, dev_c, N*sizeof(int), cudaMemcpyDeviceToHost));

  printf("\na + b = c\n");
  for(i = 0; i<N; i++){
    printf("%5d + %5d = %5d\n", a[i], b[i], c[i]);
  }

  CUDA_CALL(cudaFree(dev_a));
  CUDA_CALL(cudaFree(dev_b));
  CUDA_CALL(cudaFree(dev_c));
}

上のコードをコンパイルして実行する。

!nvcc -O2 vectorsums.cu -o vectorsums

!./vectorsums

a + b = c
    0 +     0 =     0
   -1 +     1 =     0
   -2 +     4 =     2
   -3 +     9 =     6
   -4 +    16 =    12
   -5 +    25 =    20
   -6 +    36 =    30
   -7 +    49 =    42
   -8 +    64 =    56
   -9 +    81 =    72

vectorsums.cuをPyCUDAに取り込む¶

from pycuda import autoinit
from pycuda import gpuarray
import numpy as np
from pycuda.compiler import SourceModule

mod = SourceModule(open("vectorsums.cu", "r").read())

/root/.pyenv/versions/miniconda3-4.3.30/envs/caffe2/lib/python3.6/site-packages/ipykernel_launcher.py:6: UserWarning: The CUDA compiler succeeded, but said the following:
kernel.cu(20): warning: linkage specification is not allowed

関数addを利用する¶

importedKernel = mod.get_function("add")

for i in range(5):
    a = i
    b = i*i
    a_gpu = gpuarray.to_gpu(np.array([a]))
    b_gpu = gpuarray.to_gpu(np.array([b]))
    c_gpu = gpuarray.to_gpu(np.empty_like(1))
    importedKernel(a_gpu,b_gpu,c_gpu,block=(1,1,1))
    print(a, '+',b, '=', c_gpu )

0 + 0 = 0
1 + 1 = 2
2 + 4 = 6
3 + 9 = 12
4 + 16 = 20

for i in range(10):
    a = -i
    b = i*i
    a_gpu = gpuarray.to_gpu(np.array([a]))
    b_gpu = gpuarray.to_gpu(np.array([b]))
    c_gpu = gpuarray.to_gpu(np.empty_like(1))
    importedKernel(a_gpu,b_gpu,c_gpu,block=(1,1,1))
    print(a, '+',b, '=', c_gpu )

0 + 0 = 0
-1 + 1 = 0
-2 + 4 = 2
-3 + 9 = 6
-4 + 16 = 12
-5 + 25 = 20
-6 + 36 = 30
-7 + 49 = 42
-8 + 64 = 56
-9 + 81 = 72

CUDA CとPyCUDAを組み合わせることで、コードを非常にスマートに書き換えることができるのがいい。