mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-05-12 17:26:00 +00:00
cd167e492a
* adding compilation for multiple targets * fix build * clean * update Jekinsfile * update readme * update Jenkins * use ck::half_t instead of ushort for bf16 * rename enum classes * clean * rename * clean
22 lines
337 B
C++
22 lines
337 B
C++
#ifndef CK_SYNCHRONIZATION_AMD_HPP
|
|
#define CK_SYNCHRONIZATION_AMD_HPP
|
|
|
|
#include "config.hpp"
|
|
|
|
namespace ck {
|
|
|
|
__device__ void block_sync_lds()
|
|
{
|
|
#if CK_EXPERIMENTAL_BLOCK_SYNC_LDS_WITHOUT_SYNC_VMEM
|
|
asm volatile("\
|
|
s_waitcnt lgkmcnt(0) \n \
|
|
s_barrier \
|
|
" ::);
|
|
#else
|
|
__syncthreads();
|
|
#endif
|
|
}
|
|
|
|
} // namespace ck
|
|
#endif
|