diff --git a/src/main/java/com/aparapi/Kernel.java b/src/main/java/com/aparapi/Kernel.java
index 4b9686db..67097861 100644
--- a/src/main/java/com/aparapi/Kernel.java
+++ b/src/main/java/com/aparapi/Kernel.java
@@ -2473,6 +2473,20 @@ protected final void localBarrier() {
kernelState.awaitOnLocalBarrier();
}
+ /**
+ * Wait for all kernels in the current work group to rendezvous at this call before continuing execution.
+ * This is a CUDA-style alias for {@link #localBarrier()} and maps to
+ * barrier(CLK_LOCAL_MEM_FENCE) in OpenCL.
+ *
+ * @annotion Experimental
+ * @see #localBarrier()
+ */
+ @OpenCLDelegate
+ @Experimental
+ protected final void syncThreads() {
+ kernelState.awaitOnLocalBarrier();
+ }
+
/**
* Wait for all kernels in the current work group to rendezvous at this call before continuing execution.
* It will also enforce memory ordering, such that modifications made by each thread in the work-group, to the memory,
diff --git a/src/main/java/com/aparapi/internal/writer/KernelWriter.java b/src/main/java/com/aparapi/internal/writer/KernelWriter.java
index 0a72a33f..0caecb20 100644
--- a/src/main/java/com/aparapi/internal/writer/KernelWriter.java
+++ b/src/main/java/com/aparapi/internal/writer/KernelWriter.java
@@ -159,6 +159,8 @@ public abstract class KernelWriter extends BlockWriter{
javaToCLIdentifierMap.put("localBarrier()V", "barrier(CLK_LOCAL_MEM_FENCE)");
+ javaToCLIdentifierMap.put("syncThreads()V", "barrier(CLK_LOCAL_MEM_FENCE)");
+
javaToCLIdentifierMap.put("globalBarrier()V", "barrier(CLK_GLOBAL_MEM_FENCE)");
javaToCLIdentifierMap.put("localGlobalBarrier()V", "barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE)");
diff --git a/src/test/java/com/aparapi/codegen/test/SyncThreads.java b/src/test/java/com/aparapi/codegen/test/SyncThreads.java
new file mode 100644
index 00000000..c64393cd
--- /dev/null
+++ b/src/test/java/com/aparapi/codegen/test/SyncThreads.java
@@ -0,0 +1,46 @@
+/**
+ * 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 SyncThreads extends Kernel {
+ @Override
+ public void run() {
+ syncThreads();
+ }
+}
+/**{OpenCL{
+
+ typedef struct This_s{
+ int passid;
+ }This;
+ int get_pass_id(This *this){
+ return this->passid;
+ }
+ __kernel void run(
+ int passid
+ ){
+ This thisStruct;
+ This* this=&thisStruct;
+ this->passid = passid;
+ {
+ barrier(CLK_LOCAL_MEM_FENCE);
+ return;
+ }
+ }
+
+ }OpenCL}**/
diff --git a/src/test/java/com/aparapi/codegen/test/SyncThreadsTest.java b/src/test/java/com/aparapi/codegen/test/SyncThreadsTest.java
new file mode 100644
index 00000000..941bb3e1
--- /dev/null
+++ b/src/test/java/com/aparapi/codegen/test/SyncThreadsTest.java
@@ -0,0 +1,50 @@
+/**
+ * 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 SyncThreadsTest extends com.aparapi.codegen.CodeGenJUnitBase {
+ private static final String[] expectedOpenCL = {
+ "typedef struct This_s{\n"
+ + " int passid;\n"
+ + " }This;\n"
+ + " int get_pass_id(This *this){\n"
+ + " return this->passid;\n"
+ + " }\n"
+ + " __kernel void run(\n"
+ + " int passid\n"
+ + " ){\n"
+ + " This thisStruct;\n"
+ + " This* this=&thisStruct;\n"
+ + " this->passid = passid;\n"
+ + " {\n"
+ + " barrier(CLK_LOCAL_MEM_FENCE);\n"
+ + " return;\n"
+ + " }\n"
+ + " }"};
+ private static final Class expectedException = null;
+
+ @Test
+ public void SyncThreadsTest() {
+ test(com.aparapi.codegen.test.SyncThreads.class, expectedException, expectedOpenCL);
+ }
+
+ @Test
+ public void SyncThreadsTestWorksWithCaching() {
+ test(com.aparapi.codegen.test.SyncThreads.class, expectedException, expectedOpenCL);
+ }
+}
diff --git a/src/test/java/com/aparapi/runtime/SyncThreadsSupportTest.java b/src/test/java/com/aparapi/runtime/SyncThreadsSupportTest.java
new file mode 100644
index 00000000..652c917e
--- /dev/null
+++ b/src/test/java/com/aparapi/runtime/SyncThreadsSupportTest.java
@@ -0,0 +1,98 @@
+/**
+ * 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.runtime;
+
+import com.aparapi.Kernel;
+import com.aparapi.Range;
+import com.aparapi.device.Device;
+import com.aparapi.device.JavaDevice;
+import com.aparapi.internal.kernel.KernelManager;
+
+import static org.junit.Assert.assertArrayEquals;
+import static org.junit.Assert.assertTrue;
+
+import java.util.Arrays;
+import java.util.LinkedHashSet;
+import java.util.List;
+
+import org.junit.After;
+import org.junit.Test;
+
+public class SyncThreadsSupportTest {
+ private static final int SIZE = 64;
+
+ private static class JTPKernelManager extends KernelManager {
+ private JTPKernelManager() {
+ LinkedHashSet preferredDevices = new LinkedHashSet(1);
+ preferredDevices.add(JavaDevice.THREAD_POOL);
+ setDefaultPreferredDevices(preferredDevices);
+ }
+
+ @Override
+ protected List getPreferredDeviceTypes() {
+ return Arrays.asList(Device.TYPE.JTP);
+ }
+ }
+
+ @After
+ public void tearDown() {
+ Util.resetKernelManager();
+ }
+
+ @Test
+ public void syncThreadsSynchronizesWorkItemsInJTPMode() {
+ KernelManager.setKernelManager(new JTPKernelManager());
+ Device device = KernelManager.instance().bestDevice();
+ assertTrue(device instanceof JavaDevice);
+
+ int[] target = new int[SIZE];
+ SyncThreadsKernel kernel = new SyncThreadsKernel(target);
+ try {
+ Range range = device.createRange(SIZE, SIZE);
+ kernel.execute(range);
+ assertArrayEquals(expectedValues(), target);
+ } finally {
+ kernel.dispose();
+ }
+ }
+
+ private static int[] expectedValues() {
+ int[] expected = new int[SIZE];
+ for (int i = 0; i < SIZE; i++) {
+ expected[i] = (i + 1) % SIZE;
+ }
+ return expected;
+ }
+
+ private static class SyncThreadsKernel extends Kernel {
+ private final int[] target;
+
+ @Local
+ private final int[] scratch = new int[SIZE];
+
+ private SyncThreadsKernel(int[] target) {
+ this.target = target;
+ }
+
+ @Override
+ public void run() {
+ int id = getLocalId();
+ scratch[id] = id;
+ syncThreads();
+ target[id] = scratch[(id + 1) % SIZE];
+ }
+ }
+}