not do hipify by default

aiter/jit/core.py

@@ -283,7 +283,7 @@ def build_module(
     is_python_module,
     is_standalone,
     torch_exclude,
-    hipify=True,
+    hipify=False,
     prebuild=0,
 ):
     lock_path = f"{bd_dir}/lock_{md_name}"
@@ -755,7 +755,7 @@ def wrapper(*args, custom_build_args={}, **kwargs):
                 is_python_module = d_args["is_python_module"]
                 is_standalone = d_args["is_standalone"]
                 torch_exclude = d_args["torch_exclude"]
-                hipify = d_args.get("hipify", True)
+                hipify = d_args.get("hipify", False)
                 hip_clang_path = d_args.get("hip_clang_path", None)
                 prev_hip_clang_path = None
                 if hip_clang_path is not None and os.path.exists(hip_clang_path):

aiter/ops/communication.py

@@ -1,22 +1,24 @@
 # SPDX-License-Identifier: MIT
 # Copyright (C) 2024-2025, Advanced Micro Devices, Inc. All rights reserved.
 
+import logging
+
 import torch
-from torch import Tensor
 import torch.distributed as dist
+from torch import Tensor
+
+# from ..dist.utils import get_open_port, get_distributed_init_method, get_ip
+import aiter
+
 from ..dist.parallel_state import (
+    destroy_distributed_environment,
+    destroy_model_parallel,
     ensure_model_parallel_initialized,
+    get_tp_group,
     init_distributed_environment,
     set_custom_all_reduce,
-    get_tp_group,
-    destroy_model_parallel,
-    destroy_distributed_environment,
 )
 
-# from ..dist.utils import get_open_port, get_distributed_init_method, get_ip
-import aiter
-import logging
-
 logger = logging.getLogger("aiter")
 
 
@@ -63,7 +65,7 @@ def all_reduce_asm(inp: torch.Tensor):
             return torch.empty_like(inp)
     else:
         # note: outside of cuda graph context,
-        # custom allreduce incurs a cost of cudaMemcpy, which should
+        # custom allreduce incurs a cost of hipMemcpy, which should
         # be small(<=1% of overall latency) compared to the performance
         # gains of using custom kernels
         return aiter.all_reduce_asm_(

aiter/ops/custom_all_reduce.py

@@ -1,11 +1,11 @@
 # SPDX-License-Identifier: MIT
 # Copyright (C) 2024-2025, Advanced Micro Devices, Inc. All rights reserved.
 
-import torch
 from typing import List
-from ..jit.core import (
-    compile_ops,
-)
+
+import torch
+
+from ..jit.core import compile_ops
 
 MD_NAME = "module_custom_all_reduce"
 
@@ -155,7 +155,7 @@ def register_buffer(
 
 # def gen_get_graph_buffer_ipc_meta_fake_tensors(_fa: int) -> List[torch.Tensor]:
 
-#     handle_sz = 64  # sizeof(cudaIpcMemHandle_t) is 64 byte
+#     handle_sz = 64  # sizeof(hipIpcMemHandle_t) is 64 byte
 #     num_buffers = 4  # ???
 #     handles = torch.empty((handle_sz * num_buffers,), dtype=torch.uint8, device="cuda")
 

csrc/ck_batched_gemm_a8w8/include/batched_gemm_a8w8_common.cuh

@@ -1,6 +1,6 @@
 #pragma once
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (C) 2024-2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #ifdef USE_ROCM
 
@@ -14,9 +14,9 @@
 
 #include <ATen/ATen.h>
 #include <torch/extension.h>
-#include <ATen/cuda/CUDAContext.h>
-#include <c10/cuda/CUDAGuard.h>
-#include <c10/cuda/CUDAStream.h>
+#include <ATen/hip/HIPContext.h>
+#include <ATen/hip/impl/HIPGuardImplMasqueradingAsCUDA.h>
+#include <ATen/hip/impl/HIPStreamMasqueradingAsCUDA.h>
 
 #include "ck/ck.hpp"
 #include "ck/tensor_operation/gpu/device/gemm_specialization.hpp"
@@ -272,7 +272,7 @@ __forceinline__ torch::Tensor batched_gemm_a8w8_rowwise_impl(
     int BatchStrideB = N * K;
     int BatchStrideE = M * N;
 
-    const at::cuda::OptionalCUDAGuard device_guard(device_of(XQ));
+    const at::hip::OptionalHIPGuardMasqueradingAsCUDA device_guard(device_of(XQ));
     auto device_gemm = DeviceGemmInstance{};
     auto invoker = device_gemm.MakeInvoker();
 
@@ -307,7 +307,7 @@ __forceinline__ torch::Tensor batched_gemm_a8w8_rowwise_impl(
 
     TORCH_CHECK(device_gemm.IsSupportedArgument(argument), "This GEMM is not supported!");
 
-    invoker.Run(argument, StreamConfig{at::cuda::getCurrentCUDAStream().stream()});
+    invoker.Run(argument, StreamConfig{at::hip::getCurrentHIPStream()});
     return Y;
 }
 

csrc/ck_batched_gemm_bf16/include/batched_gemm_bf16_common.cuh

@@ -1,6 +1,6 @@
 #pragma once
 // SPDX-License-Identifier: MIT
-// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (C) 2024-2025, Advanced Micro Devices, Inc. All rights reserved.
 
 #ifdef USE_ROCM
 
@@ -14,9 +14,9 @@
 
 #include <ATen/ATen.h>
 #include <torch/extension.h>
-#include <ATen/cuda/CUDAContext.h>
-#include <c10/cuda/CUDAGuard.h>
-#include <c10/cuda/CUDAStream.h>
+#include <ATen/hip/HIPContext.h>
+#include <ATen/hip/impl/HIPGuardImplMasqueradingAsCUDA.h>
+#include <ATen/hip/impl/HIPStreamMasqueradingAsCUDA.h>
 
 #include "ck/ck.hpp"
 #include "ck/tensor_operation/gpu/device/gemm_specialization.hpp"
@@ -151,7 +151,7 @@ __forceinline__ torch::Tensor batched_gemm_bf16_impl(
     int BatchStrideB = N * K;
     int BatchStrideE = M * N;
 
-    const at::cuda::OptionalCUDAGuard device_guard(device_of(XQ));
+    const at::hip::OptionalHIPGuardMasqueradingAsCUDA device_guard(device_of(XQ));
     auto device_gemm = DeviceGemmInstance{};
     auto invoker = device_gemm.MakeInvoker();
 
@@ -184,7 +184,7 @@ __forceinline__ torch::Tensor batched_gemm_bf16_impl(
 
     TORCH_CHECK(device_gemm.IsSupportedArgument(argument), "This GEMM is not supported!");
 
-    invoker.Run(argument, StreamConfig{at::cuda::getCurrentCUDAStream().stream()});
+    invoker.Run(argument, StreamConfig{at::hip::getCurrentHIPStream()});
     return Y;
 }