CUDA* と SYCL* プログラミングの比較#

このセクションでは、CUDA* と SYCL* のプログラミング・モデルを比較し、概念と API を CUDA* から SYCL* にマップする方法を示します。

実行モデル#

カーネル関数#

CUDA* では、カーネル関数は __global__ 宣言指定子で定義され、GPU 上の複数のスレッド間で同時に実行されます。CUDA* カーネルから呼び出される関数は、__device__ 指定子で修飾する必要があります。__host__ 宣言指定子は、CPU で実行されるホストコードから呼び出し可能な関数を修飾する際に使用されます。

それに対し、SYCL* カーネルは CPU、GPU、FPGA などの SYCL* 対応デバイスで実行可能な関数です。これらのカーネルは、ホストコードから起動され SYCL* デバイスで同時に実行できます。CUDA* とは異なり、SYCL* カーネル関数は特別な宣言指定子を必要とせず、標準の C++ 構文で定義できます。

次の表は、カーネル関数を定義する CUDA* と SYCL* の同等性を示しています。

CUDA*

SYCL*

__global__ void foo_kernel() {}

void foo_kernel() {}

__device__ void foo_device() {}

void foo_device() {}

__host__ void foo_host() {}

void foo_host() {}

__host__ __device__ void foo_host_device() {}

void foo_host_device() {}

__CUDA_ARCH__ マクロ使用して CUDA* __host__ __device__ 関数のコードパスを区別する場合、SYCL* で __SYCL_DEVICE_ONLY__ マクロを使用して同等の機能を実現できます。

次の表は、デバイス関数を指定する CUDA* と SYCL* の同等性を示しています。

CUDA*

SYCL*

 
__host__ __device__ int foo(int i) { 
#ifdef __CUDA_ARCH__
 return i + 1; 
#else
 return 0; 
#endif 
} 
__global__ void kernel() { 
    foo(); 
} 
int main() { 
    return foo(); 
}
 
int foo(int i) { 
#ifdef __SYCL_DEVICE_ONLY__ 
    return i + 1; 
#else 
    return 0; 
#endif 
} 

void kernel() { 
    foo(); 
} 

int main() { 
    return foo(); 
}

実行階層#

カーネル実行インスタンスは、CUDA* と SYCL* プログラミング・モデルの両方で並列処理を効率良く活用するため階層的に構成されています。CUDA* では、これらのインスタンスはスレッドと呼ばれますが、SYCL* では work-item と呼ばれます。CUDA* スレッドはブロックに構成でき、ブロックはグリッドに構成できます。SYCL* の work-item は work-group に構成でき、work-group は ND-range に構成できます。

ハードウェアの観点から見ると、CUDA* カーネルが GPU で実行されると、ストリーミング・マルチプロセッサー (SM) が、ワープと呼ばれる 32 スレッドのグループ内のスレッドを生成、管理、スケジュール、および実行します。比較すると、SYCL* では、sub-group は、同時に実行される work-group 内の関連する work-item のコレクションを表します。

CUDA* コードを SYCL* コードに移行するには、次の表に示すように CUDA* 実行階層を SYCL* 階層にマップできます。

CUDA*

SYCL*

スレッド

work-item

ワープ

sub-group

ブロック

work-group

グリッド

ND-range

スレッドのインデックス作成#

前述したように、CUDA* スレッドと SYCL* work-item は階層的に構成されています。

CUDA* には、スレッドをサポートするビルトイン変数が用意されています。

  • スレッド ID: threadIdx.x/y/z

  • ブロック ID: blockIdx.x/y/z

  • ブロック次元: blockDim.x/y/z

  • グリッド次元: gridDim.x/y/z

SYCL* には同等のビルトイン変数があります。

  • スレッド ID: sycl::nd_item.get_local_id(0/1/2)

  • Work-group ID: sycl::nd_item.get_group(0/1/2)

  • Work-group 次元: sycl::nd_item.get_local_range().get(0/1/2)

  • ND-range 次元: sycl::nd_item.get_group_range(0/1/2)

