OpenCL

[avinashkrc]

I am trying to work on this code but everytime I gets runtime error "error: field may not be qualified with an address space global uint64_t external;

Here is my data structure in hdr.h:

typedef unsigned long uint64_t;
typedef struct _elem 
{ 
  global uint64_t external;
} elem; 

typedef struct _buf
{ 
  global elem *kbuf; 
}buf;

My OpenCL kernel is:

#include "hdr.h"
__kernel void svmtest(global buf *tx) {
  if(get_global_id(0) == 0)
  printf((__constant char *)" %u %s\n", (tx[0].kbuf->external), "test");

}


My OpenCL C code in brief is as follows:

     parent =(elem*)clSVMAlloc(context, CL_MEM_READ_WRITE|CL_MEM_SVM_FINE_GRAIN_BUFFER , sizeof(elem),0);
     buf *p =(buf*)clSVMAlloc(context, CL_MEM_READ_WRITE|CL_MEM_SVM_FINE_GRAIN_BUFFER , sizeof(buf),0);

     parent->external = 9;

     p[0].kbuf = parent;

     status = clSetKernelArgSVMPointer(kernel, index++, p);
     status= clSetKernelExecInfo(
          kernel,
          CL_KERNEL_EXEC_INFO_SVM_PTRS,
          sizeof(parent),  &parent
      );
However if I change uint64_ external to uint64_t * external in hdr.h and modify the code accordingly then everything works fine. Can't I pass values in structures, is it must to pass a pointer?


[dipak · Accepted solution]

Only variables defined at program scope and static variables inside a function can be declared in the global address space. As the OpenCL spec says:

__global (or global)

--------------------------------

The __global or global address space name is used to refer to memory objects (buffer or image objects) allocated from the global memory pool.

A buffer memory object can be declared as a pointer to a scalar, vector or user-defined struct. This allows the kernel to read and/or write any location in the buffer.

..

Some examples are:

global float4 *color; // An array of float4 elements

typedef struct {

float a[3];

int b[2];

} foo_t;

global foo_t *my_info; // An array of foo_t elements.

...

Variables defined at program scope and static variables inside a function can also be declared in the global address space.

...

For more details, please refer this: The OpenCL™ C 2.0 Specification

Thanks.