OpenCLFunctionLoad["src",fun,argtypes,blockdims]
文字列 src をコンパイルし,fun をOpenCLFunctionとしてWolfram言語内で使えるようにする.
OpenCLFunctionLoad[File[srcfile],fun,argtypes,blockdim]
ソースコードファイル srcfile をコンパイルしてから,fun をOpenCLFunctionとしてロードする.
OpenCLFunctionLoad[File[libfile],fun,argtypes,blockdim]
前にコンパイルされたライブラリ libfile から fun をOpenCLFunctionとしてロードする.
OpenCLFunctionLoad
OpenCLFunctionLoad["src",fun,argtypes,blockdims]
文字列 src をコンパイルし,fun をOpenCLFunctionとしてWolfram言語内で使えるようにする.
OpenCLFunctionLoad[File[srcfile],fun,argtypes,blockdim]
ソースコードファイル srcfile をコンパイルしてから,fun をOpenCLFunctionとしてロードする.
OpenCLFunctionLoad[File[libfile],fun,argtypes,blockdim]
前にコンパイルされたライブラリ libfile から fun をOpenCLFunctionとしてロードする.
詳細とオプション
- OpenCLLink アプリケーションは,Needs["OpenCLLink`"]を使ってロードしなければならない.
- libfile がダイナミックライブラリである場合には,ダイナミックライブラリ関数 fun がロードされる.
- 使用できる引数と戻り値の型,および対応するOpenCLの型:
-
_Integer mint Wolfram言語整数 "Integer32" int 32ビット整数 "Integer64" long/long long 64ビット整数 _Real Real_t GPUの実数型 "Double" double 機械倍数 "Float" float 機械浮動小数 {base, rank, io} OpenCLMemory 特定の基底型,階数,入出力オプションのメモリ "Local" | "Shared" mint ローカルメモリあるいは共有メモリのパラメータ {"Local" | "Shared", type} mint ローカルメモリあるいは共有メモリのパラメータ - 指定{base, rank, io}では,有効な io は,"Input","Output","InputOutput"である.
- {base}が渡されると,{base,_,"InputOutput"}がデフォルトで使用される.{base,rank}が渡されると,{base,rank,"InputOutput"}が使用される.
- 引数指定{base}は{base,_,"InputOutput"}に等しく,{base,rank}は{base,rank,"InputOutput"}に等しい.
- 階数は,{base,_,io}あるいは{base,io}を使って省略することができる.
- 使用可能な基底型:
-
_Integer _Real _Complex "Byte" "Bit16" "Integer32" "Byte[2]" "Bit16[2]" "Integer32[2]" "Byte[4]" "Bit16[4]" "Integer32[4]" "Byte[8]" "Bit16[8]" "Integer32[8]" "Byte[16]" "Bit16[16]" "Integer32[16]" "UnsignedByte" "UnsignedBit16" "UnsignedInteger" "UnsignedByte[2]" "UnsignedBit16[2]" "UnsignedInteger[2]" "UnsignedByte[4]" "UnsignedBit16[4]" "UnsignedInteger[4]" "UnsignedByte[8]" "UnsignedBit16[8]" "UnsignedInteger[8]" "UnsignedByte[16]" "UnsignedBit16[16]" "UnsignedInteger[16]" "Double" "Float" "Integer64" "Double[2]" "Float[2]" "Integer64[2]" "Double[4]" "Float[4]" "Integer64[4]" "Double[8]" "Float[8]" "Integer64[8]" "Double[16]" "Float[16]" "Integer64[16]" - OpenCLFunctionLoadは,異なる引数を使って2度以上呼び出すことができる.
- OpenCLFunctionLoadでロードされた関数は,Wolfram言語カーネルと同じプロセスで実行される.
- OpenCLFunctionLoadでロードされた関数は,Wolfram言語カーネルが終了する際にアンロードされる.
- ブロック次元はリストあるいは整数でよく,起動するのに各ブロックに対していくつのスレッドが必要であるかを表す.
- ブロック次元の最大サイズは,OpenCLInformationの"Maximum Work Group Size"特性によって返される.
- 起動時に,スレッドの数が(OpenCLFunctionの追加の引数として)指定されていない場合には,最大の階数と次元を持つ要素の次元が選ばれる.画像については,階数は2に設定される.
- 起動時に,スレッドの数がブロック次元の倍数ではない場合には,スレッド数はブロック次元の倍数になるようにインクリメントされる.
- 使用できるオプション:
-
"CompileOptions" {} 直接OpenCLコンパイラに渡されるコンパイルオプション "Defines" Automatic OpenCLプリプロセッサに渡される定義 "Device" $OpenCLDevice 計算に使用されるOpenCLデバイス "IncludeDirectories" {} コンパイルに含むディレクトリ "Platform" $OpenCLPlatform 計算に使用するOpenCLプラットフォーム "ShellCommandFunction" None コンパイルに使用するシェルコマンドで呼び出す関数 "ShellOutputFunction" None コンパイルコマンドを実行したシェル出力で呼び出す関数 "TargetPrecision" Automatic 計算に使用する精度 "WorkingDirectory" Automatic 一時的なファイルが生成されるディレクトリ
例題
すべて開く すべて閉じる例 (5)
Needs["OpenCLLink`"]src = "__kernel void myKernel( __global mint * global0Id, __global mint * global1Id, mint width, mint height) {
int xIndex = get_global_id(0);
int yIndex = get_global_id(1);
int index = xIndex + yIndex*width;
if (xIndex < width && yIndex < height) {
global0Id[index] = get_local_id(0);
global1Id[index] = get_local_id(1);
}
}";fun = OpenCLFunctionLoad[src, "myKernel", {{_Integer}, {_Integer}, _Integer, _Integer}, {16, 16}]width = 64;
height = 64;
global0Id = ConstantArray[0, {width, height}];
global1Id = ConstantArray[0, {width, height}];res = fun[global0Id, global1Id, width, height];ArrayPlotを使って結果をプロットする:
ArrayPlot /@ res"SupportFiles/vectorAdd.cl"からOpenCLソースファイルへのパスを定義する:
srcf = FileNameJoin[{$OpenCLLinkPath, "SupportFiles", "vectorAdd.cl"}]vectorAdd = OpenCLFunctionLoad[File[srcf], "vectorAdd", {{_Integer, _, "Input"}, {_Integer, _, "Input"}, {_Integer, _, "Output"}, _Integer}, 16]vectorAdd[Range[32], ConstantArray[2, 32], ConstantArray[0, 32], 32]OpenCLLink ライブラリの例である"addTwo_Dobule"を見付ける:
libPath = FindLibrary["addTwo_Double"]OpenCLFunctionLoadを使ってライブラリをロードする:
libFun = OpenCLFunctionLoad[File[libPath], "oAddTwo", {{_Integer, "Input"}, {_Integer, "Output"}}, 16];libFun[ConstantArray[1, 16], ConstantArray[1, 16]]この例のソースコードは OpenCLLink にバンドルされている:
FileNameJoin[{$OpenCLLinkPath, "CSource", "addTwo.cl"}]OpenCLFunctionを呼び出す際に追加の引数を与えることができる.引数は,起動するスレッドの数(あるいは大域的な作業グループサイズ)を示す.上の例を使う:
srcf = FileNameJoin[{$OpenCLLinkPath, "SupportFiles", "vectorAdd.cl"}]vectorAdd = OpenCLFunctionLoad[File[srcf], "vectorAdd", {{_Integer, _, "Input"}, {_Integer, _, "Input"}, {_Integer, _, "Output"}, _Integer}, 16]32のスレッドで関数を呼び出す.これにより,vectorAddの最初の32個の値のみが計算される:
vectorAdd[Range[64], ConstantArray[2, 64], ConstantArray[0, 64], 256, 32]コードに構文エラーがあると,"compilation failed(コンパイルの失敗)"というエラーが返される:
OpenCLFunctionLoad["__kernel void zero( __global mint * in, mint length) {
int index = get_global_id(0);
if (index < length)
in[index] = 0z;
}", "zero", {{_Integer}, _Integer}, {10}];"ShellOutputFunction" オプションを使ってビルドログを表示することができる:
OpenCLFunctionLoad["__kernel void zero( __global mint * in, mint length) {
int index = get_global_id(0);
if (index < length)
in[index] = 0z;
}", "zero", {{_Integer}, _Integer}, {10}, "ShellOutputFunction" -> Print];上のエラーは,コード中にタイポ(コード中の0の後にz)があることを示す:
OpenCLFunctionLoad["__kernel void zero( __global mint * in, mint length) {
int index = get_global_id(0);
if (index < length)
in[index] = 0;
}", "zero", {{_Integer}, _Integer}, {10}, "ShellOutputFunction" -> Print]スコープ (2)
テンプレートされた関数 (1)
テンプレートされた関数は,マクロを使ってシミュレーションを行うことができる.未定義のマクロとして
を残す:
src = "__kernel void imageColorNegate(__global Generic_t * in, __global Generic_t * out, mint width, mint height, mint channels) {
int ii;
int xIndex = get_global_id(0);
int yIndex = get_global_id(1);
int index = channels*(xIndex + yIndex*width);
if (xIndex < width && yIndex < height) {
for (ii = 0; ii < channels; ii++)
out[index+ii] = 255 - in[index+ii];
}
}";intColorNegate = OpenCLFunctionLoad[src, "imageColorNegate", {{_Integer, _, "Input"}, {_Integer, _, "Output"}, _Integer, _Integer, _Integer}, {16, 16}, "Defines" -> {"Generic_t" -> "mint"}]intColorNegate = OpenCLFunctionLoad[src, "imageColorNegate", {{"Float", _, "Input"}, {"Float", _, "Output"}, _Integer, _Integer, _Integer}, {16, 16}, "Defines" -> {"Generic_t" -> "float"}]共有メモリとローカルメモリ (1)
OpenCLFunctionLoadを使って起動に"Local"メモリと"Shared"メモリのどちらかを指定することができる.以下のコードは,共有メモリを使って勾配計算の大域メモリを保存する:
code = "
__kernel void grad(__global mint * img, mint n, __local mint * smem) {
int tx = get_local_id(0);
int bx = get_group_id(0);
int dx = get_local_size(0);
int index = tx + bx*dx;
#define S(txOffset) smem[txOffset + 1]
S(tx) = img[index];
if (tx == 0) {
S(tx - 1) = index > 0 ? img[index - 1] : 0;
} else if (tx == dx - 1) {
S(tx + 1) = index < n-1 ? img[index + 1] : 0;
}
barrier(CLK_LOCAL_MEM_FENCE);
tx += 1;
if (index < n)
img[index] = (S(tx + 1) - S(tx-1))/2;
}";入力引数を指定し,最後の引数を共有メモリについて"Shared"とする.ブロックの大きさを256と設定する:
fun = OpenCLFunctionLoad[code, "grad", {{_Integer}, _Integer, "Shared"}, 256]n = Times@@ImageDimensions[[image]];関数を呼び出す.共有メモリの大きさを(blockSize+2)⋆sizeof (int)と設定し,起動スレッド数を画像の平坦化された長さに設定する:
fun[[image], n, (256 + 2) * 4, n]共有メモリの大きさを指定するよりよい方法は,型を使う方法である:
fun = OpenCLFunctionLoad[code, "grad", {{_Integer}, _Integer, {"Shared", _Integer}}, 256]fun[[image], n, 256 + 2, n]アプリケーション (10)
画像入力 (1)
入力は画像でもよい.画像間の線形補間を行うコードを書く(これはImageComposeを使って行うこともできる):
src = "
__kernel void linearCombine(__global mint * output, __global mint * input0, __global mint * input1, float a, float b, mint width, mint height, mint channels) {
int xIndex = get_global_id(0);
int yIndex = get_global_id(1);
if (xIndex >= width || yIndex >= height)
return ;
int pos = channels * (yIndex * width + xIndex);
for (int ii = 0; ii < channels; ii++) {
output[pos + ii] = input0[pos + ii] * a + input1[pos + ii] * b;
}
}
";上のソースコードからOpenCLFunctionをロードする:
ImageLinearCombine = OpenCLFunctionLoad[src, "linearCombine", {{_Integer, _, "Output"}, {_Integer, _, "Input"}, {_Integer, _, "Input"}, "Float", "Float", _Integer, _Integer, _Integer}, {16, 16}]height,width,channel の値を設定する.またメモリをoutputに割り当てる:
{height, width, channels} = Flatten[{ImageDimensions[[image]], ImageChannels[[image]]}];
output = OpenCLMemoryAllocate[Integer, {width, height, channels}]ImageLinearCombine[output, [image], [image], 0.5, 0.5, width, height, channels, {width, height}]Image[OpenCLMemoryGet[output], "Byte"]上を使って,関数OpenCLImageLinearCombineを作成することができる:
OpenCLImageLinearCombine[input0_Image, a_Real, input10_Image, b_Real] :=
Module[{input1, width, height, channels, output},
input1 = If[ImageDimensions[input0] === ImageDimensions[input10],
input10,
ImageResize[input10, ImageDimensions[input0]]
];
{height, width, channels} = Flatten[{ImageDimensions[input0], ImageChannels[input0]}];
output = OpenCLMemoryAllocate[Integer, {width, height, channels}];
ImageLinearCombine[output, input0, input1, a, b, width, height, channels, {width, height}];
With[{res = Image[OpenCLMemoryGet[output], "Byte"]},
OpenCLMemoryUnload[output];
res
]
]関数はImageComposeと似た構文を持つ:
OpenCLImageLinearCombine[[image], -1.0, [image], 2.0]Manipulateを使って補間係数を操作することができる:
Manipulate[OpenCLImageLinearCombine[[image], a, [image], b], {{a, 0.0}, -2.0, 1.0, 0.01}, {{b, 1.0}, -2.0, 2.0, 0.01}]効果が作成できる.以下では滑らかなアニメーションを見ることができる:
Animate[OpenCLImageLinearCombine[[image], ii, [image], 1 - ii], {ii, 0.0, 1.0}]一様な乱数生成 (1)
一様な乱数生成器は,多くのアプリケーションにおいてよく見られるシード問題である.OpenCLに一様な乱数生成器を実装する:
srcf = FileNameJoin[{$OpenCLLinkPath, "SupportFiles", "URNG_Kernels.cl"}]ソースをOpenCLFunctionとしてロードする.このアルゴリズムは画像を使って乱数の上限を提供する:
urng = OpenCLFunctionLoad[File[srcf], "noise_uniform", {{"Float[4]", _, "Input"}, {"Float[4]", _, "Output"}, _Integer}, {64, 1}]OpenCLFunctionを呼び出す.画像は,指定された適切な型を使って解釈することができる限り,OpenCLFunctionに直接渡すことができる:
res = urng[[image], [image], 1, {512, 512}]以下は通常のアヒルの画像とは異なっている.これはアルファチャンネルが1に(SetAlphaChannelを使って)設定された4チャンネル画像である:
ImageChannels[[image]]ImageAdd[[image], EdgeDetect[First[res], 15]]メルセンヌツイスタを使った乱数生成 (1)
メルセンヌ(Mersenne)ツイスタは,もう一つの一様な乱数生成器アルゴリズムである(上のアルゴリズムよりもこちらの方が高度である).実装は以下の場所にある:
srcf = FileNameJoin[{$OpenCLLinkPath, "SupportFiles", "MersenneTwister_kernel.cl"}]OpenCLFunctionをロード,つまり型_Realを指定する.これは,Realの型はCPUの機能(これが倍数精度をサポートするかどうか)に依存することを意味する:
mersenneTwister = OpenCLFunctionLoad[{srcf}, "MersenneTwister", {{_Real, _, "Output"}, {_Integer, _, "Input"}, {_Integer, _, "Input"}, {_Integer, _, "Input"}, {_Integer, _, "Input"}, _Integer}, 32]メルセンヌツイスタの入出力パラメータを設定する(詳細についてはアルゴリズムのページを参照のこと):
MTRNGCount = 4096;
PATHN = 2 ^ 25;
NPerRNG = Ceiling[PATHN / MTRNGCount];
NPerRNG = If[EvenQ[NPerRNG], NPerRNG, NPerRNG + 1];
RANDN = MTRNGCount * NPerRNG;
{hsMatrixA, hsMaskB, hsMaskC} = RandomInteger[{-2147483647, 2147483647}, {3, MTRNGCount}];
hsSeed = RandomInteger[{-2147483647, 2147483647}, MTRNGCount];
output = OpenCLMemoryAllocate[Real, RANDN]OpenCLFunctionを呼び出す:
mersenneTwister[output, hsMatrixA, hsMaskB, hsMaskC, hsSeed, NPerRNG, MTRNGCount]ListPlot[OpenCLMemoryGet[output][[ ;; 1000]]]mersenneTwister[output, hsMatrixA, hsMaskB, hsMaskC, hsSeed, NPerRNG, MTRNGCount];//AbsoluteTimingBlockRandom[SeedRandom[1, Method -> "MersenneTwister"];RandomReal[1, RANDN]];//AbsoluteTiming接頭部和のアルゴリズム (1)
スキャン,つまり接頭辞の総和のアルゴリズムは,FoldListに似ており,さまざまな場合に使えて大変便利なプリミティブアルゴリズムである.OpenCLの実装は以下の場所にある:
srcf = FileNameJoin[{$OpenCLLinkPath, "SupportFiles", "Scan.cl"}]scanExclusiveShared = OpenCLFunctionLoad[{srcf}, "scanExclusiveLocal1", {{"Integer32", "Output"}, {"Integer32", "Input"}, {"Local", "Integer32"}, "Integer32"}, 256, "Defines" -> {"WORKGROUP_SIZE" -> 256}];
scanExclusiveShared2 = OpenCLFunctionLoad[{srcf}, "scanExclusiveLocal2", {{"Integer32", "InputOutput"}, {"Integer32", "Output"}, {"Integer32", "Input"}, {"Local", "Integer32"}, "Integer32", "Integer32"}, 256, "Defines" -> {"WORKGROUP_SIZE" -> 256}];
uniformUpdate = OpenCLFunctionLoad[{srcf}, "uniformUpdate", {{"Integer32", "InputOutput"}, {"Integer32", "InputOutput"}}, 256, "Defines" -> {"WORKGROUP_SIZE" -> 256}];data = RandomInteger[10, 256];dest = OpenCLMemoryAllocate[Integer, 256];blockDim = 256;
gridDim = blockDim * Ceiling[Length[data] / (4 * blockDim)];buffer = OpenCLMemoryAllocate[Integer, 1 + (gridDim / blockDim)];scanExclusiveShared[dest, data, 512, 4 * blockDim, gridDim];
scanExclusiveShared2[buffer, dest, data, 512, 1 + (gridDim / blockDim), 1 + (gridDim / blockDim)];
uniformUpdate[dest, buffer];OpenCLMemoryGet[dest]OpenCLMemory要素を解放する:
OpenCLMemoryUnload[dest, buffer]行列操作 (1)
行列転置は,多くのアプリケーションにおける基本的なアルゴリズムである.入力を指定する:
size = 5;
input = RandomReal[1., {size, size}];
output = ConstantArray[0., {size, size}];
localWorkSize = 16;
globalworkSize = Ceiling[size / localWorkSize] * localWorkSize;OpenCLFunctionをロードする:
oclTranspose = OpenCLFunctionLoad[{FileNameJoin[{$OpenCLLinkPath, "SupportFiles", "transpose.cl"}]}, "transpose", {{"Float", _, "Output"}, {"Float", _, "Input"}, _Integer, _Integer, _Integer, {"Shared", _Integer}}, {localWorkSize, localWorkSize}]OpenCLFunctionを呼び出す:
res = oclTranspose[output, input, 0, size, size, localWorkSize * localWorkSize, {globalworkSize, globalworkSize}];結果のMatrixFormを示す:
First[res]//MatrixFormMatrixForm[Transpose[input]]行列乗算 (1)
srcf = FileNameJoin[{$OpenCLLinkPath, "SupportFiles", "matrixMul.cl"}]blockSize = 4;OpenCLFunctionをロードする.入力が階数2でなければならないことが指定されている:
MatrixMultiply = OpenCLFunctionLoad[{srcf}, "matrixMul", {{"Float", 2, "Output"}, {"Float", 2, "Input"}, {"Float", 2, "Input"}, {"Local", "Float"}, {"Local", "Float"}, _Integer, _Integer}, {blockSize, blockSize}, "Defines" -> {"BLOCK_SIZE" -> blockSize}]A = RandomReal[1.0, {8, 8}];
B = RandomReal[1.0, {8, 8}];
out = OpenCLMemoryAllocate["Float", {8, 8}]OpenCLFunctionを呼び出す:
MatrixMultiply[out, A, B, blockSize * blockSize, blockSize * blockSize, 8, 8]OpenCLMemoryGetを使って出力メモリを得る:
OpenCLMemoryGet[out]//MatrixFormDot[A, B]//MatrixForm高速フーリエ変換 (1)
一次元の離散高速フーリエ(Fourier)変換は,OpenCLLink を使って実装することができる.この実装は,入力が2のベキであると仮定する:
srcf = FileNameJoin[{$OpenCLLinkPath, "SupportFiles", "FFT_Kernels.cl"}]OpenCLFunctionLoadを使ってOpenCLFunctionをロードする:
fft = OpenCLFunctionLoad[{srcf}, "kfft", {{"Float", _, "InputOutput"}, {"Float", _, "Output"}}, 64]in = RandomReal[1.0, 1024];
out = ConstantArray[0.0, 1024];出力メモリを呼び出して複雑なリストを作成し,最初の50個の要素だけを表示する:
MapThread[Complex, fft[in, out, 64]][[ ;; 50]]上の結果は,Fourierを使った場合と同じである:
Fourier[in, FourierParameters -> {1, -1}][[ ;; 50]]金融派生商品 (1)
ブラック・ショールズ(Black–Scholes)は,金融派生商品の投資をモデル化するもので,これはOpenCLに実装されている:
srcf = FileNameJoin[{$OpenCLLinkPath, "SupportFiles", "BlackScholes.cl"}]OpenCLFunctionをロードする:
BlackScholes = OpenCLFunctionLoad[{srcf}, "BlackScholes", {{_Real, _, "Output"}, {_Real, _, "Output"}, {_Real, _, "Input"}, {_Real, _, "Input"}, {_Real, _, "Input"}, _Real, _Real, _Integer}, 128, "TargetPrecision" -> "Single"]numberOfOptions = 64;
call = OpenCLMemoryAllocate["Float", numberOfOptions];
put = OpenCLMemoryAllocate["Float", numberOfOptions];
currentPrices = RandomReal[{25.0, 35.0}, numberOfOptions];
strikePrices = RandomReal[{20.0, 40.0}, numberOfOptions];
strikeTimes = RandomReal[{0.1, 10.0}, numberOfOptions];
riskFree = 0.02;
volatility = 0.30;OpenCLFunctionを呼び出す:
BlackScholes[call, put, currentPrices, strikePrices, strikeTimes, riskFree, volatility, numberOfOptions, numberOfOptions]OpenCLMemoryGet[call]結果はFinancialDerivativeの出力に一致する:
MapThread[FinancialDerivative[{"European", "Call"}, {"StrikePrice" -> #1, "Expiration" -> #2}, {"InterestRate" -> riskFree, "Volatility" -> volatility, "CurrentPrice" -> #3}]&, {strikePrices, strikeTimes, currentPrices}]numberOfOptions = 2048;
call = OpenCLMemoryAllocate["Float", numberOfOptions];
put = OpenCLMemoryAllocate["Float", numberOfOptions];
currentPrices = RandomReal[{25.0, 35.0}, numberOfOptions];
strikePrices = RandomReal[{20.0, 40.0}, numberOfOptions];
strikeTimes = RandomReal[{0.1, 10.0}, numberOfOptions];
riskFree = 0.02;
volatility = 0.30;C2050では,2048個のオプションを評価するのに1/100秒かかる:
BlackScholes[call, put, currentPrices, strikePrices, strikeTimes, riskFree, volatility, numberOfOptions, numberOfOptions];//AbsoluteTimingCore i7 950では,, FinancialDerivative は1.13秒かかる.上は280倍の速さである.オプションの数を増やすと,速さの違いがより大きくなる:
MapThread[FinancialDerivative[{"European", "Call"}, {"StrikePrice" -> #1, "Expiration" -> #2}, {"InterestRate" -> riskFree, "Volatility" -> volatility, "CurrentPrice" -> #3}]&, {strikePrices, strikeTimes, currentPrices}];//AbsoluteTimingガウシアンフィルタ (1)
再帰的なガウス行列がガウシアンフィルタを近似するのに使われる.アルゴリズムは,ガウス行列が分離可能であるという事実に依存している:
GaussianMatrix[10]//ArrayPlotOuter[Times, GaussianMatrix[{{10}}], GaussianMatrix[{{10}}]]//ArrayPlotsrcf = FileNameJoin[{$OpenCLLinkPath, "SupportFiles", "recursiveGaussian.cl"}]OpenCLFunctionLoadを使って2つの関数をロードする:
RecusiveGaussian = OpenCLFunctionLoad[File[srcf], "RecursiveGaussian_kernel", {{"UnsignedByte[4]", _, "Input"}, {"UnsignedByte[4]", _, "Output"}, _Integer, _Integer, "Float", "Float", "Float", "Float", "Float", "Float", "Float", "Float"}, {256, 1}];
OpenCLTranspose = OpenCLFunctionLoad[File[srcf], "transpose_kernel", {{"UnsignedByte[4]", _, "Output"}, {"UnsignedByte[4]", _, "Input"}, "Local", _Integer, _Integer, _Integer}, {16, 16}];σ = 5.0;PDF[NormalDistribution[μ, σ], p]Plot[PDF[NormalDistribution[0, 5.0], x], {x, -6, 6}, Filling -> Axis]alpha = 1.695 / σ;
ema = Exp[-alpha];
ema2 = Exp[-2 * alpha];
k = (1 - ema) * (1 - ema) / (1 + 2 * alpha * ema - ema2);
a0 = k;
a1 = k * (alpha - 1) * ema;
a2 = k * (alpha + 1) * ema;
a3 = -k * ema2;
b1 = -2 * ema;
b2 = ema2;
coefp = (a0 + a1) / (1 + b1 + b2);
coefn = (a2 + a3) / (1 + b1 + b2);入力,出力,一時的なストレッジにOpenCLMemoryを割り当てる:
{width, height} = ImageDimensions[[image]];
input = OpenCLMemoryLoad[[image], "UnsignedByte[4]"];
temp = OpenCLMemoryAllocate["UnsignedByte[4]", {width, height}];
output = OpenCLMemoryAllocate["UnsignedByte[4]", {width, height}];関数を呼び出す.まず水平にガウス行列を行ってから転置し,次にガウス行列を垂直に行ってから転置し,完全なガウス行列を得る:
RecusiveGaussian[input, temp, width, height, a0, a1, a2, a3, b1, b2, coefp, coefn, {width, 1}];
OpenCLTranspose[output, temp, 16 * 16 * 4, width, height, 16, {width, height}];
RecusiveGaussian[output, temp, width, height, a0, a1, a2, a3, b1, b2, coefp, coefn, {height, 1}];
OpenCLTranspose[output, temp, 16 * 16 * 4, width, height, 16, {width, height}];Image[OpenCLMemoryGet[output], "Byte", "ColorSpace" -> "RGB"]AbsoluteTiming[
RecusiveGaussian[input, temp, width, height, a0, a1, a2, a3, b1, b2, coefp, coefn, {width, 1}];
OpenCLTranspose[output, temp, 16 * 16 * 4, width, height, 16, {width, height}];
RecusiveGaussian[output, temp, width, height, a0, a1, a2, a3, b1, b2, coefp, coefn, {height, 1}];
OpenCLTranspose[output, temp, 16 * 16 * 4, width, height, 16, {width, height}];]GaussianFilter[[image], {10, σ}];//AbsoluteTimingソート (1)
バイトニックソートは,整数の任意集合をソートする.これは原則としてマージソートに似ている.OpenCLの実装は,2のベキの長さのリストのみに使えるもので,以下の場所にある:
srcf = FileNameJoin[{$OpenCLLinkPath, "SupportFiles", "bitonicSort.cl"}]BitonicSort = OpenCLFunctionLoad[{srcf}, "bitonicSort", {{_Integer, _, "InputOutput"}, _Integer, _Integer, _Integer, _Integer}, 16]入力の長さを設定し,これをロードする.方向は,最高から最低へ,あるいは最低から最高へソートするかを示す.この場合は最低から最高へソートする:
width = 64;
list = OpenCLMemoryLoad[Reverse[Range[width]]];
direction = 1;OpenCLMemoryGet[list]マージソートに似たバイトニックソートを呼び出す.完全ソートを行うためには,複数回呼出しを行う必要がある:
For[stage = 0, stage < Log[2, width], stage++,
For[passOfStage = 0, passOfStage <= stage, passOfStage++, BitonicSort[list, stage, passOfStage, width, direction]
]
]OpenCLMemoryGet[list]考えられる問題 (5)
最大作業項目の大きさ(ブロック次元)はOpenCLInformationで返される:
OpenCLInformation[$OpenCLPlatform, $OpenCLDevice, "Maximum Work Item Sizes"]OpenCLコードで倍精度操作を使用するには,コードヘッダに以下のpramasを置く必要がある:
#ifdef USING_DOUBLE_PRECISIONQ
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
#pragma OPENCL EXTENSION cl_amd_fp64 : enable
#endif /* USING_DOUBLE_PRECISIONQ */
関数の呼出しにおけるエラーは,OpenCLLink を使用不可な状態に置くことがある.これはユーザが任意のカーネルを書くことができるようにしたことによる副次的な結果である.カーネルコード内の無限ループ,バッファオーバーフロー等は,OpenCLLink およびビデオドライバを使用不可なものにすることがある.極端な場合には,表示ドライバがクラッシュする場合もあるが,通常はOpenCLコードをさらに評価した場合に無効な結果を返すだけである.
OpenCL実装のバグの中には,"IncludeDirectories"の1つにスペースが含まれていると,カーネルをクラッシュさせるものもある.
__constant等のメモリ修飾子の使用は,OpenCLLink ではサポートされていない.OpenCLFunctionに渡されるメモリは__globalでなければならない.
インタラクティブな例題 (5)
マンデルブロ集合 (1)
マンデルブロ集合は,再帰方程式
(
は複素数)を満足するすべての点をプロットする.以下ではOpenCLにおいて集合を実装する(滑らかに色の移行が行われるように,少し複雑な色付け方法が使われている):
src = "
__kernel void mandelbrot_kernel(__global mint * set, float zoom, float bailout, mint width, mint height) {
int xIndex = get_global_id(0);
int yIndex = get_global_id(1);
int ii;
float x0 = zoom*(width/3 - xIndex);
float y0 = zoom*(height/2 - yIndex);
float tmp, x = 0, y = 0;
float c;
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;
}
c = ii - log(log(sqrt(x*x + y*y)))/log(2.0f);
if (ii == MAX_ITERATIONS) {
set[3*(xIndex + yIndex*width)] = 0;
set[3*(xIndex + yIndex*width) + 1] = 0;
set[3*(xIndex + yIndex*width) + 2] = 0;
} else {
set[3*(xIndex + yIndex*width)] = ii*c/4 + 20;
set[3*(xIndex + yIndex*width) + 1] = ii*c/4;
set[3*(xIndex + yIndex*width) + 2] = ii*c/4 + 5;
}
}
}
";MandelbrotSet = OpenCLFunctionLoad[src, "mandelbrot_kernel", {{_Integer, _, "Output"}, "Float", "Float", _Integer, _Integer}, {16, 16}, "Defines" -> {"MAX_ITERATIONS" -> 100}]width = 2048;
height = 1024;
mem = OpenCLMemoryAllocate[Integer, {height, width, 3}];res = MandelbrotSet[mem, 0.0017, 8.0, width, height, {width, height}]Image[OpenCLMemoryGet[First[res]], "Byte"]Manipulate[
MandelbrotSet[mem, zoom, 8.0, width, height, {width, height}];
Image[OpenCLMemoryGet[First[res]], "Byte"], {{zoom, 0.0017}, 0.0001, 0.003, 0.0001}]ジュリア集合 (1)
マンデルブロ集合は,ジュリア集合が制約された形である.以下はジュリア集合のコードである:
code = "
#ifdef USING_DOUBLE_PRECISIONQ
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
#pragma OPENCL EXTENSION cl_amd_fp64 : enable
#endif /* USING_DOUBLE_PRECISIONQ */
__kernel void julia_kernel(__global Real_t * set, mint width, mint height, Real_t cx, Real_t cy) {
int xIndex = get_global_id(0);
int yIndex = get_global_id(1);
int ii;
Real_t x = ZOOM_LEVEL*(width/2 - xIndex);
Real_t y = ZOOM_LEVEL*(height/2 - yIndex);
Real_t tmp;
Real_t c;
if (xIndex < width && yIndex < height) {
for (ii = 0; ii < MAX_ITERATIONS && x*x + y*y < BAILOUT; ii++) {
tmp = x*x - y*y + cx;
y = 2*x*y + cy;
x = tmp;
}
c = log(0.1f + sqrt(x*x + y*y));
set[xIndex + yIndex*width] = c;
}
}
";{width, height} = {512, 512};
jset = OpenCLMemoryAllocate[Real, {height, width}];OpenCLFunctionをロードする:
JuliaCalculate = OpenCLFunctionLoad[code, "julia_kernel", {{_Real, _, "Output"}, _Integer, _Integer, _Real, _Real}, {16, 16}, "Defines" -> {"MAX_ITERATIONS" -> 10, "ZOOM_LEVEL" -> "0.0050", "BAILOUT" -> "4.0"}];ジュリア集合を計算し,これをReliefPlotを使ってプロットする:
Manipulate[
JuliaCalculate[jset, width, height, c[[1]], c[[2]], {width, height}];
ReliefPlot[OpenCLMemoryGet[jset], DataRange -> {{-2.0, 2.0}, {-2.0, 2.0}}, ImageSize -> 256, ColorFunction -> "SunsetColors"],
{{c, {0, 1}}, {-2, -2}, {2, 2}, Locator}]ジュリア集合を計算し,これをグレイスケール画像として表示する:
Manipulate[JuliaCalculate[jset, width, height, c, d, {width, height}]; Image[OpenCLMemoryGet[jset], ImageSize -> 256], {{c, 0.0}, -2.0, 2.0, Slider}, {{d, 0.0}, -2.0, 2.0, Slider}]画像調整 (1)
ImageAdjustは,画像をスケールし直して高い値と低い値を入力する.ガンマ補正も考慮する.以下では,簡単な形のImageAdjustをOpenCLで定義する:
src = "
mint xclamp(mint val, mint low, mint high) {
return val <= low ? low : (val >= high ? high : val);
}
mint adjust(mint pixel, float lowIn, float highIn, float lowOut, float highOut, float gamma) {
float res, val;
val = xclamp(pixel, lowIn, highIn);
res = pow((val - lowIn) / (highIn - lowIn), gamma);
res = res * (highOut - lowOut) - lowOut;
return res + 0.5f;
}
__kernel void imageAdjust(__global mint * img, mint width, mint height, mint channels, float lowIn, float highIn, float lowOut, float highOut, float gamma) {
int xIndex = get_global_id(0);
int yIndex = get_global_id(1);
if (xIndex >= width || yIndex >= height)
return ;
int pos = channels * (yIndex * width + xIndex);
for (int ii = 0; ii < channels; ii++) {
img[pos + ii] = adjust(img[pos + ii], 255*lowIn, 255*highIn, 255*lowOut, 255*highOut, gamma);
}
}";OpenCLFunctionをロードする:
cOpenCLImageAdjust = OpenCLFunctionLoad[src, "imageAdjust", {{_Integer}, _Integer, _Integer, _Integer, "Float", "Float", "Float", "Float", "Float"}, {16, 16}];簡単なWolfram言語ラッパー関数を定義して,OpenCL関数がImageAdjustと似た構文を持つようにする:
OpenCLImageAdjust[img_Image, {lowIn_Real, highIn_Real}, gamma_ : 1.0] /; Head[gamma] == Real := OpenCLImageAdjust[img, {lowIn, highIn}, {0.0, 1.0}, gamma]
OpenCLImageAdjust[img_Image, {lowIn_Real, highIn_Real}, {lowOut_Real, highOut_Real}, gamma_ : 1.0] :=
Module[{width, height, channels},
{height, width, channels} = Flatten[{ImageDimensions[img], ImageChannels[img]}];
cOpenCLImageAdjust[img, width, height, channels, lowIn, highIn, lowOut, highOut, gamma, {width, height}]//First
]0.3から0.8の間の値を0.0から1.0の間の値にスケールし直すことで,画像を調節する:
OpenCLImageAdjust[[image], {0.3, 0.8}]Manipulateを使って値をスケールし直すことによって画像を調整する:
Manipulate[OpenCLImageAdjust[[image], {0.1, high}], {high, 0.11, 1.0, 0.01}]0.3から0.8の間の値を0.0から1.0の間の値にスケールし直すことで,画像を調節する:
OpenCLImageAdjust[[image], {0.3, 0.8}, {0.0, 1.0}]跳ねるボール (1)
以下の例では,さまざまな初期力で箱の中の各素粒子の位置を計算する.素粒子物理学の部分はOpenCLで行い,残りの作業はすべてWolfram言語で行う:
BallBounceEffect[bb1_] :=
Module[{tsize, fsize, z, r1, r, v, acc, device, state, BlockDim, GridDim, res, vc},
tsize = 100;
fsize = 80;
z = Table[fsize - 2 + RandomReal[2], {i, tsize * tsize}];
r1 = Table[.15 + .1 * Sin[2 * N[Pi] * (i + j) / tsize], {i, tsize}, {j, tsize}];
First[r1];
r = Flatten[r1];
v = ConstantArray[0.0, tsize * tsize];
acc = 2;
device = Automatic;
state = ConstantArray[1, tsize * tsize];
BlockDim = 256;
GridDim = Ceiling[(tsize * tsize) / BlockDim] * BlockDim;
vc = Flatten[
Table[
RGBColor[0, .5 + (.25 - r[[(i - 1) * tsize + j]]) * 2, .5 + (.25 - r[[(i - 1) * tsize + j]]) * 2 ], {i, tsize}, {j, tsize}]
];
Graphics3D[{AbsolutePointSize[0],
Point[Dynamic[Refresh[
res = bb1[v, z, r, state, acc, tsize, GridDim];
z = res[[2]];
v = res[[1]];
state = res[[3]];
Flatten[Table[{i, j, z[[(i - 1) * tsize + j]]}, {i, tsize}, {j, tsize}], 1], UpdateInterval -> 0]], VertexColors -> vc], Sphere[{tsize / 2, tsize / 2, -10}, .05], Sphere[{tsize / 2, tsize / 2, fsize}, .05],
{Black, Polygon[{{tsize, 0, 0}, {tsize, tsize, 0}, {0, tsize, 0}, {0, 0, 0}}]}
}, Boxed -> False]
]OpenCLコードを定義して,関数をWolfram言語内にロードする:
BallBouncePatternDemo[] :=
Module[{code, bb1, BlockDim},
code = "
__kernel void bb(__global float* v, __global float* z,__global float* r,__global mint *state, mint ac, mint size ) {
int i=get_global_id(0);
float acc=ac/10.0;
if(i < size * size ) {
if(v[i]<=0) {
v[i]=0;
state[i]=1;
}
if(z[i]<=0) {
z[i]=0;
state[i]=-1;
}
v[i]+=state[i]*acc;
z[i]-=state[i]*v[i]*r[i]/.25;
}
}";
BlockDim = 256;
bb1 = OpenCLFunctionLoad[code, "bb", {{"Float"}, {"Float"}, {"Float", _, "Input"}, {_Integer}, _Integer, _Integer}, {BlockDim}];
Mouseover[Graphics[{LightGray, Circle[], Inset[Style["Bring Mouse Here", Bold, Blue]]}],
BallBounceEffect[bb1]
]
]BallBouncePatternDemo[][image]n体シミュレーション (1)
n体シミュレーションは,古典的なニュートン問題である.これをOpenCLで実装する:
srcf = FileNameJoin[{$OpenCLLinkPath, "SupportFiles", "NBody.cl"}];OpenCLFunctionをロードする:
NBody = OpenCLFunctionLoad[{srcf}, "nbody_sim", {{"Float[4]", _, "Input"}, {"Float[4]", _, "Input"}, _Integer, "Float", "Float", {"Local", "Float"}, {"Float[4]", _, "Output"}, {"Float[4]", _, "Output"}}, 256]numParticles = 1024;
deltaT = 0.05;
epsSqrt = 50.0;pos = OpenCLMemoryLoad[RandomReal[512, {numParticles, 4}], "Float[4]"];
vel = OpenCLMemoryLoad[RandomReal[1, {numParticles, 4}], "Float[4]"];
newPos = OpenCLMemoryAllocate["Float[4]", {numParticles}];
newVel = OpenCLMemoryAllocate["Float[4]", {numParticles}];NBody[pos, vel, numParticles, deltaT, epsSqrt, 256 * 4, newPos, newVel, 1024];
NBody[newPos, newVel, numParticles, deltaT, epsSqrt, 256 * 4, pos, vel, 1024];Graphics3D[Point[Take[#, 3]& /@ OpenCLMemoryGet[pos]]]結果をDynamicで示す:
Graphics3D[Point[
Dynamic[Refresh[
NBody[pos, vel, numParticles, deltaT, epsSqrt, 256 * 4, newPos, newVel, 1024];
NBody[newPos, newVel, numParticles, deltaT, epsSqrt, 256 * 4, pos, vel, 1024];
Take[#, 3]& /@ OpenCLMemoryGet[pos], UpdateInterval -> 0]]]]おもしろい例題 (1)
SymbolicC (1)
OpenCLLink ではSymbolicCのコード生成機能を使用することができる.SymbolicCを使うには,ユーザはこれをロードしなければならない:
Needs["SymbolicC`"]OpenCLLink ではSymbolicCのコード生成機能を使用することができる.ここでは,Wolfram言語文を取って,これをSymbolicC 式に変換するメソッドであるtoSymbolicCを定義する(Wolfram言語のコマンドすべてを変換することはできないが,これらのコマンドをユーザが加えることはできる):
ClearAll[toSymbolicC]
SetAttributes[toSymbolicC, {HoldAll}]
toSymbolicC[x_List] := toSymbolicC /@ x
toSymbolicC[Times[-1, x_]] := "-" <> GenerateCode[toSymbolicC[x]]
toSymbolicC[(op : (Plus | Times))[args___]] := COperator[op, toSymbolicC[{args}]]
toSymbolicC[(op : (Minus | BitNot | Not | Decrement | Increment | PreDecrement | PreIncrement))[x_]] := COperator[op, toSymbolicC[x]]
toSymbolicC[(op : (Mod | Divide | Subtract | BitShiftRight | BitShiftLeft))[x_, y_]] := COperator[op, {toSymbolicC[x], toSymbolicC[y]}]
toSymbolicC[(op : (ArcCos | ArcSin | Ceiling | Cos | Cosh | Exp | Abs | Floor | Sin | Sinh | Sqrt | Tan | Tanh | Log))[x_]] := CStandardMathOperator[op, toSymbolicC[x]]
toSymbolicC[Power[x_, r : Rational[_, _]]] := CStandardMathOperator[Power, {toSymbolicC[x], toSymbolicC[r]}]
toSymbolicC[Power[x_, 2]] := COperator[Times, {toSymbolicC[x], toSymbolicC[x]}]
toSymbolicC[Power[x_, y_]] := CStandardMathOperator[Power, {toSymbolicC[x], toSymbolicC[y]}]
toSymbolicC[CompoundExpression[stmts__]] := toSymbolicC /@ stmts
toSymbolicC[If[cond_, trueStmt_]] := CIf[toSymbolicC[cond], toSymbolicC[trueStmt]]
toSymbolicC[If[cond_, trueStmt_, falseStmt_]] := CIf[toSymbolicC[cond], toSymbolicC[trueStmt], toSymbolicC[falseStmt]]
toSymbolicC[x_] := xtoSymbolicC[Cos[x] ^ 2 + Sin[x] * 2 + x ^ 8 + 3]Cに変換するには,ToCCodeStringを使う:
ToCCodeString[%]これを OpenCLLink の記号的コード生成機能に結び付けてOpenCLMapSource関数を作成することができる:
SetAttributes[OpenCLMapSource, {HoldAll}];
OpenCLMapSource[f_] := ToCCodeString[With[{fun = f[xx] /. xx -> CArray["in", "index"]}, SymbolicOpenCLFunction["map", {{CPointerType[{"__global", "mint"}], "in"}, {CPointerType[{"__global", "mint"}], "out"}, {"int", "length"}},
CBlock[{
SymbolicOpenCLDeclareIndexBlock[1],
CIf[COperator[Less, {"index", "length"}],
CAssign[CArray["out", "index"], toSymbolicC[fun]]
]
}]
]]]OpenCLMapSourceはWolfram言語の純関数を使うことができる:
OpenCLMapSource[# + Sin[#]&]コードを使って予め定義されたWolfram言語関数で作業することもできる:
myFun[x_] := Cos[x] ^ 2 + Sin[x] * 2 + x ^ 8 + 3OpenCLMapSource[myFun]上のコードをOpenCLFunctionLoadを使ってロードする:
addTwo = OpenCLFunctionLoad[OpenCLMapSource[# + 2&], "map", {{_Integer, "Input"}, {_Integer, "Output"}, _Integer}, 256]addTwo[ConstantArray[1, 100], ConstantArray[1, 100], 100]これを一般的なものにするためには,OpenCLMap関数を実装することができる:
SetAttributes[OpenCLMap, HoldFirst];
OpenCLMap[fun_, input_List] :=
Module[{len = Length[input], res, output, oclFun},
output = OpenCLMemoryAllocate[Integer, Length[input]];
oclFun = OpenCLFunctionLoad[OpenCLMapSource[fun], "map", {{_Integer, "Input"}, {_Integer, "Output"}, _Integer}, 256];
oclFun[input, output, len];
res = OpenCLMemoryGet[output];
OpenCLMemoryUnload[output];
res
]OpenCLMap[# + 2&, ConstantArray[1, 100]]BitNotの演算子を使う:
OpenCLMap[BitNot, ConstantArray[1, 100]]テクニカルノート
関連するガイド
関連リンク
テキスト
Wolfram Research (2010), OpenCLFunctionLoad, Wolfram言語関数, https://reference.wolfram.com/language/OpenCLLink/ref/OpenCLFunctionLoad.html.
CMS
Wolfram Language. 2010. "OpenCLFunctionLoad." Wolfram Language & System Documentation Center. Wolfram Research. https://reference.wolfram.com/language/OpenCLLink/ref/OpenCLFunctionLoad.html.
APA
Wolfram Language. (2010). OpenCLFunctionLoad. Wolfram Language & System Documentation Center. Retrieved from https://reference.wolfram.com/language/OpenCLLink/ref/OpenCLFunctionLoad.html
BibTeX
@misc{reference.wolfram_2026_openclfunctionload, author="Wolfram Research", title="{OpenCLFunctionLoad}", year="2010", howpublished="\url{https://reference.wolfram.com/language/OpenCLLink/ref/OpenCLFunctionLoad.html}", note=[Accessed: 18-August-2026]}
BibLaTeX
@online{reference.wolfram_2026_openclfunctionload, organization={Wolfram Research}, title={OpenCLFunctionLoad}, year={2010}, url={https://reference.wolfram.com/language/OpenCLLink/ref/OpenCLFunctionLoad.html}, note=[Accessed: 18-August-2026]}