OpenCLの概要

この記事では、OpenClプログラミングの基本について説明します。 OpenClは、 c99標準に構造が似ているGPU / CPUプログラミング言語です。 Khronos Groupは開発に取り組んでおり、完全なドキュメントはWebサイトで入手できます。 「まあ、それはささいなことで、インターネットを掘り起こすだけ」というトピックに関する論争を避けるために、私はすぐに予約をします。 ここでは、初心者プログラマの生活を可能な限り簡素化し、最初のプロジェクトのビデオカードの計算能力を使用できるようにする基本原則をまとめます。 2〜3個の深刻なOpenClプログラムを書いた人は、もう興味を示さなくなります。 この記事はある意味で私の最後の記事の続きです。

コンパイラ

まず第一に、問題は、コード自体をどこに書くかです。 私が知る限り、.NETには、スタジオでカーネルコードを直接処理するためのwhiはありません。 したがって、サードパーティのエディターを使用する必要があります。 AMD、nVidia、およびIntelは、それらのSDKをバンドルしています。 何らかの理由で、Intelovskyの方が好きです。 オプションとして、ファンによって書かれたいくつかのエディターがあります。 これらのうち、私はOpenCLTemplateに最も付属しているエディターが好きです。 これらが編集者であることは注目に値します。 コードのコンパイルは、GPU / CPUで実行する直前に発生します。

デバイスメモリモデル

言語自体を説明する前に、対話するデバイスの物理モデルについて簡単に説明します。 言語コマンドは、「ワークアイテム」と呼ばれるオブジェクトで実行されます。 各「作業項目」は互いに独立しており、残りと並行してコードを実行できます。 プロセスが、使用中のワークアイテムまたは他のワークアイテムによってすでに処理されたワークアイテムからデータを受信する場合、共有メモリを介してこれを実行できます。 合計メモリは非常に遅くなりますが、大量にあります。 計算を高速化するために、ローカルメモリがあります。 CUDAに精通している場合、「共有メモリ」と呼ばれます。 一般的な方法よりもはるかに高速ですが、すべてのプロセスがアクセスできるわけではありません。 1つのグループの作業項目のみがローカルメモリにアクセスできます。 これらのグループは、「コンピューティングユニット」または「ワークグループ」と呼ばれます（最初の名前は鉄レベルの物理パーティションを指し、2番目はプログラムレベルの論理パーティションを指します）。 デバイスに応じて、これらの各グループには異なる数の作業項目があります（たとえば、NVIDIA GT200では240、Radeon 5700シリーズでは256）。 これらのユニットの数はかなり少ない数に制限されています（NVIDIA GT200では30、Radeon 5700シリーズでは9-10）。 ワークアイテムが単独でアクセスできる超高速の「プライベートメモリ」もあります。

OpenCLデバイスドライバーは、ワークアイテムとワークグループの起動と操作を自動化します。 たとえば、100万個のプロセスを実行する必要があり、自由に使えるワークアイテムが1,000個しかない場合、ドライバーは完了後に次のタスクで各プロセスを自動的に開始します。 物理レベルの理解は、プロセス間の相互作用とメモリへのプロセスのアクセスの可能性を理解するためにのみ必要です。

基本機能

この言語はほぼ標準のC ++に基づいているため、OpenCLとそれを区別する機能のみを検討します。 前回の記事で引用した最も単純なカーネルプログラムのコードを検討してください。 このコードは、v1とv2の2つのベクトルを追加し、結果を最初のベクトルに入れます。

__kernel void

floatVectorSum(__global float * v1,

__global float * v2)

{

int i = get_global_id(0);

v1[i] = v1[i] + v2[i];

}



* This source code was highlighted with Source Code Highlighter .

手続きのお知らせ

まず、神秘的な「__kernel」が目を引きます。 このディレクティブは、外部から呼び出すプロシージャをマークする必要があります。 外部から作業するときに手順が不要な場合は、省略できます。

メモリの種類

データ型「__global」は、実行中のデバイスのグローバルアドレス空間から割り当てられたメモリを指します。 かなり遅いですが、ゆとりがあります。 最新のビデオカードの場合、ギガバイト単位で測定されます。 プロセッサで作業する場合、グローバルはRAMを指します。

