include/ck/utility/synchronization.hpp Source File

• include
• ck
• utility

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_vmcnt 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

ck.hpp

ck

Definition: ck.hpp:264

ck::s_nop

__device__ void s_nop()

Definition: synchronization.hpp:50

ck::block_sync_lds_direct_load

__device__ void block_sync_lds_direct_load()

Definition: synchronization.hpp:32

ck::block_sync_lds

__device__ void block_sync_lds()

Definition: synchronization.hpp:10