"vscode:/vscode.git/clone" did not exist on "c9f98622b1daba55477dcb330de1739fb2f02ce6"
synchronization.hpp 407 Bytes
Newer Older
Chao Liu's avatar
Chao Liu committed
1
2
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
Chao Liu's avatar
Chao Liu committed
3
#pragma once
4
#ifndef CK_NOGPU
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
21
22
23
#endif
}

} // namespace ck
#endif
24