CuPy解説

Software

ryosuke-okuta
of 21
All materials on our website are shared by users. If you have any questions about copyright issues, please report us to resolve them. We are always happy to assist you.
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