0

I have an OpenCL kernel code, which does not behave as expected. The similar C code compiled with gcc works fine.

struct data {
    short* a;
};

typedef struct data Data;

inline void foo(Data* d) {
    short b[1] = {99};
    d->a = b;
}

__kernel void bar(__global short* output) {
    Data d;
    foo(&d);
    short val = d.a[0];
    int id = get_global_id(0);
    output[id] = val;
}

Always outputs [0, 0, ..., 0].
If I initialize d.a in __kernel bar and only assign d->a[0] = 99 in foo it works as expected and outputs [99, 99, ..., 99]

Thanks in advance!

UPDATE:
I'm using Java and JOCL for the host code.

As ScottD suggested I've changed d->a = b; in function foo to *d->a = *b;.
And it works great in C version. But causes the following error for OpenCL on MacOS:

Exception in thread "main" org.jocl.CLException:
CL_BUILD_PROGRAM_FAILURE Build log for device 0:
CVMS_ERROR_COMPILER_FAILURE: CVMS compiler has crashed or hung building an element.
at org.jocl.CL.clBuildProgram(CL.java:9368)
...

Or a JVM termination on Windows with AMD CPU:

# A fatal error has been detected by the Java Runtime Environment:
# EXCEPTION_ACCESS_VIOLATION (0xc0000005) at pc=0x000007fedfeb007a, pid=3816, tid=4124
# JRE version: 7.0-b147
# Java VM: Java HotSpot(TM) 64-Bit Server VM (21.0-b17 mixed mode windows-amd64 compressed oops)
# Problematic frame:
# C  [amdocl64.dll+0x60007a]
catdog
  • 87
  • 2
  • 3
  • 9

1 Answers1

2

I believe the problem is this: Function foo sets a pointer used by the caller to the address of a local variable that goes out of scope when foo returns. When the caller accesses that pointer, data in the out of scope variable may or may not still be 99. To demonstrate, make a gcc debug build for this code. It works. Now add a printf(hello\n") after foo(&d) and before val=d.a[0]. Now it fails. This is because the printf call overites the stack memory containing the out of scope 99 value.

Probably you intended:

*d->a = *b; in place of d->a = b;

  • Changing `d->a = b;` to `*d->a = *b;` perfectly works for gcc version. But causes the following error in OpenCL kernel: `Build log for device 0: CVMS_ERROR_COMPILER_FAILURE: CVMS compiler has crashed or hung building an element.` **:(** – catdog May 08 '13 at 19:48
  • Well that is not good. All I know to do is experiment. For example, you could drop the 'inline' and see what happens. You could also try 'd->a[0] = b[0];'. Also make sure short is not a problem. I think older opencl (1.0) does not support short unless enabled as an extension. –  May 09 '13 at 00:46