複数のデバイスにおける CUDALink

複数のデバイスにおける CUDALink

はじめに	メモリ
CUDALink 関数	帯域幅
CUDALink プログラミング

Wolfram言語の中核である関数型，およびリスト志向の特性により，CUDALink は組込みの即時データ並列化が提供でき，利用できるGPUカードに自動的に計算を分配する．

はじめに

まず CUDALink アプリケーションをロードする．

Wolfram Language code: Needs["CUDALink`"]

デバイス数と同じだけワーカーカーネルを開始する．

Wolfram Language code: LaunchKernels[$CUDADeviceCount]

$CUDADeviceCountはシステム上のデバイス数である．

$CUDADeviceCount

システム上のCUDAデバイス数

$CUDADeviceCountはシステム上のCUDA GPUデバイス数を取得する

CUDALink を全ワーカーカーネル上にロードする．

Wolfram Language code: ParallelNeeds["CUDALink`"]

CUDALink は既存のWolfram言語並列計算機能を使って複数のGPU上で実行する．このセクションでは，以下の関数を使う．

ParallelNeeds	パッケージを全並列サブカーネルにロードする
DistributeDefinitions	並列計算に必要な定義を分配する
ParallelEvaluate	利用できる全並列カーネルで入力式を評価して得られた結果のリストを返す

CUDALink はWolfram言語の並列計算機能で複数のGPUを使う

すべてのカーネル上の$CUDADevice変数を設定する．

Wolfram Language code: ParallelEvaluate[$CUDADevice = $KernelID]

CUDALink 関数

画像処理，線形代数，高速フーリエ(Fourier)変換等の高レベル CUDALink 関数は他のWolfram言語関数と同様に異なるカーネルで使用できる．違いは，$CUDADeviceが計算が行われるデバイスに設定されるということだけである．

画像の名前をExampleDataのTestImagesデータ集合から取られるように設定する．

Wolfram Language code: imgNames = {"Girl", "Girl2", "Girl3", "Couple"};

変数imgNamesをワーカーカーネルに分配する．

Wolfram Language code: DistributeDefinitions[imgNames]

ExampleDataから取られた画像にCUDAErosionを適用する．

Wolfram Language code: ParallelEvaluate[CUDAErosion[ExampleData[{"TestImage", imgNames[[$KernelID]]}], 2]]//AbsoluteTiming

速度が2倍改善されたことに注目されたい．これらの画像は小さく，データは転送されなければならないので，パフォーマンスは4倍は上がらない．

