2

Can I know OpenCL's function call stack size?

I'm using NVIDIA OpenCL1.2 in Ubuntu. (NVIDIA CC=5.2) And I found some unexpected result in my testcode.

When some function invoked 64 times, the next invoked function seems like can not access the arguments.

In my thought, call stack overflow makes this progblem.

Below is my example code and result:

void testfunc(int count, int run)
{
    if(run==0) return;
    count++;
    printf("count=%d  run=%d\n", count, run);
    run = run - 1;
    testfunc(count, run);
}

__kernel void hello(__global int * in_img, __global int * out_img)
{
    int run;
    int count=0;
    run = 70;
    testfunc(count, run);   
}

And this is the result : 

count=1  run=70
count=2  run=69
count=3  run=68
count=4  run=67
count=5  run=66
count=6  run=65
count=7  run=64
   .....
count=59  run=12
count=60  run=11
count=61  run=10
count=62  run=9
count=63  run=8
count=64  run=7
count=0  run=0  // <--- Why count and run values are ZERO?
count=0  run=0
count=0  run=0
count=0  run=0
count=0  run=0
count=0  run=0
soongk
  • 259
  • 3
  • 17

1 Answers1

6

Recursion is not supported in OpenCL 1.x. From AMD's Introduction to OpenCL:

Key restrictions in the OpenCL C language are:

  • Function pointers are not supported.
  • Bit-fields are not supported.
  • Variable length arrays are not supported.
  • Recursion is not supported.
  • No C99 standard headers such as ctypes.h,errno.h,stdlib.h,etc. can be included.

AFAIK not all implementations have a call-stack like feature at all. In these cases, and possibly in your case, any function calls are inlined in the calling scope.

Aderstedt
  • 6,301
  • 5
  • 28
  • 44
  • Hello Aderstedt. Thanks for answer. Then, isn't there any call stacks in OpenCL? AFAIK, OpenCL's function call uses inlined function and all of context changing uses only registers. If it is right, call stack can not exist in OpenCL. Is it right what I told?? – soongk Oct 22 '15 at 07:40
  • Yes. Some platforms (like CPUs) might use call stacks internally, but you can't rely on having a deep stack. – Aderstedt Oct 22 '15 at 07:59
  • Thank you Aderstedt. Your help so informative for me. – soongk Oct 23 '15 at 01:15