CUDA* C++ プログラミング・ガイド (英語) および SYCL* 2020 仕様の実行空間のインデックス線形化セクションによると、形状 (dx, dy, dz) のブロックまたは work-group と id (x, y, z) の要素が与えられた場合、インデックスは x + y * Dx + z * Dx * Dy (CUDA*) および z + y * Dz + x * Dz * Dy (SYCL*) になります。この不一致の原因は、SYCL* が C++ の多次元配列インデックスと密接に連携しているためです。インデックスの計算方法の違いにより、次の表に示すように、CUDA* 実行空間の右端の次元 z は SYCL* の左端の次元 x にマップされます。

CUDA*

SYCL*

gridDim.x/y/z

sycl::nd_item.get_group_range(2/1/0)

blockIdx.x/y/z

sycl::nd_item.get_group(2/1/0)

blockDim.x/y/z

sycl::nd_item.get_local_range().get(2/1/0)

threadIdx.x/y/z

sycl::nd_item.get_local_id(2/1/0)

warpsSize

sycl::nd_item.get_sub_group().get_local_range().get(0)

カーネルの起動#

CUDA* は、<<<...>>> 実行構成構文と dim3 タイプを使用して、グリッドとブロックの次元とサイズを指定します。関数呼び出しオペレーターは、実行構成構文と組み合わせて、カーネル関数を実行するストリームの送信に使用されます。カーネルが送信されると、インデックス空間が定義され、各スレッドとブロックは一意のスレッド ID とブロック ID を受け取ります。これらの ID は、インデックス空間内のインデックス計算に使用され、ビルトイン変数を介してカーネル内でアクセスできます。

SYCL* は、メンバー関数 parallel_for (sycl::queue と sycl::range で提供される) を使用して、ND-range と work-group の次元とサイズを指定します。ユーザーは、SYCL* カーネルの属性 [[sycl::reqd_sub_group_size(dim)]] を適用して、指定された sub-group サイズでカーネルをコンパイルし実行する必要があることを示すこともできます。各デバイスは、 info::device::sub_group_sizes で定義される sub-group サイズのみをサポートします。

次の表は、SYCL* に移行されたカーネルを起動する例のオリジナル CUDA* コードを示しています。

オリジナルの CUDA* コード

移行された SYCL* コード

 
__global__ void foo() { 
   int a = threadIdx.x; 
} 
int main() { 
   dim3 size_1(100, 200, 300); 
   dim3 size_2(5, 10, 20); 
   foo<<<size_1, size_2>>>(); 
}
 
void foo(sycl::nd_item<3> item) { 
   int a = item.get_local_id(2); 
} 
int main() { 
   sycl::queue q; 
   sycl::range<3> size_1(300, 200, 100); 
   sycl::range<3> size_2(20, 10, 5); 
   q.parallel_for( 
      sycl::nd_range<3>(size_1 * size_2, size_2), [
      =](sycl::nd_item<3> item) [[sycl::reqd_sub_group_size(32)]] { 
         foo(item); 
   }); 
}

移行された SYCL* コード内の次の情報を確認してください。

  • sycl::nd_range のコンストラクターでは、最初のパラメーターのグローバルサイズは、work-group ではなく work-item です。グローバルサイズは、CUDA* に合わせて size_1 と size_2 の積にする必要があります。

  • スレッド・インデックスのセクションでは、CUDA* 実行空間の右端の次元 z を SYCL* の左端の次元 x にマップする必要があることを説明しています。つまり、この例では次元のサイズを入れ替える必要があります。

メモリーモデル#

