Loading article…
OpenCL（Open Computing Language）は、中央処理装置（CPU）、グラフィックス処理装置（GPU）、デジタル信号処理装置（ DSP）、フィールドプログラマブルゲートアレイ（FPGA）、その他のプロセッサやハードウェアアクセラレータなど、多様なプラットフォーム上で実行されるプログラムを作成するためのフレームワークです。OpenCLは、これらのデバイスをプログラミングするためのプログラミング言語（ C99ベース）と、プラットフォームを制御し、計算デバイス上でプログラムを実行するためのアプリケーションプログラミングインターフェイス（API）を規定しています。OpenCLは、タスクベースおよびデータベースの並列処理を用いた並列コンピューティングのための標準インターフェイスを提供します。
OpenCL は、非営利のオープン スタンダード組織であるKhronos Groupによって維持されているオープン スタンダードです。適合実装 (適合性テスト スイートに合格) は、AMD、Arm、Cadence、Google、Imagination、Intel、 Nvidia 、Qualcomm、Samsung、SPI 、 Verisiliconなど、さまざまな企業から提供されています。[ 6 ] [ 7 ]
OpenCL は、コンピューティング システムを、ホストプロセッサ (CPU)に接続された中央処理装置(CPU) やグラフィックス処理装置 (GPU) などの「アクセラレータ」である多数の演算装置から構成されるものと見なします。OpenCL は、プログラムを記述するためのC ライクな言語を定義します。OpenCL デバイスで実行される関数は「カーネル」と呼ばれます。[ 8 ] : 17単一の演算装置は通常、複数の演算装置で構成され、演算装置は複数の処理要素(PE) から構成されます。単一のカーネル実行は、すべての PE または多くの PE で並列に実行できます。演算装置を演算装置と PE に分割する方法はベンダー次第です。演算装置は「コア」と考えることができますが、コアの概念は OpenCL でサポートされているすべてのタイプのデバイス (または「CPU」のカテゴリ内) 全体で定義するのが難しく、[ 9 ] : 49–50演算装置の数は、ベンダーのマーケティング資料で主張されているコアの数 (実際にはSIMD レーンをカウントしている可能性がある) と一致しない場合があります。[ 10 ]
OpenCL は、C に似たプログラミング言語に加えて、ホスト上で実行されているプログラムが計算デバイス上でカーネルを起動し、ホスト メモリとは (少なくとも概念的には) 分離されているデバイス メモリを管理できるようにするアプリケーション プログラミング インターフェイス (API) を定義しています。OpenCL 言語のプログラムは実行時にコンパイルされることを想定しているため、OpenCL を使用するアプリケーションは、さまざまなホスト デバイス向けの実装間で移植可能です。[ 11 ] OpenCL 標準では、 CおよびC++用のホスト API が定義されています。Python 、[ 12 ] Java、Perl、[ 13 ] D [ 14 ]および.NET [ 9 ]などの他のプログラミング言語およびプラットフォーム用のサードパーティ API も存在します。15 OpenCL 標準の実装は、C および C++ 用のAPI を実装するライブラリと、対象となる計算デバイス用の OpenCL Cコンパイラで構成されます。
OpenCLプログラミングモデルを他の言語に開放したり、カーネルソースを検査から保護したりするために、Standard Portable Intermediate Representation (SPIR) [ 15 ]を、フロントエンドコンパイラとOpenCLバックエンド間でカーネルを転送するためのターゲット非依存の方法として使用できます。
最近では、Khronos Group がSYCL [ 16 ]を承認しました。これは、プログラミングの生産性を向上させるために、純粋なC++ 17に基づくシングルソースeDSLとして OpenCL の高レベルプログラミング モデルです。C++ カーネルに興味はあるものの SYCL シングルソース プログラミング スタイルには興味がない人は、「C++ for OpenCL」言語で書かれた計算カーネル ソースで C++ の機能を使用できます。[ 17 ]
OpenCLは、計算デバイス用に4レベルのメモリ階層を定義しています。 [ 11 ]
すべてのデバイスがこの階層の各レベルをハードウェアで実装する必要はありません。階層内のさまざまなレベル間の一貫性は緩和されており、明示的な同期構造、特にバリアによってのみ強制されます。
デバイスはホスト CPU とメモリを共有する場合としない場合があります。[ 11 ]ホスト API は、デバイスのメモリ バッファへのハンドルと、ホストとデバイス間でデータをやり取りするための関数を提供します。
計算カーネルを記述するために使用されるプログラミング言語は、カーネル言語と呼ばれます。OpenCLは、C / C++ベースの言語を採用し、デバイス上で実行されるカーネル計算を指定しますが、アクセラレータの異種ハードウェアリソースへの効率的なマッピングを容易にするために、いくつかの制限と追加が加えられています。従来、OpenCL標準ではOpenCL Cを使用してアクセラレータをプログラミングしていましたが、後にOpenCL Cのすべての機能を継承しつつ、カーネルソースでC++機能を使用できるOpenCLカーネル言語用のC++が開発されました。
OpenCL C [ 18 ]は、OpenCL のデバイス モデルに適合するように調整されたC99ベースの言語方言です。メモリ バッファはメモリ階層の特定のレベルに存在し、ポインタには、これを反映して領域修飾子__global、__local、__constant、および__privateが注釈として付けられます。デバイス プログラムにmain関数がある代わりに、OpenCL C 関数には__kernelというマークが付けられ、ホスト プログラムから呼び出されるプログラムへのエントリ ポイントであることを示します。関数ポインタ、ビット フィールド、可変長配列は省略され、再帰は禁止されています。[ 19 ] C標準ライブラリは、数学プログラミング向けのカスタム標準関数セットに置き換えられます。
OpenCL Cは、ベクトル型と演算、同期、ワークアイテムとワークグループを操作する関数を用いた並列処理を容易にするために拡張されています。 [ 19 ]特に、 C言語の対応する型と同様に動作するfloatやdoubleなどのスカラー型に加えて、OpenCLはfloat4 （単精度浮動小数点数の4次元ベクトル）などの固定長ベクトル型を提供します。このようなベクトル型は、さまざまな基本型に対して長さ2、3、4、8、16で利用可能です。[ 18 ]： §6.1.2これらの型に対するベクトル化された演算は、CPU上でOpenCLプログラムを実行する際に、 SIMD命令セット（SSEやVMXなど）にマッピングされるように設計されています。[ 11 ]その他の特殊な型には、2次元および3次元の画像型があります。[ 18 ]： 10-11
以下は、OpenCL C で記述された行列とベクトルの乗算アルゴリズムです。
// A*x を乗算し、結果を y に格納します。// A は行優先行列であり、(i,j) 要素は A[i*ncols+j] にあります。__kernel void matvec ( __global const float * A , __global const float * x ,uint ncols 、__global float * y ){size_t i = get_global_id ( 0 ); // グローバルID。行インデックスとして使用されます。__global float const * a = & A [ i * ncols ]; // i 番目の行へのポインタfloat sum = 0.f ; // ドット積のアキュムレータfor ( size_t j = 0 ; j < ncols ; j ++ ) {sum += a [ j ] * x [ j ];}y [ i ] =合計;}カーネル関数matvecは、呼び出しごとに、行列Aの1行とベクトルxの内積を計算します。
これを完全な行列ベクトル乗算に拡張するために、OpenCLランタイムはカーネルを行列の行にマッピングします。ホスト側では、 clEnqueueNDRangeKernel関数がこれを実行します。この関数は、実行するカーネル、その引数、および行列Aの行数に対応するワークアイテムの数を引数として受け取ります。
この例では、高速フーリエ変換(FFT) の実装をロードして実行します。実装は以下に示します。[ 20 ]このコードは、OpenCL ライブラリに利用可能な最初のグラフィックカードを要求し、読み書き用のメモリ バッファ (グラフィック カードの観点から) を作成し、 FFT カーネルをJIT コンパイルしてから、最後にカーネルを非同期で実行します。この例では、変換の結果は読み込まれません。これは説明のためのサンプル コードであり、本格的な使用を意図したものではないため、エラー処理は完全に省略されています。
#include <stdio.h>#include <time.h>#include "CL/opencl.h"#define NUM_ENTRIES 1024int main () // (int argc, const char* argv[]){// 定数// カーネルのソースコードは文字列として表現されます// ファイル「fft1D_1024_kernel_src.cl」内にあります。詳細は次のリストを参照してください。const char * KernelSource =#include "fft1D_1024_kernel_src.cl";// 利用可能なGPUを検索中const cl_uint num = 1 ;clGetDeviceIDs ( NULL , CL_DEVICE_TYPE_GPU , 0 , NULL , ( cl_uint * ) & num );cl_device_idデバイス[ 1 ];clGetDeviceIDs ( NULL , CL_DEVICE_TYPE_GPU , num , devices , NULL );// GPUデバイスを使用して計算コンテキストを作成しますcl_context context = clCreateContextFromType ( NULL , CL_DEVICE_TYPE_GPU , NULL , NULL , NULL );// コマンドキューを作成するclGetDeviceIDs ( NULL , CL_DEVICE_TYPE_DEFAULT , 1 , devices , NULL );cl_command_queue queue = clCreateCommandQueue ( context , devices [ 0 ], 0 , NULL );// バッファメモリオブジェクトを割り当てるcl_mem memobjs [] = { clCreateBuffer ( context , CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR , sizeof ( float ) * 2 * NUM_ENTRIES , NULL , NULL ),clCreateBuffer ( context , CL_MEM_READ_WRITE , sizeof ( float ) * 2 * NUM_ENTRIES , NULL , NULL ) };// 計算プログラムを作成する// const char* fft1D_1024_kernel_src[1] = { };cl_program program = clCreateProgramWithSource ( context , 1 , ( const char ** ) & KernelSource , NULL , NULL );// 計算プログラムの実行可能ファイルを作成するclBuildProgram ( program , 0 , NULL , NULL , NULL , NULL );// 計算カーネルを作成するcl_kernel kernel = clCreateKernel ( program , "fft1D_1024" , NULL );// 引数の値を設定しますsize_t local_work_size [ 1 ] = { 256 };clSetKernelArg ( kernel , 0 , sizeof ( cl_mem ), ( void * ) & memobjs [ 0 ]);clSetKernelArg ( kernel , 1 , sizeof ( cl_mem ), ( void * ) & memobjs [ 1 ]);clSetKernelArg ( kernel , 2 , sizeof ( float ) * ( local_work_size [ 0 ] + 1 ) * 16 , NULL );clSetKernelArg ( kernel , 3 , sizeof ( float ) * ( local_work_size [ 0 ] + 1 ) * 16 , NULL );// 作業項目の寸法を持つND範囲オブジェクトを作成し、カーネルを実行しますsize_t global_work_size [ 1 ] = { 256 };global_work_size [ 0 ] = NUM_ENTRIES ;local_work_size [ 0 ] = 64 ; //Nvidia: 192 または 256clEnqueueNDRangeKernel ( queue , kernel , 1 , NULL , global_work_size , local_work_size , 0 , NULL , NULL );}ファイル「fft1D_1024_kernel_src.cl」内の実際の計算（「G80アーキテクチャへのFFTの適合」に基づく）：[ 21 ]
R "( // このカーネルは長さ 1024 の FFT を計算します。長さ 1024 の FFT は、// 基数 16 の関数、別の基数 16 の関数、そして基数 4 の関数への呼び出しに分解されます。__kernel void fft1D_1024 ( __global float2 * in , __global float2 * out , __local float * sMemx , __local float * sMemy ) { int tid = get_local_id ( 0 ); int blockIdx = get_group_id ( 0 ) * 1024 + tid ; float2 data [ 16 ];// グローバルメモリへの/からのデータの開始インデックスin = in + blockIdx ; out = out + blockIdx ;globalLoads ( data , in , 64 ); // グローバル読み取りを統合fftRadix16Pass ( data ); // インプレース 16 基数パスtwiddleFactorMul ( data , tid , 1024 , 0 );// ローカル メモリを使用したローカル シャッフルlocalShuffle ( data , sMemx , sMemy , tid , ((( tid & 15 ) * 65 ) + ( tid >> 4 ))); fftRadix16Pass (データ); // インプレース radix-16 はtwiddleFactorMul ( data , tid , 64 , 4 ); // 回転因子の乗算localShuffle ( data , sMemx , sMemy , tid , ((( tid >> 4 ) * 64 ) + ( tid & 15 )));// 基数4の関数呼び出しが4回fftRadix4Pass ( data ); // 基数4の関数呼び出し1回目fftRadix4Pass ( data + 4 ); // 基数4の関数呼び出し2回目fftRadix4Pass ( data + 8 ); // 基数4の関数呼び出し3回目fftRadix4Pass ( data + 12 ); // 基数4の関数呼び出し4回目// 統合されたグローバルはglobalStores ( data , out , 64 );を書き込む} ) "OpenCL FFT の完全なオープンソース実装は Apple の Web サイトで見つけることができます。[ 22 ]
OpenCL C++ は、OpenCL C と C++14 を組み合わせた言語の短命な仕様です。これは、パラメータを渡すことでオンラインモードでのみビルドされることを想定していました。この言語のサポートを検出するための拡張機能は説明されていません。実際にこの言語をサポートしたドライバがあったかどうかは不明です。[ 23 ]-cl-std=c++clBuildProgram()
2020年、Khronosは[ 25 ] 、従来のOpenCL Cの機能とC++17の機能を組み合わせた、コミュニティ主導のOpenCL用C++プログラミング言語[ 26 ]への移行を発表しました。この言語では、OpenCL Cとの後方互換性を維持しながら、標準C++の豊富な言語機能を活用できます。これにより、OpenCLカーネルコード開発者は、使い慣れたプログラミングフローやツールを継続して使用でき、OpenCL Cで利用可能な既存の拡張機能やライブラリも活用できるため、C++機能へのスムーズな移行が可能になります。
言語のセマンティクスは、Khronos Group がホストする OpenCL-Docs [ 27 ]リポジトリのリリースで公開されているドキュメントに記載されていますが、現時点では Khronos Group によって承認されていません。C++ for OpenCL 言語は独立したドキュメントには記載されておらず、C++ と OpenCL C の仕様に基づいています。オープンソースのClangコンパイラは、リリース 9 以降、C++ for OpenCL をサポートしています。[ 28 ]
C++ for OpenCL は元々 Clang コンパイラ拡張機能として開発され、リリース 9 に登場しました。[ 29 ] OpenCL C と密接に結合しており、Clang 固有の機能が含まれていなかったため、そのドキュメントは他の仕様やリファレンス カードのソースとともに、Khronos Group からOpenCL-Docs リポジトリ[ 27 ]に再ホストされました。C++ for OpenCL バージョン 1.0 を説明するこのドキュメントの最初の公式リリースは 2020 年 12 月に公開されました。 [ 30 ] C++ for OpenCL 1.0 には C++17 の機能が含まれており、OpenCL C 2.0 との下位互換性があります。2021 年 12 月に、OpenCL 3.0 標準と完全に互換性のある新しい暫定版 C++ for OpenCL バージョン 2021 がリリースされました。[ 31 ]最新の C++ for OpenCL ドキュメントの作業中のドラフトは、Khronos の Web サイトにあります。[ 32 ]
C++ for OpenCL は、ネストされた並列処理とブロックを除いて、OpenCL C のほとんどの機能 (構文的にも意味的にも) をサポートしています。[ 33 ]ただし、サポートされている機能には、主に C++ と C の意味論の違いに関連するわずかな違いがあります。たとえば、C++ は暗黙的な型変換に対してより厳格であり、restrict型修飾子をサポートしていません。[ 33 ]次の C++ 機能は、C++ for OpenCL ではサポートされていません。仮想関数、dynamic_cast演算子、非配置new / delete演算子、例外、メンバ関数へのポインタ、関数への参照、C++ 標準ライブラリ。[ 33 ] C++ for OpenCL は、OpenCL C のメモリ領域 (アドレス空間) の分離の概念を C++ の機能 (関数キャスト、テンプレート、クラスメンバ、参照、ラムダ関数、演算子) に拡張しています。C++ 機能のほとんどは、カーネル関数では使用できません。たとえば、オーバーロードやテンプレート、パラメータ型の任意のクラスレイアウトなどです。[ 33 ]
以下のコードスニペットは、C++の機能を利用して、OpenCL言語で複素数演算を含むカーネルを実装する方法を示しています。
// T に異なる型 (double、float、half) が使用される場合に、さまざまな精度で複素数計算を実行できるクラス Complex を定義します。template < typename T > class complex_t { T m_re ; // 実数成分。T m_im ; // 虚数成分。public : complex_t ( T re , T im ) : m_re { re }, m_im { im } {}; // 複素数乗算の演算子を定義します。complex_t operator * ( const complex_t & other ) const { return { m_re * other . m_re - m_im * other . m_im , m_re * other . m_im + m_im * other . m_re }; } T get_re () const { return m_re ; } T get_im () const { return m_im ; } };// 入力バッファから読み取った複素数の乗算を計算し、計算結果を出力バッファに格納するヘルパー関数。template < typename T > void compute_helper ( __global T * in , __global T * out ) { auto idx = get_global_id ( 0 ); // 各ワークアイテムは、入力バッファから 4 つの連続したアイテムを使用します// - 各複素数につき 2 つ。auto offset = idx * 4 ; auto num1 = complex_t { in [ offset ], in [ offset + 1 ]}; auto num2 = complex_t { in [ offset + 2 ], in [ offset + 3 ]}; // 複素数の乗算を実行します。auto res = num1 * num2 ; // 各ワークアイテムは、2 つの連続したアイテムを出力バッファに書き込みます。out [ idx * 2 ] = res . get_re (); out [ idx * 2 + 1 ] = res . get_im (); }// このカーネルは、単精度での複素数乗算に使用されます。__kernel void compute_sp ( __global float * in , __global float * out ) { compute_helper ( in , out ); }#ifdef cl_khr_fp16 // このカーネルは、デバイスがサポートしている場合、半精度での複素数乗算に使用されます。 #pragma OPENCL EXTENSION cl_khr_fp16: enable __kernel void compute_hp ( __global half * in , __global half * out ) { compute_helper ( in , out ); } #endifOpenCL 用の C++ 言語は、OpenCL C 言語と同様のアプリケーションやライブラリに、同様の方法で使用できます。C++ 言語の豊富な機能のおかげで、OpenCL 用の C++ で記述されたアプリケーションは、OpenCL C で記述されたアプリケーションよりも複雑な機能を容易に表現できます。特に、 C++ の汎用プログラミングパラダイムは、ライブラリ開発者にとって非常に魅力的です。
C++ for OpenCL ソースは、 cl_ext_cxx_for_opencl拡張機能をサポートする OpenCL ドライバによってコンパイルできます。これにより、-cl-std=CLC++の使用が可能になりますclBuildProgram()。[ 34 ] Arm は、 2020 年 12 月にこの拡張機能のサポートを発表しました。[ 35 ]しかし、OpenCL デバイスで高速化されるアルゴリズムの複雑さが増すにつれて、より多くのアプリケーションが、Clang [ 36 ]などのスタンドアロン コンパイラを使用して、C++ for OpenCL カーネルをオフラインでコンパイルし、実行可能なバイナリ形式または SPIR-V [ 37 ]などのポータブル バイナリ形式に変換することが予想されます。このような実行可能ファイルは、専用の OpenCL API を使用して OpenCL アプリケーションの実行中にロードできます。[ 38 ]
OpenCL 1.0向けにC++でコンパイルされたバイナリは、OpenCL 2.0準拠デバイスで実行できます。また、そのようなカーネルソースで使用されている言語機能によっては、以前のOpenCLバージョンやOpenCL 3.0をサポートするデバイスでも実行可能です。
OpenCLドライバとは別に、OpenCL用にC++で書かれたカーネルは、OpenCL Cカーネルと同様に、 clspv [ 39 ]コンパイラとclvk [ 40 ]ランタイムレイヤーを使用してVulkanデバイス上で実行できるようにコンパイルできます。
OpenCL 用の C++ は、ドキュメントに記載されている貢献者のコミュニティによって開発されたオープン言語です。[ 32 ]言語のセマンティック定義やオープンソースツールのサポートへの新しい貢献は、主要な設計思想に合致し、経験豊富な貢献者によってレビューおよび承認されれば、関心のある人なら誰でも受け入れます。[ 17 ]
OpenCL は、当初は商標権を保有するApple Inc.によって開発され、 AMD、IBM、Qualcomm、Intel、Nvidiaの技術チームとの協力により最初の提案へと改良されました。Appleはこの最初の提案をKhronos Groupに提出しました。2008 年 6 月 16 日、CPU、GPU、組み込みプロセッサ、ソフトウェア企業の代表者からなる Khronos Compute Working Group が結成されました[ 41 ]。このグループは 5 か月かけて OpenCL 1.0 の仕様の技術的な詳細を 2008 年 11 月 18 日までに完成させました[ 42 ]。この技術仕様は Khronos メンバーによってレビューされ、2008 年 12 月 8 日に公開が承認されました[ 43 ] 。
OpenCL 1.0 は、2009 年 8 月 28 日にMac OS X Snow Leopardと共にリリースされました。Apple のプレスリリースによると: [ 44 ]
Snow Leopardは、Open Computing Language（OpenCL）によって最新のハードウェアのサポートをさらに拡張しました。これにより、これまでグラフィックスアプリケーションのみが利用できたGPUの膨大なギガフロップス演算能力を、あらゆるアプリケーションが活用できるようになります。OpenCLはC言語をベースとしており、オープンスタンダードとして提案されています。
AMD は、Stream フレームワークで、現在非推奨となっているClose to Metal の代わりに OpenCL をサポートすることに決定しました。[ 45 ] [ 46 ] RapidMind は、複数のベンダーの GPU を単一のインターフェースでサポートするために、開発プラットフォームで OpenCL を採用したことを発表しました。[ 47 ] 2008 年 12 月 9 日、Nvidia は、GPU Computing Toolkit に OpenCL 1.0 仕様の完全なサポートを追加する意向を発表しました。[ 48 ] 2009 年 10 月 30 日、IBM は、 XL コンパイラの一部として最初の OpenCL 実装をリリースしました。[ 49 ]
グラフィックカード上のOpenCLでは、通常のCPUと比較して計算を最大1000倍高速化することが可能。 次期バージョンのOpenCLの重要な機能の一部は、倍精度演算や半精度演算など、1.0ではオプションとなっている。[ 50 ]
OpenCL 1.1は2010年6月14日にKhronosグループによって承認され[ 51 ]、並列プログラミングの柔軟性、機能性、パフォーマンスを向上させるための重要な機能が追加されました。
2011年11月15日、Khronos GroupはOpenCL 1.2仕様を発表しました[ 52 ]。この仕様では、パフォーマンスと並列プログラミング機能に関して、以前のバージョンに比べて大幅な機能追加が行われました。最も注目すべき機能は以下のとおりです。
2013年11月18日、Khronos Groupは、最終版のOpenCL 2.0仕様の承認と一般公開を発表しました。[ 54 ] OpenCL 2.0の更新と追加には、次のものが含まれます。
OpenCL 2.1 暫定仕様の承認とリリースは、2015 年 3 月 3 日にサンフランシスコで開催されたゲーム開発者会議で発表されました。リリースは 2015 年 11 月 16 日に行われました。[ 55 ]この仕様では、 C++14のサブセットに基づく OpenCL C++ カーネル言語が導入され、既存の OpenCL C カーネル言語のサポートも維持されています。Vulkanと OpenCL 2.1 は、中間表現として SPIR-V を共有しており、これにより高レベル言語のフロントエンドが共通のコンパイルターゲットを共有できます。OpenCL API の更新内容は次のとおりです。
AMD、ARM、Intel、HPC、YetiWareはOpenCL 2.1のサポートを表明している。[ 56 ] [ 57 ]
OpenCL 2.2 では、並列プログラミングの生産性を大幅に向上させるために、OpenCL C++ カーネル言語がコア仕様に組み込まれています。[ 58 ] [ 59 ] [ 60 ] 2017 年 5 月 16 日にリリースされました。 [ 61 ] 2018 年 5 月にバグ修正を含むメンテナンス アップデートがリリースされました。[ 62 ]
OpenCL 3.0 仕様は、2020 年 4 月からプレビュー状態であった後、2020 年 9 月 30 日にリリースされました。OpenCL 1.2 の機能は必須のベースラインとなり、OpenCL 2.x および OpenCL 3.0 のすべての機能はオプションとなりました。仕様では OpenCL C 言語が維持され、OpenCL C++ カーネル言語は非推奨となり、C++17とSPIR-V中間コードのサブセットを実装するClang / LLVMコンパイラに基づくC++ for OpenCL 言語[ 17 ]に置き換えられました。[ 63 ] [ 64 ] [ 65 ] Khronos openCL 拡張機能がいくつか含まれた C++ for OpenCL のバージョン 3.0.7 が IWOCL 21 で発表されました。[ 66 ]実際のバージョンは、いくつかの新しい拡張機能と修正を含む 3.0.11 です。NVIDIA は、Khronos OpenCL ワーキング グループと緊密に連携し、セマフォとメモリ共有を使用して Vulkan の相互運用性を改善しました。[ 67 ]最後のマイナーアップデートは、バグ修正と複数デバイス用の新しい拡張機能を含む 3.0.14 でした。[ 68 ]
OpenCL 2.2 のリリース時に、Khronos Group は、OpenCL が可能な限りVulkanと統合し、両方の API で OpenCL ソフトウェアの展開の柔軟性を実現すると発表しました。[ 69 ] [ 70 ]これは、Adobe の Premiere Rush が clspv [ 39 ]オープンソース コンパイラを使用して大量の OpenCL C カーネル コードをコンパイルし、Android への展開のために Vulkan ランタイムで実行することで実証されています。[ 71 ] OpenCL は、Vulkan とは独立した将来を見据えたロードマップを持っており、「OpenCL Next」が開発中で、2020 年のリリースを目指しています。OpenCL Next には、Vulkan / OpenCL 相互運用、スクラッチ パッド メモリ管理、拡張サブグループ、SPIR-V 1.4 取り込み、SPIR-V 拡張デバッグ情報などの拡張機能が統合される可能性があります。OpenCL はまた、複数のアクセラレータ タイプでの展開の柔軟性のために、Vulkan のようなローダーとレイヤー、および「柔軟なプロファイル」を検討しています。[ 72 ]
clinfoOpenCL情報を表示するためのコマンドラインツールOpenCL は、実行時にロードされる一連のヘッダーと共有オブジェクトで構成されています。ランタイムがサポートする必要のあるベンダーのクラスごとに、インストール可能なクライアントドライバ (ICD) をプラットフォームにインストールする必要があります。つまり、たとえば Linux プラットフォームで Nvidia デバイスをサポートするには、OpenCL ランタイム (ICD ローダー) がベンダーの ICD を見つけて呼び出しを適切にリダイレクトできるように、Nvidia ICD をインストールする必要があります。標準の OpenCL ヘッダーはコンシューマー アプリケーションで使用され、各関数への呼び出しは、OpenCL ランタイムによって ICD を使用して適切なドライバにプロキシされます。各ベンダーは、ドライバで各 OpenCL 呼び出しを実装する必要があります。[ 73 ]
Apple [ 74 ]、Nvidia [ 75 ] 、 ROCm、RapidMind [ 76 ]、Gallium3D [ 77 ]のOpenCL実装はすべてLLVMコンパイラ技術に基づいており、フロントエンドとしてClangコンパイラを使用しています。
2016年現在、OpenCLはグラフィックス処理ユニット（GPU）、SIMD命令を備えたCPU 、 FPGA、Movidius Myriad 2、Adapteva Epiphany、およびDSP上で動作します。
公式に準拠するためには、実装は Khronos 適合性テスト スイート (CTS) に合格し、その結果を Khronos 採用プログラムに提出する必要があります。[ 174 ]すべての OpenCL バージョンの Khronos CTS コードは、2017 年以降オープンソースで利用可能です。[ 175 ]
Khronos Groupは、OpenCLに準拠した製品の拡張リストを維持しています。[ 4 ]
標準に準拠した実装はすべて、clinfo ツールのいずれかを使用して照会できます (同じ名前で同様の機能セットを持つツールが複数あります)。[ 186 ] [ 187 ] [ 188 ]
製品とそのOpenCLサポートのバージョンは以下のとおりです。[ 189 ]
OpenCL 1.2+ のすべてのハードウェアが使用可能、OpenCL 2.x はオプションのみ、Khronos テスト スイートは 2020 年 10 月より利用可能[ 190 ] [ 191 ]
まだありません：Khronosテストスイートは準備完了、ドライバーアップデートにより2.0および2.1をサポートするすべてのハードウェアが利用可能
OpenCLの重要な特徴は、抽象化されたメモリおよび実行モデルによる移植性です。プログラマは、他のプラットフォームでの直接的な移植性を放棄しない限り、Nvidia GPU向けのインライン並列スレッド実行（PTX）などのハードウェア固有の技術を直接使用することはできません。OpenCLカーネルは、準拠する実装であればどれでも実行可能です。
しかし、カーネルのパフォーマンスは必ずしもプラットフォーム間で移植可能ではありません。既存の実装は、カーネルコードが適切に調整されていれば競争力があることが示されており、パフォーマンスの移植性の問題に対する解決策として自動調整が提案されています[ 194 ]。実験的な線形代数カーネルでは「許容できるレベルのパフォーマンス」が得られています[ 195 ] 。動作が異なる複数のカーネルを含むアプリケーション全体の移植性も研究されており、移植性には限られたトレードオフしか必要ないことが示されています[ 196 ] 。
2011年にデルフト工科大学で行われた研究では、 CUDAプログラムと、それらを直接OpenCL Cに変換したものを比較したところ、Nvidiaの実装ではCUDAがOpenCLを最大30%上回ることがわかりました。研究者らは、OpenCLプログラムに手動で最適化を適用することで比較をより公平にすることができ、その場合「OpenCLがCUDAよりもパフォーマンスが悪くなる理由はない」と指摘しました。パフォーマンスの違いは、主にプログラミングモデル（特にメモリモデル）の違いと、OpenCLに対するNVIDIAのCUDAコンパイラの最適化の違いに起因すると考えられます。[ 194 ]
D-Wave Systems Inc. の別の研究では、「OpenCL カーネルのパフォーマンスは CUDA のパフォーマンスよりも約 13% ～ 63% 遅く、エンドツーエンドの時間は約 16% ～ 67% 遅い」ことが判明しました。[ 197 ]
OpenCL では、同じプログラムを実行することで CPU と GPU 間でワークロードを共有できるため、プログラマーはデバイス間で作業を分割することで両方を活用できます。[ 198 ]これにより、デバイス間で操作の相対的な速度が異なるため、作業をどのように分割するかを決定するという問題が生じます。この問題を解決するために機械学習が提案されています。Grewe と O'Boyle は、プログラムのコンパイル時の特徴に基づいてトレーニングされたサポートベクターマシンのシステムについて説明しています。このシステムは、実際にプログラムを実行してパフォーマンスを測定することなく、静的にデバイス分割問題を決定できます。[ 199 ]
AMD RDNA 2とNvidia RTXシリーズの実際のグラフィックカードを比較したところ、OpenCLテストでは結果が確定しませんでした。Nvidia CUDAまたはOptiXの使用によるパフォーマンス向上の可能性はテストされていません。[ 200 ]
{{cite web}}: CS1メンテナンス: アーカイブサービスは非推奨になりました (リンク){{cite web}}: CS1メンテナンス: アーカイブサービスは非推奨になりました (リンク)