0 Replies Latest reply on Mar 16, 2018 2:24 PM by oostap

    half_float image scaler issue

    oostap

      Hi

      While developing Neural Network inference engine met weird issue with AMD OpenCL implementation.

      OS: Ubuntu 16.04.4

      AMD Driver: "amdgpu-pro-17.50-511655"

      The same behaviour on Windows 10 with whatever drivers are installed on system for RX 580 GPU.

       

       

      Made just simple test to: fill image2d (RGBA, half_float) with values 0.1 (converted to half_float) and scale down with scaler.

       

      Kernel:

      #pragma OPENCL EXTENSION cl_khr_fp16 : enable\

      constant sampler_t sampler =

            CLK_NORMALIZED_COORDS_TRUE

          | CLK_ADDRESS_CLAMP

          | CLK_FILTER_NEAREST;

       

      __kernel void testKernel(

          __read_only image2d_t input,

          __global float4 *out,

          unsigned int outWidth,

          unsigned int outHeight

      )

      {

       

          float inX = (float)get_global_id(0)/ outWidth;

          float inY = (float)get_global_id(1)/ outHeight;

          float2 posIn = (float2) (inX, inY);

          half4 px = read_imageh(input, sampler, posIn);

          float4 pixel = (float4)(px.x,px.y,px.z,px.w);

          out[get_global_id(1)*outWidth+get_global_id(0)]=pixel;

          printf(\"[%f,%f] %f,%f,%f,%f\",inX, inY, (float)px.x,(float)px.y,(float)px.z,(float)px.w);

      }"

       

      Test works great on mobile ARM Mali GPU (as well my entire NN inference engine) but not on AMD GPU (RX 580).

      Is FP16 image scaler working in AMD OpenCL driver?