CUDA* と SYCL* メモリーモデルは階層構造です。CUDA* では、共有メモリー、グローバルメモリー、定数メモリー、統合メモリーなど、複数のメモリー空間があります。共有メモリーによって、スレッドブロック内での効率良い通信が可能になります。グローバルメモリーは、デバイス全体のすべてのスレッドがアクセスできるため、大きな記憶領域が提供されますが、共有メモリーに比べるとアクセスは低速です。定数メモリーは読み取り専用の空間であり、定数やルックアップ・テーブルなどの不変データが保存されます。統合メモリーには、ホストとデバイスの両方からアクセスできます。

同様に、SYCL* では、ローカルメモリーは work-group 内で共有され、グローバルメモリーはすべての work-item からアクセスでき、共有メモリーはホストとデバイスでアクセスできます。SYCL* 2020 仕様では、定数メモリーは SYCL* デバイスのメモリーモデルに含まれなくなりました。

共有メモリー#

CUDA* 共有メモリーは、SYCL* ローカルメモリーにマップできます。これには、アクセスターゲットを sycl::access::target::local に設定してアクセサーを宣言します。例:

オリジナルの CUDA* コード

移行された SYCL* コード

 
__global__ void foo() { 
   __shared__ int shm[16]; 
   shm[0] = 2; 
} 
int main() { 
   foo<<<1, 1>>>(); 
}
 
void foo(int *shm) { 
   shm[0] = 2; 
} 
int main() { 
   sycl::queue q; q.submit([&](sycl::handler &cgh) { 

      sycl::local_accessor<int> shm_acc(sycl::range<1>(16), cgh); 
      cgh.parallel_for( 
         sycl::nd_range<3>(sycl::range<3>(1, 1, 1), 
         sycl::range<3>(1, 1, 1)), 
         [=](sycl::nd_item<3> item_ct1) { 
            foo(shm_acc.get_pointer()); 
      }); 
   }); 
}

グローバル、定数、および統合メモリー#

CUDA* グローバルメモリーと定数メモリーは、SYCL* グローバルメモリーにマップできます。CUDA* 統合メモリーは、SYCL* 共有メモリーにマップできます。これには、sycl::malloc_device または sycl::malloc_shared を使用してメモリーを割り当てます。例:

オリジナルの CUDA* コード

移行された SYCL* コード

 
void foo() { 
   int *mem1, *mem2; 

   cudaMalloc(&mem1, 10); 
   cudaMallocManaged(&mem2, 10); 
}
 
void foo() { 
   sycl::queue q; 
   int *mem1, *mem2; 

   mem1 = sycl::malloc_device<int>(10, q); 
   mem2 = sycl::malloc_shared<int>(10, q); 
}

CUDA* では、タイプ指定子 __device__、__constant__、および __managed__ を使用して、グローバルメモリーと統合メモリーに常駐する変数を宣言できることに注意してください。SYCL* には同等のものはありませんが、ヘルパークラスを使用して同等の機能を実現できます。dpct::global_memory、dpct::constant_memory、および dpct::shared_memory のリファレンス実装は、SYCLomatic プロジェクトで提供されています。

CUDA* デバイス API マッピング#

同期 API#

CUDA* では、CUDA* カーネル内のスレッドの実行を同期するため同期関数が使用されます。__syncthreads() は、スレッドブロック内のすべてのスレッドが関数に到達するまで、そのブロックのすべてのスレッドの実行をブロックします。さらに、__syncthreads() 以前にこれらのスレッドでアクセスされたすべてのグローバルメモリーと共有メモリーは、ブロック内のすべてのスレッドから参照できます。CUDA* の __syncthreads 関数は、sycl::group_barrier にマップできます (sycl::group オブジェクトが渡されます)。CUDA* の __syncthreads_and、__syncthreads_or、および __syncthreads_count 関数の移行には、sycl::group_barrier の後に追加のグループ・アルゴリズムが必要です。CUDA* の __syncwarp 関数は、sycl::group_barrier にマップできます (sycl::sub_group オブジェクトが渡されます)。

次の表は、同期関数の CUDA* から SYCL* へのマッピングを示しています。

CUDA*

SYCL*

