"vscode:/vscode.git/clone" did not exist on "2145a8b6d256616ba27435106d9285fc5d09152b"
synchronization.hpp 983 Bytes
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
13
14
15
16
17
18
19
#ifdef __gfx12__
    asm volatile("\
    s_wait_idle lgkmcnt(0) \n \
    s_barrier_signal \n \
    s_barrier_wait \
    " ::);
#else
Chao Liu's avatar
Chao Liu committed
20
21
22
23
    asm volatile("\
    s_waitcnt lgkmcnt(0) \n \
    s_barrier \
    " ::);
24
#endif
Chao Liu's avatar
Chao Liu committed
25
#else
26
    __syncthreads();
Chao Liu's avatar
Chao Liu committed
27
28
#endif
}
29

30
31
__device__ void block_sync_lds_direct_load()
{
32
33
34
35
36
37
38
39
#ifdef __gfx12__
    asm volatile("\
    s_wait_idle vmcnt(0) \n \
    s_wait_idle lgkmcnt(0) \n \
    s_barrier_signal \n \
    s_barrier_wait \
    " ::);
#else
40
41
42
43
44
    asm volatile("\
    s_waitcnt vmcnt(0) \n \
    s_waitcnt lgkmcnt(0) \n \
    s_barrier \
    " ::);
45
#endif
46
47
}

ltqin's avatar
ltqin committed
48
49
50
51
52
53
54
55
56
57
__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
58
59

} // namespace ck