From 2c0e170dc45daa575c238385b9e187b6f4b4d8bf Mon Sep 17 00:00:00 2001
From: Dmitriy Shabanov <shabanovd@gmail.com>
Date: Wed, 18 Oct 2017 14:19:31 +0300
Subject: [PATCH] support inline array

---
 CHANGELOG.md                                  |  3 +-
 CONTRIBUTORS.md                               |  2 +
 pom.xml                                       |  2 +-
 src/main/java/com/aparapi/Config.java         |  2 +
 .../exception/ClassParseException.java        |  2 +
 .../aparapi/internal/model/MethodModel.java   | 19 +++--
 .../aparapi/internal/writer/BlockWriter.java  | 64 +++++++++-----
 .../aparapi/internal/writer/KernelWriter.java | 49 +++++++----
 .../aparapi/codegen/test/ArrayCreation.java   | 84 +++++++++++++++++++
 .../codegen/test/ArrayCreationTest.java       | 76 +++++++++++++++++
 .../codegen/test/ObjectRefCopyTest.java       | 27 +++++-
 .../codegen/test/ReturnBooleanNewArray.java   |  2 +-
 .../test/ReturnBooleanNewArrayTest.java       |  4 +-
 .../codegen/test/ReturnBooleanVarArray.java   | 25 +++++-
 .../test/ReturnBooleanVarArrayTest.java       | 27 +++++-
 .../codegen/test/ReturnByteArrayNew.java      |  2 +-
 .../codegen/test/ReturnByteArrayNewTest.java  |  4 +-
 .../codegen/test/ReturnByteArrayVar.java      | 25 +++++-
 .../codegen/test/ReturnByteArrayVarTest.java  | 27 +++++-
 .../codegen/test/ReturnDoubleArrayNew.java    |  2 +-
 .../test/ReturnDoubleArrayNewTest.java        |  4 +-
 .../codegen/test/ReturnDoubleArrayVar.java    | 25 +++++-
 .../test/ReturnDoubleArrayVarTest.java        | 27 +++++-
 .../codegen/test/ReturnFloatArrayNew.java     |  2 +-
 .../codegen/test/ReturnFloatArrayNewTest.java |  4 +-
 .../codegen/test/ReturnFloatArrayVar.java     | 25 +++++-
 .../codegen/test/ReturnFloatArrayVarTest.java | 27 +++++-
 .../codegen/test/ReturnIntArrayNew.java       |  2 +-
 .../codegen/test/ReturnIntArrayNewTest.java   |  4 +-
 .../codegen/test/ReturnIntArrayVar.java       | 25 +++++-
 .../codegen/test/ReturnIntArrayVarTest.java   | 27 +++++-
 .../codegen/test/ReturnLongArrayNew.java      |  2 +-
 .../codegen/test/ReturnLongArrayNewTest.java  |  4 +-
 .../codegen/test/ReturnLongArrayVar.java      |  4 +-
 .../codegen/test/ReturnLongArrayVarTest.java  | 27 +++++-
 .../codegen/test/ReturnShortArrayNew.java     |  2 +-
 .../codegen/test/ReturnShortArrayNewTest.java |  4 +-
 .../codegen/test/ReturnShortArrayVar.java     | 24 +++++-
 .../codegen/test/ReturnShortArrayVarTest.java | 28 ++++++-
 .../java/com/aparapi/runtime/ArrayTest.java   | 82 ++++++++++++++++++
 40 files changed, 700 insertions(+), 97 deletions(-)
 create mode 100644 src/test/java/com/aparapi/codegen/test/ArrayCreation.java
 create mode 100644 src/test/java/com/aparapi/codegen/test/ArrayCreationTest.java
 create mode 100644 src/test/java/com/aparapi/runtime/ArrayTest.java

diff --git a/CHANGELOG.md b/CHANGELOG.md
index b2cf89f5..bd3318fb 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -1,7 +1,8 @@
 # Aparapi Changelog
 
-## 1.4.2
+## 1.5.0
 
+* Support inline array creation in the kernel, which is implemented in the GPU in private memory.
 * Updated parent pon to v6.
 * createProgram had the wrong signature producing a unsatisfied link exception that is now fixed.
 * Build now requires version 3.5.0 of maven due to changes in surefire plugin.