__syncthreads()

sycl::group_barrier(Group)

__syncthreads_and()

sycl::group_barrier(Group) と sycl::all_of_group(Group, predicate)

__syncthreads_or()

sycl::group_barrier(Group) と sycl::any_of_group(Group, predicate)

__syncthreads_count()

sycl::group_barrier(Group) と sycl::reduce_over_group(Group, predicate?1:0, sycl::ext::oneapi::plus<>())

__syncwarp()

sycl::group_barrier(Sub_group)

メモリーフェンス API#

メモリーフェンス関数を使用すると、メモリーアクセスに順序付けを強制できます。メモリーフェンス関数は、順序付けが適用されるスコープが異なります。CUDA* のメモリーフェンス関数は、異なるメモリースコープの sycl::atomic_fence にマップできます。

次の表は、フェンス関数の CUDA* から SYCL* のマッピングを示しています。

CUDA*

SYCL*

__threadfence_block()

sycl::atomic_fence(sycl::memory_order::acq_rel, sycl::memory_scope::work_group)

__threadfence()

sycl::atomic_fence(sycl::memory_order::acq_rel, sycl::memory_scope::device)

__threadfence_system()

sycl::atomic_fence(sycl::memory_order::acq_rel, sycl::memory_scope::system)

ワープ組込み API#

ワープ投票やシャッフル関数などの CUDA* ワープ組込み関数は、SYCL* のグループ・アルゴリズム API にマッピングできます。

ワープ投票 API#

次の表に示すように、CUDA* ワープ投票 API は SYCL* のグループ・アルゴリズム API にマップできます。

CUDA*

SYCL*

__all()/__all_sync()

sycl::all_of_group()

__any()/__any_sync()

sycl::any_of_group()

__ballot()/__ballot_sync()

sycl::reduce_over_group()

CUDA* ワープ組込み関数の同期バージョンには、呼び出しに参加するスレッドを指定するマスクが渡されます。同等の SYCL* API では、マスク機能を直接サポートしていません。SYCLomatic プロジェクトのマスクバージョン API のリファレンス実装を参照してください。

次の表は、SYCL* に移行されたワープ投票サンプルのオリジナル CUDA* コードを示しています。

オリジナルの CUDA* コード

移行された SYCL* コード

 
__device__ void foo(){ 
    __all_sync(mask, predicate); 
}
 
void foo(sycl::nd_item<3> item) { 
    auto g = item.get_sub_group(); 
    sycl::all_of_group(g, (~mask & (0x1 << g.get_local_linear_id())) || predicate); 
}

ワープシャッフル API#

CUDA* ワープシャッフル関数は、次の SYCL* グループ・アルゴリズムにマップできます。

CUDA*

SYCL*

__shfl()/__shfl_sync()

sycl::select_from_group()

__shfl_up()/__shfl_up_sync()

sycl::shift_group_right()

__shfl_down()/__shfl_down_sync()

sycl::shift_group_left()

__shfl_xor()/__shfl_xor_sync()

sycl::permute_group_by_xor()

CUDA* シャッフル関数は、ワープのサブセットのスレッドの動作をサポートします。同等の SYCL* API は、sub_group のサブセットの操作を直接サポートしていません。SYCLomatic プロジェクトのヘルパー実装を参照してください。

次の表は、SYCL* に移行されたワープシャッフルの例のオリジナル CUDA* コードを示しています。

オリジナルの CUDA* コード

移行された SYCL* コード

 
__device__ void foo(){ 
   __shfl_sync(mask, val, r_id, 16); 
}
 
void foo(sycl::nd_item<3> item) { 
   auto g = item.get_sub_group(); 
   unsigned int start_index = 
      (g.get_local_linear_id() / 16) * 16; 
   sycl::select_from_group(g, val, start_index + r_id % 16); 
}

CUDA* ホスト API マッピング#

デバイス管理#

