/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/utility/flush_icache.hpp Source File

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/utility/flush_icache.hpp Source File#

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/utility/flush_icache.hpp Source File
flush_icache.hpp
Go to the documentation of this file.
1 // SPDX-License-Identifier: MIT
2 // Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
3 
4 #pragma once
5 
6 #include <hip/hip_runtime.h>
7 
8 namespace ck {
9 static __global__ void flush_icache()
10 {
11  asm __volatile__("s_icache_inv \n\t"
12  "s_nop 0 \n\t"
13  "s_nop 0 \n\t"
14  "s_nop 0 \n\t"
15  "s_nop 0 \n\t"
16  "s_nop 0 \n\t"
17  "s_nop 0 \n\t"
18  "s_nop 0 \n\t"
19  "s_nop 0 \n\t"
20  "s_nop 0 \n\t"
21  "s_nop 0 \n\t"
22  "s_nop 0 \n\t"
23  "s_nop 0 \n\t"
24  "s_nop 0 \n\t"
25  "s_nop 0 \n\t"
26  "s_nop 0 \n\t"
27  "s_nop 0 \n\t" ::
28  :);
29 }
30 } // namespace ck
void flush_icache()
Definition: flush_cache.hpp:216
Definition: ck.hpp:267