diff --git a/CONTRIBUTORS.md b/CONTRIBUTORS.md
index 4fee1845..42337edf 100644
--- a/CONTRIBUTORS.md
+++ b/CONTRIBUTORS.md
@@ -19,6 +19,7 @@
 * Lorenzo Gallucci
 * Fernando Marino`
 * AMD Corporation
+* Dmitriy Shabanov <shabanovd@gmail.com>
 
 # Details
 
@@ -37,3 +38,4 @@ Below are some of the specific details of various contributions.
 * Paul Miner issue #61 and #115 (JTP Speed up and fixes to explicit puts) June 13th 2013
 * lgalluci for his fix for issue #121 (incorrect toString for 3D ranges) July 6th 2013
 * George Vinokhodov submited a fix for a bug regarding forward references.
+* Dmitriy Shabanov submited PR for inline array feature.
diff --git a/pom.xml b/pom.xml
index 8b002640..df61c926 100644
--- a/pom.xml
+++ b/pom.xml
@@ -11,7 +11,7 @@
 
     <groupId>com.aparapi</groupId>
     <artifactId>aparapi</artifactId>
-    <version>1.4.2-SNAPSHOT</version>
+    <version>1.5.0-SNAPSHOT</version>
     <packaging>jar</packaging>
 
     <prerequisites>
diff --git a/src/main/java/com/aparapi/Config.java b/src/main/java/com/aparapi/Config.java
index 06f369a8..0986dcfd 100644
--- a/src/main/java/com/aparapi/Config.java
+++ b/src/main/java/com/aparapi/Config.java
@@ -157,6 +157,8 @@ public class Config extends ConfigJNI{
 
    public static final boolean enableMONITOR = Boolean.getBoolean(propPkgName + ".enable.MONITOR");
 
+   public static final boolean enableARRAY = !Boolean.getBoolean(propPkgName + ".disable.ARRAY");
+
    public static final boolean enableNEW = Boolean.getBoolean(propPkgName + ".enable.NEW");
 
    public static final boolean enableATHROW = Boolean.getBoolean(propPkgName + ".enable.ATHROW");
diff --git a/src/main/java/com/aparapi/internal/exception/ClassParseException.java b/src/main/java/com/aparapi/internal/exception/ClassParseException.java
index fcdac3f9..39aac1fa 100644
--- a/src/main/java/com/aparapi/internal/exception/ClassParseException.java
+++ b/src/main/java/com/aparapi/internal/exception/ClassParseException.java
@@ -74,6 +74,8 @@ import com.aparapi.internal.instruction.Instruction;
       ATHROW("We don't support athrow instructions"), //
       SYNCHRONIZE("We don't support monitorenter or monitorexit instructions"), //
       NEW("We don't support new instructions"), //
+      NEWARRAY("We don't support new array instructions"), //
+      NEWMULTIARRAY("We don't support new multi array instructions"), //
       ARRAYALIAS("We don't support copying refs in kernels"), //
       SWITCH("We don't support lookupswitch or tableswitch instructions"), //
       METHODARRAYARG("We don't support passing arrays as method args"), //
diff --git a/src/main/java/com/aparapi/internal/model/MethodModel.java b/src/main/java/com/aparapi/internal/model/MethodModel.java
index 4a16a868..1766a13c 100644
--- a/src/main/java/com/aparapi/internal/model/MethodModel.java
+++ b/src/main/java/com/aparapi/internal/model/MethodModel.java
@@ -54,6 +54,7 @@ package com.aparapi.internal.model;
 
 import com.aparapi.*;
 import com.aparapi.internal.exception.*;
+import com.aparapi.internal.exception.ClassParseException.TYPE;
 import com.aparapi.internal.instruction.*;
 import com.aparapi.internal.instruction.InstructionPattern.*;
 import com.aparapi.internal.instruction.InstructionSet.*;
@@ -243,12 +244,16 @@ public class MethodModel{
             throw new ClassParseException(instruction, ClassParseException.TYPE.SYNCHRONIZE);
          }
 
-         if ((!Config.enableNEW) && (instruction instanceof New)) {
-            throw new ClassParseException(instruction, ClassParseException.TYPE.NEW);
-         }
-
-         if (instruction instanceof I_AASTORE) {
-            throw new ClassParseException(instruction, ClassParseException.TYPE.ARRAYALIAS);
+         if (instruction instanceof New) {
+             if ((instruction instanceof I_NEWARRAY)) {
+                 if (!Config.enableARRAY) {
+                     throw new ClassParseException(instruction, TYPE.NEWARRAY);
+                 }
+             } else if ((instruction instanceof I_MULTIANEWARRAY)) {
+                 throw new ClassParseException(instruction, ClassParseException.TYPE.NEWMULTIARRAY);
+             } else if (!Config.enableNEW) {
+                 throw new ClassParseException(instruction, ClassParseException.TYPE.NEW);
+             }
          }
 
          if ((!Config.enableSWITCH) && ((instruction instanceof I_LOOKUPSWITCH) || (instruction instanceof I_TABLESWITCH))) {
@@ -444,7 +449,7 @@ public class MethodModel{
          }
 
          _expressionList.add(new CloneInstruction(this, e));
-         System.out.println("clone of " + e);
+         //System.out.println("clone of " + e);
       } else if (_instruction instanceof I_DUP2) {
          Instruction e = _expressionList.getTail();
          while (!e.producesStack()) {
diff --git a/src/main/java/com/aparapi/internal/writer/BlockWriter.java b/src/main/java/com/aparapi/internal/writer/BlockWriter.java
index b43fa98a..0712ffc4 100644
--- a/src/main/java/com/aparapi/internal/writer/BlockWriter.java
+++ b/src/main/java/com/aparapi/internal/writer/BlockWriter.java
@@ -366,14 +366,14 @@ public abstract class BlockWriter{
       }
    }
 
-   public String convertType(String _typeDesc, boolean useClassModel) {
+   public String convertType(String _typeDesc, boolean useClassModel, boolean isLocal) {
       return (_typeDesc);
    }
 
    public String convertCast(String _cast) {
       // Strip parens off cast
       //System.out.println("cast = " + _cast);
-      final String raw = convertType(_cast.substring(1, _cast.length() - 1), false);
+      final String raw = convertType(_cast.substring(1, _cast.length() - 1), false, false);
       return ("(" + raw + ")");
    }
 
@@ -393,24 +393,46 @@ public abstract class BlockWriter{
              write(")");
          }
       } else if (_instruction instanceof CompositeInstruction) {
-         writeComposite((CompositeInstruction) _instruction);
+          writeComposite((CompositeInstruction) _instruction);
+
+      } else if (_instruction instanceof I_NEWARRAY) {
+          if (_instruction.getParentExpr() instanceof Return) {
+              throw new CodeGenException("'newarray' is not allowed after 'return'");
+          }
+
+          for (Instruction operand = _instruction.getFirstChild(); operand != null; operand = operand.getNextExpr()) {
+              write("[");
+              writeInstruction(operand);
+              write("]");
+          }
 
       } else if (_instruction instanceof AssignToLocalVariable) {
          final AssignToLocalVariable assignToLocalVariable = (AssignToLocalVariable) _instruction;
 
          final LocalVariableInfo localVariableInfo = assignToLocalVariable.getLocalVariableInfo();
-         if (assignToLocalVariable.isDeclaration()) {
-            final String descriptor = localVariableInfo.getVariableDescriptor();
-            // Arrays always map to __global arrays
-            if (descriptor.startsWith("[")) {
-               write(" __global ");
-            }
-            write(convertType(descriptor, true));
-         }
-         if (localVariableInfo == null) {
-            throw new CodeGenException("outOfScope" + _instruction.getThisPC() + " = ");
+
+         if ((_instruction.getFirstChild() instanceof I_NEWARRAY)) {
+             if (localVariableInfo == null) {
+                 throw new CodeGenException("outOfScope" + _instruction.getThisPC() + " = ");
+             } else {
+                 final String descriptor = localVariableInfo.getVariableDescriptor();
+                 write(convertType(descriptor, true, true));
+                 write(localVariableInfo.getVariableName());
+             }
          } else {
-            write(localVariableInfo.getVariableName() + " = ");
+             if (assignToLocalVariable.isDeclaration()) {
+                 final String descriptor = localVariableInfo.getVariableDescriptor();
+                 // Arrays always map to __global arrays
+                 if (descriptor.startsWith("[")) {
+                     write(" __global ");
+                 }
+                 write(convertType(descriptor, true, false));
+             }
+             if (localVariableInfo == null) {
+                 throw new CodeGenException("outOfScope" + _instruction.getThisPC() + " = ");
+             } else {
+                 write(localVariableInfo.getVariableName() + " = ");
+             }
          }
 
          for (Instruction operand = _instruction.getFirstChild(); operand != null; operand = operand.getNextExpr()) {
@@ -457,7 +479,7 @@ public abstract class BlockWriter{
                dim++;
             }
 
-            NameAndTypeEntry nameAndTypeEntry = ((AccessInstanceField) load).getConstantPoolFieldEntry().getNameAndTypeEntry();
+            NameAndTypeEntry nameAndTypeEntry = ((AccessField) load).getConstantPoolFieldEntry().getNameAndTypeEntry();
             if (isMultiDimensionalArray(nameAndTypeEntry)) {
                String arrayName = nameAndTypeEntry.getNameUTF8Entry().getUTF8();
                write(" * this->" + arrayName + arrayDimMangleSuffix + dim);
@@ -661,7 +683,7 @@ public abstract class BlockWriter{
 
             final LocalVariableInfo localVariableInfo = alv.getLocalVariableInfo();
             if (alv.isDeclaration()) {
-               write(convertType(localVariableInfo.getVariableDescriptor(), true));
+               write(convertType(localVariableInfo.getVariableDescriptor(), true, false));
             }
             if (localVariableInfo == null) {
                throw new CodeGenException("outOfScope" + _instruction.getThisPC() + " = ");
@@ -678,7 +700,7 @@ public abstract class BlockWriter{
          final LocalVariableInfo localVariableInfo = assignToLocalVariable.getLocalVariableInfo();
          if (assignToLocalVariable.isDeclaration()) {
             // this is bad! we need a general way to hoist up a required declaration
-            throw new CodeGenException("/* we can't declare this " + convertType(localVariableInfo.getVariableDescriptor(), true)
+            throw new CodeGenException("/* we can't declare this " + convertType(localVariableInfo.getVariableDescriptor(), true, false)
                   + " here */");
          }
          write(localVariableInfo.getVariableName());
@@ -766,22 +788,22 @@ public abstract class BlockWriter{
    }
 
    private boolean isMultiDimensionalArray(final AccessArrayElement arrayLoadInstruction) {
-      AccessInstanceField accessInstanceField = getUltimateInstanceFieldAccess(arrayLoadInstruction);
+       AccessField accessInstanceField = getUltimateInstanceFieldAccess(arrayLoadInstruction);
       return isMultiDimensionalArray(accessInstanceField.getConstantPoolFieldEntry().getNameAndTypeEntry());
    }
 
    private boolean isObjectArray(final AccessArrayElement arrayLoadInstruction) {
-      AccessInstanceField accessInstanceField = getUltimateInstanceFieldAccess(arrayLoadInstruction);
+       AccessField accessInstanceField = getUltimateInstanceFieldAccess(arrayLoadInstruction);
       return isObjectArray(accessInstanceField.getConstantPoolFieldEntry().getNameAndTypeEntry());
    }
 
-   private AccessInstanceField getUltimateInstanceFieldAccess(final AccessArrayElement arrayLoadInstruction) {
+   private AccessField getUltimateInstanceFieldAccess(final AccessArrayElement arrayLoadInstruction) {
       Instruction load = arrayLoadInstruction.getArrayRef();
       while (load instanceof I_AALOAD) {
          load = load.getFirstChild();
       }
 
-      return (AccessInstanceField) load;
+      return (AccessField) load;
    }
 
    public void writeMethod(MethodCall _methodCall, MethodEntry _methodEntry) throws CodeGenException {
diff --git a/src/main/java/com/aparapi/internal/writer/KernelWriter.java b/src/main/java/com/aparapi/internal/writer/KernelWriter.java
index 4e5de609..8099c591 100644
--- a/src/main/java/com/aparapi/internal/writer/KernelWriter.java
+++ b/src/main/java/com/aparapi/internal/writer/KernelWriter.java
@@ -52,6 +52,8 @@ under those regulations, please refer to the U.S. Bureau of Industry and Securit
  */
 package com.aparapi.internal.writer;
 
+import static com.aparapi.Kernel.PRIVATE_SUFFIX;
+
 import com.aparapi.*;
 import com.aparapi.internal.exception.*;
 import com.aparapi.internal.instruction.*;
@@ -70,24 +72,40 @@ public abstract class KernelWriter extends BlockWriter{
 
    private final String cvtBooleanArrayToCharStar = "char* ";
 
+   private final String cvtBooleanArrayToChar = "char ";
+
    private final String cvtByteToChar = "char ";
 
    private final String cvtByteArrayToCharStar = "char* ";
 
+   private final String cvtByteArrayToChar = "char ";
+
    private final String cvtCharToShort = "unsigned short ";
 
    private final String cvtCharArrayToShortStar = "unsigned short* ";
 
+   private final String cvtCharArrayToShort = "unsigned short ";
+
    private final String cvtIntArrayToIntStar = "int* ";
 
+   private final String cvtIntArrayToInt = "int ";
+
    private final String cvtFloatArrayToFloatStar = "float* ";
 
+   private final String cvtFloatArrayToFloat = "float ";
+
    private final String cvtDoubleArrayToDoubleStar = "double* ";
 
+   private final String cvtDoubleArrayToDouble = "double ";
+
    private final String cvtLongArrayToLongStar = "long* ";
 
+   private final String cvtLongArrayToLong = "long ";
+
    private final String cvtShortArrayToShortStar = "short* ";
 
+   private final String cvtShortArrayToShort = "short ";
+
    /** When declaring a __private struct pointer field, we always omit the "__private" qualifier. This is because the NVidia OpenCL compiler, at time of writing
     * erroneously complains about explicitly qualifying pointers with __private ("error: field may not be qualified with an address space").
     */
@@ -151,29 +169,29 @@ public abstract class KernelWriter extends BlockWriter{
     *          String in the Java JNI notation, [I, etc
     * @return Suitably converted string, "char*", etc
     */
-   @Override public String convertType(String _typeDesc, boolean useClassModel) {
+   @Override public String convertType(String _typeDesc, boolean useClassModel, boolean isLocal) {
       if (_typeDesc.equals("Z") || _typeDesc.equals("boolean")) {
          return (cvtBooleanToChar);
       } else if (_typeDesc.equals("[Z") || _typeDesc.equals("boolean[]")) {
-         return (cvtBooleanArrayToCharStar);
+         return isLocal ? (cvtBooleanArrayToChar) : (cvtBooleanArrayToCharStar);
       } else if (_typeDesc.equals("B") || _typeDesc.equals("byte")) {
          return (cvtByteToChar);
       } else if (_typeDesc.equals("[B") || _typeDesc.equals("byte[]")) {
-         return (cvtByteArrayToCharStar);
+         return isLocal ? (cvtByteArrayToChar) : (cvtByteArrayToCharStar);
       } else if (_typeDesc.equals("C") || _typeDesc.equals("char")) {
          return (cvtCharToShort);
       } else if (_typeDesc.equals("[C") || _typeDesc.equals("char[]")) {
-         return (cvtCharArrayToShortStar);
+         return isLocal ? (cvtCharArrayToShort) : (cvtCharArrayToShortStar);
       } else if (_typeDesc.equals("[I") || _typeDesc.equals("int[]")) {
-         return (cvtIntArrayToIntStar);
+         return isLocal ? (cvtIntArrayToInt) : (cvtIntArrayToIntStar);
       } else if (_typeDesc.equals("[F") || _typeDesc.equals("float[]")) {
-         return (cvtFloatArrayToFloatStar);
+         return isLocal ? (cvtFloatArrayToFloat) : (cvtFloatArrayToFloatStar);
       } else if (_typeDesc.equals("[D") || _typeDesc.equals("double[]")) {
-         return (cvtDoubleArrayToDoubleStar);
+         return isLocal ? (cvtDoubleArrayToDouble) : (cvtDoubleArrayToDoubleStar);
       } else if (_typeDesc.equals("[J") || _typeDesc.equals("long[]")) {
-         return (cvtLongArrayToLongStar);
+         return isLocal ? (cvtLongArrayToLong) : (cvtLongArrayToLongStar);
       } else if (_typeDesc.equals("[S") || _typeDesc.equals("short[]")) {
-         return (cvtShortArrayToShortStar);
+         return isLocal ? (cvtShortArrayToShort) : (cvtShortArrayToShortStar);
       }
       // if we get this far, we haven't matched anything yet
       if (useClassModel) {
@@ -373,8 +391,8 @@ public abstract class KernelWriter extends BlockWriter{
             argLine.append(className);
             thisStructLine.append(className);
          } else {
-            argLine.append(convertType(ClassModel.typeName(signature.charAt(0)), false));
-            thisStructLine.append(convertType(ClassModel.typeName(signature.charAt(0)), false));
+            argLine.append(convertType(ClassModel.typeName(signature.charAt(0)), false, false));
+            thisStructLine.append(convertType(ClassModel.typeName(signature.charAt(0)), false, false));
          }
 
          argLine.append(" ");
@@ -520,7 +538,7 @@ public abstract class KernelWriter extends BlockWriter{
                }
                totalSize += fSize;
 
-               final String cType = convertType(field.getNameAndTypeEntry().getDescriptorUTF8Entry().getUTF8(), true);
+               final String cType = convertType(field.getNameAndTypeEntry().getDescriptorUTF8Entry().getUTF8(), true, false);
                assert cType != null : "could not find type for " + field.getNameAndTypeEntry().getDescriptorUTF8Entry().getUTF8();
                writeln(cType + " " + field.getNameAndTypeEntry().getNameUTF8Entry().getUTF8() + ";");
             }
@@ -585,7 +603,7 @@ public abstract class KernelWriter extends BlockWriter{
          if (returnType.startsWith("[")) {
             write(" __global ");
          }
-         write(convertType(returnType, true));
+         write(convertType(returnType, true, false));
 
          write(mm.getName() + "(");
 
@@ -618,12 +636,11 @@ public abstract class KernelWriter extends BlockWriter{
                   write(", ");
                }
 
-               // Arrays always map to __global arrays
-               if (descriptor.startsWith("[")) {
+               if (descriptor.startsWith("[") && !lvi.getVariableName().endsWith(PRIVATE_SUFFIX)) {
                   write(" __global ");
                }
 
-               write(convertType(descriptor, true));
+               write(convertType(descriptor, true, false));
                write(lvi.getVariableName());
                alreadyHasFirstArg = true;
             }
diff --git a/src/test/java/com/aparapi/codegen/test/ArrayCreation.java b/src/test/java/com/aparapi/codegen/test/ArrayCreation.java
new file mode 100644
index 00000000..4d47d48c
--- /dev/null
+++ b/src/test/java/com/aparapi/codegen/test/ArrayCreation.java
@@ -0,0 +1,84 @@
+/**
+ * Copyright (c) 2016 - 2017 Syncleus, Inc.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+package com.aparapi.codegen.test;
+
+public class ArrayCreation {
+
+    public float[][] results = new float[128][2];
+    int y = 2;
+
+    float[] method(float[] a) {
+        a[0] = a[0] + 1;
+        return a;
+    }
+
+    public void run() {
+        float[] a = new float[2];
+
+        float[] b = new float[16];
+
+        //float[][] c = new float[16][16];
+
+        float[] d = method(a);
+
+        results[y] = b;
+    }
+}
+/**{OpenCL{
+ typedef struct This_s{
+ __global float *results;
+ int results__javaArrayLength0;
+ int results__javaArrayDimension0;
+ int results__javaArrayLength1;
+ int results__javaArrayDimension1;
+ int y;
+ int passid;
+ }This;
+ int get_pass_id(This *this){
+ return this->passid;
+ }
+ __global float* com_aparapi_codegen_test_VectorCreation__method(This *this,  __global float* a){
+ a[0]  = a[0] + 1.0f;
+ return(a);
+ }
+ __kernel void run(
+ __global float *results,
+ int results__javaArrayLength0,
+ int results__javaArrayDimension0,
+ int results__javaArrayLength1,
+ int results__javaArrayDimension1,
+ int y,
+ int passid
+ ){
+ This thisStruct;
+ This* this=&thisStruct;
+ this->results = results;
+ this->results__javaArrayLength0 = results__javaArrayLength0;
+ this->results__javaArrayDimension0 = results__javaArrayDimension0;
+ this->results__javaArrayLength1 = results__javaArrayLength1;
+ this->results__javaArrayDimension1 = results__javaArrayDimension1;
+ this->y = y;
+ this->passid = passid;
+ {
+ float a[2];
+ float b[16];
+ __global float* d = com_aparapi_codegen_test_VectorCreation__method(this, a);
+ this->results[this->y]  = b;
+ return;
+ }
+ }
+
+ }OpenCL}**/
diff --git a/src/test/java/com/aparapi/codegen/test/ArrayCreationTest.java b/src/test/java/com/aparapi/codegen/test/ArrayCreationTest.java
new file mode 100644
index 00000000..56de08ff
--- /dev/null
+++ b/src/test/java/com/aparapi/codegen/test/ArrayCreationTest.java
@@ -0,0 +1,76 @@
+/**
+ * Copyright (c) 2016 - 2017 Syncleus, Inc.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+package com.aparapi.codegen.test;
+
+import org.junit.Test;
+
+public class ArrayCreationTest extends com.aparapi.codegen.CodeGenJUnitBase {
+    private static final String[] expectedOpenCL = {
+        "typedef struct This_s{\n"
+            + "   __global float *results;\n"
+            + "   int results__javaArrayLength0;\n"
+            + "   int results__javaArrayDimension0;\n"
+            + "   int results__javaArrayLength1;\n"
+            + "   int results__javaArrayDimension1;\n"
+            + "   int y;\n"
+            + "   int passid;\n"
+            + "}This;\n"
+            + "int get_pass_id(This *this){\n"
+            + "   return this->passid;\n"
+            + "}\n"
+            + " __global float* com_aparapi_codegen_test_ArrayCreation__method(This *this,  __global float* a){\n"
+            + "   a[0]  = a[0] + 1.0f;\n"
+            + "   return(a);\n"
+            + "}\n"
+            + "__kernel void run(\n"
+            + "   __global float *results, \n"
+            + "   int results__javaArrayLength0, \n"
+            + "   int results__javaArrayDimension0, \n"
+            + "   int results__javaArrayLength1, \n"
+            + "   int results__javaArrayDimension1, \n"
+            + "   int y, \n"
+            + "   int passid\n"
+            + "){\n"
+            + "   This thisStruct;\n"
+            + "   This* this=&thisStruct;\n"
+            + "   this->results = results;\n"
+            + "   this->results__javaArrayLength0 = results__javaArrayLength0;\n"
+            + "   this->results__javaArrayDimension0 = results__javaArrayDimension0;\n"
+            + "   this->results__javaArrayLength1 = results__javaArrayLength1;\n"
+            + "   this->results__javaArrayDimension1 = results__javaArrayDimension1;\n"
+            + "   this->y = y;\n"
+            + "   this->passid = passid;\n"
+            + "   {\n"
+            + "      float a[2];\n"
+            + "      float b[16];\n"
+            + "       __global float* d = com_aparapi_codegen_test_ArrayCreation__method(this, a);\n"
+            + "      this->results[this->y]  = b;\n"
+            + "      return;\n"
+            + "   }\n"
+            + "}\n"
+    };
+    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = null;
+
+    @Test
+    public void ArrayCreationTest() {
+        test(ArrayCreation.class, expectedException, expectedOpenCL);
+    }
+
+    @Test
+    public void ArrayCreationTestWorksWithCaching() {
+        test(ArrayCreation.class, expectedException, expectedOpenCL);
+    }
+}
diff --git a/src/test/java/com/aparapi/codegen/test/ObjectRefCopyTest.java b/src/test/java/com/aparapi/codegen/test/ObjectRefCopyTest.java
index e3978c7d..a5cff184 100644
--- a/src/test/java/com/aparapi/codegen/test/ObjectRefCopyTest.java
+++ b/src/test/java/com/aparapi/codegen/test/ObjectRefCopyTest.java
@@ -16,11 +16,34 @@
 package com.aparapi.codegen.test;
 
 import com.aparapi.internal.exception.ClassParseException;
+import org.junit.Ignore;
 import org.junit.Test;
 
 public class ObjectRefCopyTest extends com.aparapi.codegen.CodeGenJUnitBase {
-    private static final String[] expectedOpenCL = null;
-    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = ClassParseException.class;
+    private static final String[] expectedOpenCL = {
+        "typedef struct This_s{\n"
+            + "   __global com_aparapi_codegen_test_ObjectRefCopy$DummyOOA *dummy;\n"
+            + "   int passid;\n"
+            + "}This;\n"
+            + "int get_pass_id(This *this){\n"
+            + "   return this->passid;\n"
+            + "}\n"
+            + "__kernel void run(\n"
+            + "   __global com_aparapi_codegen_test_ObjectRefCopy$DummyOOA *dummy, \n"
+            + "   int passid\n"
+            + "){\n"
+            + "   This thisStruct;\n"
+            + "   This* this=&thisStruct;\n"
+            + "   this->dummy = dummy;\n"
+            + "   this->passid = passid;\n"
+            + "   {\n"
+            + "      int myId = get_global_id(0);\n"
+            + "      this->dummy[myId]  = this->dummy[(myId + 1)];\n"
+            + "      return;\n"
+            + "   }\n"
+            + "}"
+    };
+    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = null;
 
     @Test
     public void ObjectRefCopyTest() {
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnBooleanNewArray.java b/src/test/java/com/aparapi/codegen/test/ReturnBooleanNewArray.java
index 6271a596..1d7facb4 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnBooleanNewArray.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnBooleanNewArray.java
@@ -26,4 +26,4 @@ public class ReturnBooleanNewArray {
         returnBooleanNewArray();
     }
 }
-/**{Throws{ClassParseException}Throws}**/
+/**{Throws{CodeGenException}Throws}**/
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnBooleanNewArrayTest.java b/src/test/java/com/aparapi/codegen/test/ReturnBooleanNewArrayTest.java
index e37330fe..c750ae9f 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnBooleanNewArrayTest.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnBooleanNewArrayTest.java
@@ -15,13 +15,13 @@
  */
 package com.aparapi.codegen.test;
 
-import com.aparapi.internal.exception.ClassParseException;
+import com.aparapi.internal.exception.CodeGenException;
 import org.junit.Test;
 
 public class ReturnBooleanNewArrayTest extends com.aparapi.codegen.CodeGenJUnitBase {
 
     private static final String[] expectedOpenCL = null;
-    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = ClassParseException.class;
+    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = CodeGenException.class;
 
     @Test
     public void ReturnBooleanNewArrayTest() {
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnBooleanVarArray.java b/src/test/java/com/aparapi/codegen/test/ReturnBooleanVarArray.java
index 757a8aa7..85f00fcd 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnBooleanVarArray.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnBooleanVarArray.java
@@ -28,4 +28,27 @@ public class ReturnBooleanVarArray {
         returnBooleanVarArray();
     }
 }
-/**{Throws{ClassParseException}Throws}**/
+/**{OpenCL{
+ typedef struct This_s{
+ int passid;
+ }This;
+ int get_pass_id(This *this){
+ return this->passid;
+ }
+ __global char* com_aparapi_codegen_test_ReturnBooleanVarArray__returnBooleanVarArray(This *this){
+ char ba[1024];
+ return(ba);
+ }
+ __kernel void run(
+ int passid
+ ){
+ This thisStruct;
+ This* this=&thisStruct;
+ this->passid = passid;
+ {
+ com_aparapi_codegen_test_ReturnBooleanVarArray__returnBooleanVarArray(this);
+ return;
+ }
+ }
+
+ }OpenCL}**/
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnBooleanVarArrayTest.java b/src/test/java/com/aparapi/codegen/test/ReturnBooleanVarArrayTest.java
index f1439e39..46f45c4b 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnBooleanVarArrayTest.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnBooleanVarArrayTest.java
@@ -15,13 +15,34 @@
  */
 package com.aparapi.codegen.test;
 
-import com.aparapi.internal.exception.ClassParseException;
 import org.junit.Test;
 
 public class ReturnBooleanVarArrayTest extends com.aparapi.codegen.CodeGenJUnitBase {
 
-    private static final String[] expectedOpenCL = null;
-    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = ClassParseException.class;
+    private static final String[] expectedOpenCL = {
+        "typedef struct This_s{\n"
+            + "   int passid;\n"
+            + "}This;\n"
+            + "int get_pass_id(This *this){\n"
+            + "   return this->passid;\n"
+            + "}\n"
+            + " __global char* com_aparapi_codegen_test_ReturnBooleanVarArray__returnBooleanVarArray(This *this){\n"
+            + "   char ba[1024];\n"
+            + "   return(ba);\n"
+            + "}\n"
+            + "__kernel void run(\n"
+            + "   int passid\n"
+            + "){\n"
+            + "   This thisStruct;\n"
+            + "   This* this=&thisStruct;\n"
+            + "   this->passid = passid;\n"
+            + "   {\n"
+            + "      com_aparapi_codegen_test_ReturnBooleanVarArray__returnBooleanVarArray(this);\n"
+            + "      return;\n"
+            + "   }\n"
+            + "}"
+    };
+    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = null;
 
     @Test
     public void ReturnBooleanVarArrayTest() {
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnByteArrayNew.java b/src/test/java/com/aparapi/codegen/test/ReturnByteArrayNew.java
index fe08bdc9..e47b8d12 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnByteArrayNew.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnByteArrayNew.java
@@ -25,4 +25,4 @@ public class ReturnByteArrayNew {
         returnByteArrayNew();
     }
 }
-/**{Throws{ClassParseException}Throws}**/
+/**{Throws{CodeGenException}Throws}**/
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnByteArrayNewTest.java b/src/test/java/com/aparapi/codegen/test/ReturnByteArrayNewTest.java
index a0134a0f..d322a6b4 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnByteArrayNewTest.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnByteArrayNewTest.java
@@ -15,12 +15,12 @@
  */
 package com.aparapi.codegen.test;
 
-import com.aparapi.internal.exception.ClassParseException;
+import com.aparapi.internal.exception.CodeGenException;
 import org.junit.Test;
 
 public class ReturnByteArrayNewTest extends com.aparapi.codegen.CodeGenJUnitBase {
     private static final String[] expectedOpenCL = null;
-    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = ClassParseException.class;
+    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = CodeGenException.class;
 
     
     @Test
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnByteArrayVar.java b/src/test/java/com/aparapi/codegen/test/ReturnByteArrayVar.java
index 8751bfb6..45b9cc9a 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnByteArrayVar.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnByteArrayVar.java
@@ -27,4 +27,27 @@ public class ReturnByteArrayVar {
         returnByteArrayVar();
     }
 }
-/**{Throws{ClassParseException}Throws}**/
+/**{OpenCL{
+ typedef struct This_s{
+ int passid;
+ }This;
+ int get_pass_id(This *this){
+ return this->passid;
+ }
+ __global char* com_aparapi_codegen_test_ReturnByteArrayVar__returnByteArrayVar(This *this){
+ char bytes[1024];
+ return(bytes);
+ }
+ __kernel void run(
+ int passid
+ ){
+ This thisStruct;
+ This* this=&thisStruct;
+ this->passid = passid;
+ {
+ com_aparapi_codegen_test_ReturnByteArrayVar__returnByteArrayVar(this);
+ return;
+ }
+ }
+
+ }OpenCL}**/
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnByteArrayVarTest.java b/src/test/java/com/aparapi/codegen/test/ReturnByteArrayVarTest.java
index 2b83d4d7..49cf10aa 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnByteArrayVarTest.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnByteArrayVarTest.java
@@ -15,13 +15,34 @@
  */
 package com.aparapi.codegen.test;
 
-import com.aparapi.internal.exception.ClassParseException;
 import org.junit.Test;
 
 public class ReturnByteArrayVarTest extends com.aparapi.codegen.CodeGenJUnitBase {
 
-    private static final String[] expectedOpenCL = null;
-    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = ClassParseException.class;
+    private static final String[] expectedOpenCL = {
+        "typedef struct This_s{\n"
+            + "   int passid;\n"
+            + "}This;\n"
+            + "int get_pass_id(This *this){\n"
+            + "   return this->passid;\n"
+            + "}\n"
+            + " __global char* com_aparapi_codegen_test_ReturnByteArrayVar__returnByteArrayVar(This *this){\n"
+            + "   char bytes[1024];\n"
+            + "   return(bytes);\n"
+            + "}\n"
+            + "__kernel void run(\n"
+            + "   int passid\n"
+            + "){\n"
+            + "   This thisStruct;\n"
+            + "   This* this=&thisStruct;\n"
+            + "   this->passid = passid;\n"
+            + "   {\n"
+            + "      com_aparapi_codegen_test_ReturnByteArrayVar__returnByteArrayVar(this);\n"
+            + "      return;\n"
+            + "   }\n"
+            + "}"
+    };
+    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = null;
 
     @Test
     public void ReturnByteArrayVarTest() {
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnDoubleArrayNew.java b/src/test/java/com/aparapi/codegen/test/ReturnDoubleArrayNew.java
index 64e3d29a..c7770ac5 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnDoubleArrayNew.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnDoubleArrayNew.java
@@ -25,4 +25,4 @@ public class ReturnDoubleArrayNew {
         returnDoubleArrayNew();
     }
 }
-/**{Throws{ClassParseException}Throws}**/
+/**{Throws{CodeGenException}Throws}**/
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnDoubleArrayNewTest.java b/src/test/java/com/aparapi/codegen/test/ReturnDoubleArrayNewTest.java
index 9c2433f5..d60e3c92 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnDoubleArrayNewTest.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnDoubleArrayNewTest.java
@@ -15,13 +15,13 @@
  */
 package com.aparapi.codegen.test;
 
-import com.aparapi.internal.exception.ClassParseException;
+import com.aparapi.internal.exception.CodeGenException;
 import org.junit.Test;
 
 public class ReturnDoubleArrayNewTest extends com.aparapi.codegen.CodeGenJUnitBase {
 
     private static final String[] expectedOpenCL = null;
-    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = ClassParseException.class;
+    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = CodeGenException.class;
 
     @Test
     public void ReturnDoubleArrayNewTest() {
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnDoubleArrayVar.java b/src/test/java/com/aparapi/codegen/test/ReturnDoubleArrayVar.java
index 4be6a58e..609aa988 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnDoubleArrayVar.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnDoubleArrayVar.java
@@ -27,4 +27,27 @@ public class ReturnDoubleArrayVar {
         returnDoubleArrayVar();
     }
 }
-/**{Throws{ClassParseException}Throws}**/
+/**{OpenCL{
+ typedef struct This_s{
+ int passid;
+ }This;
+ int get_pass_id(This *this){
+ return this->passid;
+ }
+ __global double* com_aparapi_codegen_test_ReturnDoubleArrayVar__returnDoubleArrayVar(This *this){
+ double doubles[1024];
+ return(doubles);
+ }
+ __kernel void run(
+ int passid
+ ){
+ This thisStruct;
+ This* this=&thisStruct;
+ this->passid = passid;
+ {
+ com_aparapi_codegen_test_ReturnDoubleArrayVar__returnDoubleArrayVar(this);
+ return;
+ }
+ }
+
+ }OpenCL}**/
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnDoubleArrayVarTest.java b/src/test/java/com/aparapi/codegen/test/ReturnDoubleArrayVarTest.java
index 1517e608..330b735d 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnDoubleArrayVarTest.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnDoubleArrayVarTest.java
@@ -15,13 +15,34 @@
  */
 package com.aparapi.codegen.test;
 
-import com.aparapi.internal.exception.ClassParseException;
 import org.junit.Test;
 
 public class ReturnDoubleArrayVarTest extends com.aparapi.codegen.CodeGenJUnitBase {
 
-    private static final String[] expectedOpenCL = null;
-    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = ClassParseException.class;
+    private static final String[] expectedOpenCL = {
+        "typedef struct This_s{\n"
+            + "   int passid;\n"
+            + "}This;\n"
+            + "int get_pass_id(This *this){\n"
+            + "   return this->passid;\n"
+            + "}\n"
+            + " __global double* com_aparapi_codegen_test_ReturnDoubleArrayVar__returnDoubleArrayVar(This *this){\n"
+            + "   double doubles[1024];\n"
+            + "   return(doubles);\n"
+            + "}\n"
+            + "__kernel void run(\n"
+            + "   int passid\n"
+            + "){\n"
+            + "   This thisStruct;\n"
+            + "   This* this=&thisStruct;\n"
+            + "   this->passid = passid;\n"
+            + "   {\n"
+            + "      com_aparapi_codegen_test_ReturnDoubleArrayVar__returnDoubleArrayVar(this);\n"
+            + "      return;\n"
+            + "   }\n"
+            + "}"
+    };
+    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = null;
 
     @Test
     public void ReturnDoubleArrayVarTest() {
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnFloatArrayNew.java b/src/test/java/com/aparapi/codegen/test/ReturnFloatArrayNew.java
index 9dd36a7c..211d3e45 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnFloatArrayNew.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnFloatArrayNew.java
@@ -25,4 +25,4 @@ public class ReturnFloatArrayNew {
         returnFloatArrayNew();
     }
 }
-/**{Throws{ClassParseException}Throws}**/
+/**{Throws{CodeGenException}Throws}**/
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnFloatArrayNewTest.java b/src/test/java/com/aparapi/codegen/test/ReturnFloatArrayNewTest.java
index 431d54a6..2437e467 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnFloatArrayNewTest.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnFloatArrayNewTest.java
@@ -15,13 +15,13 @@
  */
 package com.aparapi.codegen.test;
 
-import com.aparapi.internal.exception.ClassParseException;
+import com.aparapi.internal.exception.CodeGenException;
 import org.junit.Test;
 
 public class ReturnFloatArrayNewTest extends com.aparapi.codegen.CodeGenJUnitBase {
 
     private static final String[] expectedOpenCL = null;
-    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = ClassParseException.class;
+    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = CodeGenException.class;
 
     @Test
     public void ReturnFloatArrayNewTest() {
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnFloatArrayVar.java b/src/test/java/com/aparapi/codegen/test/ReturnFloatArrayVar.java
index f843d65b..289b0050 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnFloatArrayVar.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnFloatArrayVar.java
@@ -27,4 +27,27 @@ public class ReturnFloatArrayVar {
         returnFloatArrayVar();
     }
 }
-/**{Throws{ClassParseException}Throws}**/
+/**{OpenCL{
+ typedef struct This_s{
+ int passid;
+ }This;
+ int get_pass_id(This *this){
+ return this->passid;
+ }
+ __global float* com_aparapi_codegen_test_ReturnFloatArrayVar__returnFloatArrayVar(This *this){
+ float floats[1024];
+ return(floats);
+ }
+ __kernel void run(
+ int passid
+ ){
+ This thisStruct;
+ This* this=&thisStruct;
+ this->passid = passid;
+ {
+ com_aparapi_codegen_test_ReturnFloatArrayVar__returnFloatArrayVar(this);
+ return;
+ }
+ }
+
+ }OpenCL}**/
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnFloatArrayVarTest.java b/src/test/java/com/aparapi/codegen/test/ReturnFloatArrayVarTest.java
index 48fd6b0d..6fe511ff 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnFloatArrayVarTest.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnFloatArrayVarTest.java
@@ -15,13 +15,34 @@
  */
 package com.aparapi.codegen.test;
 
-import com.aparapi.internal.exception.ClassParseException;
 import org.junit.Test;
 
 public class ReturnFloatArrayVarTest extends com.aparapi.codegen.CodeGenJUnitBase {
 
-    private static final String[] expectedOpenCL = null;
-    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = ClassParseException.class;
+    private static final String[] expectedOpenCL = {
+        "typedef struct This_s{\n"
+            + "   int passid;\n"
+            + "}This;\n"
+            + "int get_pass_id(This *this){\n"
+            + "   return this->passid;\n"
+            + "}\n"
+            + " __global float* com_aparapi_codegen_test_ReturnFloatArrayVar__returnFloatArrayVar(This *this){\n"
+            + "   float floats[1024];\n"
+            + "   return(floats);\n"
+            + "}\n"
+            + "__kernel void run(\n"
+            + "   int passid\n"
+            + "){\n"
+            + "   This thisStruct;\n"
+            + "   This* this=&thisStruct;\n"
+            + "   this->passid = passid;\n"
+            + "   {\n"
+            + "      com_aparapi_codegen_test_ReturnFloatArrayVar__returnFloatArrayVar(this);\n"
+            + "      return;\n"
+            + "   }\n"
+            + "}"
+    };
+    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = null;
 
     @Test
     public void ReturnFloatArrayVarTest() {
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnIntArrayNew.java b/src/test/java/com/aparapi/codegen/test/ReturnIntArrayNew.java
index c8b8b1f0..abd3ac2b 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnIntArrayNew.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnIntArrayNew.java
@@ -25,4 +25,4 @@ public class ReturnIntArrayNew {
         returnIntArrayNew();
     }
 }
-/**{Throws{ClassParseException}Throws}**/
+/**{Throws{CodeGenException}Throws}**/
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnIntArrayNewTest.java b/src/test/java/com/aparapi/codegen/test/ReturnIntArrayNewTest.java
index fe067d9d..cd7858f1 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnIntArrayNewTest.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnIntArrayNewTest.java
@@ -15,13 +15,13 @@
  */
 package com.aparapi.codegen.test;
 
-import com.aparapi.internal.exception.ClassParseException;
+import com.aparapi.internal.exception.CodeGenException;
 import org.junit.Test;
 
 public class ReturnIntArrayNewTest extends com.aparapi.codegen.CodeGenJUnitBase {
 
     private static final String[] expectedOpenCL = null;
-    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = ClassParseException.class;
+    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = CodeGenException.class;
 
     @Test
     public void ReturnIntArrayNewTest() {
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnIntArrayVar.java b/src/test/java/com/aparapi/codegen/test/ReturnIntArrayVar.java
index 172757e2..9f1d0306 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnIntArrayVar.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnIntArrayVar.java
@@ -27,4 +27,27 @@ public class ReturnIntArrayVar {
         returnIntArrayVar();
     }
 }
-/**{Throws{ClassParseException}Throws}**/
+/**{OpenCL{
+ typedef struct This_s{
+ int passid;
+ }This;
+ int get_pass_id(This *this){
+ return this->passid;
+ }
+ __global int* com_aparapi_codegen_test_ReturnIntArrayVar__returnIntArrayVar(This *this){
+ int ints[1024];
+ return(ints);
+ }
+ __kernel void run(
+ int passid
+ ){
+ This thisStruct;
+ This* this=&thisStruct;
+ this->passid = passid;
+ {
+ com_aparapi_codegen_test_ReturnIntArrayVar__returnIntArrayVar(this);
+ return;
+ }
+ }
+
+ }OpenCL}**/
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnIntArrayVarTest.java b/src/test/java/com/aparapi/codegen/test/ReturnIntArrayVarTest.java
index a032b2c1..f891df6a 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnIntArrayVarTest.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnIntArrayVarTest.java
@@ -15,13 +15,34 @@
  */
 package com.aparapi.codegen.test;
 
-import com.aparapi.internal.exception.ClassParseException;
 import org.junit.Test;
 
 public class ReturnIntArrayVarTest extends com.aparapi.codegen.CodeGenJUnitBase {
 
-    private static final String[] expectedOpenCL = null;
-    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = ClassParseException.class;
+    private static final String[] expectedOpenCL = {
+        "typedef struct This_s{\n"
+            + "   int passid;\n"
+            + "}This;\n"
+            + "int get_pass_id(This *this){\n"
+            + "   return this->passid;\n"
+            + "}\n"
+            + " __global int* com_aparapi_codegen_test_ReturnIntArrayVar__returnIntArrayVar(This *this){\n"
+            + "   int ints[1024];\n"
+            + "   return(ints);\n"
+            + "}\n"
+            + "__kernel void run(\n"
+            + "   int passid\n"
+            + "){\n"
+            + "   This thisStruct;\n"
+            + "   This* this=&thisStruct;\n"
+            + "   this->passid = passid;\n"
+            + "   {\n"
+            + "      com_aparapi_codegen_test_ReturnIntArrayVar__returnIntArrayVar(this);\n"
+            + "      return;\n"
+            + "   }\n"
+            + "}"
+    };
+    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = null;
 
     @Test
     public void ReturnIntArrayVarTest() {
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnLongArrayNew.java b/src/test/java/com/aparapi/codegen/test/ReturnLongArrayNew.java
index e9143097..9a11f9a5 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnLongArrayNew.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnLongArrayNew.java
@@ -25,4 +25,4 @@ public class ReturnLongArrayNew {
         returnLongArrayNew();
     }
 }
-/**{Throws{ClassParseException}Throws}**/
+/**{Throws{CodeGenException}Throws}**/
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnLongArrayNewTest.java b/src/test/java/com/aparapi/codegen/test/ReturnLongArrayNewTest.java
index b9c5c2a5..d437ba4b 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnLongArrayNewTest.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnLongArrayNewTest.java
@@ -15,13 +15,13 @@
  */
 package com.aparapi.codegen.test;
 
-import com.aparapi.internal.exception.ClassParseException;
+import com.aparapi.internal.exception.CodeGenException;
 import org.junit.Test;
 
 public class ReturnLongArrayNewTest extends com.aparapi.codegen.CodeGenJUnitBase {
 
     private static final String[] expectedOpenCL = null;
-    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = ClassParseException.class;
+    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = CodeGenException.class;
 
     @Test
     public void ReturnLongArrayNewTest() {
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnLongArrayVar.java b/src/test/java/com/aparapi/codegen/test/ReturnLongArrayVar.java
index 31edd946..0cb75614 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnLongArrayVar.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnLongArrayVar.java
@@ -27,4 +27,6 @@ public class ReturnLongArrayVar {
         returnLongArrayVar();
     }
 }
-/**{Throws{ClassParseException}Throws}**/
+/**{OpenCL{
+
+ }OpenCL}**/
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnLongArrayVarTest.java b/src/test/java/com/aparapi/codegen/test/ReturnLongArrayVarTest.java
index f71c402d..66e77425 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnLongArrayVarTest.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnLongArrayVarTest.java
@@ -15,13 +15,34 @@
  */
 package com.aparapi.codegen.test;
 
-import com.aparapi.internal.exception.ClassParseException;
 import org.junit.Test;
 
 public class ReturnLongArrayVarTest extends com.aparapi.codegen.CodeGenJUnitBase {
 
-    private static final String[] expectedOpenCL = null;
-    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = ClassParseException.class;
+    private static final String[] expectedOpenCL = {
+        "typedef struct This_s{\n"
+            + "   int passid;\n"
+            + "}This;\n"
+            + "int get_pass_id(This *this){\n"
+            + "   return this->passid;\n"
+            + "}\n"
+            + " __global long* com_aparapi_codegen_test_ReturnLongArrayVar__returnLongArrayVar(This *this){\n"
+            + "   long longs[1024];\n"
+            + "   return(longs);\n"
+            + "}\n"
+            + "__kernel void run(\n"
+            + "   int passid\n"
+            + "){\n"
+            + "   This thisStruct;\n"
+            + "   This* this=&thisStruct;\n"
+            + "   this->passid = passid;\n"
+            + "   {\n"
+            + "      com_aparapi_codegen_test_ReturnLongArrayVar__returnLongArrayVar(this);\n"
+            + "      return;\n"
+            + "   }\n"
+            + "}"
+    };
+    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = null;
 
     @Test
     public void ReturnLongArrayVarTest() {
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnShortArrayNew.java b/src/test/java/com/aparapi/codegen/test/ReturnShortArrayNew.java
index c3abc039..443033f9 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnShortArrayNew.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnShortArrayNew.java
@@ -25,4 +25,4 @@ public class ReturnShortArrayNew {
         returnShortArrayNew();
     }
 }
-/**{Throws{ClassParseException}Throws}**/
+/**{Throws{CodeGenException}Throws}**/
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnShortArrayNewTest.java b/src/test/java/com/aparapi/codegen/test/ReturnShortArrayNewTest.java
index 4b4e5f8f..f720a785 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnShortArrayNewTest.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnShortArrayNewTest.java
@@ -15,13 +15,13 @@
  */
 package com.aparapi.codegen.test;
 
-import com.aparapi.internal.exception.ClassParseException;
+import com.aparapi.internal.exception.CodeGenException;
 import org.junit.Test;
 
 public class ReturnShortArrayNewTest extends com.aparapi.codegen.CodeGenJUnitBase {
 
     private static final String[] expectedOpenCL = null;
-    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = ClassParseException.class;
+    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = CodeGenException.class;
 
     @Test
     public void ReturnShortArrayNewTest() {
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnShortArrayVar.java b/src/test/java/com/aparapi/codegen/test/ReturnShortArrayVar.java
index 07d590bd..8efb31c0 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnShortArrayVar.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnShortArrayVar.java
@@ -27,4 +27,26 @@ public class ReturnShortArrayVar {
         returnShortArrayVar();
     }
 }
-/**{Throws{ClassParseException}Throws}**/
+/**{OpenCL{
+ typedef struct This_s{
+ int passid;
+ }This;
+ int get_pass_id(This *this){
+ return this->passid;
+ }
+ __global short* com_aparapi_codegen_test_ReturnShortArrayVar__returnShortArrayVar(This *this){
+ short shorts[1024];
+ return(shorts);
+ }
+ __kernel void run(
+ int passid
+ ){
+ This thisStruct;
+ This* this=&thisStruct;
+ this->passid = passid;
+ {
+ com_aparapi_codegen_test_ReturnShortArrayVar__returnShortArrayVar(this);
+ return;
+ }
+ }
+}OpenCL}**/
diff --git a/src/test/java/com/aparapi/codegen/test/ReturnShortArrayVarTest.java b/src/test/java/com/aparapi/codegen/test/ReturnShortArrayVarTest.java
index 42bdefbb..b703e37b 100644
--- a/src/test/java/com/aparapi/codegen/test/ReturnShortArrayVarTest.java
+++ b/src/test/java/com/aparapi/codegen/test/ReturnShortArrayVarTest.java
@@ -15,13 +15,35 @@
  */
 package com.aparapi.codegen.test;
 
-import com.aparapi.internal.exception.ClassParseException;
+
 import org.junit.Test;
 
 public class ReturnShortArrayVarTest extends com.aparapi.codegen.CodeGenJUnitBase {
 
-    private static final String[] expectedOpenCL = null;
-    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = ClassParseException.class;
+    private static final String[] expectedOpenCL = {
+        "typedef struct This_s{\n"
+            + "   int passid;\n"
+            + "}This;\n"
+            + "int get_pass_id(This *this){\n"
+            + "   return this->passid;\n"
+            + "}\n"
+            + " __global short* com_aparapi_codegen_test_ReturnShortArrayVar__returnShortArrayVar(This *this){\n"
+            + "   short shorts[1024];\n"
+            + "   return(shorts);\n"
+            + "}\n"
+            + "__kernel void run(\n"
+            + "   int passid\n"
+            + "){\n"
+            + "   This thisStruct;\n"
+            + "   This* this=&thisStruct;\n"
+            + "   this->passid = passid;\n"
+            + "   {\n"
+            + "      com_aparapi_codegen_test_ReturnShortArrayVar__returnShortArrayVar(this);\n"
+            + "      return;\n"
+            + "   }\n"
+            + "}"
+    };
+    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = null;
 
     @Test
     public void ReturnShortArrayVarTest() {
diff --git a/src/test/java/com/aparapi/runtime/ArrayTest.java b/src/test/java/com/aparapi/runtime/ArrayTest.java
new file mode 100644
index 00000000..a21af6e2
--- /dev/null
+++ b/src/test/java/com/aparapi/runtime/ArrayTest.java
@@ -0,0 +1,82 @@
+/**
+ * Copyright (c) 2016 - 2017 Syncleus, Inc.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+package com.aparapi.runtime;
+
+import static org.junit.Assert.assertArrayEquals;
+
+import com.aparapi.Kernel;
+import java.util.Arrays;
+import org.junit.Test;
+
+
+public class ArrayTest {
+    @Test
+    public void test() {
+        VectorKernel b = new VectorKernel();
+        b.test();
+    }
+
+    public static class VectorKernel extends Kernel {
+        private static final int SIZE = 32;
+
+        static int[][] target = new int[SIZE][SIZE];
+
+        public VectorKernel() {
+            for (int i = 0; i < SIZE; ++i) {
+                int[] ints = new int[SIZE];
+                Arrays.fill(ints, 99);
+
+                target[i] = ints;
+            }
+        }
+
+        private static void fillArray(int[] ints_$private$, int id) {
+            for (int i = 0; i < SIZE; i++) {
+                ints_$private$[i] = i + id;
+            }
+        }
+
+        @Override
+        public void run() {
+            int id = getGlobalId();
+
+            int[] ints = new int[SIZE];
+
+            fillArray(ints, id);
+
+            for (int i = 0; i < SIZE; i++) {
+                target[id][i] = ints[i];
+            }
+        }
+
+        void validate() {
+            for (int j = 0; j < SIZE; j++) {
+
+                int[] expected = new int[SIZE];
+                for (int i = 0; i < SIZE; i++) {
+                    expected[i] = i + j;
+                }
+
+                assertArrayEquals("target["+j+"]", expected, target[j]);
+            }
+        }
+
+        public void test() {
+            execute(SIZE);
+            validate();
+        }
+    }
+}
-- 
GitLab