Add syncThreads barrier alias

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
+    * <code>barrier(CLK_LOCAL_MEM_FENCE)</code> 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.<br> 
     * It will also enforce memory ordering, such that modifications made by each thread in the work-group, to the memory,

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)");

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}**/

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<? extends com.aparapi.internal.exception.AparapiException> 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);
+    }
+}

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<Device> preferredDevices = new LinkedHashSet<Device>(1);
+            preferredDevices.add(JavaDevice.THREAD_POOL);
+            setDefaultPreferredDevices(preferredDevices);
+        }
+
+        @Override
+        protected List<Device.TYPE> 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];
+        }
+    }
+}