diff --git a/src/test/java/com/aparapi/runtime/BufferTransfer.java b/src/test/java/com/aparapi/runtime/BufferTransfer.java
new file mode 100644
index 0000000000000000000000000000000000000000..00fdc3d91c2a7b43558d6745ad1724c63ec2bdbd
--- /dev/null
+++ b/src/test/java/com/aparapi/runtime/BufferTransfer.java
@@ -0,0 +1,268 @@
+/**
+ * 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 com.aparapi.Kernel;
+import com.aparapi.Range;
+import com.aparapi.device.Device;
+import com.aparapi.device.OpenCLDevice;
+import com.aparapi.internal.kernel.KernelManager;
+import org.junit.BeforeClass;
+import org.junit.Test;
+
+import java.util.Arrays;
+
+import static org.junit.Assert.assertTrue;
+import static org.junit.Assert.fail;
+
+public class BufferTransfer {
+
+ static OpenCLDevice openCLDevice = null;
+
+ @BeforeClass
+ public static void setUpBeforeClass() throws Exception {
+
+ Device device = KernelManager.instance().bestDevice();
+ if (device == null || !(device instanceof OpenCLDevice)) {
+ fail("no opencl device!");
+ }
+ openCLDevice = (OpenCLDevice) device;
+ }
+
+ @Test
+ public void inOutOnce() {
+
+ final int SIZE = 1024;
+ final InOutKernel kernel = new InOutKernel();
+ final Range range = openCLDevice.createRange(SIZE);
+
+ kernel.in = new int[SIZE];
+ kernel.out = new int[SIZE];
+
+ Util.fill(kernel.in, new Util.Filler() {
+ public void fill(int[] array, int index) {
+ array[index] = index;
+ }
+ });
+ kernel.execute(range);
+
+ assertTrue("in == out", Util.same(kernel.in, kernel.out));
+
+ }
+
+ @Test
+ public void auto() {
+
+ final int SIZE = 1024;
+ final AddKernel kernel = new AddKernel();
+ final Range range = openCLDevice.createRange(SIZE);
+
+ kernel.values = new int[SIZE];
+ kernel.result = new int[SIZE];
+ Util.zero(kernel.result);
+ Util.fill(kernel.values, new Util.Filler() {
+ public void fill(int[] array, int index) {
+ array[index] = index;
+ }
+ });
+
+ int[] expectedResult = Arrays.copyOf(kernel.result, kernel.result.length);
+
+ Util.apply(expectedResult, kernel.values, new Util.Operator() {
+
+ @Override
+ public void apply(int[] lhs, int[] rhs, int index) {
+ lhs[index] = lhs[index] + rhs[index];
+
+ }
+ });
+ kernel.execute(range);
+
+ assertTrue("expectedResult == result", Util.same(expectedResult, kernel.result));
+
+ kernel.execute(range);
+
+ Util.apply(expectedResult, kernel.values, new Util.Operator() {
+
+ @Override
+ public void apply(int[] lhs, int[] rhs, int index) {
+ lhs[index] = lhs[index] + rhs[index];
+
+ }
+ });
+ assertTrue("expectedResult == result", Util.same(expectedResult, kernel.result));
+
+ Util.zero(kernel.values);
+ kernel.execute(range);
+ assertTrue("expectedResult == result", Util.same(expectedResult, kernel.result));
+
+ }
+
+ @Test
+ public void explicit() {
+
+ final int SIZE = 1024;
+ final AddKernel kernel = new AddKernel();
+ kernel.setExplicit(true);
+ final Range range = openCLDevice.createRange(SIZE);
+
+ kernel.values = new int[SIZE];
+ kernel.result = new int[SIZE];
+ Util.zero(kernel.result);
+ Util.fill(kernel.values, new Util.Filler() {
+ public void fill(int[] array, int index) {
+ array[index] = index;
+ }
+ });
+
+ int[] expectedResult = Arrays.copyOf(kernel.result, kernel.result.length);
+
+ Util.apply(expectedResult, kernel.values, new Util.Operator() {
+
+ @Override
+ public void apply(int[] lhs, int[] rhs, int index) {
+ lhs[index] = lhs[index] + rhs[index];
+
+ }
+ });
+
+ kernel.execute(range).get(kernel.result);
+
+ assertTrue("after first explicit add expectedResult == result", Util.same(expectedResult, kernel.result));
+
+ kernel.execute(range).get(kernel.result);
+
+ Util.apply(expectedResult, kernel.values, new Util.Operator() {
+ @Override
+ public void apply(int[] lhs, int[] rhs, int index) {
+ lhs[index] = lhs[index] + rhs[index];
+
+ }
+ });
+ assertTrue("after second explicit add expectedResult == result", Util.same(expectedResult, kernel.result));
+
+ Util.zero(kernel.values);
+
+ kernel.put(kernel.values).execute(range).get(kernel.result);
+
+ assertTrue("after zeroing values and third explici add expectedResult == result", Util.same(expectedResult, kernel.result));
+
+ Util.zero(kernel.result);
+
+ kernel.put(kernel.result).execute(range).get(kernel.result);
+
+ Util.zero(expectedResult);
+
+ assertTrue("after zeroing values and result and forth explicit add expectedResult == result",
+ Util.same(expectedResult, kernel.result));
+
+ }
+
+ @Test
+ public void issue60Explicit() {
+
+ TestKernel kernel = new TestKernel();
+ kernel.setExplicit(true);
+ kernel.step();
+
+ }
+
+ @Test
+ public void issue60Auto() {
+ TestKernel kernel = new TestKernel();
+ kernel.step();
+
+ }
+
+ public static class InOutKernel extends Kernel {
+
+ int[] in;
+
+ int[] out;
+
+ @Override
+ public void run() {
+ int gid = getGlobalId(0);
+ in[gid] = out[gid];
+
+ }
+
+ }
+
+ public static class AddKernel extends Kernel {
+
+ int[] values;
+
+ int[] result;
+
+ @Override
+ public void run() {
+ int gid = getGlobalId(0);
+ result[gid] = result[gid] + values[gid];
+
+ }
+
+ }
+
+ private class TestKernel extends Kernel {
+ int[] simStep = new int[1];
+
+ int[] neuronOutputs = new int[3];
+
+ int[] expected = new int[]{
+ 3,
+ 0,
+ 0,
+ 0,
+ 3,
+ 0,
+ 0,
+ 0,
+ 3,
+ 0,
+ 0,
+ 0,
+ 3,
+ 0,
+ 0,
+ 0
+ };
+
+ public void step() {
+ int simSteps = 16;
+ int[][] log = new int[neuronOutputs.length][simSteps];
+ put(neuronOutputs);
+ for (simStep[0] = 0; simStep[0] < simSteps; simStep[0]++) {
+ put(simStep).execute(neuronOutputs.length).get(neuronOutputs);
+ for (int n = 0; n < neuronOutputs.length; n++)
+ log[n][simStep[0]] = neuronOutputs[n];
+ }
+ System.out.println(getTargetDevice().getShortDescription() + (isExplicit() ? ", explicit" : ", auto"));
+
+ for (int n = 0; n < neuronOutputs.length; n++)
+ System.out.println(Arrays.toString(log[n]));
+
+ assertTrue("log[2] == expected", Util.same(log[2], expected));
+ }
+
+ @Override
+ public void run() {
+ int neuronID = getGlobalId();
+ neuronOutputs[neuronID] = (simStep[0] % (neuronID + 2) == 0) ? (neuronID + 1) : 0;
+ }
+ }
+
+}
diff --git a/src/test/java/com/aparapi/runtime/CallStaticFromAnonymousKernel.java b/src/test/java/com/aparapi/runtime/CallStaticFromAnonymousKernel.java
new file mode 100644
index 0000000000000000000000000000000000000000..854f3edb788cb7bd62ad3822a8be7e9552f81f45
--- /dev/null
+++ b/src/test/java/com/aparapi/runtime/CallStaticFromAnonymousKernel.java
@@ -0,0 +1,75 @@
+/**
+ * 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 com.aparapi.Kernel;
+import com.aparapi.device.Device;
+import org.junit.Test;
+
+import static org.junit.Assert.assertTrue;
+
+class AnotherClass {
+ static public int foo(int i) {
+ return i + 42;
+ }
+};
+
+public class CallStaticFromAnonymousKernel {
+
+ static final int size = 256;
+
+ // This method is a static target in the anonymous
+ // kernel's containing class
+ public static int fooBar(int i) {
+ return i + 20;
+ }
+
+ public static void main(String args[]) {
+ CallStaticFromAnonymousKernel k = new CallStaticFromAnonymousKernel();
+ k.test();
+ }
+
+ @Test
+ public void test() {
+ final int[] values = new int[size];
+ final int[] results = new int[size];
+ for (int i = 0; i < size; i++) {
+ values[i] = i;
+ results[i] = 0;
+ }
+ Kernel kernel = new Kernel() {
+
+ // Verify codegen for resolving static call from run's callees
+ public int doodoo(int i) {
+ return AnotherClass.foo(i);
+ }
+
+ @Override
+ public void run() {
+ int gid = getGlobalId();
+ // Call a static in the containing class and call a kernel method
+ // that calls a static in another class
+ results[gid] = CallStaticFromAnonymousKernel.fooBar(values[gid]) + doodoo(gid);
+ }
+ };
+ kernel.execute(size);
+ assertTrue("ran on GPU", kernel.getTargetDevice().getType() == Device.TYPE.GPU);
+
+ for (int i = 0; i < size; i++) {
+ assertTrue("results == fooBar", results[i] == (fooBar(values[i]) + AnotherClass.foo(i)));
+ }
+ }
+}
diff --git a/src/test/java/com/aparapi/runtime/ExplicitBoolean.java b/src/test/java/com/aparapi/runtime/ExplicitBoolean.java
new file mode 100644
index 0000000000000000000000000000000000000000..6626162398b8970574708e941bd40a751c71c9ed
--- /dev/null
+++ b/src/test/java/com/aparapi/runtime/ExplicitBoolean.java
@@ -0,0 +1,88 @@
+/**
+ * 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 com.aparapi.Kernel;
+import org.junit.Test;
+
+import static org.junit.Assert.assertTrue;
+
+public class ExplicitBoolean {
+
+ private static void printArray(boolean[] a) {
+ for (int i = 0; i < a.length; i++) {
+ System.out.print((a[i] ? 1 : 0) + "\t");
+ }
+ System.out.println();
+ }
+
+ @Test
+ public void test() {
+ int size = 16;
+ ExplicitBooleanTestKernel k1 = new ExplicitBooleanTestKernel(size);
+ ExplicitBooleanTestKernel k2 = new ExplicitBooleanTestKernel(size);
+ k2.input = k1.output;
+
+ for (int i = 0; i < size; i++) {
+ k1.input[i] = Math.random() > 0.5;
+ }
+
+ if (size <= 32)
+ printArray(k1.input);
+
+ k1.go();
+
+ if (size <= 32)
+ printArray(k1.output);
+
+ assertTrue("k1.input == k1.output ", Util.same(k1.output, k1.output));
+
+ k2.go();
+
+ if (size <= 32)
+ printArray(k2.output);
+
+ assertTrue("k1.input == k2.input", Util.same(k1.output, k1.output));
+ System.out.println(k1.getTargetDevice().getShortDescription());
+ }
+
+ class ExplicitBooleanTestKernel extends Kernel {
+ public boolean[] input, output;
+ int size; // Number of work items.
+ int iterations; // Number of times to execute kernel.
+
+ public ExplicitBooleanTestKernel(int _size) {
+ size = _size;
+ input = new boolean[size];
+ output = new boolean[size];
+ setExplicit(true);
+ put(output);
+ }
+
+ public void go() {
+ put(input);
+ execute(size);
+ get(output);
+ }
+
+ @Override
+ public void run() {
+ int id = getGlobalId();
+ output[id] = input[id];
+ }
+ }
+
+}
diff --git a/src/test/java/com/aparapi/runtime/Issue102.java b/src/test/java/com/aparapi/runtime/Issue102.java
new file mode 100644
index 0000000000000000000000000000000000000000..88b573132536dabea239edd39b1231c302fb6fbf
--- /dev/null
+++ b/src/test/java/com/aparapi/runtime/Issue102.java
@@ -0,0 +1,73 @@
+/**
+ * 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 com.aparapi.Kernel;
+import org.junit.Test;
+
+import static org.junit.Assert.assertTrue;
+
+
+final class BugDataObject {
+ int value = 7;
+
+ public int getValue() {
+ return value;
+ }
+
+ public void setValue(int value) {
+ this.value = value;
+ }
+}
+
+
+public class Issue102 extends Kernel {
+ static final int size = 32;
+
+ static BugDataObject[] objects = new BugDataObject[size];
+ int[] target = new int[size];
+
+ public Issue102() {
+ for (int i = 0; i < size; ++i) {
+ objects[i] = new BugDataObject();
+ target[i] = 99;
+ }
+ }
+
+ public static void main(String[] args) {
+ Issue102 b = new Issue102();
+ b.test();
+ }
+
+ @Override
+ public void run() {
+ int id = getGlobalId();
+ target[id] = objects[id].getValue();
+ }
+
+ void validate() {
+ for (int i = 0; i < size; i++) {
+ System.out.println(target[i] + " ... " + objects[i].getValue());
+ assertTrue("target == objects", target[i] == objects[i].getValue());
+ }
+ }
+
+ @Test
+ public void test() {
+ execute(size);
+ validate();
+ }
+}
diff --git a/src/test/java/com/aparapi/runtime/Issue103.java b/src/test/java/com/aparapi/runtime/Issue103.java
new file mode 100644
index 0000000000000000000000000000000000000000..cf1ebad3361667ab6ed93871571bbb18463e3b2f
--- /dev/null
+++ b/src/test/java/com/aparapi/runtime/Issue103.java
@@ -0,0 +1,61 @@
+/**
+ * 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 com.aparapi.Kernel;
+import org.junit.Test;
+
+import static org.junit.Assert.assertArrayEquals;
+
+
+public class Issue103 extends Kernel {
+ static final int size = 32;
+
+ static int[] source = new int[size];
+ static int[] target = new int[size];
+
+ public Issue103() {
+ for (int i = 0; i < size; ++i) {
+ source[i] = 7;
+ target[i] = 99;
+ }
+ }
+
+ public static void main(String[] args) {
+ Issue103 b = new Issue103();
+ b.test();
+ }
+
+ @Override
+ public void run() {
+ int id = getGlobalId();
+ target[id] = source[id];
+ }
+
+ void validate() {
+ assertArrayEquals("target == source", target, source);
+// for (int i = 0; i < size; i++) {
+// System.out.println(target[i] + " ... " + source[i]);
+// assertTrue("target == source", target[i] == source[i]);
+// }
+ }
+
+ @Test
+ public void test() {
+ execute(size);
+ validate();
+ }
+}
diff --git a/src/test/java/com/aparapi/runtime/Issue68.java b/src/test/java/com/aparapi/runtime/Issue68.java
new file mode 100644
index 0000000000000000000000000000000000000000..ac302127a8948051a006df544f7a279dfc24b4d1
--- /dev/null
+++ b/src/test/java/com/aparapi/runtime/Issue68.java
@@ -0,0 +1,234 @@
+/**
+ * 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 com.aparapi.Kernel;
+
+abstract class ArrayAccess {
+ private final int offset;
+ private final int length;
+
+ protected ArrayAccess(int offset, int length) {
+ this.offset = offset;
+ this.length = length;
+ }
+
+ public abstract int[] getIntData();
+
+ public int getOffset() {
+ return offset;
+ }
+
+ public int getLength() {
+ return length;
+ }
+}
+
+class IntMemoryArrayAccess extends ArrayAccess {
+ private final int[] data;
+
+ public IntMemoryArrayAccess(int[] data, int offset, int length) {
+ super(offset, length);
+ this.data = data;
+ }
+
+ @Override
+ public int[] getIntData() {
+ return data;
+ }
+}
+
+public class Issue68 {
+ public static void main(String[] args) {
+ final int SQRT_LENGTH = 1024;
+ final int LENGTH = SQRT_LENGTH * SQRT_LENGTH;
+ final ArrayAccess arrayAccess = new IntMemoryArrayAccess(new int[LENGTH], 0, LENGTH);
+ new Issue68().transformColumns(SQRT_LENGTH, SQRT_LENGTH, false, arrayAccess, new int[SQRT_LENGTH], null);
+ }
+
+ private void transformColumns(final int length, final int count, final boolean isInverse, final ArrayAccess arrayAccess,
+ final int[] wTable, final int[] permutationTable) {
+ final Kernel kernel = new ColumnTableFNTRunnable(length, isInverse, arrayAccess, wTable, permutationTable, getModulus());
+ kernel.execute(count);
+ }
+
+ private int getModulus() {
+ return 2113929217;
+ }
+
+ // Runnable for calculating the column transforms in parallel
+ private class ColumnTableFNTRunnable extends Kernel {
+ private final int stride;
+ private final int length;
+ private final boolean isInverse;
+ private final int[] data;
+ private final int offset;
+ @Constant
+ private final int[] wTable;
+ @Constant
+ private final int[] permutationTable;
+ private final int permutationTableLength;
+ private int modulus;
+ private float inverseModulus;
+
+ public ColumnTableFNTRunnable(int length, boolean isInverse, ArrayAccess arrayAccess, int[] wTable, int[] permutationTable,
+ int modulus) {
+ stride = arrayAccess.getLength() / length;
+ this.length = length; // Transform length
+ this.isInverse = isInverse;
+ data = arrayAccess.getIntData();
+ offset = arrayAccess.getOffset();
+ this.wTable = wTable;
+ this.permutationTable = permutationTable;
+ permutationTableLength = (permutationTable == null ? 0 : permutationTable.length);
+ setModulus(modulus);
+ }
+
+ @Override
+ public void run() {
+ if (isInverse) {
+ inverseColumnTableFNT();
+ } else {
+ columnTableFNT();
+ }
+ }
+
+ private void columnTableFNT() {
+ int nn, istep, mmax, r;
+
+ final int offset = this.offset + getGlobalId();
+ nn = length;
+
+ if (nn < 2) {
+ return;
+ }
+
+ r = 1;
+ mmax = nn >> 1;
+ while (mmax > 0) {
+ istep = mmax << 1;
+
+ // Optimize first step when wr = 1
+
+ for (int i = offset; i < (offset + (nn * stride)); i += istep * stride) {
+ final int j = i + (mmax * stride);
+ final int a = data[i];
+ final int b = data[j];
+ data[i] = modAdd(a, b);
+ data[j] = modSubtract(a, b);
+ }
+
+ int t = r;
+
+ for (int m = 1; m < mmax; m++) {
+ for (int i = offset + (m * stride); i < (offset + (nn * stride)); i += istep * stride) {
+ final int j = i + (mmax * stride);
+ final int a = data[i];
+ final int b = data[j];
+ data[i] = modAdd(a, b);
+ data[j] = modMultiply(wTable[t], modSubtract(a, b));
+ }
+ t += r;
+ }
+ r <<= 1;
+ mmax >>= 1;
+ }
+
+ //if (permutationTable != null)
+ // {
+ columnScramble(offset);
+ // }
+ }
+
+ private void inverseColumnTableFNT() {
+ int nn, istep, mmax, r;
+
+ final int offset = this.offset + getGlobalId();
+ nn = length;
+
+ if (nn < 2) {
+ return;
+ }
+
+ // if (permutationTable != null)
+ // {
+ columnScramble(offset);
+ // }
+
+ r = nn;
+ mmax = 1;
+ istep = 0;
+ while (nn > mmax) {
+ istep = mmax << 1;
+ r >>= 1;
+
+ // Optimize first step when w = 1
+
+ for (int i = offset; i < (offset + (nn * stride)); i += istep * stride) {
+ final int j = i + (mmax * stride);
+ final int wTemp = data[j];
+ data[j] = modSubtract(data[i], wTemp);
+ data[i] = modAdd(data[i], wTemp);
+ }
+
+ int t = r;
+
+ for (int m = 1; m < mmax; m++) {
+ for (int i = offset + (m * stride); i < (offset + (nn * stride)); i += istep * stride) {
+ final int j = i + (mmax * stride);
+ final int wTemp = modMultiply(wTable[t], data[j]);
+ data[j] = modSubtract(data[i], wTemp);
+ data[i] = modAdd(data[i], wTemp);
+ }
+ t += r;
+ }
+ mmax = istep;
+ }
+ }
+
+ private void columnScramble(int offset) {
+ for (int k = 0; k < permutationTableLength; k += 2) {
+ final int i = offset + (permutationTable[k] * stride), j = offset + (permutationTable[k + 1] * stride);
+ final int tmp = data[i];
+ data[i] = data[j];
+ data[j] = tmp;
+ }
+ }
+
+ public final int modMultiply(int a, int b) {
+ final int r1 = (a * b) - ((int) (inverseModulus * a * b) * modulus), r2 = r1 - modulus;
+
+ return (r2 < 0 ? r1 : r2);
+ }
+
+ private int modAdd(int a, int b) {
+ final int r1 = a + b, r2 = r1 - modulus;
+
+ return (r2 < 0 ? r1 : r2);
+ }
+
+ private int modSubtract(int a, int b) {
+ final int r1 = a - b, r2 = r1 + modulus;
+
+ return (r1 < 0 ? r2 : r1);
+ }
+
+ private void setModulus(int modulus) {
+ inverseModulus = 1.0f / (modulus + 0.5f); // Round down
+ this.modulus = modulus;
+ }
+ }
+}
diff --git a/src/test/java/com/aparapi/runtime/Issue69.java b/src/test/java/com/aparapi/runtime/Issue69.java
new file mode 100644
index 0000000000000000000000000000000000000000..3c12deb015404873dad921ba2b7c93095bb178d2
--- /dev/null
+++ b/src/test/java/com/aparapi/runtime/Issue69.java
@@ -0,0 +1,51 @@
+/**
+ * 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 com.aparapi.Kernel;
+import com.aparapi.Range;
+
+public class Issue69 {
+
+ public static void main(String[] args) {
+ final int globalArray[] = new int[512];
+ Kernel kernel = new Kernel() {
+ @Override
+ public void run() {
+ globalArray[getGlobalId()] = getGlobalId();
+ }
+ };
+ for (int loop = 0; loop < 100; loop++) {
+
+ System.out.printf("%3d free = %10d\n", loop, Runtime.getRuntime().freeMemory());
+ kernel.execute(Range.create(512, 64), 1);
+ for (int i = 0; i < globalArray.length; ++i) {
+ if (globalArray[i] != i)
+ System.err.println("Wrong!");
+ }
+ }
+ for (int loop = 0; loop < 100; loop++) {
+
+ System.out.printf("%3d free = %10d\n", loop, Runtime.getRuntime().freeMemory());
+ kernel.execute(Range.create(512, 64), 2);
+ for (int i = 0; i < globalArray.length; ++i) {
+ if (globalArray[i] != i)
+ System.err.println("Wrong!");
+ }
+ }
+ }
+
+}
diff --git a/src/test/java/com/aparapi/runtime/LoadCL.java b/src/test/java/com/aparapi/runtime/LoadCL.java
new file mode 100644
index 0000000000000000000000000000000000000000..49fa4066844386f2460c0a8322754cd5e04f1176
--- /dev/null
+++ b/src/test/java/com/aparapi/runtime/LoadCL.java
@@ -0,0 +1,71 @@
+/**
+ * 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 com.aparapi.Range;
+import com.aparapi.device.Device;
+import com.aparapi.device.OpenCLDevice;
+import com.aparapi.internal.kernel.KernelManager;
+import com.aparapi.opencl.OpenCL;
+import com.aparapi.opencl.OpenCL.Resource;
+import org.junit.Test;
+
+import static org.junit.Assert.assertTrue;
+
+public class LoadCL {
+
+ @Test
+ public void test() {
+ final int size = 32;
+ final float[] in = new float[size];
+
+ for (int i = 0; i < size; i++) {
+ in[i] = i;
+ }
+
+ final float[] squares = new float[size];
+ final float[] quads = new float[size];
+ final Range range = Range.create(size);
+
+ final Device device = KernelManager.instance().bestDevice();
+
+ if (device instanceof OpenCLDevice) {
+ final OpenCLDevice openclDevice = (OpenCLDevice) device;
+
+ final Squarer squarer = openclDevice.bind(Squarer.class);
+ squarer.square(range, in, squares);
+
+ for (int i = 0; i < size; i++) {
+ assertTrue("in[" + i + "] * in[" + i + "] = in[" + i + "]^2", in[i] * in[i] == squares[i]);
+ }
+
+ squarer.square(range, squares, quads);
+
+ for (int i = 0; i < size; i++) {
+ assertTrue("in[" + i + "]^2 * in[" + i + "]^2 = in[" + i + "]^4", in[i] * in[i] * in[i] * in[i] == quads[i]);
+ }
+ }
+ }
+
+ @Resource("com/aparapi/test/runtime/squarer.cl")
+ interface Squarer extends OpenCL<Squarer> {
+ public Squarer square(//
+ Range _range,//
+ @GlobalReadWrite("in") float[] in,//
+ @GlobalReadWrite("out") float[] out);
+ }
+}
+
diff --git a/src/test/java/com/aparapi/runtime/RangeSize.java b/src/test/java/com/aparapi/runtime/RangeSize.java
new file mode 100644
index 0000000000000000000000000000000000000000..0cac50f3c99125c78025d01f36ae0cc00d6a8d8f
--- /dev/null
+++ b/src/test/java/com/aparapi/runtime/RangeSize.java
@@ -0,0 +1,43 @@
+/**
+ * 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 com.aparapi.Range;
+import org.junit.Test;
+
+import static org.junit.Assert.assertTrue;
+
+public class RangeSize {
+
+ @Test
+ public void test384x384() {
+ Range range = Range.create2D(384, 384);
+ System.out.println("local[0] " + range.getLocalSize(0));
+ System.out.println("local[1] " + range.getLocalSize(1));
+ System.out.println("workGroupSize " + range.getWorkGroupSize());
+ assertTrue("Range > max work size", range.getLocalSize(0) * range.getLocalSize(1) <= range.getWorkGroupSize());
+ }
+
+ @Test
+ public void test384x320() {
+ Range range = Range.create2D(384, 320);
+ System.out.println("local[0] " + range.getLocalSize(0));
+ System.out.println("local[1] " + range.getLocalSize(1));
+ System.out.println("workGroupSize " + range.getWorkGroupSize());
+ assertTrue("Range > max work size", range.getLocalSize(0) * range.getLocalSize(1) <= range.getWorkGroupSize());
+ }
+
+}
diff --git a/src/test/java/com/aparapi/runtime/Test12x4_4x2.java b/src/test/java/com/aparapi/runtime/Test12x4_4x2.java
new file mode 100644
index 0000000000000000000000000000000000000000..a1a72032cf9354788598bd29a6c312a4225ce277
--- /dev/null
+++ b/src/test/java/com/aparapi/runtime/Test12x4_4x2.java
@@ -0,0 +1,521 @@
+/**
+ * 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 com.aparapi.Kernel;
+import com.aparapi.Range;
+import com.aparapi.device.Device;
+import org.junit.Test;
+
+public class Test12x4_4x2 {
+ @SuppressWarnings("deprecation")
+ @Test
+ public void test() {
+ // globalThreadId, threadId, globalX, globalY, localX, localY
+ final int[][] test = new int[][]{
+ {
+ 0, //globalThreadId
+ 0,//threadId
+ 0,//globalX
+ 0,//globalY
+ 0,//localX
+ 0
+ //localY
+ },
+ {
+ 1,//globalThreadId
+ 1,//threadId
+ 1,//globalX
+ 0,//globalY
+ 1,//localX
+ 0
+ //localY
+ },
+ {
+ 2,//globalThreadId
+ 2,//threadId
+ 2,//globalX
+ 0,//globalY
+ 2,//localX
+ 0
+ //localY
+ },
+ {
+ 3,//globalThreadId
+ 3,//threadId
+ 3,//globalX
+ 0,//globalY
+ 3,//localX
+ 0
+ //localY
+ },
+ {
+ 4,//globalThreadId
+ 4,//threadId
+ 0,//globalX
+ 1,//globalY
+ 0,//localX
+ 1
+ //localY
+ },
+ {
+ 5,//globalThreadId
+ 5,//threadId
+ 1,//globalX
+ 1,//globalY
+ 1,//localX
+ 1
+ //localY
+ },
+ {
+ 6,//globalThreadId
+ 6,//threadId
+ 2,//globalX
+ 1,//globalY
+ 2,//localX
+ 1
+ //localY
+ },
+ {
+ 7,//globalThreadId
+ 7,//threadId
+ 3,//globalX
+ 1,//globalY
+ 3,//localX
+ 1
+ //localY
+ },
+ {
+ 8,//globalThreadId
+ 0,//threadId
+ 4,//globalX
+ 0,//globalY
+ 0,//localX
+ 0
+ //localY
+ },
+ {
+ 9,//globalThreadId
+ 1,//threadId
+ 5,//globalX
+ 0,//globalY
+ 1,//localX
+ 0
+ //localY
+ },
+ {
+ 10,//globalThreadId
+ 2,//threadId
+ 6,//globalX
+ 0,//globalY
+ 2,//localX
+ 0
+ //localY
+ },
+ {
+ 11,//globalThreadId
+ 3,//threadId
+ 7,//globalX
+ 0,//globalY
+ 3,//localX
+ 0
+ //localY
+ },
+ {
+ 12,//globalThreadId
+ 4,//threadId
+ 4,//globalX
+ 1,//globalY
+ 0,//localX
+ 1
+ //localY
+ },
+ {
+ 13,//globalThreadId
+ 5,//threadId
+ 5,//globalX
+ 1,//globalY
+ 1,//localX
+ 1
+ //localY
+ },
+ {
+ 14,//globalThreadId
+ 6,//threadId
+ 6,//globalX
+ 1,//globalY
+ 2,//localX
+ 1
+ //localY
+ },
+ {
+ 15,//globalThreadId
+ 7,//threadId
+ 7,//globalX
+ 1,//globalY
+ 3,//localX
+ 1
+ //localY
+ },
+ {
+ 16,//globalThreadId
+ 0,//threadId
+ 8,//globalX
+ 0,//globalY
+ 0,//localX
+ 0
+ //localY
+ },
+ {
+ 17,//globalThreadId
+ 1,//threadId
+ 9,//globalX
+ 0,//globalY
+ 1,//localX
+ 0
+ //localY
+ },
+ {
+ 18,//globalThreadId
+ 2,//threadId
+ 10,//globalX
+ 0,//globalY
+ 2,//localX
+ 0
+ //localY
+ },
+ {
+ 19,//globalThreadId
+ 3,//threadId
+ 11,//globalX
+ 0,//globalY
+ 3,//localX
+ 0
+ //localY
+ },
+
+ {
+ 20,//globalThreadId
+ 4,//threadId
+ 8,//globalX
+ 1,//globalY
+ 0,//localX
+ 1
+ //localY
+ },
+ {
+ 21,//globalThreadId
+ 5,//threadId
+ 9,//globalX
+ 1,//globalY
+ 1,//localX
+ 1
+ //localY
+ },
+ {
+ 22,//globalThreadId
+ 6,//threadId
+ 10,//globalX
+ 1,
+ 2,//localX
+ 1
+ //localY
+ },
+ {
+ 23,//globalThreadId
+ 7,//threadId
+ 11,//globalX
+ 1,//globalY
+ 3,//localX
+ 1
+ //localY
+ },
+ {
+ 24,//globalThreadId
+ 0,//threadId
+ 0,//globalX
+ 2,//globalY
+ 0,//localX
+ 0
+ //localY
+ },
+ {
+ 25,//globalThreadId
+ 1,//threadId
+ 1,//globalX
+ 2,//globalY
+ 1,//localX
+ 0
+ //localY
+ },
+ {
+ 26,//globalThreadId
+ 2,//threadId
+ 2,//globalX
+ 2,//globalY
+ 2,//localX
+ 0
+ //localY
+ },
+ {
+ 27,//globalThreadId
+ 3,//threadId
+ 3,//globalX
+ 2,//globalY
+ 3,//localX
+ 0
+ //localY
+ },
+ {
+ 28,//globalThreadId
+ 4,//threadId
+ 0,//globalX
+ 3,//globalY
+ 0,//localX
+ 1
+ //localY
+ },
+ {
+ 29,//globalThreadId
+ 5,//threadId
+ 1,//globalX
+ 3,//globalY
+ 1,//localX
+ 1
+ //localY
+ },
+ {
+ 30,//globalThreadId
+ 6,//threadId
+ 2,//globalX
+ 3,//globalY
+ 2,//localX
+ 1
+ //localY
+ },
+ {
+ 31,//globalThreadId
+ 7,//threadId
+ 3,//globalX
+ 3,//globalY
+ 3,//localX
+ 1
+ //localY
+ },
+ {
+ 32,//globalThreadId
+ 0,//threadId
+ 4,//globalX
+ 2,//globalY
+ 0,//localX
+ 0
+ //localY
+ },
+ {
+ 33,//globalThreadId
+ 1,//threadId
+ 5,//globalX
+ 2,//globalY
+ 1,//localX
+ 0
+ //localY
+ },
+ {
+ 34,//globalThreadId
+ 2,//threadId
+ 6,//globalX
+ 2,//globalY
+ 2,//localX
+ 0
+ //localY
+ },
+ {
+ 35,//globalThreadId
+ 3,//threadId
+ 7,//globalX
+ 2,//globalY
+ 3,//localX
+ 0
+ //localY
+ },
+ {
+ 36,//globalThreadId
+ 4,//threadId
+ 4,//globalX
+ 3,//globalY
+ 0,//localX
+ 1
+ //localY
+ },
+ {
+ 37,//globalThreadId
+ 5,//threadId
+ 5,//globalX
+ 3,//globalY
+ 1,//localX
+ 1
+ //localY
+ },
+ {
+ 38,//globalThreadId
+ 6,//threadId
+ 6,//globalX
+ 3,//globalY
+ 2,//localX
+ 1
+ //localY
+ },
+ {
+ 39,//globalThreadId
+ 7,//threadId
+ 7,//globalX
+ 3,//globalY
+ 3,//localX
+ 1
+ //localY
+ },
+ {
+ 40,//globalThreadId
+ 0,//threadId
+ 8,//globalX
+ 2,//globalY
+ 0,//localX
+ 0
+ //localY
+ },
+ {
+ 41,//globalThreadId
+ 1,//threadId
+ 9,//globalX
+ 2,//globalY
+ 1,//localX
+ 0
+ //localY
+ },
+ {
+ 42,//globalThreadId
+ 2,//threadId
+ 10,//globalX
+ 2,//globalY
+ 2,//localX
+ 0
+ //localY
+ },
+ {
+ 43,//globalThreadId
+ 3,//threadId
+ 11,//globalX
+ 2,//globalY
+ 3,//localX
+ 0
+ //localY
+ },
+
+ {
+ 44,//globalThreadId
+ 4,//threadId
+ 8,//globalX
+ 3,//globalY
+ 0,//localX
+ 1
+ //localY
+ },
+ {
+ 45,//globalThreadId
+ 5,//threadId
+ 9,//globalX
+ 3,//globalY
+ 1,//localX
+ 1
+ //localY
+ },
+ {
+ 46,//globalThreadId
+ 6,//threadId
+ 10,//globalX
+ 3,//globalY
+ 2,//localX
+ 1
+ //localY
+ },
+ {
+ 47,//globalThreadId
+ 7,//threadId
+ 11,//globalX
+ 3,//globalY
+ 3,//localX
+ 1
+ //localY
+ },
+ };
+ Kernel kernel = new Kernel() {
+
+ @Override
+ public boolean isAllowDevice(Device _device) {
+ return _device.getType() == Device.TYPE.JTP;
+ }
+
+ @Override
+ public void run() {
+ int x = getGlobalId(0);
+ int y = getGlobalId(1);
+ int lx = getLocalId(0);
+ int ly = getLocalId(1);
+ int w = getGlobalSize(0);
+ int h = getGlobalSize(1);
+ int globalThreadId = getGlobalId(1) * getGlobalSize(0) + getGlobalId(0);
+ int threadId = getLocalId(1) * getLocalSize(0) + getLocalId(0);
+ synchronized (test) {
+ boolean show = false;
+ if (globalThreadId != test[globalThreadId][0]) {
+ System.out.println("bad globalThreadId");
+ show = true;
+ }
+ if (threadId != test[globalThreadId][1]) {
+ System.out.println("bad threadId");
+ show = true;
+ }
+ if (x != test[globalThreadId][2]) {
+ System.out.println("bad globalx");
+ show = true;
+ }
+ if (y != test[globalThreadId][3]) {
+ System.out.println("bad globaly");
+ show = true;
+ }
+ if (lx != test[globalThreadId][4]) {
+ System.out.println("bad localx");
+ show = true;
+ }
+ if (ly != test[globalThreadId][5]) {
+ System.out.println("bad localy");
+ show = true;
+ }
+ if (show) {
+ System.out.println("derived =>" + globalThreadId + " " + threadId + " " + x + "," + y + " " + lx + "," + ly + " "
+ + w + "," + h);
+ System.out.println("data =>" + test[globalThreadId][0] + " " + test[globalThreadId][1] + " "
+ + test[globalThreadId][2] + "," + test[globalThreadId][3] + " " + test[globalThreadId][4] + ","
+ + test[globalThreadId][5] + " " + w + "," + h);
+ }
+ }
+ }
+
+ };
+ kernel.execute(Range.create2D(12, 4, 4, 2));
+
+ }
+}
diff --git a/src/test/java/com/aparapi/runtime/UseStaticArray.java b/src/test/java/com/aparapi/runtime/UseStaticArray.java
new file mode 100644
index 0000000000000000000000000000000000000000..2b0f631e1958bd1349b872495b07f9b4bf9e2b13
--- /dev/null
+++ b/src/test/java/com/aparapi/runtime/UseStaticArray.java
@@ -0,0 +1,61 @@
+/**
+ * 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 com.aparapi.Kernel;
+import com.aparapi.device.Device;
+import org.junit.Test;
+
+import static org.junit.Assert.assertArrayEquals;
+import static org.junit.Assert.assertTrue;
+
+public class UseStaticArray extends Kernel {
+
+ static final int size = 256;
+
+ static final int[] values = new int[size];
+
+ static final int[] results = new int[size];
+
+ public static void main(String args[]) {
+ UseStaticArray k = new UseStaticArray();
+ k.test();
+ }
+
+ @Override
+ public void run() {
+ int gid = getGlobalId();
+ results[gid] = values[gid];
+ }
+
+ @Test
+ public void test() {
+
+ for (int i = 0; i < size; i++) {
+ values[i] = i;
+ results[i] = 0;
+ }
+
+ execute(size);
+
+ assertTrue("ran on GPU", getTargetDevice().getType() == Device.TYPE.GPU);
+
+ assertArrayEquals("results == fooBar", results, values);
+// for (int i = 0; i < size; i++) {
+// assertTrue("results == fooBar", results[i] == values[i]);
+// }
+ }
+}
diff --git a/src/test/java/com/aparapi/runtime/Util.java b/src/test/java/com/aparapi/runtime/Util.java
new file mode 100644
index 0000000000000000000000000000000000000000..c354f1ed379ebad2d9460efba9e75b658da0eafd
--- /dev/null
+++ b/src/test/java/com/aparapi/runtime/Util.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.runtime;
+
+import java.util.Arrays;
+
+public class Util {
+ static void fill(int[] array, Filler _filler) {
+ for (int i = 0; i < array.length; i++) {
+ _filler.fill(array, i);
+ }
+ }
+
+ static boolean same(int[] lhs, int[] rhs, Comparer _comparer) {
+ boolean same = lhs != null && rhs != null && lhs.length == rhs.length;
+ for (int i = 0; same && i < lhs.length; i++) {
+ same = _comparer.same(lhs, rhs, i);
+ }
+ return (same);
+ }
+
+ static void zero(int[] array) {
+ Arrays.fill(array, 0);
+ }
+
+ static boolean same(int[] lhs, int[] rhs) {
+ return (same(lhs, rhs, new Comparer() {
+
+ @Override
+ public boolean same(int[] lhs, int[] rhs, int index) {
+
+ return lhs[index] == rhs[index];
+ }
+ }));
+ }
+
+ static boolean same(boolean[] lhs, boolean[] rhs) {
+ boolean same = lhs != null && rhs != null && lhs.length == rhs.length;
+ for (int i = 0; same && i < lhs.length; i++) {
+ same = lhs[i] == rhs[i];
+ }
+ return (same);
+ }
+
+ static void apply(int[] lhs, int[] rhs, Operator _operator) {
+ for (int i = 0; i < lhs.length; i++) {
+ _operator.apply(lhs, rhs, i);
+ }
+ }
+
+ interface Filler {
+ void fill(int[] array, int index);
+ }
+
+ interface Comparer {
+ boolean same(int[] lhs, int[] rhs, int index);
+ }
+
+ interface Operator {
+ void apply(int[] lhs, int[] rhs, int index);
+ }
+
+}
diff --git a/src/test/java/com/aparapi/runtime/squarer.cl b/src/test/java/com/aparapi/runtime/squarer.cl
new file mode 100644
index 0000000000000000000000000000000000000000..2ccc8b07158c0613c44d2818af082d017c3df475
--- /dev/null
+++ b/src/test/java/com/aparapi/runtime/squarer.cl
@@ -0,0 +1,6 @@
+
+__kernel void square( __global float *in, __global float *out){
+ const size_t id = get_global_id(0);
+ out[id] = in[id]*in[id];
+}
+