gpucoder.atomicExch

Write an entry-point function, anyCols, that accepts an M-by-N matrix, A. The function returns a 1-by-N vector, B, where each element of B is equal to 1 if the corresponding column of A contains a nonzero element, and the element is equal to 0 otherwise.

function B = anyCols(A) %#codegen
    coder.gpu.kernelfun();
    rows = size(A,1);
    cols = size(A,2);
    B = zeros([1 cols]);
    coder.gpu.kernel();
    for i=1:cols
        coder.gpu.kernel();
        for j=1:rows
            if (A(j,i) ~= 0)
                B(i) = 1;
            end
        end
    end
end

If you parallelize the for-loops across multiple threads, it is possible for multiple threads to read from and write to B(i) at once. To write to B(i) without interference from other threads, use gpucoder.atomicExch to exchange the value of B(i) with 1. If the loop sets B(i) equal to 1, the value of B(i) does not change for the rest of the loop execution, so parallelizing the loop across multiple threads does not change the result.

function B = anyCols(A) %#codegen
    coder.gpu.kernelfun();
    rows = size(A,1);
    cols = size(A,2);
    B = zeros([1 cols]);
    coder.gpu.kernel();
    for i=1:cols
        coder.gpu.kernel();
        for j=1:rows
            if (A(j,i) ~= 0)
                B(i) = gpucoder.atomicExch(B(i),1);
            end
        end
    end
end

Create a GPU code configuration object and generate code for anyCols.

A = zeros(1024);
cfg = coder.gpuConfig("mex");
codegen anyCols -config cfg -args {A}

The generated code contains the kernel anyCols_kernel2, which calculates each column of B by using the gpu_atomicExch function.

static __global__
    __launch_bounds__(256, 1) void anyCols_kernel2(const real_T A[1048576],
                                                   real_T B[1024])
{
  uint64_T gThreadId;
  int32_T i;
  int32_T j;
  gThreadId = mwGetGlobalThreadIndex();
  j = static_cast<int32_T>(gThreadId % 1024ULL);
  i = static_cast<int32_T>((gThreadId - static_cast<uint64_T>(j)) / 1024ULL);
  if ((i < 1024) && (A[j + (i << 10)] != 0.0)) {
    gpu_atomicExch(&B[i], 1.0);
  }
}