CUDA* デバイス管理 API は、次に示すように SYCL* デバイスクラスとそのメンバー関数にマップできます。

CUDA*

SYCL*

cudaGetDeviceCount()

sycl::device::get_devices()

cudaSetDevice()

sycl::device dev { device_selector } //Select sycl::device and make it ready for creating sycl::queue

cudaGetDevice()

sycl::queue.get_device() //Get active device from sycl::queue created

cudaGetDeviceProperties()/cudaDeviceGetAttribute()

sycl::device.get_info<info type>()

次の表は、SYCL* に移行されたデバイス管理の例のオリジナル CUDA* コードを示しています。

オリジナルの CUDA* コード

移行された SYCL* コード

 
int device_count; 
cudaGetDeviceCount(&device_count); 

for(int i = 0; i < device_count; i++) { 
   cudaSetDevice(i); 
   cudaDeviceProp prop; 
   cudaGetDeviceProperties(&prop, i); 
   int warp_size = prop.warpSize; 
   … 
   kernel<<<size_1, size_2>>>(); 
}
 
auto devices = 
      sycl::device::get_devices(sycl::info::device_type::gpu); 

for(auto &device : devices) { 
   sycl::queue q(device); 
   auto sub_group_sizes = 
      device.get_info<sycl::info::device::sub_group_sizes>(); 
   ... 
   q.parallel_for(sycl::nd_range<3>(
      size_1 * size_2, size_2), 
      [=](sycl::nd_item<3> item){ 
   kernel(item); 
   }); 
}

ストリーム管理#

CUDA* ストリーム管理 API は、次に示すように SYCL* キュークラスとそのメンバー関数にマップできます。

CUDA*

SYCL*

cudaStreamCreate()

sycl::queue のコンストラクター

cudaStreamDestroy()

sycl::queue のデストラクター

cudaStreamAddCallback()

std::async()

cudaStreamSynchronize()

sycl::queue.wait()

cudaStreamWaitEvent()

sycl::queue.ext_oneapi_submit_barrier()

次の表は、SYCL* に移行されたストリーム管理の例のオリジナル CUDA* コードを示しています。

オリジナルの CUDA* コード

移行された SYCL* コード

 
void callback(cudaStream_t st, 
                    cudaError_t status, 
                    void *vp) {…} 

void test() { 
   cudaStream_t stream; 
   cudaEvent_t event; 

   cudaStreamCreate(&stream); 
   cudaStreamAddCallback(stream, callback, 0, 0); 
   cudaStreamSynchronize(stream); 
   cudaStreamWaitEvent(stream, event, 0); 
   cudaStreamDestroy(stream); 
}
 
void callback(sycl::queue st, int status, void *vp) {…} 

void test() { sycl::queue q; sycl::event event; std::async([&]() { 
   q.wait(); callback(q, 0, 0); }); 
   q.wait(); 
   q.ext_oneapi_submit_barrier({event}); 
}

sycl::queue のコンストラクターは、デフォルトでアウトオブオーダーのプロパティーを持つキューを作成することに注意してください。インオーダー・キューを作成するには、sycl::property::queue::in_order::in_order() を使用します。

メモリー管理#

CUDA* メモリー管理 API は、次に示すように SYCL* USM ポインターベースのメモリー管理 API にマップできます。

CUDA*

SYCL*

cudaMalloc()

sycl::malloc_device()

cudaMallocHost()

sycl::malloc_host()

cudaMallocManaged()

sycl::malloc_shared()

cudaMemcpy()

sycl::queue.memcpy()

cudaMemset()

sycl::queue.memset()

cudaFree()/cudaFreeHost()

sycl::free()

次の表は、SYCL* に移行されたメモリー管理の例のオリジナル CUDA* コードを示しています。

オリジナルの CUDA* コード

