gpucoder.atomicCAS

グローバル メモリまたは共有メモリの変数の値のアトミックな比較交換

R2021b 以降

ページ内をすべて折りたたむ

    構文

    A = gpucoder.atomicCAS(A,B,C)
    [A,oldA] = gpucoder.atomicCAS(A,B,C)

    説明

    gpucoder.atomicCAS 関数は、グローバルまたは共有の GPU メモリ位置から値を読み取り、その値を 2 番目の値と比較し、2 つの値が等しい場合は、3 番目の値をメモリ位置に書き戻します。生成された GPU コードでは、操作は "アトミック" となります。これは GPU スレッドが他のスレッドからの干渉を受けることなく、読み取り・変更・書き込み操作を実行することを意味します。

    A = gpucoder.atomicCAS(A,B,C) は、B を A の値と比較し、値が同じであれば、C の値を A に書き込みます。

    生成された CUDA® コード内にある GPU で実行したい for ループ内で、gpucoder.atomicCAS 関数を直接呼び出します。ループの各反復処理では、A への読み書きが可能でなければなりません。gpucoder.atomicCAS を含むループの前に coder.gpu.kernel プラグマを使用します。

    例

    [A,oldA] = gpucoder.atomicCAS(A,B,C) は、A の以前の値を oldA として返します。

    例

    すべて折りたたむ

    関数 gpucoder.atomicCAS を使用してシンプルなアトミック比較交換演算を実行し、対応する CUDA atomicCAS() API を呼び出す CUDA コードを生成します。

    1 つのファイル内に、行列入力 a、b、および c を受け入れるエントリポイント関数 myAtomicCAS を記述します。

    function a = myAtomicCAS(a,b,c)

coder.gpu.kernel;
for i =1:numel(a)
    a(i) = gpucoder.atomicCAS(a(i), b, c);
end

end

    コード生成で使用する uint32 行列の型を作成するには、coder.newtype 関数を使用します。

    A = coder.newtype('uint32', [1 30], [0 1]);
B = coder.newtype('uint32', [1 1], [0 0]);
C = coder.newtype('uint32', [1 1], [0 0]);
inputArgs = {A,B,C};

    CUDA ライブラリを生成するには、関数 codegen を使用します。

    cfg = coder.gpuConfig('lib');
cfg.GenerateReport = true;

codegen -config cfg -args inputArgs myAtomicCAS -d myAtomicCAS

    生成された CUDA コードには、atomicCAS() CUDA API の呼び出しを使用する myAtomicCAS_kernel1 カーネルが含まれます。

    //
// File: myAtomicCAS.cu
//
...

static __global__ __launch_bounds__(1024, 1) void myAtomicCAS_kernel1(
    const uint32_T c, const uint32_T b, const int32_T i, uint32_T a_data[])
{
  uint64_T loopEnd;
  uint64_T threadId;

...

  for (uint64_T idx{threadId}; idx <= loopEnd; idx += threadStride) {
    int32_T b_i;
    b_i = static_cast<int32_T>(idx);
    atomicCAS(&a_data[b_i], b, c);
  }
}
...

void myAtomicCAS(uint32_T a_data[], int32_T a_size[2], uint32_T b, uint32_T c)
{
  dim3 block;
  dim3 grid;
...

  if (validLaunchParams) {
    cudaMemcpy(gpu_a_data, a_data, a_size[1] * sizeof(uint32_T),
               cudaMemcpyHostToDevice);
    myAtomicCAS_kernel1<<<grid, block>>>(c, b, i, gpu_a_data);
    cudaMemcpy(a_data, gpu_a_data, a_size[1] * sizeof(uint32_T),
               cudaMemcpyDeviceToHost);
...

}

    入力引数

    すべて折りたたむ

    共有またはグローバルの GPU メモリ位置への参照。スカラーとして指定します。参照を gpucoder.atomicCAS の入力引数および出力引数として使用します。

    データ型: int32 | uint32 | uint64

    オペランド。スカラーとして指定します。

    データ型: int32 | uint32 | uint64

    拡張機能

    すべて展開する

    C/C++ コード生成
    MATLAB® Coder™ を使用して C および C++ コードを生成します。

    GPU コード生成
    GPU Coder™ を使用して NVIDIA® GPU のための CUDA® コードを生成します。

    バージョン履歴

    R2021b で導入

    参考

    関数

    トピック