diff --git a/src/main/java/com/aparapi/internal/writer/KernelWriter.java b/src/main/java/com/aparapi/internal/writer/KernelWriter.java index 0a72a33f..c914ba85 100644 --- 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 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 index c968ae3f..eb87de42 100644 --- 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 new file mode 100644 index 00000000..8249a24b --- /dev/null +++ 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 new file mode 100644 index 00000000..a2bb2c93 --- /dev/null +++ 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 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 new file mode 100644 index 00000000..b1135a69 --- /dev/null +++ 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 new file mode 100644 index 00000000..565a54ab --- /dev/null +++ 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 index 6d26bab1..25a77324 100644 --- 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 index 1876a41e..82ab651b 100644 --- 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 index 62c1b98a..430aa77f 100644 --- 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 index 66587d10..e218f032 100644 --- 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 index ca86a118..cdf8aefb 100644 --- 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 index 945598ac..da4f31ef 100644 --- 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;