gpucoder.atomicInc

Atomically increment variable in global or shared memory within upper bound

Since R2021b

Syntax

A = gpucoder.atomicInc(A,B)

[A,oldA] = gpucoder.atomicInc(A,B)

Description

The gpucoder.atomicInc function reads a value from a global or shared GPU memory location, increments the value, and writes the result 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.atomicInc(A,B) increments the value of A within the upper bound B. If the value of A is greater than or equal to B, it is reset to zero.

Call the gpucoder.atomicInc 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.atomicInc.

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

Examples

Increment within a Upper Bound Using CUDA atomicInc

Perform a simple atomic warp around increment operation by using the gpucoder.atomicInc function and generate CUDA code that calls corresponding CUDA atomicInc() APIs.

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

function a = myAtomicInc(a,b)

coder.gpu.kernel;
for i =1:numel(a)
    a(i) = gpucoder.atomicInc(a(i), b);
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]);
inputArgs = {A,B};

To generate a CUDA library, use the codegen function.

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

codegen -config cfg -args inputArgs myAtomicInc -d myAtomicInc

The generated CUDA code contains the myAtomicInc_kernel1 kernel with calls to the atomicInc() CUDA APIs.

//
// File: myAtomicInc.cu
//
...

static __global__ __launch_bounds__(1024, 1) void myAtomicInc_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);
    atomicInc(&a_data[b_i], b);
  }
}
...

void myAtomicInc(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);
    myAtomicInc_kernel1<<<grid, block>>>(b, i, gpu_a_data);
    cudaMemcpy(a_data, gpu_a_data, a_size[1] * sizeof(uint32_T),
               cudaMemcpyDeviceToHost);
...

}

Input Arguments

A — Reference to shared or global memory location

scalar

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

Data Types: uint32

B — Operand

scalar

Operand, specified as a scalar.

Data Types: uint32

Extended Capabilities

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