PyCUDAを使ってみよう
PyCUDAについて
PyCUDAを使用すると、PythonからNvidiaのCUDA並列計算APIにアクセスできます。
PyCUDAの主な機能には次のものがあります。
完全：すべてのCUDAの機能をPythonにマップ
柔軟で高速な、自動的に調整されたコードの実行時コード生成（RTCG）を有効にできる
追加された堅牢性：
オブジェクトの有効期間の自動管理
自動エラーチェック
すべてのCUDAエラーは自動的にPython例外に変換
追加された利便性：
既製のオンGPU線形代数、リダクション、スキャンが付属
FFTおよびLAPACKのアドオンパッケージが利用可能
速い。 ラッピングオーバーヘッドがほぼゼロ。
完全で役立つドキュメント。
パフォーマンス上の利点：
ライブラリ、デバッグサポート、プロファイリング、および複数のGPUに関する説明については、PyCUDAFAQを参照してください。
インストール
PyCUDAは次のようにインストールします。
code: bash
$ pip install -U pycuda
GoogleColab を利用している場合は、まずランタイムのタイプでGPUを選択してから、
コードセルに次のように入力して実行します。
code: ipyton
!pip install -U pycuda
利用方法
code: examples/hello_gpu.py
import pycuda.autoinit
import pycuda.driver as drv
import numpy
from pycuda.compiler import SourceModule
mod = SourceModule("""
__global__ void multiply_them(float *dest, float *a, float *b)
{
const int i = threadIdx.x;
}
""")
multiply_them = mod.get_function("multiply_them")
a = numpy.random.randn(400).astype(numpy.float32)
b = numpy.random.randn(400).astype(numpy.float32)
dest = numpy.zeros_like(a)
multiply_them(
drv.Out(dest), drv.In(a), drv.In(b),
block=(400,1,1), grid=(1,1))
print(dest-a*b)
このプログラムは画面をゼロで埋め尽くします。裏側ではもっと興味深いことが起こっています。
まず、PyCUDAがCUDAソースコードをコンパイルしてデバイスに転送します。
このCUDAソースコードは定数である必要はありません。コンパイルするコードをPythonで動的に生成させることができます。( メタプログラミングを参照) PyCUDAのnumpyインタラクションコードは、デバイスに領域を自動的に割り当て、numpy配列aとbをコピーし、400x1x1のシングル・ブロック・グリッドを起動し、転送先にコピーしました。
カーネルの呼び出し間でデータをデバイスに保持することもできます。常にデータをコピーする必要はありません。
この例にはクリーンアップ処理がないことに注意してください。怠けてコードしなかったのではなく、単に必要がありません。 PyCUDAは、必要なクリーンアップを自動的に推測して、それを実行してくれます。
もう少し詳しく説明することにしましょう。
初期化
まず、PyCudaを使用する前に、インポートして初期化する必要があります。
code: python
import pycuda.driver as cuda
import pycuda.autoinit
from pycuda.compiler import SourceModule
pycuda.autoinit のインポートは省略することができます。その場合、必要に応じて、初期化、コンテキストの作成、およびクリーンアップを手動で実行します。
データの転送
プログラムに必要なことのひとつに、データをデバイスに転送することがあります。 PyCudaでは、ほとんどの場合、ホスト上のnumpy配列からデータを転送します。 実際には、文字列を含めPythonバッファーインターフェイスを利用しているすべてのオブジェクトを転送することができます。
まず。4x4配列を乱数を作成しましょう。
code: python
import numpy
a = numpy.random.randn(4,4)
ここで、配列a は倍精度の数値で構成されていますが、ほとんどのNVIDIAデバイスは単精度だけをサポートしているので、次のようにfloat32型に変換しておきます。
code: python
a = a.astype(numpy.float32)
次に、データを転送する場所が必要なので、デバイスにメモリを割り当てる必要があります。
code: python
a_gpu = cuda.mem_alloc(a.nbytes)
このデータをデバイスに転送します。
code: python
cuda.memcpy_htod(a_gpu, a)
カーネルの実行
ここでは、理解しやすくするために単純な処理で説明します。a_gpuの各エントリを2倍にするコードを記述しましょう。 これには、対応するCUDA Cコードを記述して、それをpycuda.compiler.SourceModuleのコンストラクターに与えます。
code: python
mod = SourceModule("""
__global__ void doublify(float *a)
{
int idx = threadIdx.x + threadIdx.y*4;
}
""")
エラーがない場合は、コードがコンパイルされてデバイスに展開されます。 pycuda.driver.Functionへの参照を見つけて呼び出し、a_gpuを引数として与えて、ブロックサイズを 4x4 にします。
code: python
func = mod.get_function("doublify")
func(a_gpu, block=(4,4,1))
最後に、GPUからデータをフェッチして、元のデータと一緒に表示します。
code: python
a_doubled = numpy.empty_like(a)
cuda.memcpy_dtoh(a_doubled, a_gpu)
print(a_doubled)
print(a)
https://gyazo.com/f386f63f3b10c2e0bb0c9a9e7846319e
動きましたね！
便利なことにPyCudaがこのコードを引き継ぎ、すべてのクリーンアップを実行してくれるので、これで終了です。
コードを示しておきましょう。
code: pytohn
import pycuda.driver as cuda
import pycuda.autoinit
from pycuda.compiler import SourceModule
import numpy
a = numpy.random.randn(4,4)
a_gpu = cuda.mem_alloc(a.nbytes)
cuda.memcpy_htod(a_gpu, a)
mod = SourceModule("""
__global__ void doublify(float *a)
{
int idx = threadIdx.x + threadIdx.y*4;
}
""")
func = mod.get_function("doublify")
func(a_gpu, block=(4,4,1))
a_doubled = numpy.empty_like(a)
cuda.memcpy_dtoh(a_doubled, a_gpu)
print(a_doubled)
print(a)
明示的メモリコピーのショートカット
pycuda.driver.In、pycuda.driver.Out、およびpycuda.driver.InOutのハンドラー引数を使用すると、一部のメモリ転送を簡略化できます。 たとえば、a_gpuを作成する代わりに、aを置き換えても問題がない場合は、次のようにすることができます。
code: python
func(cuda.InOut(a), block=(4, 4, 1))
準備された呼び出し
組み込みのpycuda.driver.Function.__call__()メソッドを使用した関数呼び出しでは、型の識別にオーバーヘッドが発生します（デバイスインターフェイス を参照）。 このオーバーヘッドなしで上記と同じ効果を達成するために、関数は引数の型にバインドされてから呼び出されます。（Pythonの標準ライブラリstructモジュールを参照） これにより、numpy.numberクラスを使用して明示的な引数サイズを割り当てる必要もなくなります。 code: python
grid = (1, 1)
block = (4, 4, 1)
func.prepare("P") # Python: integer, C:void *
func.prepared_call(grid, block, a_gpu)
おまけ：もっと簡単にしてみる
pycuda.gpuarray.GPUArrayを使用すると、同じ効果をはるかに少ないコードで実現できます。
code: python
import pycuda.gpuarray as gpuarray
import pycuda.driver as cuda
import pycuda.autoinit
import numpy
a_gpu = gpuarray.to_gpu(numpy.random.randn(4,4).astype(numpy.float32))
a_doubled = (2*a_gpu).get()
print(a_doubled)
print(a_gpu)
より高度な使用方法
構造体
いくつかの可変長配列を2倍にするために、次の構造体(struct) があるとします。
code: python
mod = SourceModule("""
struct DoubleOperation {
int datalen, __padding; // 64ビットptrを整列
float *ptr;
};
__global__ void double_array(DoubleOperation *a) {
for (int idx = threadIdx.x; idx < a->datalen; idx += blockDim.x) {
}
}
""")
グリッド内の各ブロックは、配列の1つを2倍にします。 forループでは、スレッドよりも多くのデータ要素を2倍にすることができますが、十分な数のスレッドがあることを保証できる場合は効率的ではありません。 次に、構造体のラッパークラスが作成され、2つの配列がインスタンス化されます。
code: python
class DoubleOpStruct:
mem_size = 8 + numpy.intp(0).nbytes
def __init__(self, array, struct_arr_ptr):
self.data = cuda.to_device(array)
self.shape, self.dtype = array.shape, array.dtype
cuda.memcpy_htod(int(struct_arr_ptr), numpy.getbuffer(numpy.int32(array.size)))
cuda.memcpy_htod(int(struct_arr_ptr) + 8, numpy.getbuffer(numpy.intp(int(self.data))))
def __str__(self):
return str(cuda.from_device(self.data, self.shape, self.dtype))
struct_arr = cuda.mem_alloc(2 * DoubleOpStruct.mem_size)
do2_ptr = int(struct_arr) + DoubleOpStruct.mem_size
array1 = DoubleOpStruct(numpy.array(1, 2, 3, dtype=numpy.float32), struct_arr) array2 = DoubleOpStruct(numpy.array(0, 4, dtype=numpy.float32), do2_ptr) print("original arrays", array1, array2)
このコードは、pycuda.driver.to_device()関数とpycuda.driver.from_device()関数を使用して値を割り当ておよびコピーし、割り当てられたメモリブロックへのオフセットを使用する方法を示します。 最後に、コードを実行できます。 以下は、両方の配列を2倍にし、次に2番目の配列のみを2倍にすることを示しています。
code: python
func = mod.get_function("double_array")
func(struct_arr, block = (32, 1, 1), grid=(2, 1))
print("doubled arrays", array1, array2)
func(numpy.intp(do2_ptr), block = (32, 1, 1), grid=(1, 1))
print("doubled second only", array1, array2, "\n")
メタプログラミング
従来型のプログラミングでは、タスクを実行するプログラムを作成します。メタプログラミング(metaprograming)では、タスクを実行するプログラムを生成するプログラムを作成します。
これはかなり難しく思えるかもしれません。
メタプログラミングが重要な理由
まずはじめに、なぜメタプログラミングが重要であるかを考えてみましょう。
自動チューニング
多くの場合、CUDAプログラマーの時間のかなりの部分は、コードチューニングに費やされます。このチューニングでは、次のような項目が重要になります。
ブロックあたりの最適なスレッド数はいくつにするか？
一度にどのくらいのデータを処理する必要があるか？
どのデータを共有メモリにロードするか？ また、その必要があるか？
対応するブロックの大きさをどうするか？
試行錯誤の中で、コードの実行時間のパターンを運良く見つけることができ、最速のバージョンを確実に選択できることがあります。残念ながら、こうしたチューニングは根拠が薄く信頼性が低くなるだけでなく、新しいハードウェア世代で完全に失敗する可能性があります。
PyCUDAが提供しようとしているこの問題の解決策は次のとおりです。
試行錯誤はしない：実行時にベンチマークを行い、最も速く機能するものを採用する
これは、CUDAランタイムAPIに対するPyCUDAの重要な利点です。コードの実行中にこれらの決定を行うことができます。 ATLASやFFTWなど、多くの著名なコンピューティングパッケージが同様の手法を利用していますが、それらはかなり複雑な最適化ドライバールーチンを必要とします。それに対して、PyCUDAでは快適なPythonを使いながらチューニングを行うことができます。
データ型
コードは、実行時にさまざまなデータ型を処理する必要があるあります。たとえば、単精度と倍精度の両方の浮動小数点数で機能する必要があるようなケースです。両方のバージョンをプリコンパイルすることもできますが、それはコードのサイズが増大するばかりか、複雑さ増して保守性が劣化してしまいます。必要なときに必要なコードを生成する方がはるかに簡単です。
与えられた問題に特化したコード
ライブラリを作成しているとして、ユーザーはライブラリにいくつかのタスクを実行するように要求するはずです。コードを必要以上に汎用化して遅くするのではなく、解決を求められている問題だけを意図的にコードを生成できれば、それがどれほど単純になるかか想像してみてください。 PyCUDAはこれを実現してくれます。
定数は変数よりも高速
問題のサイズが実行ごとに異なるけれど、同じサイズのデータに対してより多くのカーネル呼び出しを実行する場合は、データサイズを定数としてハードコーディングしてコンパイルする方が高速になります。これは、主にフェッチ時間の短縮とレジスタ圧力の低下により、パフォーマンスに大きなメリットをもたらす可能性があります。特に、定数による乗算は、一般的な変数-変数の乗算よりもはるかに効率的に実行されます。
ループ展開(Loop Unrolling)
ループ展開は高速化の基本テクニックのひとつです。CUDAプログラミングガイドには、nvccの優れた点と、ループを展開する方法が記載されています。ただし、バージョン2.1の時点では、それは真実ではなく、少なくとも私の経験によれば、#pragmaunrollは何もしてくれません。メタプログラミングを使用すると、Pythonで必要なサイズにループを動的に展開することができます。
ランタイムコード生成
PyCUDAの、pycuda.compiler.SourceModuleに与えるCUDA Cコードは定数である必要はありません。PyCUDAは、処理するコードの出所について何も想定していません。このため、コード生成に関連するロジックをアプリケーションのニーズに合わせて柔軟に設計することができます。
実行時にコードを生成するランタイムコード生成(RTCG: Run Time Code Generation) では、次の３つの方法があります。
単純なテキストのキーワード置換
テキストテンプレート
構文ツリーの構築
単純なテキストのキーワード置換
この単純な手法は、ソースコードで検索と置換を処理してコードを生成します。実行時にソースコードに型や定数を代入するなど、驚くほど広範囲のユースケースに適用することができます。 Cプリプロセッサマクロと組み合わせることで、その技術的範囲がさらに広がります。
Pythonの標準ライブラリstring などを使用すると、外部ソフトウェアに依存することなくキーワード置換を処理することができます。 テキストテンプレート
制御フローと条件分岐が必要だけれど、すべての入出力がテキストであれば、Webページの生成などで一般的に使用されているテンプレートエンジンを使用すると、キーワード置換の機能がスマートに処理することができます。
Python用のテンプレートエンジンは数多く存在しますが、最も有名なものの2つはJinja2とCheetahです。
code: python
from jinja2 import Template
tpl = Template(”””
__global__ void twice({{ type_name }} *tgt)
{
int idx = threadIdx.x +
{{ thread_block_size }} * {{ block_size }} * blockIdx.x;
{% for i in range( block_size ) %}
{% set offset = i * thread_block_size %}
{% endfor %}
}”””)
rendered tpl = tpl.render(
type_name = "float",
block_size = block_size,
thread_block_size = thread_block_size )
smod = SourceModule(rendered tpl)
構文ツリーの構築
生成されるコードがテキスト以外の場合、テンプレートエンジンの使用には限界があります。
この場合、ホスト言語でターゲットコードの完全な表現を導入することが適切になります。 最も一般的なそのような表現は、構文ツリーの形式です。 構文ツリーの構築により、ホスト言語のすべての機能を使用してコードを生成できます。 特に、テンプレート作成はほとんど「フラット」であり、出力の行に沿って方向付けられますが、構文ツリーの構築により、ユーザーは、たとえば関数の階層を使用して目的のコードを生成できます。
次のコードは、前の例と同じ展開されたベクトル加算コードを生成するために、
CodePyパッケージを使用して抽象構文木(AST: Abstract Syntax Trees)を構築する例です。
code: python
from codepy.cgen import (
FunctionBody, FunctionDeclaration,
Typedef, POD, Value,
Pointer, Module, Block, Initializer, Assign
)
from codepy.cgen.cuda import CudaGlobal
mod = Module([
FunctionBody(
CudaGlobal(FunctionDeclaration(
Value("void", "add"),
arg decls =[Pointer(POD(dtype, name))
Block([
Initializer (
POD(numpy.int32, "idx"),
"threadIdx .x + %d∗blockIdx.x"
% ( thread_block_size ∗ block_size )),
] + [
Assign(
"tgt idx+%d" % (o∗thread_block_size), o * thread_block_size ,
o * thread_block_size ))
for o in range( block size )]))])
smod = SourceModule(mod)
補足説明：
CodePyは、Python用のC / C ++メタプログラミングツールキットです。
ソースコードをコンパイルし、Pythonインタープリターに動的にロードすることができます。
参考