I have this code :
[codebox]#include <stdio.h>
#include <math.h>
#include <stdlib.h>
device float p[12];
device int counter=0;
device int array_counter[5+1];
// Operation on a single element of a 1D array
global void test0D_cu(int j )
{
int jj,ix;
ix= threadIdx.x+0; // from 0 to 4
jj = j + ix - 2; // " 2 6
atomicAdd(&counter,1); // increment the (saving)array index
atomicExch(&array_counter[counter],jj); // save for each thread, this count in an array
// printf ( " ix %d array_counter[counter] %d counter %d \n",ix , array_counter[counter] , counter );
}
int main(void)
{
int j,h_array_counter[5+1];
float p1[12];
unsigned char *d_addr = NULL;
cudaGetSymbolAddress( (void**)&d_addr,“p”);
cudaMemset( d_addr, 0.0f , sizeof(p)); // reset the device array to 0.0
cudaGetSymbolAddress( (void**)&d_addr,"array_counter");
cudaMemset( d_addr, 0 , sizeof(array_counter)); // reset the device array to 0
dim3 gridsize (1);
dim3 blocksize ( 5);
j = 4; // one test element indices
test0D_cu <<< gridsize,blocksize >>> (j);
cudaMemcpyFromSymbol(p1,p,sizeof(p)); // recover the main array and array_counter
cudaMemcpyFromSymbol(&h_array_counter,array_counter,sizeof(array_counter));
int ii=1;
for ( int m = 2 ; m <= 6; m ++){
printf ( "ii %d m %d p1[jj] %e %d \n",ii, m, p1[m],h_array_counter[ii]);
ii++;
}
}
[/codebox]
Which get these results with :
- emulation (what I would like to see)
[codebox]
ii 1 m 2 p1[jj] 0.000000e+00 2
ii 2 m 3 p1[jj] 0.000000e+00 3
ii 3 m 4 p1[jj] 0.000000e+00 4
ii 4 m 5 p1[jj] 0.000000e+00 5
ii 5 m 6 p1[jj] 0.000000e+00 6
[/codebox]
- release ( the last collumn is wrong)
[codebox]
ii 1 m 2 p1[jj] 0.000000e+00 0
ii 2 m 3 p1[jj] 0.000000e+00 0
ii 3 m 4 p1[jj] 0.000000e+00 0
ii 4 m 5 p1[jj] 0.000000e+00 0
ii 5 m 6 p1[jj] 0.000000e+00 6
[/codebox]
I would like to get the same result for the release case. Is this possible? It seems that only the last iteration is save in the array, why is that?
Do I have to use __threadfence() ,syncthreads() or anything else ? I already tried __syncthreads() and it did not change anything.
Thank you for your help
Haw can I edit the topic title?