gpucoder.atomicCAS - Atomically compare and swap the value of a variable in global or shared

  memory - MATLAB ([original](https://in.mathworks.com/help/gpucoder/ref/gpucoder.atomiccas.html)) ([raw](?raw))

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

Since R2021b

Syntax

Description

[A,oldA] = gpucoder.atomicCAS([A](#mw%5Fe9d9ebed-1d2e-4a4a-b334-74d763ce88a4),[B](#mw%5Fe9d9ebed-1d2e-4a4a-b334-74d763ce88a4),[C](#mw%5Fe9d9ebed-1d2e-4a4a-b334-74d763ce88a4)) compares B to the value of A in global or shared memory and if the values are the same writes the value of C into A. 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.

example

Examples

collapse all

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

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

function a = myAtomicCAS(a,b,c)

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

end

To create a type for a matrix of doubles for use in code generation, use thecoder.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(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

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

Data Types: int32 | uint32 | uint64

Version History

Introduced in R2021b

See Also

Functions

Topics