synchronization.hpp Source File

synchronization.hpp Source File#

Composable Kernel: synchronization.hpp Source File
synchronization.hpp
Go to the documentation of this file.
1// SPDX-License-Identifier: MIT
2// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
3
4#pragma once
5
6#include "ck/ck.hpp"
7
8namespace ck {
9
10#if CK_EXPERIMENTAL_BLOCK_SYNC_LDS_WITHOUT_SYNC_VMEM
11#ifdef __gfx12__
12__device__ void llvm_amdgcn_s_wait_dscnt(short cnt) __asm("llvm.amdgcn.s.wait.dscnt");
13#endif
14#endif
15
16__device__ void block_sync_lds()
17{
18#if CK_EXPERIMENTAL_BLOCK_SYNC_LDS_WITHOUT_SYNC_VMEM
19#if defined(__gfx12__)
20 llvm_amdgcn_s_wait_dscnt(0);
21 asm volatile("s_barrier_signal -1\n\t"
22 "s_barrier_wait -1");
23#elif defined(__gfx11__)
24 // asm volatile("\
25 // s_waitcnt lgkmcnt(0) \n \
26 // s_barrier \
27 // " ::);
28 __builtin_amdgcn_s_waitcnt(0xfc07);
29 __builtin_amdgcn_s_barrier();
30#else
31 // asm volatile("\
32 // s_waitcnt lgkmcnt(0) \n \
33 // s_barrier \
34 // " ::);
35 __builtin_amdgcn_s_waitcnt(0xc07f);
36 __builtin_amdgcn_s_barrier();
37#endif
38#else
39 __syncthreads();
40#endif
41}
42
43__device__ void block_sync_lds_direct_load()
44{
45#ifdef __gfx12__
46 asm volatile("\
47 s_wait_loadcnt 0x0 \n \
48 s_wait_dscnt 0x0 \n \
49 s_barrier_signal -1 \n \
50 s_barrier_wait -1 \
51 " ::);
52#else
53 asm volatile("\
54 s_waitcnt vmcnt(0) \n \
55 s_waitcnt lgkmcnt(0) \n \
56 s_barrier \
57 " ::);
58#endif
59}
60
61__device__ void s_nop()
62{
63#if 1
64 asm volatile("\
65 s_nop 0 \n \
66 " ::);
67#else
68 __builtin_amdgcn_sched_barrier(0);
69#endif
70}
71
72} // namespace ck
ck
Definition ck.hpp:268
ck::s_nop
__device__ void s_nop()
Definition synchronization.hpp:61
ck::block_sync_lds_direct_load
__device__ void block_sync_lds_direct_load()
Definition synchronization.hpp:43
ck::block_sync_lds
__device__ void block_sync_lds()
Definition synchronization.hpp:16