グローバルに加えて、「__ local」があります。 ワークグループのみがアクセスできます。 そのようなグループごとに、約8キロバイトが割り当てられます。

また、高速メモリは「__privat」です。 これは、別のスレッド（ワークアイテム）のみがアクセスできるメモリです。 合計で、このメモリの32個のレジスタがストリームに割り当てられます。

カーネルの作成時に宣言できる残りのタイプのメモリは、__ globalタイプに基づいています。 1つは「__constant」で、読み取り専用で使用できます。 次に、これらは「__read_only」、「__ write_only」、および「__read_write」です。これらの構造の使用は画像に対してのみ許可されています。

プロセス識別子

ビデオカードで起動した後、すべてのプロセスは同等であり、同等のコードを実行します。 しかし、明らかに、同じアクションを複数回繰り返す必要はありません。各プロセスはタスクの独自の部分を実行する必要があります。 それらを取り巻く世界で彼らの場所を特定することは、プロセス識別子です。 最も単純な識別子は「get_global_id（0）」です。 上記の例の場合、このプロセスが追加するベクトルのi番号を指します。 1次元のベクトルではなく2次元の画像を処理する場合、2つの軸に沿ったプロセスの位置を知る必要があります。 もちろん、この値は計算できます。 しかし、これらは不必要な操作です。 そのため、便宜上、起動時に2次元空間が必要であることを指定できます。 その後、プロセスで両方の位置識別子を取得できます：「get_global_id（0）」、「get_global_id（1）」。 3次元空間でも同じことができます。 多くの場合、作業するスペースの寸法も必要になる場合があります。 たとえば、ほとんどすべての処理中の画像の場合、幅と高さが必要です。 スペースの次元を取得するには、識別子「get_global_size（i）」が使用されます。 さらに、ワーキンググループ内のプロセス識別子-「get_local_id（i）」、「get_local_size（i）」、およびグループ自体の識別子-「get_group_id（i）」、「get_num_groups（i）」があります。 これらの関係のほとんどは互いに関連しています：num_groups * local_size = global_size、local_id + group_id * local_size = global_id、global_size％local_size = 0

課金の最適化

OpenCLおよびビデオカードの開発者は、彼らの発案の主な目標は複雑な計算を加速することであることに気付きました。 これを行うために、いくつかの特殊な機能が言語に追加されました。これらの機能を使用すると、数学的な問題を迅速に処理できます。

インラインベクトル

最初の機能は、ベクトルとベクトル演算です。 OpenClでは、2、4、8、および16次元のベクトルを変数として宣言できます。 これはそれに応じて行われます：int2、int4、int8、int16。 また、double、byte、および他のすべてのタイプを宣言できます。 対応する次元のベクトルは、加算/減算/除算/乗算することができます。同様に、任意のベクトルを数値で除算/乗算することができます。

uint4 sumall = (uint4)(1,1,1,1);

small += (uint4)(1,1,1,1);

sumall = sumall/2;



* This source code was highlighted with Source Code Highlighter .

さらに、ベクター用に最適化された多数の関数があり、それらを直接操作できます。 このような関数には、距離計算関数、ベクトル積関数が含まれます。 例：

float4 dir1 = (float4)(1, 1, 1, 0);

float4 dir2 = (float4)(1, 2, 3, 0);

float4 normal = cross(dir1, dir2);



* This source code was highlighted with Source Code Highlighter .

また、ベクトルをマージして、一方から他方に部品を取り、それらをより大きなものに接着することもできます。

int4 vi0 = (int4) -7 ;

int4 vi1 = (int4) ( 0, 1, 2, 3 ) ;

vi0.lo = vi1.hi; //

int8 v8 = (int8)(vi0.s0123, vi1.s0123); //



* This source code was highlighted with Source Code Highlighter .

シンプルな機能

OpenClのもう1つの機能は、組み込み関数ライブラリです。 OpenClのmath.libの標準セットに加えて、いわゆるネイティブ関数があります。 これらは、ビデオカードの特定の機能の使用と失礼な数学に直接基づく機能です。 それらを超精密な計算で使用することはお勧めできませんが、画像フィルタリングの場合、違いに気付くことはできません。 このような関数には、たとえば、「native_sin」、「native_cos」、「native_powr」が含まれます。 これらの機能の詳細な説明は行いませんが、多くの機能があり、原則は異なります。 必要な場合は、ドキュメントを参照してください。

