Referencing length of a primitive array from kernel code causes error "clSetKernelArg (array length) failed invalid arg size" or "clEnqueueNDRangeKernel() failed invalid kernel args"

[GoogleCodeExporter]

Using the pre-built jar+dll, the error is "clSetKernelArg (array length) failed 
invalid arg size". Compiling from trunk results in "clEnqueueNDRangeKernel() 
failed invalid kernel args".

import com.amd.aparapi.Kernel;
import com.amd.aparapi.Range;

public class ArrayLengthBug
{
    public static void main(String... args)
    {
        TestKernel kernel = new TestKernel();
        kernel.execute(1);
        kernel.dispose();
        return;
    }

    public static class TestKernel extends Kernel
    {
        public final int[] buffer = new int[100];

        public void run()
        {
            buffer[0] = buffer[buffer.length - 1];
            return;
        }
    }
}

There was a bug in the old sources that has been in trunk where argPos was 
incremented after setting the length parameter instead of before it. That's 
what causes the first error message.

In the sources from trunk, if the array is more than one-dimensional, it also 
passes an extra javaArrayDimension parameter to calculate offsets. This 
parameter is also used for object arrays. It is not used for primitive arrays, 
but it is present in the generated OpenCL code, and causes an error when the 
kernel is executed because the parameter had not been set. Generated OpenCL 
from trunk code:

typedef struct This_s{
   __global int *buffer;
   int buffer__javaArrayLength0;
   int buffer__javaArrayDimension0;
   int passid;
}This;
int get_pass_id(This *this){
   return this->passid;
}
__kernel void run(
   __global int *buffer,
   int buffer__javaArrayLength0,
   int buffer__javaArrayDimension0,
   int passid
){
   This thisStruct;
   This* this=&thisStruct;
   this->buffer = buffer;
   this->buffer__javaArrayLength0 = buffer__javaArrayLength0;
   this->buffer__javaArrayDimension0 = buffer__javaArrayDimension0;
   this->passid = passid;
   {
      this->buffer[0]  = this->buffer[(this->buffer__javaArrayLength0 - 1)];
      return;
   }
}

The fix I'm using is to add a conditional around the generation of this 
parameter in KernelWriter.java. In the write() method, there is a loop that 
begins:
            for(int i = 0; i < numDimensions; i++) {

The body of the loop writes the length and dimension parameters. The second 
half of the body, starting with the line constructing "dimName" should be 
conditional upon the parameter either being an object or multi-dimensional. I 
added this condition around that block which from what I can tell has fixed the 
issue for me:
               if ((className != null) || (numDimensions > 1)) {

(I also moved the corresponding StringBuilder declarations inside)

The workaround would be handle the array length yourself: set it as a field, 
and don't reference ".length" directly.

Unrelated, but what I was investigating at the time: there's a mistake in 
Aparapi.cpp that due to in-order execution, has no effect:

      if (passid == 0) {
         int writeCount = writeEventCount;
         if(writeEventCount > 0) {
            cl_event* writeEvents = jniContext->writeEvents;
         }
         //... comments omitted ...
      } else {

The variables writeCount and cl_event are being redeclared in the if() block, 
so when the call to clEnqueueNDRangeKernel() is made below, they still have 
their original values. However, the call to clEnqueueNDRangeKernel doesn't 
actually need to wait for the write events because the command queue is not 
created with out-of-order execution (the event count/list should be explicitly 
passed as 0/NULL to clEnqueueNDRangeKernel). Were this changed though, this 
would be a problem (race condition between writes and execution).

Original issue reported on code.google.com by paul.mi...@gmail.com on 13 Jun 2013 at 6:56

[GoogleCodeExporter]

Patches are attached to issue 115 since they overlap.

Original comment by paul.mi...@gmail.com on 13 Jun 2013 at 7:54

[GoogleCodeExporter]

What is the status of this issue ticket?

I am also curious about the investigation of the in-order vs. out-of-order 
execution race condition...

Original comment by pnnl.edg...@gmail.com on 14 Jul 2013 at 10:58

[GoogleCodeExporter]

This is still failing for me on Trunk. Maybe the patch has not been delivered 
yet?

Original comment by jytoumit@gmail.com on 23 Jul 2014 at 8:55