AMD Logo AMD Developer Central
AMD Developer Forums
Decrease font size
Increase font size
Topic Title: static arrays in kernels
Topic Summary:
Created On: 11/06/2009 02:36 PM
Status: Post and Reply
Linear : Threading : Single : Branch
Search Topic Search Topic
Topic Tools Topic Tools
View similar topics View similar topics
View topic in raw text format. Print this topic.
 11/06/2009 02:36 PM
User is offline View Users Profile Print this message

Author Icon
twiig

Posts: 18
Joined: 10/21/2009

I have been having a bit of trouble compiling a kernel with a static array of integers.  When compiling the kernel (see code below) on the GPU, I receive the following:

 

For test only: Expires on Sun Feb 28 00:00:00 2010

CL Error: clBuildProgram (-11)

Build Log:

Link failed

Press any key to continue . . .

And, when compiled for the CPU, I receive:
For test only: Expires on Sun Feb 28 00:00:00 2010
C:\Users\tyler\AppData\Local\Temp\OCLB365.tmp.obj:fake.text+0x2c): undefined reference to `__vla_alloc'
C:\Users\tyler\AppData\Local\Temp\OCLB365.tmp.obj:fake.text+0xba): undefined reference to `__vla_dealloc'
C:\Users\tyler\AppData\Local\Temp\OCLB365.tmp.obj:fake.text+0x18e): undefined reference to `__vla_alloc'
C:\Users\tyler\AppData\Local\Temp\OCLB365.tmp.obj:fake.text+0x21d): undefined reference to `__vla_dealloc'
CL Error: clBuildProgram (-11)
Build Log:
Compilation failed
Press any key to continue . . .
Any ideas?  Note that if the only line of code in the kernel is the array declaration, the same result is obtained.

Code:
__kernel void combination(const int n, const int k, __global int *out)
{
   size_t gid = get_global_id(0);
   int a[k];
   int index = gid*100;

   for(int i = 0; i < k; i++)
      a[i] = gid+i;

   while(index != -1)
   {
      for(int i = 0; i < k; i++)
         out[index++] = a[i];

      int j = k - 1;
      if(a[j] < n - 1)
      {
         ++a[j];
         continue;
      }

      while(a[j] - a[j-1] == 1)
         --j;

      if(j == 1)
      {
         index = -1;
         continue;
      }

      int z = ++a[j-1];

      while(j < k)
      {
         a[j] = ++z;
         ++j;
      }
   }
}
 11/06/2009 02:45 PM
User is offline View Users Profile Print this message

Author Icon
MicahVillmow

Posts: 525
Joined: 02/05/2008

twiig,
This is not possible in OpenCL and violates the spec, section 6.8.d. "Variable length arrays and structures with flexible(or unsized) arrays are not supported."

-------------------------
Micah Villmow
Advanced Micro Devices Inc.
--------------------------------
The information presented in this document is for informational purposes only and may contain technical inaccuracies, omissions and typographical errors. Links to third party sites are for convenience only, and no endorsement is implied.

 11/07/2009 01:01 PM
User is offline View Users Profile Print this message

Author Icon
twiig

Posts: 18
Joined: 10/21/2009

Ok, thanks!  I guess I missed that.

Just to follow up, if I wanted some form of addressable ray in which I don't know the size beforehand, I should use a pre-allocated local memory array passed in?

 11/09/2009 01:42 PM
User is offline View Users Profile Print this message

Author Icon
MicahVillmow

Posts: 525
Joined: 02/05/2008

Twiig,
In this case you want to pass it in as an argument to the kernel.

-------------------------
Micah Villmow
Advanced Micro Devices Inc.
--------------------------------
The information presented in this document is for informational purposes only and may contain technical inaccuracies, omissions and typographical errors. Links to third party sites are for convenience only, and no endorsement is implied.

Statistics
6123 users are registered to the AMD Developer Forums forum.
There are currently 0 users logged in.

FuseTalk Hosting Executive Plan v3.2 - © 1999-2009 FuseTalk Inc. All rights reserved.

Contact AMD | Terms and Conditions | Forum Rules | ©2009 Advanced Micro Devices, Inc. | Privacy | Trademark information