synchronization.hpp 1 KB
Newer Older
Chao Liu's avatar
Chao Liu committed
1
// SPDX-License-Identifier: MIT
Illia Silin's avatar
Illia Silin committed
2
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
Chao Liu's avatar
Chao Liu committed
3

Chao Liu's avatar
Chao Liu committed
4
#pragma once
Chao Liu's avatar
Chao Liu committed
5

Chao Liu's avatar
Chao Liu committed
6
#include "ck/ck.hpp"
Chao Liu's avatar
Chao Liu committed
7
8
9
10
11

namespace ck {

__device__ void block_sync_lds()
{
12
#if CK_EXPERIMENTAL_BLOCK_SYNC_LDS_WITHOUT_SYNC_VMEM
Chao Liu's avatar
Chao Liu committed
13
14
15
16
17
    asm volatile("\
    s_waitcnt lgkmcnt(0) \n \
    s_barrier \
    " ::);
#else
18
    __syncthreads();
Chao Liu's avatar
Chao Liu committed
19
20
#endif
}
21

22
23
24
25
26
27
28
__device__ void block_sync_lds_direct_load()
{
    asm volatile("\
    s_waitcnt vmcnt(0) \n \
    s_waitcnt lgkmcnt(0) \n \
    s_barrier \
    " ::);
29
30
31
32
33
34
#if CK_USE_AMD_LDS_DIRECT_LOAD_INLINE_ASM
    // When direct loads and `waitcnt` instructions are submitted using inline asm, the usage of
    // `sched_barrier` is necessary to make sure that no instructions that use the loaded memory
    // are scheduled by the compiler before the `waitcnt` instruction.
    __builtin_amdgcn_sched_barrier(0);
#endif
35
36
}

ltqin's avatar
ltqin committed
37
38
39
40
41
42
43
44
45
46
__device__ void s_nop()
{
#if 1
    asm volatile("\
    s_nop 0 \n \
    " ::);
#else
    __builtin_amdgcn_sched_barrier(0);
#endif
}
Chao Liu's avatar
Chao Liu committed
47
48

} // namespace ck