From 19aa1c5c673953c05fbafcba229643403876835a Mon Sep 17 00:00:00 2001 From: gshaowei6 Date: Sun, 17 May 2026 02:04:33 +0800 Subject: [PATCH] Add syncThreads barrier alias --- src/main/java/com/aparapi/Kernel.java | 14 +++ .../aparapi/internal/writer/KernelWriter.java | 2 + .../com/aparapi/codegen/test/SyncThreads.java | 46 +++++++++ .../aparapi/codegen/test/SyncThreadsTest.java | 50 ++++++++++ .../runtime/SyncThreadsSupportTest.java | 98 +++++++++++++++++++ 5 files changed, 210 insertions(+) create mode 100644 src/test/java/com/aparapi/codegen/test/SyncThreads.java create mode 100644 src/test/java/com/aparapi/codegen/test/SyncThreadsTest.java create mode 100644 src/test/java/com/aparapi/runtime/SyncThreadsSupportTest.java 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]; + } + } +}