移行された SYCL* コード

 
void test() { 
   int *dev_ptr, *host_ptr, *shared_ptr; 
   int size; 

   ... 
   cudaMalloc(&dev_ptr, size); 
   cudaMallocHost(&host_ptr, size); 
   cudaMallocManaged(&shared_ptr, size); 
   cudaMemset(dev_ptr, size, 0); 
   cudaMemcpy(host_ptr, 
               dev_ptr, size, cudaMemcpyHostToDevice); 
   cudaMemcpy(shared_ptr, 
               host_ptr, size, cudaMemcpyHostToDevice); 
   ... 
   int a = shared_ptr[0]; 
   ... 
   cudaFree(dev_ptr); 
   cudaFree(host_ptr); 
   cudaFree(shared_ptr); 
}
 
void test() { 
   sycl::queue q; 
   int *dev_ptr, *host_ptr, *shared_ptr; 
   int size; 
   ... 
   dev_ptr = (int *)sycl::malloc_device(size, q); 
   host_ptr = (int *)sycl::malloc_host(size, q); 
   shared_ptr = (int *)sycl::malloc_shared(size, q); 
   q.memset(dev_ptr, size, 0).wait(); 
   q.memcpy(host_ptr, dev_ptr, size).wait(); 
   q.memcpy(shared_ptr, host_ptr, size).wait(); 
   ... 
   int a = shared_ptr[0]; 
   ... 
   sycl::free(dev_ptr, q); 
   sycl::free(host_ptr, q); 
   sycl::free(shared_ptr, q); 
}

エラー処理#

CUDA* ランタイム・ライブラリーのエラー処理は、主に API 呼び出しが返すエラーコードに依存します。SYCL* では、同期エラーはランタイムが例外をスローすると即座に報告されます。この例外をキャッチして処理するには、try-catch 文を使用します。例:

オリジナルの CUDA* コード

移行された SYCL* コード

 
void test() { 
   int *ptr; 
   if (cudaMalloc(&ptr, sizeof(int))) { 
      std::cout << "error" << std::endl; 
   } 
}
 
void test() try { 
   int *ptr; 
   sycl::queue q; ptr = sycl::malloc_device<int>(1, q); 
} 
catch (sycl::exception const &exc) { 
   std::cerr << exc.what() 
             << "Exception caught at file:"<< __FILE__ << ", line:"<< __LINE__ << std::endl; 
   std::exit(1); 
}

非同期エラーは発生してもすぐには報告されません。キューはオプションですが、構築時に exception_list をパラメーターとして非同期ハンドラーを受け取ることができます。async_handler 呼び出しは、キューのメンバー関数 queue::wait_and_throw()、queue::throw_asynchronous() によってトリガーされるか、未処理の非同期エラーを含むキューの破棄時に自動的にトリガーされます。呼び出されると、async_handler が起動され、キューまたはコンテキストに関連付けられた未処理の非同期エラーを表す例外オブジェクトのリストを含む exception_list 引数を受け取ります。次の例は、SYCL* に移行された非同期例外ハンドラーの実装の 1 つを示しています。

オリジナルの CUDA* コード

移行された SYCL* コード

 
void test() { 
   int *ptr; 
   kernel<<<1, 1>>>(); 
   if (cudaDeviceSynchronize()) { 
      std::cout << "error" << std::endl; 
      } 
}
 
auto exception_handler = [](
      cl::sycl::exception_list exceptions) { 

   for (std::exception_ptr const &e : exceptions) { 
      try { 
         std::rethrow_exception(e); 
      } catch (cl::sycl::exception const &e) { 
         std::cerr << "Caught asynchronous SYCL exception:"<< std::endl 
         << e.what() 
         << std::endl 
         << "Exception caught at file:"<< __FILE__ << ", line:"<< __LINE__ << std::endl; 
      } 
   } 
}; 
void test() { 
   sycl::queue q{exception_handler}; 
   q.parallel_for( 
      sycl::nd_range<3>(size_1 * size_2, size_2),
      [=](sycl::nd_item<3> item){ 
         kernel(item); 
   }).wait_and_throw(); 
}