flush_icache.hpp Source File

• include
• ck
• utility

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

8namespace ck {

9static __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

Definition ck.hpp:268