Why does CUDA atomic function not work?

Question

The problem is I want to apply atomicMin to an int array. My array has 6 elements and the code is:

#include <stdlib.h>
#include <stdio.h>

__global__ void kernel(int * globalmax, int * r){

    int val = r[threadIdx.x];
    *globalmax = atomicMin(globalmax, val);
}


int main(int argc, char **argv) {

    int m = 0;
    int * d_i;
    int r[] = {-3, 150, 11, -1, 15, -4};
    int * d_r;
    cudaMalloc((void **) &d_r, sizeof( int) * 6);
    cudaMalloc((void **) &d_i, sizeof(int));
    cudaMemcpy(d_r, r, sizeof( int) * 6, cudaMemcpyHostToDevice);
    cudaMemcpy(d_i, &m, sizeof(int) , cudaMemcpyHostToDevice);

    kernel<<<1, 7>>>(d_i, d_r);

    cudaMemcpy(&m, d_i, sizeof(int), cudaMemcpyDeviceToHost);
    printf("%d", m);
    cudaFree(d_i);
    cudaFree(d_r);
    return 0;
}

When I call the kernel code with 6 threads if the minimum element position is 6 the minimum element can not be returned and if I call a kernel with 7 thread I get the right answer. What is wrong with this code?

Answer 1

You shouldn't be trying to store the old value back to the atomic variable. Change:

*globalmax = atomicMin(globalmax, val);

to:

atomicMin(globalmax, val);