include/ck_tile/host/flush_icache.hpp Source File

• include
• ck_tile
• host

flush_icache.hpp

Go to the documentation of this file.

1 // SPDX-License-Identifier: MIT

2 // Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.

3 

4 #pragma once

5 

6 #include <hip/hip_runtime.h>

7 

8 namespace ck_tile {

9 static __global__ void flush_cache()

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_tile

ck_tile

Definition: cluster_descriptor.hpp:13