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?
You shouldn't be trying to store the old value back to the atomic variable. Change:
*globalmax = atomicMin(globalmax, val);
to:
atomicMin(globalmax, val);