CuPy解説

Software

ryosuke-okuta
of 21
Description
Text
  • CuPy解説 2015/12/19 Chainer Meetup #1@スマートニュース (株)Preferred Networks 奥田 遼介
  • 自己紹介 奥田 遼介  -2014東北大学 修士  文字列処理など  2014 (株)プリファードインフラストラクチャー  2014- (株)プリファードネットワークス  映像解析系、製造業系にかかわる研究開発  ChainerやCuPyの開発 2
  • 何を話すか?  CuPyの簡単な解説  Cupyでどのくらい早くなるか?  利用上の注意点  ElementwiseKenrnel, ReductionKernelの使い方解説  CuPyの実装のすごーくざっくーりした全体概要 3
  • CuPyとは何か?  CUDA上で計算を行うNumPyサブセットのライブラリ  関数・メソッドのサブセットを実装  Chainer v1.5.0では 約174個の関数が実装済み  行列積などはcuBLASを利用(高速)  配列のスライス、転置、reshape 等が自由にできる  カスタムカーネルの記述も可能  elementwise, reduction  とにかく簡単にGPUが使える事を追求  Python上で簡単に多次元配列といえばNumPy  PC上で簡単にGPUといえばCUDA  CUDA +NumPy =CuPy 4
  • CuPy の使い方は?  numpy の代わりに cupy を使う(だいたい動く)  CPU/GPU の両方で動く関数の書き方  chainer.cuda.get_array_module()を使うと、引数に cupy.ndarray があるかないかで numpy / cupy のどちらかを返してくれ ます  例えば下は NumPy と CuPy の両方で動く logsumexp の実装例  (より省メモリな実装を考えてみてください) def logsumexp(x, axis=None): xp = cuda.get_array_module(x) x_max = x.max(axis=axis) return x_max + xp.log( xp.exp(x – x_max).sum(axis=axis)) 5
  • CuPyってどのくらい早くなるの?  状況によりけりですが、最大数十倍程度速くなります def test(xp): a = xp.arange(1000000).reshape(1000, -1) return a.T * 2 test(numpy) t1 = datetime.datetime.now() for i in range(1000): test(numpy) t2 = datetime.datetime.now() print(t2 -t1) test(cupy) t1 = datetime.datetime.now() for i in range(1000): test(cupy) t2 = datetime.datetime.now() print(t2 -t1) 時間[ms] 倍率 NumPy 2929 1 CuPy 585 5 6
  • 5倍しかなってない!もっと早くならないの?  なります  メモリープールを有効にしましょう  cupy.cuda.set_allocator(cupy.cuda.MemoryPool().malloc)  面倒であれば「import chainer」でもOK  なぜ早くなるか?  CUDAではmalloc, freeの実行時にCPUとGPUが同期する  CuPyでは一度mallocした領域を使いまわすことで同期を回避 cupy.cuda.set_allocator(cupy.cuda.MemoryPool().malloc) test(cupy) t1 = datetime.datetime.now() for i in range(cnt): test(cupy) t2 = datetime.datetime.now() print(t2 -t1) 時間[ms] 倍率 NumPy 2929 1 CuPy 585 5 CuPy メモリプール 123 24 7
  • Chainer(CuPy)が遅い・・・・?  v1.5からものすごく速くなりました!  v1.4まではCuPyのコードがPythonだったので遅かった  なので、Cythonにしました  例えば、Chainerから見るてどのくらい早くなったか? 速度 mnist ptb v1.4.1 1.0 1.0 v1.2 (pyCUDAの頃) 2.3 1.6 v1.5 (cython) 2.6 1.2※ ※ GradientClippingが行われるようになってちょっと計算量が増えています 8
  • おわり(その1)  ここより下のスライドはよりCuPyを活用したい人のため の情報です  Elementwise,Reductionを作りたい人以外には不要な情報 です。 9
  • Elementwiseな操作をしたい時は?  要素ごとの2乗誤差をとる関数を定義  使い方  ブロードキャストや型チェックは全自動 squared_diff = cupy.ElementwiseKernel( ‘float32 x, float32 y’, #入力 ‘float32 z’, #出力 ‘z = (x - y) * (x - y)’, #計算 ‘squared_diff’) #名前 squared_diff(cupy.arange(10), 10)) 10
  • 型をジェネリックに扱いたい時は?  1文字の型名はプレースホルダーになる squared_diff = cupy.ElementwiseKernel( ‘T x, T y’, //入力 ‘T z’, //出力 ‘z = (x - y) * (x - y)’, //計算 ‘squared_diff’) //名前 11
  • 配列に自由にアクセスしたい時は?  例えば、bincountがしたいときは?  「raw」をつけると配列の添え字アクセスができる for x in arr_x: bin[x] += 1 cupy.ElementwiseKernel( 'S x', 'raw U bin', 'atomicAdd(&bin[x], 1)', 'bincount_kernel') 12
  • Reduceしたい時は?  例えばL2距離  使い方  「axis」や「keepdims」などのオプションも使えます l2norm_kernel = cupy.ReductionKernel( 'T x', # 入力 'T y', # 出力 'x * x', # 前処理 'a + b’, # リデュース 'y = sqrt(a)', # 後処理 '0', # 初期値 'l2norm') # 名前 x = cupy.arange(10, dtype='f').reshape(2, 5) L2norm_kernel(x, axis=1, keepdims=True) 13
  • おわり(その2)  ここより下のスライドはよりCuPyを作りたい人のための 情報です  NumPyのかわりにCuPyを使いたいだけの人には完全に不 要な情報です 14
  • CuPyの実装が知りたい時は?  コードを読んでください  今回はとてもざっくりと全体の概要だけ説明します 15
  • CuPyの全体構造は?  大まかに3層構造 CUDA(cuBLAS, cuRNAD, cuDNN) ndarray ufunc, elementwise, reduction CUDA Python wrapper cupy.cuda cupy.core 関数群(演算、テンソル操作) cupy 16
  • CUDA Python wrapperとは?  薄いラッパー  driver.pyx, runtime.pyx, cublas.pyx, curand.pyx, cudnn.pyx  C関数をPython関数にする  cupy_cuda.h, cupy_cudnn.h  CUDAが無い環境でCuPyをインストールするためのヘッダ  少し厚めのラッパー  memory.pyx GPUのポインタをくるむ  function.pyx CUDAのkernelとmoduleをくるむ  compiler.pyx cuファイルのコンパイルを支援  stream.pyx streamとeventをくるむ 17
  • ndarray の実体ってどうなっているの?  Cython側の実装:cupy/core/core.pyx  CUDA側の実装:cupy/core/carray.cuh  メンバとしては以下のようなものを持っている  詳しく説明しません template class CArray { private: T* data_; int size_; int shape_[ndim]; int strides_[ndim]; 18
  • cupy.add(x, y)を実行したら何が起こるのか?  とても複雑なことが起きます  ufunc の__call__が呼び出される  引数のオプションを解決(args, kwdargs)  引数の個数チェック  引数のdeviceをチェック  引数の型を正規化  ブロードキャストを実行  add をどの型の組み合わせで実行するかを解決  戻り値の割り当て実行  計算が高速に実行できるように引数を最適化  CUDAのカーネルコードを生成(キャッシュ付)  CUDAのカーネル呼び出しのための引数を構築  カーネル呼び出しを実行 19
  • CuPyの機能が足りないと思った時は?  もう一度よく考える  NumPy/CuPyでは別の方法で同じことができる場合も多いです  複数の関数とIndexingの組み合わせで結構どうにかなります  ぜひPRを!  NumPyの機能をCuPyに移植したPRは基本マージされます  効率は特に考えなくて良いので、まずは実装してみてください 20
  • おわり(まとめ)  ぜひCuPyを活用してください  Preferred Networksでは人材採用に力を入れています  フルタイム、アルバイト、インターンなどに興味のある 方はホームページからコンタクトお願いします  https://www.preferred-networks.jp/job_ja 21
Comments
Top