From c7be5c160f49b47ca28153c361d76f022cd9e69b Mon Sep 17 00:00:00 2001 From: patrickmaub Date: Tue, 12 May 2026 20:22:51 -0500 Subject: [PATCH] Fix final null array kernel refs --- .../aparapi/internal/kernel/KernelRunner.java | 3 + .../aparapi/internal/model/Entrypoint.java | 22 ++++++++ .../aparapi/internal/writer/KernelWriter.java | 20 +++++-- .../aparapi/codegen/test/FinalNullArray.java | 29 ++++++++++ .../codegen/test/FinalNullArrayTest.java | 56 +++++++++++++++++++ .../java/com/aparapi/runtime/NullRefTest.java | 39 +++++++++++-- 6 files changed, 159 insertions(+), 10 deletions(-) create mode 100644 src/test/java/com/aparapi/codegen/test/FinalNullArray.java create mode 100644 src/test/java/com/aparapi/codegen/test/FinalNullArrayTest.java diff --git a/src/main/java/com/aparapi/internal/kernel/KernelRunner.java b/src/main/java/com/aparapi/internal/kernel/KernelRunner.java index 25975a2a..e9454c62 100644 --- a/src/main/java/com/aparapi/internal/kernel/KernelRunner.java +++ b/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(); diff --git a/src/main/java/com/aparapi/internal/model/Entrypoint.java b/src/main/java/com/aparapi/internal/model/Entrypoint.java index 24ec1a14..95735e13 100644 --- a/src/main/java/com/aparapi/internal/model/Entrypoint.java +++ b/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 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); } diff --git a/src/main/java/com/aparapi/internal/writer/KernelWriter.java b/src/main/java/com/aparapi/internal/writer/KernelWriter.java index 0a72a33f..9d5d1f57 100644 --- a/src/main/java/com/aparapi/internal/writer/KernelWriter.java +++ b/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,7 +427,7 @@ 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()); @@ -433,7 +435,9 @@ public void writePragma(String _name, boolean _enable) { 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()); } } diff --git a/src/test/java/com/aparapi/codegen/test/FinalNullArray.java b/src/test/java/com/aparapi/codegen/test/FinalNullArray.java new file mode 100644 index 00000000..a084f7e9 --- /dev/null +++ b/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; + } + } +} diff --git a/src/test/java/com/aparapi/codegen/test/FinalNullArrayTest.java b/src/test/java/com/aparapi/codegen/test/FinalNullArrayTest.java new file mode 100644 index 00000000..27c6142a --- /dev/null +++ b/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 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); + } +} diff --git a/src/test/java/com/aparapi/runtime/NullRefTest.java b/src/test/java/com/aparapi/runtime/NullRefTest.java index 8d312509..927bddec 100644 --- a/src/test/java/com/aparapi/runtime/NullRefTest.java +++ b/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 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 {