Main Content

gpucoder.atomicAnd

指定された値とグローバル メモリまたは共有メモリの変数の間でのビット単位 AND のアトミックな実行

R2021b 以降

    説明

    [A,oldA] = gpucoder.atomicAnd(A,B) は、B とグローバル メモリまたは共有メモリの A の値の間でビット単位の AND を実行し、結果を A に書き戻します。この演算は、読み取り、変更、書き込みの操作全体が他のスレッドからの干渉なしに実行されることが保証されているという意味でアトミックです。入力引数と出力引数の順序は、構文の規定と一致しなければなりません。

    すべて折りたたむ

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

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

    function a = myAtomicAnd(a,b)
    
    coder.gpu.kernelfun;
    for i =1:numel(a)
        [a(i),~] = gpucoder.atomicAnd(a(i), b);
    end
    
    end
    

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

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

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

    cfg = coder.gpuConfig('lib');
    cfg.GenerateReport = true;
    
    codegen -config cfg -args inputArgs myAtomicAnd -d myAtomicAnd
    

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

    //
    // File: myAtomicAnd.cu
    //
    ...
    
    static __global__ __launch_bounds__(1024, 1) void myAtomicAnd_kernel1(
        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);
        atomicAnd(&a_data[b_i], b);
      }
    }
    ...
    
    void myAtomicAnd(uint32_T a_data[], int32_T a_size[2], uint32_T b)
    {
      dim3 block;
      dim3 grid;
    ...
    
        cudaMemcpy(gpu_a_data, a_data, a_size[1] * sizeof(uint32_T),
                   cudaMemcpyHostToDevice);
        myAtomicAnd_kernel1<<<grid, block>>>(b, i, gpu_a_data);
        cudaMemcpy(a_data, gpu_a_data, a_size[1] * sizeof(uint32_T),
                   cudaMemcpyDeviceToHost);
    ...
    
    }
    

    入力引数

    すべて折りたたむ

    オペランド。スカラー、ベクトル、行列、または多次元配列として指定します。入力 A と入力 B は以下の要件を満たさなければなりません。

    • 同じデータ型。

    • 同じサイズまたは互換性のあるサイズ。たとえば、AMN 列の行列で、B がスカラーまたは 1N 列の行ベクトル。

    データ型: int32 | uint32 | uint64

    制限

    • gpucoder.stencilKernel プラグマの関数ハンドル入力にアトミック関数の呼び出しを含めることはできません。次に例を示します。

      out1 = gpucoder.stencilKernel(@myAtomicAnd,A,[3 3],'same',B);
      

    バージョン履歴

    R2021b で導入