mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-05-14 02:02:46 +00:00
2ca6f22fab
* Add rotating buffer feature for universal gemm
* adding changes in tile_engine
* Updated code to merge kernel_launch
* removing comments
* Enable rotating buffer changes to flatmm
* Created diff launch_kernel function for rotating buffer
* Simplfied calculation using macros
* merge code with new changes in tile_engine
* clang formatted
* Redefine macros
[ROCm/composable_kernel commit: 99857e10e6]
31 lines
883 B
C++
31 lines
883 B
C++
// SPDX-License-Identifier: MIT
|
|
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
|
|
|
|
#pragma once
|
|
|
|
#include <hip/hip_runtime.h>
|
|
|
|
namespace ck_tile {
|
|
static __global__ void flush_cache()
|
|
{
|
|
asm __volatile__("s_icache_inv \n\t"
|
|
"s_nop 0 \n\t"
|
|
"s_nop 0 \n\t"
|
|
"s_nop 0 \n\t"
|
|
"s_nop 0 \n\t"
|
|
"s_nop 0 \n\t"
|
|
"s_nop 0 \n\t"
|
|
"s_nop 0 \n\t"
|
|
"s_nop 0 \n\t"
|
|
"s_nop 0 \n\t"
|
|
"s_nop 0 \n\t"
|
|
"s_nop 0 \n\t"
|
|
"s_nop 0 \n\t"
|
|
"s_nop 0 \n\t"
|
|
"s_nop 0 \n\t"
|
|
"s_nop 0 \n\t"
|
|
"s_nop 0 \n\t" ::
|
|
:);
|
|
}
|
|
} // namespace ck_tile
|