Fix final null array refs in kernels

src/main/java/com/aparapi/internal/kernel/KernelRunner.java

@@ -1727,6 +1727,9 @@ else if (Config.enableShowGeneratedOpenCL) {
                   int i = 0;
 
                   for (final Field field : entryPoint.getReferencedFields()) {
+                     if (entryPoint.isFinalNullArrayField(field.getName())) {
+                        continue;
+                     }
                      try {
                         field.setAccessible(true);
                         args[i] = new KernelArg();

src/main/java/com/aparapi/internal/model/Entrypoint.java

@@ -810,6 +810,14 @@ private boolean noCL(ClassModelMethod m) {
       return found;
    }
 
+   private boolean isFinalNullArrayField(Field field) throws IllegalAccessException {
+      if (kernelInstance == null || !field.getType().isArray() || !Modifier.isFinal(field.getModifiers())) {
+         return false;
+      }
+      field.setAccessible(true);
+      return field.get(kernelInstance) == null;
+   }
+
    private FieldEntry getSimpleGetterField(MethodModel method) {
       return method.getAccessorVariableFieldEntry();
    }
@@ -842,6 +850,20 @@ public Set<String> getArrayFieldArrayLengthUsed() {
       return (arrayFieldArrayLengthUsed);
    }
 
+   public boolean isFinalNullArrayField(String fieldName) {
+      for (final Field field : referencedFields) {
+         if (field.getName().equals(fieldName)) {
+            try {
+               return isFinalNullArrayField(field);
+            } catch (final SecurityException | IllegalAccessException e) {
+               e.printStackTrace();
+               return false;
+            }
+         }
+      }
+      return false;
+   }
+
    public MethodModel getMethodModel() {
       return (methodModel);
    }

src/main/java/com/aparapi/internal/writer/KernelWriter.java

@@ -392,6 +392,8 @@ public void writePragma(String _name, boolean _enable) {
             numDimensions++;
             signature = signature.substring(1);
          }
+         final boolean literalNullArray = isPointer && privateMemorySize == null
+               && _entryPoint.isFinalNullArrayField(field.getName());
 
          // If it is a converted array of objects, emit the struct param
          String className = null;
@@ -425,15 +427,17 @@ public void writePragma(String _name, boolean _enable) {
             assignLine.append("this->");
             assignLine.append(field.getName());
             assignLine.append(" = ");
-            assignLine.append(field.getName());
+            assignLine.append(literalNullArray ? "NULL" : field.getName());
          }
 
          argLine.append(field.getName());
          thisStructLine.append(field.getName());
          if (privateMemorySize == null) {
             assigns.add(assignLine.toString());
          }
-         argLines.add(argLine.toString());
+         if (!literalNullArray) {
+            argLines.add(argLine.toString());
+         }
          if (privateMemorySize != null) {
             thisStructLine.append("[").append(privateMemorySize).append("]");
          }
@@ -456,12 +460,14 @@ public void writePragma(String _name, boolean _enable) {
                lenAssignLine.append("this->");
                lenAssignLine.append(lenName);
                lenAssignLine.append(" = ");
-               lenAssignLine.append(lenName);
+               lenAssignLine.append(literalNullArray ? "0" : lenName);
 
                lenArgLine.append("int " + lenName);
 
                assigns.add(lenAssignLine.toString());
-               argLines.add(lenArgLine.toString());
+               if (!literalNullArray) {
+                  argLines.add(lenArgLine.toString());
+               }
                thisStruct.add(lenStructLine.toString());
 
                if (numDimensions > 1) {
@@ -475,12 +481,14 @@ public void writePragma(String _name, boolean _enable) {
                   dimAssignLine.append("this->");
                   dimAssignLine.append(dimName);
                   dimAssignLine.append(" = ");
-                  dimAssignLine.append(dimName);
+                  dimAssignLine.append(literalNullArray ? "0" : dimName);
 
                   dimArgLine.append("int " + dimName);
 
                   assigns.add(dimAssignLine.toString());
-                  argLines.add(dimArgLine.toString());
+                  if (!literalNullArray) {
+                     argLines.add(dimArgLine.toString());
+                  }
                   thisStruct.add(dimStructLine.toString());
                }
             }

src/test/java/com/aparapi/codegen/test/FinalNullArray.java

@@ -0,0 +1,29 @@
+/**
+ * Copyright (c) 2016 - 2018 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 com.aparapi.Kernel;
+
+public class FinalNullArray extends Kernel {
+    private final int[] nullArray = null;
+
+    @Override
+    public void run() {
+        if (nullArray == null) {
+            return;
+        }
+    }
+}

src/test/java/com/aparapi/codegen/test/FinalNullArrayTest.java

@@ -0,0 +1,56 @@
+/**
+ * Copyright (c) 2016 - 2018 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 FinalNullArrayTest extends com.aparapi.codegen.CodeGenJUnitBase {
+    private static final String[] expectedOpenCL = {
+    "typedef struct This_s{\n" +
+" __global int *nullArray;\n" +
+" int passid;\n" +
+" }This;\n" +
+" int get_pass_id(This *this){\n" +
+" return this->passid;\n" +
+" }\n" +
+"\n" +
+" __kernel void run(\n" +
+" int passid\n" +
+" ){\n" +
+" This thisStruct;\n" +
+" This* this=&thisStruct;\n" +
+" this->nullArray = NULL;\n" +
+" this->passid = passid;\n" +
+" {\n" +
+" if (this->nullArray == NULL){\n" +
+" return;\n" +
+" }\n" +
+" return;\n" +
+" }\n" +
+" }\n" +
+" "};
+    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = null;
+
+    @Test
+    public void FinalNullArrayTest() {
+        test(com.aparapi.codegen.test.FinalNullArray.class, expectedException, expectedOpenCL);
+    }
+
+    @Test
+    public void FinalNullArrayTestWorksWithCaching() {
+        test(com.aparapi.codegen.test.FinalNullArray.class, expectedException, expectedOpenCL);
+    }
+}

src/test/java/com/aparapi/runtime/NullRefTest.java

@@ -16,19 +16,50 @@
 package com.aparapi.runtime;
 
 import com.aparapi.Kernel;
-import org.junit.Ignore;
+import com.aparapi.Range;
+import com.aparapi.device.Device;
+import com.aparapi.device.OpenCLDevice;
+import com.aparapi.internal.kernel.KernelManager;
+import org.junit.After;
+import org.junit.Before;
 import org.junit.Test;
 
+import java.util.Arrays;
+import java.util.List;
+
+import static org.junit.Assume.assumeTrue;
+
 public class NullRefTest {
-    @Ignore("Known bug, runs on CPU but not GPU.")
+    private static OpenCLDevice openCLDevice = null;
+
+    private class CLKernelManager extends KernelManager {
+        @Override
+        protected List<Device.TYPE> getPreferredDeviceTypes() {
+            return Arrays.asList(Device.TYPE.ACC, Device.TYPE.GPU, Device.TYPE.CPU);
+        }
+    }
+
+    @Before
+    public void setUpBeforeClass() throws Exception {
+        KernelManager.setKernelManager(new CLKernelManager());
+        Device device = KernelManager.instance().bestDevice();
+        assumeTrue(device != null && device instanceof OpenCLDevice);
+        openCLDevice = (OpenCLDevice) device;
+    }
+
+    @After
+    public void classTeardown() {
+        Util.resetKernelManager();
+    }
+
     @Test
     public void test() {
-        new NullRefTest().doTest();
+        doTest();
     }
 
     private void doTest() {
         final Kernel kernel = new NullRefKernel();
-        kernel.execute(1);
+        kernel.execute(Range.create(openCLDevice, 1, 1));
     }
 
     private class NullRefKernel extends Kernel {