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

       else

       value += ptr[i];

  }

  if (type == 2)

       value /= ratio;

  out[index] = value;

}

Outcomes