AnsweredAssumed Answered

Optimizing OpenCL 1D downsampler

Question asked by gales on Feb 29, 2016
Latest reply on Mar 31, 2016 by ekondis

I would like a very fast 1D array resampler but haven't been able to find one so I'm trying a two stage approach. Linear interpolate in an upsampler to some integer factor greater that what I really want and downsample to the value I need.


Example: incomming data set is 1024 samples long and I want 801 samples in my final array. I then upsample to 1602 then downsample by a factor of 2 to get 801.


Problem is it's just not fast enough even on my W9100 GPU. I'm trying to process 1G samples/sec. Funny thing is the clFFT has no problem processing at these rates but a simple downsampler chokes.

I'm doing a maximum, minimum, and average output in parallel on seperate queues. I tried doing them all together in a float_3 array output but that was even slower.


Does anyone have any ideas on how to optimize this?


__kernel void Downsample(__global float *in, __global float *out, int type, int ratio)


  int gx = get_global_id(0);

  int gw = get_global_size(0);

  int gy = get_global_id(1);

  int gh = get_global_size(1);


  int index = gy * gw + gx;


  __global float *ptr = in + index * ratio;


  float value = ptr[0];

  for (int i = 1; i < ratio; i++)


       if (type == 0)

       value = max(value, ptr[i]);

       else if (type == 1)

            value = min(value, ptr[i]);


       value += ptr[i];


  if (type == 2)

       value /= ratio;

  out[index] = value;