title: OpenCL bicubic interpolation kernel failed with error CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST
body: <p>Bicubic interpolation is one of the common interpolation method, but I can not find any working implementation on OpenCL. I was decided to write bicubic interpolation on OpenCL myself, but ...</p>
  
  <p>I have some problem with kernel programm.  </p>
  
  <p>When I run kernel execution, program failed with error CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST. No any other information about cause of  error. I am using javacl binding form google code: <a href="http://code.google.com/p/javacl" rel="nofollow">http://code.google.com/p/javacl</a>, AMD Accelerated Parallel Processing SDK 2.3 on Ubuntu linux 10.10, hardware AMD Radeon 5xxxHD</p>
  
  <p>I haven`t opencl debugger on ubuntu for AMD APP SDK (</p>
  
  <p>If I uncomment float4 val=read_imagef(signal, sampler, (float2)(x+iX,y+iY)); and comment  calculation of bicubic interpolation  "float4 val=..." all work without any error(but using bilinear interpolation). I think that this error because of invalid memory access or register memory overflow.</p>
  
  <pre><code>const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_FILTER_LINEAR | CLK_ADDRESS_CLAMP_TO_EDGE;
  const float CATMULL_ROM[16]={-0.5F,1.5F,-1.5F,0.5F,1.0F,-2.5F,2.0F,-0.5F,-0.5F,0.0F,0.5F,0.0F,0.0F,1.0F,0.0F,0.0F};
  __kernel void bicubicUpscale(int scale,read_only image2d_t signal, write_only image2d_t upscale) {
  
  int x = get_global_id(0)-2, y = get_global_id(1)-2;
  
  float C[16];
  float T[16];
  
  
  for (int i = 0; i &lt; 16; i++)
  {
      C[i]=0.0F;
      T[i]=0.0F;
  }
  
  for (int i = 0; i &lt; 4; i++)
      for (int j = 0; j &lt; 4; j++)
          for (int k = 0; k &lt; 4; k++)
          {
              T[4*i+j] += read_imagef(signal, sampler, (int2)(x+k,y+i)).x * CATMULL_ROM[4*j+k];
          }
  for (int i = 0; i &lt; 4; i++)
      for (int j = 0; j &lt; 4; j++)
          for (int k = 0; k &lt; 4; k++)
          {
              C[4*i+j] += CATMULL_ROM[4*i+k] * T[4*k+j];
          }
  
  for (int i = 0; i &lt; scale; i++)
  {
      for (int j = 0; j &lt; scale; j++)
      {
          float iX=(float)j/(float) scale;
          float iY=(float)i/(float) scale;
          //float4 val=read_imagef(signal, sampler, (float2)(x+iX,y+iY));
          float val= iX * (iX * (iX * (iY * (iY * (iY * C[0] + C[1]) + C[2]) + C[3])
          + (iY * (iY * (iY * C[4] + C[5]) + C[6]) + C[7]))
          + (iY * (iY * (iY * C[8] + C[9]) + C[10]) + C[11]))
          + (iY * (iY * (iY * C[12] + C[13]) + C[14]) + C[15]);
          write_imagef(upscale, (int2)(x*scale+j, y*scale+i), val);
      }
  }
  }
  </code></pre>
  
  <hr>
  
  <p>I rewrite this program for using local memory, but it still not working correctly</p>
  
  <pre><code>const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_FILTER_LINEAR | CLK_ADDRESS_CLAMP_TO_EDGE;
  const float CATMULL_ROM[]={-0.5F,1.5F,-1.5F,0.5F,1.0F,-2.5F,2.0F,-0.5F,-0.5F,0.0F,0.5F,0.0F,0.0F,1.0F,0.0F,0.0F};
  __kernel void bicubicUpscale(local float* sharedBuffer,int scale,read_only image2d_t signal, write_only image2d_t upscale) {
  int x = get_global_id(0)-2, y = get_global_id(1)-2;
  //int locX=get_local_id(0);
  
  int offsetT = (y+2)*512+(x+2)*32+16;
  int offsetC = (y+2)*512+(x+2)*32;
  
  global float* C=&amp;sharedBuffer[offsetT];
  global float* T=&amp;sharedBuffer[offsetT];
  
  for (int i = 0; i &lt; 32; i++){
      sharedBuffer[offsetC+ i]=0.0F;
  }
  
  for (int i = 0; i &lt; 4; i++)
      for (int j = 0; j &lt; 4; j++)
          for (int k = 0; k &lt; 4; k++){
              //T[4*i+j] = mad(read_imagef(signal, sampler, (int2)(x+k,y+i)).x,CATMULL_ROM[4*j+k],T[4*i+j]);
              T[i+j] += read_imagef(signal, sampler, (int2)(x+k,y+i)).x * CATMULL_ROM[4*j+k];
          }
  for (int i = 0; i &lt; 4; i++)
      for (int j = 0; j &lt; 4; j++)
          for (int k = 0; k &lt; 4; k++){
              //C[4*i+j] = mad(CATMULL_ROM[4*i+k],T[4*k+j],C[4*i+j]);
              sharedBuffer[offsetC +4*i+j] += CATMULL_ROM[4*i+k] * sharedBuffer[offsetT + 4*k+j];
          }
  
  
  barrier (CLK_GLOBAL_MEM_FENCE);
  
  
  for (int i = 0; i &lt; scale; i++)
      for (int j = 0; j &lt; scale; j++)
          {
              float iX=(float)j/(float) scale;
              float iY=(float)i/(float) scale;
              float4 val= iX * (iX * (iX * (iY * (iY * (iY * C[0] + C[1]) + C[2]) + C[3])
              + (iY * (iY * (iY * C[4] + C[5]) + C[6]) + C[7]))
              + (iY * (iY * (iY * C[8] + C[9]) + C[10]) + C[11]))
              + (iY * (iY * (iY * C[12] + C[13]) + C[14]) + C[15]);
              write_imagef(upscale, (int2)(x*scale+j, y*scale+i), val);
          }
  }
  </code></pre>
  
  <p>Do you know any decision for this problem.</p>
  
  <p><a href="https://docs.google.com/leaf?id=0B_sJf8pL_3q_NDE3ZjliNGEtZDY5NS00MzI1LWI2NzAtNGQ3ZDNlYzIxZWMy&amp;hl=en" rel="nofollow">Java sources + maven2 build</a>. Use command "mvn clean compile exec:java" to compile and run demo.</p>
  
  <p>Regards, 
     Igor</p>