一般的な機能

「単純な関数」に加えて、開発者は多くの呼ばれる共通関数を作成しました。 これらは、画像処理で一般的な機能です。 例：mad（a、b、c）= a * b + c、mix（a、b、c）= a +（ba）* c。 これらの関数は、対応する数学アクションよりも高速です。

例

サイトwww.cmsoft.com.brには、ネイティブ関数と共通関数を使用してコードを最適化する可能性を示す素晴らしい例があります。

kernel void regularFuncs()

{

for ( int i=0; i<5000; i++)

{

float a=1, b=2, c=3, d=4;

float e = a*b+c;

e = a*b+c*d;

e = sin(a);

e = cos(b);

e = a*b+c*d;

e = sin(a);

e = cos(b);

e = a*b+c*d;

e = sin(a);

e = cos(b);

float4 vec1 = (float4)(1, 2, 3, 0);

float4 vec2 = (float4)(-1, 3, 1, 0);

float4 vec = distance(vec1, vec2);

double x=1, y=2, z=3;

double resp = x*y+z;

}

}

kernel void nativeFuncs()

{

for ( int i=0; i<5000; i++)

{

float a=1, b=2, c=3, d=4;

float e = mad(a,b,c);

e = mad(a,b,c*d);

e = native_sin(a);

e = native_cos(b);

e = mad(a,b,c*d);

e = native_sin(a);

e = native_cos(b);

e = mad(a,b,c*d);

e = native_sin(a);

e = native_cos(b);

float4 vec1 = (float4)(1, 2, 3, 0);

float4 vec2 = (float4)(-1, 3, 1, 0);

float4 vec = fast_distance(vec1, vec2);

double x=1, y=2, z=3;

double resp = mad(x,y,z);

}

}



* This source code was highlighted with Source Code Highlighter .

2番目の手順（最適化を使用）は35倍高速です。

許可

OpenClには、さまざまな追加機能を含めることができる多くのディレクティブがあることに注意してください。 これには2つの理由があります。 最初-歴史的に、これらの機能のすべてがサポートされていませんでした。 第二に、これらの機能はパフォーマンスに影響を与える可能性があります。 通常、機能は次のコマンドによって有効になります。

#pragma OPENCL EXTENSION extension name : behavior



* This source code was highlighted with Source Code Highlighter .

例として。 次のコマンドが含まれます：バイトタイプ、計算の倍精度、およびすべての数学関数を使用する機能

#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable

#pragma OPENCL EXTENSION cl_khr_fp64 : enable



* This source code was highlighted with Source Code Highlighter .

同期する

障壁

多くの場合、コンピューティングでは、同期が必要です。 これはいくつかの方法で実現されます。 最初は障壁です。 バリアは、他のすべてのプロセスまたはそのワーキンググループのプロセスに到達するまでプロセスが停止するようなチームです。 2つの例を示します。

kernel void localVarExample()

{

int i = get_global_id(0);

__local int x[10];

x[i] = i;

barrier(CLK_LOCAL_MEM_FENCE);

if (i>0) int y = x[i-1];

}

kernel void globalVarExample()

{

int i = get_global_id(0);

__global int x[10];

x[i] = i;

barrier(CLK_GLOBAL_MEM_FENCE);

if (i>0) int y = x[i-1];

}



* This source code was highlighted with Source Code Highlighter .

最初の例では、バリアコマンドでワークグループのすべてのプロセスが予想され、2番目では、OpenCLデバイスのすべてのプロセスが予想されます。

この例の特徴であるコマンド「__local int x [10];」に注目する価値があります。 および「__global int x [10];」。 これらを使用すると、プロセスのグループおよび実行中のすべてのプロセスでグローバル変数を選択できます。

単一の操作

スレッド間で同期するための2番目のオプションはアトミックです。 これらは、メモリへの同時アクセスを防ぐ機能です。 それらを使用する前に、次のディレクティブを含める必要があります。

#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable

#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable

#pragma OPENCL EXTENSION cl_khr_global_int32_extended_atomics : enable

#pragma OPENCL EXTENSION cl_khr_local_int32_extended_atomics : enable

