title: OpenCL program doesn't build when a function is called from a kernel
body: <p>This question follows from <a href="https://stackoverflow.com/questions/44856649/code-terminates-after-saying-could-not-create-kernel-on-eclipse/44881295#44881295">this</a> question.</p>
  
  <p>I have a kernel which calls a regular function. When I build and run my code, I get the following output:</p>
  
  <pre><code>Number of devices: 2
  building program failed
  -----COULD NOT CREATE KERNEL!!---
  </code></pre>
  
  <p>The problematic portion of my <code>.cl</code> file is as follows:</p>
  
  <pre><code>    #define  IDCT_INT_MIN   (- IDCT_INT_MAX - 1)
      #define  IDCT_INT_MAX   2147483647
        ....
        ....
        ....
        ....
        ....
        ....
  
      #define SCALE(x,n)      ((x) &lt;&lt; (n))
  
      #define but(a,b,x,y)    { x = SUB(a,b); y = ADD(a,b); }
  
      static  int DESCALE (int x, int n)
      {
          return (x + (1 &lt;&lt; (n - 1)) - (x &lt; 0)) &gt;&gt; n;
      }
  
      static  int ADD(int x, int y)
      {
          int r = x + y;
  
          return r;       
      }
  
      static  int SUB(int x, int y)
      {
          int r = x - y;
  
          return r;        
      }
  
      static  int CMUL(int c, int x)
      {
          int r = c * x;
          r = (r + (1 &lt;&lt; (C_BITS - 1))) &gt;&gt; C_BITS;
          return r;
      }
  
      static  void rot(int f, int k, int x, int y, int *rx, int *ry) {
          int COS[2][8] = {
              {c0_1, c1_1, c2_1, c3_1, c4_1, c5_1, c6_1, c7_1},
              {c0_s2, c1_s2, c2_s2, c3_s2, c4_s2, c5_s2, c6_s2, c7_s2}
          };
  
          *rx = SUB(CMUL(COS[f][k], x), CMUL(COS[f][8 - k], y));
          *ry = ADD(CMUL(COS[f][8 - k], x), CMUL(COS[f][k], y));
      }
  
  void idct_1D(__private int *Y);
  
  __kernel void IDCT(__global int *input, __global uchar *output) 
  {
      int Y[64];
      int k, l;
      int Yc[8];
  
      for (k = 0; k &lt; 8; k++)
      {
          for (l = 0; l &lt; 8; l++) Y(k, l) = SCALE(input[(k &lt;&lt; 3) + l], S_BITS);
          idct_1d(&amp;Y(k, 0));
      }
  
      for (l = 0; l &lt; 8; l++)
      {
  
          for (k = 0; k &lt; 8; k++)
      {
              Yc[k] = Y(k, l);
      }
  
          idct_1d(Yc);
  
          for (k = 0; k &lt; 8; k++)
          {
              int r = 128 + DESCALE(Yc[k], S_BITS + 3);
              r = r &gt; 0 ? (r &lt; 255 ? r : 255) : 0;
              X(k, l) = r;
          }
  
      }
  }
  
  void idct_1D(__private int *Y) 
  {
  
  int z1[8], z2[8], z3[8];
  
  
      but(Y[0], Y[4], z1[1], z1[0]);
      rot(1, 6, Y[2], Y[6], &amp;z1[2], &amp;z1[3]);
      but(Y[1], Y[7], z1[4], z1[7]);
      z1[5] = CMUL(sqrt2, Y[3]);
      z1[6] = CMUL(sqrt2, Y[5]);
  
      but(z1[0], z1[3], z2[3], z2[0]);
      but(z1[1], z1[2], z2[2], z2[1]);
      but(z1[4], z1[6], z2[6], z2[4]);
      but(z1[7], z1[5], z2[5], z2[7]);
  
      z3[0] = z2[0];
      z3[1] = z2[1];
      z3[2] = z2[2];
      z3[3] = z2[3];
      rot(0, 3, z2[4], z2[7], &amp;z3[4], &amp;z3[7]);
      rot(0, 1, z2[5], z2[6], &amp;z3[5], &amp;z3[6]);
  
      but(z3[0], z3[7], Y[7], Y[0]);
      but(z3[1], z3[6], Y[6], Y[1]);
      but(z3[2], z3[5], Y[5], Y[2]);
      but(z3[3], z3[4], Y[4], Y[3]);
  }   
  </code></pre>
  
  <p>The error is resulting from the function <code>idct_1D</code> which I am calling from the kernel <code>IDCT</code>.</p>
  
  <p><strong>How can I resolve the error with my kernel function?</strong> </p>
  
  <p><strong>Is there any way to call a function from a kernel?</strong></p>
  
  <p>EDIT:</p>
  
  <p>After following pmdj's answer, I wrote my build statements in the following way:</p>
  
  <pre><code>ret= clBuildProgram(program, 1, &amp;device_id, NULL, NULL, NULL);
  
  if (ret != CL_SUCCESS) 
  {
   printf("building program failed\n");
  
    size_t log_size;
    char buffer[2048];
  clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, 0, NULL, &amp;log_size);
  clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &amp;log_size);
  printf("%s\n", buffer);
  
   }
  </code></pre>
  
  <p>I got the following response after building and running my code:</p>
  
  <pre><code>Number of devices: 2
  building program failed
  ptxas application ptx input, line 71; error   : Call has wrong number of parameters
  ptxas application ptx input, line 112; error   : Call has wrong number of parameters
  ptxas application ptx input, line 153; error   : Call has wrong number of parameters
  ptxas application ptx input, line 194; error   : Call has wrong number of parameters
  ptxas application ptx input, line 235; error   : Call has wrong number of parameters
  ptxas application ptx input, line 276; error   : Call has wrong number of parameters
  ptxas application ptx input, line 317; error   : Call has wrong number of parameters
  ptxas application ptx input, line 358; error   : Call has wrong number of parameters
  ptxas application ptx input, line 392; error   : Call has wrong number of parameters
  ptxas application ptx input, line 520; error   : Call has wrong number of parameters
  ptxas application ptx input, line 648; error   : Call has wrong number of parameters
  ptxas application ptx input, line 776; error   : Call has wrong number of parameters
  ptxas application ptx input, line 904; error   : Call has wrong number of parameters
  ptxas application ptx input, line 1032; error   : Call has wrong number of parameters
  ptxas application ptx input, line 1160; error   : Call has wrong number of parameters
  ptxas application ptx input, line 1288; error   : Call has wrong number of parameters
  ptxas fatal   : Ptx assembly aborted due to errors
  
  -----COULD NOT CREATE KERNEL!!---
  </code></pre>
  
  <p>I checked this <a href="https://github.com/HPCE/hpce-2016-cw6/issues/32" rel="nofollow noreferrer">github link</a> and removed all the comments and printf statements from my <code>.cl</code> file. Despite that, I keep getting the same errors after building and running the code.</p>
  
  <p>EDIT:</p>
  
  <p>The code reflects the change I tried to implement in the code after following pmdj's suggestion. Now I am only getting syntax errors. </p>
  
  <p>The output I am now getting is as follows:</p>
  
  <pre><code>Number of devices: 2
  building program failed
  &lt;kernel&gt;:98:70: error: invalid address space for pointee of pointer argument to __kernel function
  __kernel void IDCT(__global int *input, __global uchar *output, int *Yc, int *Yin) 
                                                                       ^
  &lt;kernel&gt;:98:79: error: invalid address space for pointee of pointer argument to __kernel function
  __kernel void IDCT(__global int *input, __global uchar *output, int *Yc, int *Yin) 
                                                                                ^
  &lt;kernel&gt;:105:30: error: called object type '__attribute__((address_space(16776963))) int *' is not a function or function pointer
                  for (l = 0; l &lt; 8; l++) Yin(k, l) = SCALE(input[(k &lt;&lt; 3) + l], S_BITS);
                                          ~~~^
  &lt;kernel&gt;:106:3: warning: implicit declaration of function 'idct_1d' is invalid in C99
                  idct_1d(&amp;Yin(k, 0));
                  ^
  &lt;kernel&gt;:106:15: error: called object type '__attribute__((address_space(16776963))) int *' is not a function or function pointer
                  idct_1d(&amp;Yin(k, 0));
                           ~~~^
  &lt;kernel&gt;:114:15: error: called object type '__attribute__((address_space(16776963))) int *' is not a function or function pointer
                          Yc[k] = Yin(k, l);
                                  ~~~^
  
  -----COULD NOT CREATE KERNEL!!---
  </code></pre>
  
  <p>EDIT:</p>
  
  <p>I put <code>__private</code> before the function parameters and I am again getting ptx errors.</p>
  
  <p>EDIT:</p>
  
  <p>The code now works. It turns out I was calling the function with the wrong name (<code>idct_1d</code> instead of <code>idct_1D</code>).</p>
  
  <p>The headings of my functions are now as follows:</p>
  
  <pre><code>void idct_1D(int *Y);
  
  
  __kernel void IDCT(__global int *input, __global uchar *output);
  </code></pre>