【GPU】メモリについて「cudaMemcpy」

pytorchの並列化のレスポンスの調査のため、gpuメモリについて調べた軌跡をメモ。
この記事では、もしかしたらってこうかなーってのしかわかってない。

これらのサイトを参考にした。非常に勉強になった。
http://d.hatena.ne.jp/interleave/20091103/1257259065
http://news.mynavi.jp/column/architecture/338/

cudaでhello world!!

実装は見なくていいという場合は飛ばしてください。
いや　本当に動くの？と思いながらこの記事を元に動かした。
http://www.gdep.jp/page/view/255

環境：ubuntu14, cuda8.0

main.cu

# include <stdio.h>

__global__ void
    kernel( void ) {
}

int main( void ) {
    kernel<<<1,1>>>();
    printf( "Hello, World!\n" );
    return 0;
}

nvcc main.cu
./a.out
# Hello, World!って出た。

GPUで文字列を形成し、それをホスト側に移して表示する

ホストがcpu側、デバイスはgpu側。

main.cu

# include <stdio.h>

//デバイス側で文字を表示
__device__ void
device_strcpy(char *dst, const char *src)
{
    while (*dst++ = *src++);
}

__global__ void kernel(char *A)
{
    device_strcpy(A, "Hello, World!");
}

int main() {
   char *d_hello;
   char hello[32]; //cpu側メモリを確保

   cudaMalloc((void**)&d_hello, 32);  //gpu側メモリを確保

   kernel<<<1,1>>>(d_hello);

   cudaMemcpy(hello, d_hello, 32, cudaMemcpyDeviceToHost); //gpuからcpuにコピー

   cudaFree(d_hello);

   puts(hello);
}

gdepさんの記事は非常に勉強になる。

こちらの記事も試して見た。
http://d.hatena.ne.jp/interleave/20091103/1257259065

• 1KB転送

cudaMalloc : 932.438000
cudaMemcpyHostToDevice : 5247.837000
cudaMemcpyDeviceToHost : 7321.464000
cudaFree : 0.365000

• 1MB転送

cudaMalloc : 938.814000
cudaMemcpyHostToDevice : 359.169000
cudaMemcpyDeviceToHost : 365.630000
cudaFree : 0.141000

• 512MB転送

cudaMalloc : 924.903000
cudaMemcpyHostToDevice : 150.473000
cudaMemcpyDeviceToHost : 136.740000
cudaFree : 0.274000

やっぱり転送はデータをでっかくして送った方がスピードが出てる。

cudaMallocについて

整理すると
cudaMalloc：gpu(デバイス)メモリ確保
cudaMemcpyHostToDevice：メモリ（ホスト）からgpuメモリに転送
cudaMemcpyDeviceToHost：gpuメモリからメモリに転送

cudaMalloc(&d_tmp, N);
cudaMemcpy(d_tmp, input, N, cudaMemcpyHostToDevice);
cudaMemcpy(output, d_tmp, N, cudaMemcpyDeviceToHost);

で、何となくcudaに慣れてきたところで、pytorchの中身へ。
pytorchはcpuだとcとかc++でgpuはcudaファイルが動いてる。
今回見るのはcuファイル。
今回目をつけたのはcudaMemcpyとcudamemcpyasync。
いかにもって名前でcudamemcpyasyncは非同期だけどcudaMemcpyって同期だよね。
だからcudaMemcpyやってる間はコピーを待っちゃうよね？gpuが大量にあってもcudaMemcpyで処理待ちしちゃうよね？って想定でコードを解析。

PReLU

preluの実装の中では至る所でTHCTensor_クラスが使われてる。

void THNN_(PReLU_updateOutput)(
           THCState *state,
           THCTensor *input,
           THCTensor *output,
           THCTensor *weight,
           long nOutputPlane)
{
  THCTensor_(resizeAs)(state, output, input);

  weight = THCTensor_(newContiguous)(state, weight);
  real *w = THCTensor_(data)(state, weight);
...
略

THCTensor

thcって頭についてるのはtorchから使ってるluaの名残の名前？な気がする。
THCStorage_を呼んでる。

略
...
THC_API int THCTensor_(getDevice)(THCState* state, const THCTensor* tensor) {
  if (!tensor->storage) return -1;
  return THCStorage_(getDevice)(state, tensor->storage);
}

THCStorageCopy

ここでcudaMemcpy使ってるんですね。いやでもなぁ。
活性化関数のとこでメモリ転送って、いやぁー。ここレスポンスに関係ないようなぁ。
そもそもTHCudaCheckがわからん。調べるとみんなここでエラー出たーとか言ってて、解決策はcudaの環境下　実装上の失敗。だからなんかチェックしてくれる機能である。中身が知りたい。

void THCStorage_(copyCPU)(THCState *state, THCStorage *self, struct THStorage *src)
{
  THArgCheck(self->size == src->size, 2, "size does not match");
  THCudaCheck(cudaMemcpy(self->data, src->data, self->size * sizeof(real), cudaMemcpyHostToDevice));
}

その他メモ

GPUは超低速プロセサ
...人を運ぶタスクにおいてgpuはバスでcpuはスポーツカー
http://news.mynavi.jp/column/architecture/338/
【コラム】コンピュータアーキテクチャの話 (336) CPUと異なるGPUにおけるコアと2次キャッシュの関係性
...cpuとgpuではl2キャッシュのついてる場所が違う。
...cpuはl2スイッチで矛盾が起こる共有の仕方をしているので、コヒーレンシプロトコルで矛盾を無くしている。gpuはメモリ側にl2がついてるので、l2の目的が違う。
http://news.mynavi.jp/photo/column/architecture/336/images/011l.jpg
骨まで理解するPCアーキテクチャ(GPU編) 第4回
...gpuメモリのデータ転送時間かかるので、でっかいデータをまとめてやれよ
http://pc.watch.impress.co.jp/docs/column/1month-kouza/646073.html

• 確かに動くが理解はしてない。

gcc -Wall -c cuda_main.c
nvcc -c cuda_kernel.cu
gcc -I/usr/local/cuda/include -L/usr/local/cuda/lib64 cuda_main.o cuda_kernel.o -lcudart
./a.out
http://qiita.com/gunn/items/34075251ec6687e06800

• 使ってないgpuを使用するスクリプト

# !/usr/bin/env bash

# Number of free GPUs on a machine
export n_gpus=`lspci | grep -i "nvidia" | wc -l`

# Return -1 if there are no GPUs on the machine
if [ $n_gpus -eq 0 ]; then
  echo "-1"
  exit -1
fi

f_gpu=`nvidia-smi | sed -e '1,/Processes/d' \
  | tail -n+3 | head -n-1 | awk '{print $2}'\
  | awk -v ng=$n_gpus 'BEGIN{for (n=0;n<ng;++n){g[n] = 1}} {delete g[$1];} END{for (i in g) print i}' \
  | tail -n 1`

# return -1 if no free GPU was found
if [ `echo $f_gpu | grep -v '^$' | wc -l` -eq 0 ]; then
  echo "-1"
  exit -1
else
  echo $f_gpu
fi

CUDA_VISIBLE_DEVICES=`scripts/free-gpu` python

あとtensorflowには2017/9/1時点では並列gpuを使用するときに、
エラーなっちゃうバグがまだ残ってるぽい
https://github.com/tensorflow/tensorflow/issues/152

cuda toolkitの全バージョン

cudaのソース