@Local annotations on array fields seem ignored, causing crash at runtime

What steps will reproduce the problem?

1. Run TestLocal class (see attached file TestLocal.java) on APARAPI (r1076 at 
least, tested with KernelWriter class at r1155)

What is the expected output? What do you see instead?

I expect:

[496.0, 496.0, 496.0, 496.0]

(since: 496 = 32 * 31 / 2)

However, on my machine JVM crashes producing a message like this:

#
# A fatal error has been detected by the Java Runtime Environment:
#
#  EXCEPTION_ACCESS_VIOLATION (0xc0000005) at pc=0x0000000077783290, pid=16084, 
tid=11240
#
# JRE version: Java(TM) SE Runtime Environment (7.0_55-b13) (build 1.7.0_55-b13)
# Java VM: Java HotSpot(TM) 64-Bit Server VM (24.55-b03 mixed mode 
windows-amd64 compressed oops)
# Problematic frame:
# C  [ntdll.dll+0x53290]
#
# Failed to write core dump. Minidumps are not enabled by default on client 
versions of Windows
#
# An error report file with more information is saved as:
# D:\work\com.amd.aparapi\hs_err_pid16084.log
#
# If you would like to submit a bug report, please visit:
#   http://bugreport.sun.com/bugreport/crash.jsp
# The crash happened outside the Java Virtual Machine in native code.
# See problematic frame for where to report the bug.
#

Which is no surprise since OpenCL code tags "temp" variable as __global instead 
of __local:

typedef struct This_s{
   __global float *temp;
   __global float *val$input;
   __global float *val$output;
   int passid;
}This;
int get_pass_id(This *this){
   return this->passid;
}
__kernel void run(
   __global float *temp, 
   __global float *val$input, 
   __global float *val$output, 
   int passid
){
   This thisStruct;
   This* this=&thisStruct;
   this->temp = temp;
   this->val$input = val$input;
   this->val$output = val$output;
   this->passid = passid;
   {
      int localId = get_local_id(0);
      this->temp[localId]  = this->val$input[get_global_id(0)];
      barrier(CLK_LOCAL_MEM_FENCE);
      for (int offset = get_local_size(0) / 2; offset>0; offset = offset >> 1){
         if (localId<offset){
            this->temp[localId]  = this->temp[localId] + this->temp[(localId + offset)];
         }
         barrier(CLK_LOCAL_MEM_FENCE);
      }
      if (localId==0){
         this->val$output[get_group_id(0)]  = this->temp[0];
      }
      return;
   }
}

What version of the product are you using? On what operating system?

APARAPI Win64 (Aparapi_2014_04_29_Win64.zip, got from subversion)

Windows 7 Professional

GPU environment is a Nvidia Quadro FX880M with 1 GiB, driver: ForceWare: 327.02 

Please provide any additional information below.

I suspect that the problem is in KernelWriter's imports; those for @Local and 
@Constant (used implicitly in constants LOCAL_ANNOTATION_NAME and 
CONSTANT_ANNOTATION_NAME, respectively) point to @OpenCL.Local and 
@OpenCL.Constant instead of 
@com.amd.aparapi.Kernel.Local and @com.amd.aparapi.Kernel.Constant, so 
KernelWriter.write(Entrypoint) searches for different XXX_ANNOTATION_NAME than 
those actually present in code. Moreover, only parameters (not fields) can be 
annotated with @OpenCL.Local and @OpenCL.Constant, so one should not be able to 
find them on any field.

In fact, changing imports in KernelWriter (see file KernelWriter.java.patch) 
solves this issue, allowing attached code to complete.
Original issue reported on code.google.com by lgallucci@gmail.com on 29 May 2014 at 4:23
Attachments:
rzel / aparapi

@Local annotations on array fields seem ignored, causing crash at runtime #145
