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

• develop
• include
• ck
• utility

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

ck::utility::flush_icache

void flush_icache()

Definition: flush_cache.hpp:361

ck

Definition: ck.hpp:268