cubecl_cpp/hip/
dialect.rs

1macro_rules! hip_op {
2    ($ty: ty, $impl: expr) => {
3        #[pliron::derive::op_interface_impl]
4        impl $crate::shared::operation::OpToCPP<$crate::target::Hip> for $ty {
5            fn to_cpp(&self, ctx: &pliron::context::Context) -> String {
6                $crate::shared::closure_inference_hack::<$ty, String>(self, ctx, $impl)
7            }
8        }
9    };
10}
11use cubecl_core::ir::dialect::synchronization::{SyncOp, SyncScope};
12pub(super) use hip_op;
13
14macro_rules! hip_op_with_out {
15    ($ty: ty, $impl: expr) => {
16        #[pliron::derive::op_interface_impl]
17        impl $crate::shared::operation::OpToCPP<$crate::target::Hip> for $ty {
18            fn to_cpp(&self, ctx: &pliron::context::Context) -> String {
19                use cubecl_core::ir::prelude::*;
20                use $crate::shared::CppValue;
21                let op = $crate::shared::closure_inference_hack::<$ty, String>(self, ctx, $impl);
22                let out = self.get_result(ctx).fmt_left(ctx);
23                format!("{out} = {op};\n")
24            }
25        }
26    };
27}
28pub(super) use hip_op_with_out;
29
30hip_op!(SyncOp, |op, ctx| {
31    match op.scope(ctx).0 {
32        SyncScope::Plane => {
33            // HIP has no `__syncwarp`. AMD wavefronts execute in lockstep, so the
34            // execution half of the sync is a compiler scheduling barrier; the
35            // wavefront-scope fence supplies the memory ordering `__syncwarp`
36            // carries on CUDA (LDS/global writes by other lanes of the wave are
37            // visible past the sync).
38
39            "
40__builtin_amdgcn_fence(__ATOMIC_ACQ_REL, \"wavefront\");
41__builtin_amdgcn_wave_barrier();\n"
42        }
43        SyncScope::Cube => "__syncthreads();\n",
44        // As on CUDA: `__syncthreads` is the block's own ordering, and the fence is what makes
45        // this block's writes visible to the others.
46        SyncScope::Device => "__threadfence();\n__syncthreads();\n",
47        SyncScope::Unit => "",
48    }
49    .into()
50});
cubecl_cpp/hip/dialect.rs

cubecl_cpp/hip/
dialect.rs
