gpucoder.atomicDec

Atomically decrement a variable in global or shared memory within a specified upper bound

Since R2021b

Syntax

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

Description

[A,oldA] = gpucoder.atomicDec(A,B) decrements the value of A in global or shared memory within the upper bound B. The operation is atomic in a sense that the entire read-modify-write operation is guaranteed to be performed without interference from other threads. The order of the input and output arguments must match the syntax provided.

Examples

Decrement within a Upper Bound Using CUDA atomicDec

Perform a simple atomic wrap around decrement operation by using the gpucoder.atomicDec function and generate CUDA^® code that calls corresponding CUDA atomicDec() APIs.

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

function a = myAtomicDec(a,b)

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

end

To create a type for a matrix of doubles 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 myAtomicDec -d myAtomicDec

The generated CUDA code contains the myAtomicDec_kernel1 kernel with calls to the atomicDec() CUDA APIs.

//
// File: myAtomicDec.cu
//
...

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

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

}

Input Arguments

A, B — Operands

scalars | vectors | matrices | multidimensional arrays

Operands, specified as scalars, vectors, matrices, or multidimensional arrays. Inputs A and B must satisfy the following requirements:

• Have the same data type.

• Have the same size or have sizes that are compatible. For example, A is an M-by-N matrix and B is a scalar or 1-by-N row vector.

Data Types: uint32

Limitations

• Function handle input to the gpucoder.stencilKernel pragma cannot contain calls to atomic functions. For example,

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

Version History

Introduced in R2021b