Wolfram Language code: (CUDAErosion[ExampleData[{"TestImage", #}], 2]& /@ imgNames)//AbsoluteTiming

その他の場合は，データの転送にかかる時間は計算に必要な時間に比べると多くない．ここでは2000のランダムな整数ベクトルを割り当てる．

Wolfram Language code: lsts = Table[RandomReal[1, 100], {ii, 2000}];

各デバイスにCUDAFoldをマップする．

Wolfram Language code: ParallelMap[CUDAFold[Plus, 0, #]&, lsts];//AbsoluteTiming

より速度が上がった．

Wolfram Language code: Map[CUDAFold[Plus, 0, #]&, lsts];//AbsoluteTiming

CUDALink プログラミング

CUDAFunctionは最適化されており，GPUに局所的であるため，DistributeDefinitionsを使ってワーカーカーネルと共有することはできない．このセクションではGPUをプログラムする別の方法を述べる．

2を足す

ベクトルに2を足す基本的なCUDAコードをロードする．

Wolfram Language code:

code = "
__global__ void addTwo(mint * in, mint * out, mint length) {
	mint index = threadIdx.x + blockIdx.x*blockDim.x;
	if (index < length)
		out[index] = in[index] + 2;
}";

CUDAFunctionをロードする．割当てにSetDelayedが使われている．これにより，DistributeDefinitionsはCUDAFunctionLoad呼出しのすべての従属変数を分配することができる．

Wolfram Language code: cudaFun := CUDAFunctionLoad[code, "addTwo", {{_Integer, _, "Input"}, {_Integer, _, "Output"}, _Integer}, 256]

入力パラメータを設定する．

Wolfram Language code:

listSize = 1000;
A = ConstantArray[1, {listSize}];
B = ConstantArray[1, {listSize}];

ワーカーカーネル間に定義を分配する．

Wolfram Language code: DistributeDefinitions[cudaFun, listSize, A, B]

異なるCUDAデバイスを使って各ワーカーカーネルでCUDAFunctionを実行する．

Wolfram Language code: ParallelEvaluate[res = cudaFun[A, B, listSize];]

最初の20要素を示す結果を集める．

Wolfram Language code: ParallelEvaluate[Take[First@res,20]]

マンデルブロ(Mandelbrot)集合

これは CUDALink ドキュメントの別のセクションで定義されたものと同じCUDAコードである．

Wolfram Language code:

src = "
__global__ void mandelbrot_kernel(char * set, float zoom, float bailout, mint width, mint height) {
   int xIndex = threadIdx.x + blockIdx.x*blockDim.x;
   int yIndex = threadIdx.y + blockIdx.y*blockDim.y;
   mint ii;

   Real_t x0 = zoom*(width/3 - xIndex);
   Real_t y0 = zoom*(height/2 - yIndex);
   Real_t tmp, x = 0, y = 0;

   if (xIndex < width && yIndex < height) {
       for (ii = 0; (x*x+y*y <= bailout) && (ii < MAX_ITERATIONS); ii++) {
            tmp = x*x - y*y +x0;
            y = 2*x*y + y0;
            x = tmp;
        }
        if (ii == MAX_ITERATIONS) {
            set[xIndex + yIndex*width] = 0;
        } else {
            set[xIndex + yIndex*width] = 1;
        }
    }
}
";

CUDAFunctionをロードする．

Wolfram Language code:

mfun := CUDAFunctionLoad[src, "mandelbrot_kernel", {{"Byte", _, "Output"}, "Float", "Float", _Integer, _Integer}, {16, 16}, "Defines" -> {"MAX_ITERATIONS" -> 1000}]

Wolfram Language code: {width, height} = {2048, 1024};

ワーカーカーネルと変数を共有する．

Wolfram Language code: DistributeDefinitions[mfun, width, height]

それぞれ異なるズームレベルでカーネルを開始し，"Bit"画像を返す．

Wolfram Language code:

ParallelEvaluate[
	mset = CUDAMemoryAllocate["Byte", {height, width}];
	res = mfun[mset, 0.001 * $KernelID, 8.0, width, height, {width, height}];
	Image[CUDAMemoryGet[First[res]], "Bit"]
	]

乱数生成器

以下のファイルにメルセンヌツイスタ(Mersenne Twister)が実装されている．

Wolfram Language code: srcf = FileNameJoin[{$CUDALinkPath, "SupportFiles", "random.cu"}]

関数をWolfram言語にロードする．

Wolfram Language code:

mersenneTwister := CUDAFunctionLoad[{srcf}, "MersenneTwister", {{_Real, _, "Output"}, {_Integer, _, "Input"}, {_Integer, _, "Input"}, {_Integer, _, "Input"}, {_Integer, _, "Input"}, _Integer}, 32];

CUDAFunctionのための入力変数を設定する．

Wolfram Language code:

MTRNGCount = 4096;
PATHN = 2 ^ 14;
NPerRNG = Ceiling[PATHN / MTRNGCount];
NPerRNG = If[EvenQ[NPerRNG], NPerRNG, NPerRNG + 1];
RANDN = MTRNGCount * NPerRNG;

mersenneTwister関数と入力パラメータをを分配する．

Wolfram Language code: DistributeDefinitions[mersenneTwister, MTRNGCount, PATHN, NPerRNG, RANDN]

シードの値を割り当てる．乱数が相関性を持たないようにするために，シードの評価は各ワーカーカーネルで行われなければならないことに注意．出力メモリも割り当てられ，計算が行われ，結果が可視化される．

Wolfram Language code:

ParallelEvaluate[
	{hsMatrixA, hsMaskB, hsMaskC} = RandomInteger[{-Developer`$MaxMachineInteger, Developer`$MaxMachineInteger}, {3, MTRNGCount}];
	hsSeed = RandomInteger[{-Developer`$MaxMachineInteger, Developer`$MaxMachineInteger}, MTRNGCount];output = CUDAMemoryAllocate[Real, RANDN];
	mersenneTwister[output, hsMatrixA, hsMaskB, hsMaskC, hsSeed, NPerRNG, MTRNGCount];
	ListPlot[CUDAMemoryGet[output]]
	]

メモリ

CUDAMemoryはカーネルとそれがロードされているデバイスの両方に結び付けられているため，ワーカーカーネルには分配できない．

マスターカーネルにメモリをロードする．

Wolfram Language code: x = CUDAMemoryLoad[{1, 2, 3}]

それから定義を分配する．

Wolfram Language code: DistributeDefinitions[x]

分配されたCUDAMemoryはワーカーカーネルによって操作することはできない．

Wolfram Language code: ParallelEvaluate[CUDAMemoryGet[x]]

メモリはParallelEvaluateを使ってワーカーカーネルにロードできる．

Wolfram Language code: ParallelEvaluate[mem = CUDAMemoryLoad[{1, 2, 3}]]

ParallelEvaluateを使ってメモリ上でさらに処理を行う．

Wolfram Language code: ParallelEvaluate[CUDAMemoryGet[mem]]

帯域幅

場合によっては，データの転送に費やす時間が計算に必要な時間より多くなることがある．

大きなリストをロードする．

Wolfram Language code: lsts = Table[RandomReal[1, 1000000], {ii, 10}];

並列処理では大きなリストをワーカーカーネルと共有しなければならないので，直列処理を行う場合より大幅に時間がかかる．

Wolfram Language code: ParallelMap[CUDASort, lsts];//AbsoluteTiming

データを転送する必要がない直列処理の方が速い．

Wolfram Language code: Map[CUDASort, lsts];//AbsoluteTiming

Top

その他のラーニングリソース

テクニカルサポート

Wolframソリューション

教育のためのWolframソリューション

使い始める

Grow Your Skills

Wolframと繋がる

大人用の教育プログラム

若者のための教育プログラム

読む

複数のデバイスにおける CUDALink

はじめに

CUDALink 関数

CUDALink プログラミング

2を足す

マンデルブロ(Mandelbrot)集合

乱数生成器

メモリ

帯域幅

複数のデバイスにおける CUDALink

はじめに

CUDALink 関数

CUDALink プログラミング

2を足す

マンデルブロ(Mandelbrot)集合

乱数生成器

メモリ

帯域幅

関連するガイド

関連するテクニカルノート