mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-05-15 10:37:44 +00:00
21a146fb2f
* wavelet gemm programming model support for CK
* GEMM pipeline update for wavelet progrmmaing model
* Updated wavelet programming pipeline
* fixes for global-write for math-wave
* fixed bug in global writes
* Updated comments for better readability
* fixed clang format errors
* added block_lds without barrier sync
* clean
* clean
* clean
* clean
* refactor
* prototype
4 layouts
fix default stride
all problem sizes
tidy
move file
update build script
restore old file
fix build
* refactor standalone test to use gemm test harness
* simplify gemm test
* update build script
* remove redundant
* early return when cmd arg doesn't match
* tidy
* report failure when result not validated
* tidy
* Add comment depicting B2C mapping pattern.
* Formatting & comments.
* Comparison with custom B2C mapping pattern.
* Example for wavelet gemm.
* Add wavelet to Gemm standalone test.
* Remove debug code.
* Remove dangling #endif directive.
Co-authored-by: root <Raman Jana>
Co-authored-by: Chao Liu <chao.liu2@amd.com>
Co-authored-by: Adam Osewski <aosewski@amd.com>
Co-authored-by: Anthony Chang <ac.chang@outlook.com>
Co-authored-by: Adam Osewski <19374865+aosewski@users.noreply.github.com>
[ROCm/composable_kernel commit: 1cfa87608a]
34 lines
518 B
C++
34 lines
518 B
C++
// SPDX-License-Identifier: MIT
|
|
// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
|
|
|
|
#pragma once
|
|
|
|
#include "ck/ck.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
|
|
}
|
|
|
|
__device__ void s_nop()
|
|
{
|
|
#if 1
|
|
asm volatile("\
|
|
s_nop 0 \n \
|
|
" ::);
|
|
#else
|
|
__builtin_amdgcn_sched_barrier(0);
|
|
#endif
|
|
}
|
|
|
|
} // namespace ck
|