Min Max reduce kernels

Accelerated Computing CUDA CUDA Programming and Performance
1

hello !

how to modify reduce example for sorting arrays ?

i wanna know the min/max value in array with its position in array.

i have to try expand float array to float2 array where f.x is value and f.y is position

__global__ void ExpandData(float* indata,int shift,int n,float2* expdata)

{	

	const unsigned int index=(blockIdx.x*blockDim.x+threadIdx.x);

	

	expdata[index].x=indata[index];

	expdata[index].y=index;

}

and sort it with single reduce, but results are wrong …

__global__ void FindSpikeReduce(float2* indata,float* sumdata,int n,int shift,float4* outdata)

{

	const unsigned int tid=(blockIdx.x*threadIdx.x);

	float2 temp={0,0};

	

	for(unsigned int stride=n/2; stride>0; stride>>=1) //0918728093

   {

       if (tid < stride) 

    {

     //sumdata[tid].x+=sumdata[tid+stride].x;

           if (indata[tid].x<indata[tid+stride].x)

     {

      temp=indata[tid];

      indata[tid]=indata[tid+stride];

      indata[tid+stride]=temp;

     }

       }

       __syncthreads();

   }

   outdata[shift].x=indata[0].x;

   outdata[shift].y=indata[0].y;

   outdata[shift].z=sumdata[0];

   outdata[shift].w=shift;

}

classic loop finding gives correct results :

for (j = 1; j < 8; j++) {

    sum+=fp_PowerSpectrum[index+j];	

       if (fp_PowerSpectrum[index+j] > temp) {

          temp = fp_PowerSpectrum[index+j];

    tpos=j;

      }
2

Have you tried CUDPP?

CUDPP: CUDA Data Parallel Primitives Library
[url=“http://www.gpgpu.org/developer/cudpp/”]http://www.gpgpu.org/developer/cudpp/[/url]

3

CUDPP sort is working only for positive floats but i need sort negatives too …

any hint to modify reduce sample from sdk or what is wrong in my kernels ???

4

hmm,

__shared__ float value_max[NUMTHREADS];

__shared__ float value_min[NUMTHREADS];

__shared__ int index_max[NUM_THREADS];

__shared__ int index_min[NUM_THREADS];

and then later in the code, instead of

value_min[index] = fminf(value_min[index], value_min[index+xx]);

value_max[index] = fmaxf(value_max[index], value_max[index+xx]);

you do something like:

if (value_min[index] < value_min[index+xx])

{

  value_min[index] = value_min[index+xx];

  index_min[index] = index_min[index+xx];

}

if (value_max[index] < value_max[index+xx])

{

  value_max[index] = value_max[index+xx];

  index_max[index] = index_max[index+xx];

}