/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/utility/synchronization.hpp Source File

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/utility/synchronization.hpp Source File#

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/utility/synchronization.hpp Source File
synchronization.hpp
Go to the documentation of this file.
1 // SPDX-License-Identifier: MIT
2 // Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
3 
4 #pragma once
5 
6 #include "ck/ck.hpp"
7 
8 namespace ck {
9 
10 __device__ void block_sync_lds()
11 {
12 #if CK_EXPERIMENTAL_BLOCK_SYNC_LDS_WITHOUT_SYNC_VMEM
13 #ifdef __gfx12__
14  asm volatile("\
15  s_wait_dscnt 0x0 \n \
16  s_barrier_signal -1 \n \
17  s_barrier_wait -1 \
18  " ::);
19 #else
20  // asm volatile("\
21  // s_waitcnt lgkmcnt(0) \n \
22  // s_barrier \
23  // " ::);
24  __builtin_amdgcn_s_waitcnt(0xc07f);
25  __builtin_amdgcn_s_barrier();
26 #endif
27 #else
28  __syncthreads();
29 #endif
30 }
31 
32 __device__ void block_sync_lds_direct_load()
33 {
34 #ifdef __gfx12__
35  asm volatile("\
36  s_wait_loadcnt 0x0 \n \
37  s_wait_dscnt 0x0 \n \
38  s_barrier_signal -1 \n \
39  s_barrier_wait -1 \
40  " ::);
41 #else
42  asm volatile("\
43  s_waitcnt vmcnt(0) \n \
44  s_waitcnt lgkmcnt(0) \n \
45  s_barrier \
46  " ::);
47 #endif
48 }
49 
50 __device__ void s_nop()
51 {
52 #if 1
53  asm volatile("\
54  s_nop 0 \n \
55  " ::);
56 #else
57  __builtin_amdgcn_sched_barrier(0);
58 #endif
59 }
60 
61 } // namespace ck
Definition: ck.hpp:267
__device__ void s_nop()
Definition: synchronization.hpp:50
__device__ void block_sync_lds_direct_load()
Definition: synchronization.hpp:32
__device__ void block_sync_lds()
Definition: synchronization.hpp:10