Skip to content

GitLab

Switch branch/tag
Find file Normal viewHistoryPermalink
synchronization.hpp 361 Bytes
Newer Older
Chao Liu's avatar
Code clean up (#20)  
Chao Liu committed
1
2 
#ifndef CK_SYNCHRONIZATION_AMD_HPP
#define CK_SYNCHRONIZATION_AMD_HPP
carlushuang's avatar
support gcc with cpu only compile  
carlushuang committed
3 
#ifndef CK_NOGPU
Chao Liu's avatar
Code clean up (#20)  
Chao Liu committed
4
5
6
7
8
9
10 

#include "config.hpp"

namespace ck {

__device__ void block_sync_lds()
{
Chao Liu's avatar
Compile for gfx908 and gfx90a (#130)  
Chao Liu committed
11 
#if CK_EXPERIMENTAL_BLOCK_SYNC_LDS_WITHOUT_SYNC_VMEM
Chao Liu's avatar
Code clean up (#20)  
Chao Liu committed
12
13
14
15
16 
    asm volatile("\
    s_waitcnt lgkmcnt(0) \n \
    s_barrier \
    " ::);
#else
Chao Liu's avatar
[MIOpen Downstream] Initial MIOpen integration (#52)  
Chao Liu committed
17 
    __syncthreads();
Chao Liu's avatar
Code clean up (#20)  
Chao Liu committed
18
19
20
21
22 
#endif
}

} // namespace ck
#endif
carlushuang's avatar
support gcc with cpu only compile  
carlushuang committed
23 
#endif