Syncleus · ItsOtherMauridian · May 12, 2026
diff --git a/src/main/java/com/aparapi/internal/writer/KernelWriter.java b/src/main/java/com/aparapi/internal/writer/KernelWriter.java
@@ -218,7 +218,12 @@ public abstract class KernelWriter extends BlockWriter{
 
       if (barrierAndGetterMappings != null) {
          // this is one of the OpenCL barrier or size getter methods
-         // write the mapping and exit
+         // write the mapping and exit. OpenCL work-item query functions return size_t,
+         // but the corresponding Kernel methods return Java int; preserve the Java
+         // narrowing semantics before any enclosing cast/arithmetic is applied.
+         if (methodSignature.endsWith(")I") && !methodName.equals("getPassId")) {
+            write("(int)");
+         }
          if (argc > 0) {
             write(barrierAndGetterMappings);
             write("(");
@@ -313,6 +318,37 @@ private boolean isThis(Instruction instruction) {
       return instruction instanceof I_ALOAD_0;
    }
 
+   private static boolean usesLongMultiply(MethodModel methodModel) {
+      if (methodModel == null) {
+         return false;
+      }
+
+      for (Instruction instruction = methodModel.getPCHead(); instruction != null; instruction = instruction.getNextPC()) {
+         if (instruction instanceof I_LMUL) {
+            return true;
+         }
+      }
+      return false;
+   }
+
+   private static boolean usesLongMultiply(Entrypoint entryPoint) {
+      if (entryPoint == null) {
+         return false;
+      }
+
+      if (usesLongMultiply(entryPoint.getMethodModel())) {
+         return true;
+      }
+
+      for (MethodModel calledMethod : entryPoint.getCalledMethods()) {
+         if (usesLongMultiply(calledMethod)) {
+            return true;
+         }
+      }
+
+      return false;
+   }
+
    public void writePragma(String _name, boolean _enable) {
       write("#pragma OPENCL EXTENSION " + _name + " : " + (_enable ? "en" : "dis") + "able");
       newLine();
@@ -536,6 +572,37 @@ public void writePragma(String _name, boolean _enable) {
          newLine();
       }
 
+      if (usesLongMultiply(_entryPoint)) {
+         // Some OpenCL drivers miscompile native 64-bit integer multiplication.
+         // Compute the low 64 bits with 32-bit operations to preserve Java long multiply semantics.
+         write("inline ulong aparapi_umul64_lo(ulong a, ulong b){");
+         in();
+         {
+            newLine();
+            write("uint a0 = (uint)a;");
+            newLine();
+            write("uint a1 = (uint)(a >> 32);");
+            newLine();
+            write("uint b0 = (uint)b;");
+            newLine();
+            write("uint b1 = (uint)(b >> 32);");
+            newLine();
+            write("uint lo = a0 * b0;");
+            newLine();
+            write("uint hi = mul_hi(a0, b0);");
+            newLine();
+            write("uint cross = hi + (a0 * b1) + (a1 * b0);");
+            newLine();
+            write("return (((ulong)cross) << 32) | (ulong)lo;");
+            out();
+            newLine();
+         }
+         write("}");
+         newLine();
+         write("inline long aparapi_lmul(long a, long b){ return (long)aparapi_umul64_lo((ulong)a, (ulong)b); }");
+         newLine();
+      }
+
       // Emit structs for oop transformation accessors
       for (final ClassModel cm : _entryPoint.getObjectArrayFieldsClasses().values()) {
          final ArrayList<FieldEntry> fieldSet = cm.getStructMembers();
@@ -773,6 +840,32 @@ public void writePragma(String _name, boolean _enable) {
          write(" >> ");
          writeInstruction(binaryInstruction.getRhs());
 
+         if (needsParenthesis) {
+            write(")");
+         }
+      } else if (_instruction instanceof I_LMUL) {
+         final BinaryOperator binaryInstruction = (BinaryOperator) _instruction;
+         final Instruction parent = binaryInstruction.getParentExpr();
+         boolean needsParenthesis = true;
+
+         if (parent instanceof AssignToLocalVariable) {
+            needsParenthesis = false;
+         } else if (parent instanceof AssignToField) {
+            needsParenthesis = false;
+         } else if (parent instanceof AssignToArrayElement) {
+            needsParenthesis = false;
+         }
+
+         if (needsParenthesis) {
+            write("(");
+         }
+
+         write("aparapi_lmul(");
+         writeInstruction(binaryInstruction.getLhs());
+         write(", ");
+         writeInstruction(binaryInstruction.getRhs());
+         write(")");
+
          if (needsParenthesis) {
             write(")");
          }

diff --git a/src/test/java/com/aparapi/codegen/test/CompositeArbitraryScopeTest.java b/src/test/java/com/aparapi/codegen/test/CompositeArbitraryScopeTest.java
@@ -28,7 +28,7 @@ public class CompositeArbitraryScopeTest extends com.aparapi.codegen.CodeGenJUni
 " }\n" +
 "\n" +
 " void com_aparapi_codegen_test_CompositeArbitraryScope__t5(This *this){\n" +
-" int gid = get_global_id(0);\n" +
+" int gid = (int)get_global_id(0);\n" +
 " int numRemaining = 1;\n" +
 " int thisCount = 0;\n" +
 " for (; numRemaining>0 && gid>0; numRemaining++){\n" +
@@ -40,7 +40,7 @@ public class CompositeArbitraryScopeTest extends com.aparapi.codegen.CodeGenJUni
 " return;\n" +
 " }\n" +
 " void com_aparapi_codegen_test_CompositeArbitraryScope__t4(This *this){\n" +
-" int gid = get_global_id(0);\n" +
+" int gid = (int)get_global_id(0);\n" +
 " int numRemaining = 1;\n" +
 " while (numRemaining>0 && gid>0){\n" +
 " numRemaining++;\n" +
@@ -54,7 +54,7 @@ public class CompositeArbitraryScopeTest extends com.aparapi.codegen.CodeGenJUni
 " return;\n" +
 " }\n" +
 " void com_aparapi_codegen_test_CompositeArbitraryScope__t3(This *this){\n" +
-" int gid = get_global_id(0);\n" +
+" int gid = (int)get_global_id(0);\n" +
 " int numRemaining = 1;\n" +
 " while (numRemaining>0){\n" +
 " numRemaining++;\n" +
@@ -67,7 +67,7 @@ public class CompositeArbitraryScopeTest extends com.aparapi.codegen.CodeGenJUni
 " return;\n" +
 " }\n" +
 " void com_aparapi_codegen_test_CompositeArbitraryScope__t2(This *this){\n" +
-" int gid = get_global_id(0);\n" +
+" int gid = (int)get_global_id(0);\n" +
 " int numRemaining = 1;\n" +
 " for (; numRemaining>0; numRemaining){\n" +
 " {\n" +
@@ -78,7 +78,7 @@ public class CompositeArbitraryScopeTest extends com.aparapi.codegen.CodeGenJUni
 " return;\n" +
 " }\n" +
 " void com_aparapi_codegen_test_CompositeArbitraryScope__t1(This *this){\n" +
-" int gid = get_global_id(0);\n" +
+" int gid = (int)get_global_id(0);\n" +
 " int numRemaining = 1;\n" +
 " while (numRemaining>0){\n" +
 " numRemaining++;\n" +
@@ -96,7 +96,7 @@ public class CompositeArbitraryScopeTest extends com.aparapi.codegen.CodeGenJUni
 " This* this=&thisStruct;\n" +
 " this->passid = passid;\n" +
 " {\n" +
-" int gid = get_global_id(0);\n" +
+" int gid = (int)get_global_id(0);\n" +
 " int numRemaining = 1;\n" +
 " com_aparapi_codegen_test_CompositeArbitraryScope__t1(this);\n" +
 " com_aparapi_codegen_test_CompositeArbitraryScope__t2(this);\n" +

diff --git a/src/test/java/com/aparapi/codegen/test/LongCastMultiply.java b/src/test/java/com/aparapi/codegen/test/LongCastMultiply.java
@@ -0,0 +1,31 @@
+/**
+ * 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 LongCastMultiply extends Kernel {
+    long[] values = new long[1];
+
+    @Override
+    public void run() {
+        values[0] = calculate(Integer.MAX_VALUE + getGlobalId());
+    }
+
+    long calculate(int value) {
+        return (long) value * 100;
+    }
+}
diff --git a/src/test/java/com/aparapi/codegen/test/LongCastMultiplyTest.java b/src/test/java/com/aparapi/codegen/test/LongCastMultiplyTest.java
@@ -0,0 +1,62 @@
+/**
+ * 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 LongCastMultiplyTest extends com.aparapi.codegen.CodeGenJUnitBase {
+    private static final String[] expectedOpenCL = {"inline ulong aparapi_umul64_lo(ulong a, ulong b){\n" +
+        "   uint a0 = (uint)a;\n" +
+        "   uint a1 = (uint)(a >> 32);\n" +
+        "   uint b0 = (uint)b;\n" +
+        "   uint b1 = (uint)(b >> 32);\n" +
+        "   uint lo = a0 * b0;\n" +
+        "   uint hi = mul_hi(a0, b0);\n" +
+        "   uint cross = hi + (a0 * b1) + (a1 * b0);\n" +
+        "   return (((ulong)cross) << 32) | (ulong)lo;\n" +
+        "}\n" +
+        "inline long aparapi_lmul(long a, long b){ return (long)aparapi_umul64_lo((ulong)a, (ulong)b); }\n" +
+        "typedef struct This_s{\n" +
+        "   __global long *values;\n" +
+        "   int passid;\n" +
+        "}This;\n" +
+        "int get_pass_id(This *this){\n" +
+        "   return this->passid;\n" +
+        "}\n" +
+        "long com_aparapi_codegen_test_LongCastMultiply__calculate(This *this, int value){\n" +
+        "   return((aparapi_lmul((long)value, 100L)));\n" +
+        "}\n" +
+        "__kernel void run(\n" +
+        "   __global long *values, \n" +
+        "   int passid\n" +
+        "){\n" +
+        "   This thisStruct;\n" +
+        "   This* this=&thisStruct;\n" +
+        "   this->values = values;\n" +
+        "   this->passid = passid;\n" +
+        "   {\n" +
+        "      this->values[0]  = com_aparapi_codegen_test_LongCastMultiply__calculate(this, (2147483647 + (int)get_global_id(0)));\n" +
+        "      return;\n" +
+        "   }\n" +
+        "}\n" +
+        "\n"};
+    private static final Class<? extends com.aparapi.internal.exception.AparapiException> expectedException = null;
+
+    @Test
+    public void LongCastMultiplyTest() {
+        test(com.aparapi.codegen.test.LongCastMultiply.class, expectedException, expectedOpenCL);
+    }
+}
diff --git a/src/test/java/com/aparapi/codegen/test/LongMultiplyCastOverflow.java b/src/test/java/com/aparapi/codegen/test/LongMultiplyCastOverflow.java
@@ -0,0 +1,31 @@
+/**
+ * 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 LongMultiplyCastOverflow extends Kernel {
+   int[] values = new int[1];
+   long[] results = new long[1];
+
+   @Override public void run() {
+      results[0] = calculate(values[0]);
+   }
+
+   long calculate(int value) {
+      return (long) value * 100;
+   }
+}
diff --git a/src/test/java/com/aparapi/codegen/test/LongMultiplyCastOverflowTest.java b/src/test/java/com/aparapi/codegen/test/LongMultiplyCastOverflowTest.java
@@ -0,0 +1,36 @@
+/**
+ * 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 static org.junit.Assert.assertFalse;
+import static org.junit.Assert.assertTrue;
+
+import com.aparapi.internal.model.ClassModel;
+import com.aparapi.internal.model.Entrypoint;
+import com.aparapi.internal.writer.KernelWriter;
+import org.junit.Test;
+
+public class LongMultiplyCastOverflowTest {
+   @Test public void LongMultiplyCastOverflowTest() throws Exception {
+      ClassModel classModel = ClassModel.createClassModel(LongMultiplyCastOverflow.class);
+      Entrypoint entrypoint = classModel.getEntrypoint("run", new LongMultiplyCastOverflow());
+      String opencl = KernelWriter.writeToString(entrypoint);
+
+      assertTrue(opencl.contains("inline ulong aparapi_umul64_lo"));
+      assertTrue(opencl.contains("aparapi_lmul((long)value, 100L)"));
+      assertFalse(opencl.contains(" * 100L"));
+   }
+}
diff --git a/src/test/java/com/aparapi/codegen/test/ObjectArrayMemberAccessTest.java b/src/test/java/com/aparapi/codegen/test/ObjectArrayMemberAccessTest.java
@@ -42,7 +42,7 @@ public class ObjectArrayMemberAccessTest extends com.aparapi.codegen.CodeGenJUni
 " this->dummy = dummy;\n" +
 " this->passid = passid;\n" +
 " {\n" +
-" int myId = get_global_id(0);\n" +
+" int myId = (int)get_global_id(0);\n" +
 " this->dummy[myId].mem=this->dummy[myId].mem + 2;\n" +
 " this->dummy[myId].floatField=this->dummy[myId].floatField + 2.0f;\n" +
 " return;\n" +

diff --git a/src/test/java/com/aparapi/codegen/test/ObjectArrayMemberCallTest.java b/src/test/java/com/aparapi/codegen/test/ObjectArrayMemberCallTest.java
@@ -52,7 +52,7 @@ public class ObjectArrayMemberCallTest extends com.aparapi.codegen.CodeGenJUnitB
         + "   this->dummy = dummy;\n"
         + "   this->passid = passid;\n"
         + "   {\n"
-        + "      int myId = get_global_id(0);\n"
+        + "      int myId = (int)get_global_id(0);\n"
         + "      this->dummy[myId].mem=com_aparapi_codegen_test_ObjectArrayMemberCall$DummyOOA__addEmUp( &(this->dummy[myId]), this->dummy[myId].mem, 2);\n"
         + "      int tmp = com_aparapi_codegen_test_ObjectArrayMemberCall$DummyOOA__addToMem( &(this->dummy[myId]), 2);\n"
         + "      int tmp2 = com_aparapi_codegen_test_ObjectArrayMemberCall$DummyOOA__addEmUpPlusOne( &(this->dummy[myId]), 2, tmp);\n"

diff --git a/src/test/java/com/aparapi/codegen/test/ObjectArrayMemberGetterSetterTest.java b/src/test/java/com/aparapi/codegen/test/ObjectArrayMemberGetterSetterTest.java
@@ -105,7 +105,7 @@ public class ObjectArrayMemberGetterSetterTest extends com.aparapi.codegen.CodeG
         + "   this->out = out;\n"
         + "   this->passid = passid;\n"
         + "   {\n"
-        + "      int myId = get_global_id(0);\n"
+        + "      int myId = (int)get_global_id(0);\n"
         + "      int tmp = com_aparapi_codegen_test_DummyOOA__getMem( &(this->dummy[myId]));\n"
         + "      com_aparapi_codegen_test_DummyOOA__setMem( &(this->dummy[myId]), (com_aparapi_codegen_test_DummyOOA__getMem( &(this->dummy[myId])) + 2));\n"
         + "      com_aparapi_codegen_test_DummyOOA__setMem( &(this->dummy[myId]), (com_aparapi_codegen_test_TheOtherOne__getMem( &(this->other[myId])) + com_aparapi_codegen_test_ObjectArrayMemberGetterSetter__getSomething(this)));\n"

diff --git a/src/test/java/com/aparapi/codegen/test/ObjectArrayMemberHierarchyTest.java b/src/test/java/com/aparapi/codegen/test/ObjectArrayMemberHierarchyTest.java
@@ -54,7 +54,7 @@ public class ObjectArrayMemberHierarchyTest extends com.aparapi.codegen.CodeGenJ
         + "   this->dummy = dummy;\n"
         + "   this->passid = passid;\n"
         + "   {\n"
-        + "      int myId = get_global_id(0);\n"
+        + "      int myId = (int)get_global_id(0);\n"
         + "      this->dummy[myId].intField=(com_aparapi_codegen_test_ObjectArrayMemberHierarchy$DummyParent__getIntField( &(this->dummy[myId])) + 2) + com_aparapi_codegen_test_ObjectArrayMemberHierarchy__getSomething(this);\n"
         + "      com_aparapi_codegen_test_ObjectArrayMemberHierarchy$DummyOOA__setFloatField( &(this->dummy[myId]), (this->dummy[myId].floatField + 2.0f));\n"
         + "      return;\n"

diff --git a/src/test/java/com/aparapi/codegen/test/ObjectRefCopyTest.java b/src/test/java/com/aparapi/codegen/test/ObjectRefCopyTest.java
@@ -37,7 +37,7 @@ public class ObjectRefCopyTest extends com.aparapi.codegen.CodeGenJUnitBase {
             + "   this->dummy = dummy;\n"
             + "   this->passid = passid;\n"
             + "   {\n"
-            + "      int myId = get_global_id(0);\n"
+            + "      int myId = (int)get_global_id(0);\n"
             + "      this->dummy[myId]  = this->dummy[(myId + 1)];\n"
             + "      return;\n"
             + "   }\n"

diff --git a/src/test/scala/com/aparapi/SimpleScalaTest.scala b/src/test/scala/com/aparapi/SimpleScalaTest.scala
@@ -75,7 +75,7 @@ class SimpleScalaTest {
                        |   this->passid = passid;
                        |   {
                        |      {
-                       |         int i = get_global_id(0);
+                       |         int i = (int)get_global_id(0);
                        |         this->result$2[i]  = (((this->inA$2[i] + this->inB$2[i]) / (this->inA$2[i] / this->inB$2[i])) * ((this->inA$2[i] - this->inB$2[i]) / (this->inA$2[i] * this->inB$2[i]))) - (((this->inB$2[i] - this->inA$2[i]) * (this->inB$2[i] + this->inA$2[i])) * ((this->inB$2[i] - this->inA$2[i]) / (this->inB$2[i] * this->inA$2[i])));
                        |      }
                        |      return;