OpenCLLink プログラミング
This feature is not supported on the Wolfram Cloud.

ユーザがCラッパーコードを書く必要がないので,Wolfram言語でOpenCLをプログラミングするのは簡単である.(Cラッパーコードは,かなり冗長で,理解しにくく,デバックしにくいことがある.)また,OpenCLLink を使うことによって,新しいバージョンの標準がリリースされた場合に互換性があることが保証される.

このセクションでは,OpenCLプログラミングについて簡単に説明する.ユーザがOpenCLコードをロードしてWolfram言語内から使用できるようにするOpenCLFunctionLoadを使う.

OpenCLFunctionLoadOpenCL関数をソースからWolfram言語にロードする

Wolfram言語のOpenCLプログラミング.

詳しくは CUDALink プログラミング」を読まれたい.

OpenCLプログラミング

OpenCLプログラムは,入力リストの各要素について計算を行う小さなコードである.最初のプログラムは各要素に2を加える.

__kernel void addTwo_kernel(__global mint * arry, mint len) {
    
    int index = get_global_id(0);

    if (index >= len) return;

arry[index] += 2;
}

以下は,上のプログラムをセクションに分解したものである.

_kernel void addTwo_kernel (_global int arry, int len) {

_kernelコンストラクトは,関数がOpenCL GPU上で実行されるように宣言する.残りは,_global接頭辞を持つポインタ付きの関数引数である.

int index = get_global_id (0);

これで関数を実行するスレッドのインデックス値を得る.インデックス値は,0から起動されるスレッドの数までである.

if (index >= len) return;

これで,プログラムが入力配列の長さを超えるメモリを書かないことを確実にする.スレッドの数はブロックサイズの倍数で起動されるので,入力配列の大きさがブロックサイズの倍数ではない場合には,この条件文が必要である.

arry[index] += 2;

これで各要素に2を加える.

プログラムをWolfram言語にロードする

一旦プログラムが書かれると,OpenCLLink を使ってWolfram言語にロードすることができる.これはOpenCLFunctionLoadを使って行う.

OpenCLLink アプリケーションをロードする.

In[1]:=
Click for copyable input

First, assign the program to a string.

In[2]:=
Click for copyable input

関数をロードする.OpenCLFunctionLoadの引数は,ソースコード,ロードする関数の名前,関数の署名,およびブロック次元である.

In[3]:=
Click for copyable input
Out[3]=

関数を実行する.

In[4]:=
Click for copyable input
Out[4]=

結果は,出力リスト要素の集合である.

CUDAをOpenCLに移植する

OpenCLLink はOpenCLプログラミングで必要なCラッパーを処理するので,ユーザはOpenCLカーネルコードだけに注目することができ,これが唯一CUDAから移植する必要のあるコードである.

OpenCLプログラムを移植することについては,CUDAとOpenCLの間で1対1で対応する形で関数名が変更されている.以下にその対応表を挙げる.

このセクションでは,上の表を使って以下のCUDAコードをOpenCLに移植する.

__global__ void myKernel(mint * global0Id, mint * global1Id, mint width, mint height) {
int xIndex = threadIdx.x + blockDim.x * blockIdx.x;
int yIndex = threadIdx.y + blockDim.y * blockIdx.y;
int index = xIndex + yIndex*width;
    if (xIndex < width && yIndex < height) {
     global0Id[index] = threadIdx.x;
global1Id[index] = threadIdx.y;
}
}

以下はCUDAコードをOpenCLに変換したものである.

In[5]:=
Click for copyable input

以下が移植で変更された部分である.

  • _global_voidmyKernelが,_kernel voidmyKernelになった
  • int global0Id, int global1Idが,_global int global0Id, _global int global1Idになった
  • threadIdx.x + blockDim.x blockIdx.xが,get_global_id (0)になった
  • threadIdx.y + blockDim.y blockIdx.yが,get_global_id (1)になった
  • threadIdx.xが,get_local_id (0)になった
    • threadIdx.yが,get_local_id (1)になった

    コードのロードについて唯一必要な変更は,CUDAFunctionLoadOpenCLFunctionLoadで置換することだけである.

    In[6]:=
    Click for copyable input
    Out[6]=

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

    In[7]:=
    Click for copyable input

    関数を起動する.

    In[11]:=
    Click for copyable input

    結果を可視化する.

    In[12]:=
    Click for copyable input
    Out[12]=

用語

CUDAとOpenCLの間で,使われる用語が少し異なることに注意する必要がある.Wolfram言語では,両方のいいところを組み合せている.

以下は用語の対応表である.

CUDAOpenCL
streaming multiprocessor(ストリーミング・マルチプロセッサ)device(デバイス)
multiprocessor(マルチプロセッサ)compute unit(計算ユニット)
global memory(グローバルメモリ)global memory(グローバルメモリ)
shared memory(共有メモリ)local memory(ローカルメモリ)
local memory(ローカルメモリ)private memory(プライベートメモリ)
kernel(カーネル)program(プログラム)
block(ブロック)work group(ワークグループ)
thread(スレッド)work item(ワークアイテム)

メモリ

メモリマネージャの動作は,CUDALinkOpenCLLink で同じである.メモリマネージャについては,CUDALink のメモリ」に詳しく説明されている.

メモリは1つのリンクに結び付いているので,CUDALinkOpenCLLink のメモリを互いに交換することはできない.