Min Max reduce kernels

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);





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) 



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














classic loop finding gives correct results :

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


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

          temp = fp_PowerSpectrum[index+j];



Have you tried CUDPP?

CUDPP: CUDA Data Parallel Primitives Library

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 ???


__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];