Hi, i wrote a simple kernel to perform odd even transposition sort. I have to sort an array of floats and I also need to know the new positions of the original indexes.
This is the kernel:
const unsigned int SIZE = 8;
global void sort_device (float* vec_in, float* vec_out, float* indices)
{
unsigned int tid = threadIdx.x;
__shared__ int i;
__shared__ float input;
__shared__ float inde1;
__shared__ float tmp1;
__shared__ float tmp2;
input[tid]= vec_in[tid];
inde1[tid]= indices[tid];
__syncthreads();
for (i=0; i<SIZE; i++)
{ if(tid%2==0)
{
if (input[tid]<input[tid+1])
{
tmp1 = input[tid];
input[tid] = input [tid+1];
input [tid+1] =tmp1;
/*tmp2 = inde1[tid];
inde1[tid] = inde1[tid+1];
inde1[tid+1] =tmp2;
*/
}
}
__syncthreads();
if((tid%2!=0)&&(tid!= SIZE))
{ if (input[tid]<input[tid+1])
{ tmp1 = input[tid];
input[tid] = input [tid+1];
input [tid+1] =tmp1;
/*tmp2 = inde1[tid];
inde1[tid] = inde1 [tid+1];
inde1 [tid+1] =tmp2;
*/
}
}
__syncthreads();
}
vec_out [tid]= input[tid];
indices[tid]=inde1[tid];
__syncthreads();
}
…
unsigned int block_size = SIZE;
dim3 dimGrid(1, 1);
dim3 dimBlock(block_size, 1);
sort_device <<< dimGrid, dimBlock >>> (vec_in_d, vec_out_d, indices_d);
…
This does it fine but if i uncomment the commented part (to have the position of the indexes) this makes problems to both vectors…
I don´t know why this happens. I don´t care about timing considerations because i have no much time in real life to do this…
Why does it happens? Is there a solution to this particular kernel?
