1 Reply Latest reply on Mar 31, 2016 3:51 PM by ekondis

    Optimizing OpenCL 1D downsampler

    gales

      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;

      }