#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable

#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable



* This source code was highlighted with Source Code Highlighter .

この関数の最も簡単な例：

__kernel void test(global int * num)

{

atom_inc(&num[0]);

}



* This source code was highlighted with Source Code Highlighter .

「atom_inc（＆num [0]）;」の代わりにnum ++が書き込まれた場合、すべてのプロセスが同時にメモリにアクセスして同じ値を読み取るため、プログラム実行の結果は予測不能でした。 合計で、11のユニット操作関数があります：「add、sub、xchg、inc、dec、cmp_xchg、min、max、および、or、xor」。

これらの関数を使用してセマフォを作成することは難しくありません。

void GetSemaphor(__global int * semaphor) {

int occupied = atom_xchg(semaphor, 1);

while (occupied > 0)

{

occupied = atom_xchg(semaphor, 1);

}

}



void ReleaseSemaphor(__global int * semaphor)

{

int prevVal = atom_xchg(semaphor, 0);

}



* This source code was highlighted with Source Code Highlighter .

画像を操作する

このガイドに最後に含めたいのは、OpenCLを介して画像を操作することです。 クリエーターは、ユーザーの脳を最小限に抑える必要がある画像を使用して作品を作成しようとしました。 とてもいいです。 タイプimage2d_tおよびimage3d_tの画像のダウンロードが可能です。 前者は普通の画像で、後者は三次元です。 また、ロードされるイメージは、「__ read_only」、「__ write_only」、「__ read_write」のいずれかの形式である必要があります。 イメージからのデータの読み取りと書き込みは、特別な手順によってのみ可能です：value = read_imageui（イメージ、サンプラー、位置）、write_imageui（イメージ、位置、値）。

私の意見では、「サンプラー」の概念を除いて、ここのすべてが明確です。 サンプラーは、画像での作業を最適化するものです。 「正規化された座標」、「アドレスモード」、「フィルターモード」の3つのパラメーターがあります。 最初の2つの意味があります：「CLK_NORMALIZED_COORDS_TRUE、CLK_NORMALIZED_COORDS_FALSE」。 名前に応じて、入力座標が正規化されているかどうかを示す必要があります。 2番目は、画像の境界の外側から座標を読み取ろうとしている場合の対処方法を示しています。 可能なオプション：イメージ（CLK_ADDRESS_MIRRORED_REPEAT）をミラーリングするには、最も近い境界値（CLK_ADDRESS_CLAMP_TO_EDGE）を取り、ベースカラー（CLK_ADDRESS_CLAMP）を取り、何もしない（ユーザーはこれがCLK_ADDRESS_NONEが起こらないことを保証します）。 3番目は、入力に整数座標がない場合の処理​​を示しています。 可能なオプション：最も近い値に近似（CLK_FILTER_NEAREST）、線形補間（CLK_FILTER_LINEAR）。

簡単な例。 エリア内の平均値で画像を吸収します。

__kernel void ImageDiff(__read_only image2d_t bmp1, __write_only image2d_t bmpOut)



{

const sampler_t smp = CLK_NORMALIZED_COORDS_FALSE |

CLK_ADDRESS_CLAMP_TO_EDGE |

CLK_FILTER_NEAREST;

int2 coords = (int2)(get_global_id(0), get_global_id(1));

uint4 sumall = (uint4)(0,0,0,0);

int sum = 0;

for ( int i=-10;i<11;i++)

for ( int j=-10;j<11;j++)

{

int2 newpol = (int2)(i,j)+coords;

sumall+= read_imageui(bmp1, smp, newpol);

sum++;

}

sumall = sumall/sum;

write_imageui(bmpOut, coords, sumall);

}



* This source code was highlighted with Source Code Highlighter .

有用性

まあ、私は簡単な説明を管理したと思います。 誰かがそれを必要とするならば、今、より詳細な研究のためのいくつかのリンク。

ドキュメントを含む公式サイト。

例と明確な説明があるサイト。

良いpdfnichek、そこにはOpenClデバイスの構造がよく描かれています。

ロシア語のOpenCLに関する2つのプレゼンテーションもあります。 それらにはかなりの情報があり、関連するテキストはありません。 確かに、良い例があります。 最初のもの 。 二番目 。