gpucoder.atomicCAS

Atomically compare and swap value of variable in global or shared memory

Since R2021b

collapse all in page

    Syntax

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

    Description

    The gpucoder.atomicCAS function reads a value from a global or shared GPU memory location, compares the value to a second value, and, if the two values are equal, writes a third value back to the memory location. In generated GPU code, the operation is atomic, which means that a GPU thread performs the read-modify-write operation without interference from other threads.

    A = gpucoder.atomicCAS(A,B,C) compares B to the value of A and, if the values are the same, writes the value of C into A.

    Call the gpucoder.atomicCAS function directly inside a for-loop that you want to execute on the GPU in the generated CUDA® code. Each iteration of the loop must be able to read and write to A. Use the coder.gpu.kernel pragma before the loop that contains gpucoder.atomicCAS.

    example

    [A,oldA] = gpucoder.atomicCAS(A,B,C) returns the previous value of A as oldA.

    Examples

    collapse all

    Perform a simple atomic compare and swap operation by using the gpucoder.atomicCAS function and generate CUDA code that calls corresponding CUDA atomicCAS() APIs.

    In one file, write an entry-point function myAtomicCAS that accepts matrix inputs a,b, and c.

    function a = myAtomicCAS(a,b,c)

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

end

    To create a type for a uint32 matrix for use in code generation, use the coder.newtype function.

    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};

    To generate a CUDA library, use the codegen function.

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

codegen -config cfg -args inputArgs myAtomicCAS -d myAtomicCAS

    The generated CUDA code contains the myAtomicCAS_kernel1 kernel with calls to the atomicCAS() CUDA APIs.

    //
// 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);
...

}

    Input Arguments

    collapse all

    Reference to a shared or global GPU memory location, specified as a scalar. Use the reference as an input and output argument of gpucoder.atomicCAS.

    Data Types: int32 | uint32 | uint64

    Operands, specified as scalars.

    Data Types: int32 | uint32 | uint64

    Extended Capabilities

    expand all

    C/C++ Code Generation
    Generate C and C++ code using MATLAB® Coder™.

    GPU Code Generation
    Generate CUDA® code for NVIDIA® GPUs using GPU Coder™.

    Version History

    Introduced in R2021b

    See Also

    Functions

    Topics