[clang] Connect the HIP device parser into the interpreter (PR #218848)
Aditya Sinha via cfe-commits
cfe-commits at lists.llvm.org
Mon Sep 7 23:03:45 PDT 2026
https://github.com/AdityaSinha149 updated https://github.com/llvm/llvm-project/pull/218848
>From c8360cd7546c6b339903d1d47f0638e4080c4768 Mon Sep 17 00:00:00 2001
From: AdityaSinha149 <adsinha at amd.com>
Date: Wed, 19 Aug 2026 22:51:56 +0530
Subject: [PATCH] [clang-repl] Hip environment initialized
---
clang/include/clang/Interpreter/Interpreter.h | 39 +++++++++--
clang/lib/Interpreter/Interpreter.cpp | 57 +++++++++-------
.../test/Interpreter/HIP/hip-environment.hip | 11 ++++
clang/test/Interpreter/HIP/lit.local.cfg | 2 +
clang/test/lit.cfg.py | 54 ++++++++++++++-
clang/tools/clang-repl/ClangRepl.cpp | 65 +++++++++++++------
6 files changed, 178 insertions(+), 50 deletions(-)
create mode 100644 clang/test/Interpreter/HIP/hip-environment.hip
create mode 100644 clang/test/Interpreter/HIP/lit.local.cfg
diff --git a/clang/include/clang/Interpreter/Interpreter.h b/clang/include/clang/Interpreter/Interpreter.h
index c2622b23d5d9c..4504b679504e0 100644
--- a/clang/include/clang/Interpreter/Interpreter.h
+++ b/clang/include/clang/Interpreter/Interpreter.h
@@ -46,6 +46,8 @@ class Decl;
class IncrementalParser;
class IncrementalCUDADeviceParser;
+enum class OffloadType { CUDA, HIP };
+
/// Create a pre-configured \c CompilerInstance for incremental processing.
class IncrementalCompilerBuilder {
using DriverCompilationFn = llvm::Error(const driver::Compilation &);
@@ -65,29 +67,52 @@ class IncrementalCompilerBuilder {
// Offload options
void SetOffloadArch(llvm::StringRef Arch) { OffloadArch = Arch; };
- // CUDA specific
- void SetCudaSDK(llvm::StringRef path) { CudaSDKPath = path; };
+ void SetDeviceSDK(OffloadType Type, llvm::StringRef Path) {
+ if (Type == OffloadType::HIP)
+ RocmSDKPath = Path;
+ else
+ CudaSDKPath = Path;
+ }
+
+ // Retained for compatibility with existing CUDA callers.
+ void SetCudaSDK(llvm::StringRef Path) {
+ SetDeviceSDK(OffloadType::CUDA, Path);
+ }
// Hand over the compilation.
void SetDriverCompilationCallback(std::function<DriverCompilationFn> C) {
CompilationCB = C;
}
- llvm::Expected<std::unique_ptr<CompilerInstance>> CreateCudaHost();
- llvm::Expected<std::unique_ptr<CompilerInstance>> CreateCudaDevice();
+ llvm::Expected<std::unique_ptr<CompilerInstance>>
+ CreateHost(OffloadType Type);
+ llvm::Expected<std::unique_ptr<CompilerInstance>>
+ CreateDevice(OffloadType Type);
+
+ // Retained for compatibility with existing CUDA callers.
+ llvm::Expected<std::unique_ptr<CompilerInstance>> CreateCudaHost() {
+ return CreateHost(OffloadType::CUDA);
+ }
+ llvm::Expected<std::unique_ptr<CompilerInstance>> CreateCudaDevice() {
+ return CreateDevice(OffloadType::CUDA);
+ }
private:
llvm::Expected<std::unique_ptr<CompilerInstance>>
create(std::string TT, std::vector<const char *> &ClangArgv);
- llvm::Expected<std::unique_ptr<CompilerInstance>> createCuda(bool device);
+ llvm::Expected<std::unique_ptr<CompilerInstance>>
+ createOffload(OffloadType Type, bool device);
std::vector<const char *> UserArgs;
std::optional<std::string> TargetTriple;
llvm::StringRef OffloadArch;
+ llvm::StringRef RocmSDKPath;
llvm::StringRef CudaSDKPath;
+ std::string OffloadCUID;
+
std::optional<std::function<DriverCompilationFn>> CompilationCB;
};
@@ -147,8 +172,8 @@ class Interpreter {
create(std::unique_ptr<CompilerInstance> CI,
std::unique_ptr<IncrementalExecutorBuilder> IEB = nullptr);
static llvm::Expected<std::unique_ptr<Interpreter>>
- createWithCUDA(std::unique_ptr<CompilerInstance> CI,
- std::unique_ptr<CompilerInstance> DCI);
+ createWithDevice(OffloadType Type, std::unique_ptr<CompilerInstance> CI,
+ std::unique_ptr<CompilerInstance> DCI);
const ASTContext &getASTContext() const;
ASTContext &getASTContext();
diff --git a/clang/lib/Interpreter/Interpreter.cpp b/clang/lib/Interpreter/Interpreter.cpp
index 092f3ede771f6..684e256f20830 100644
--- a/clang/lib/Interpreter/Interpreter.cpp
+++ b/clang/lib/Interpreter/Interpreter.cpp
@@ -45,12 +45,14 @@
#include "clang/Serialization/ASTReader.h"
#include "clang/Serialization/ModuleCache.h"
#include "clang/Serialization/ObjectFilePCHContainerReader.h"
+#include "llvm/ADT/StringExtras.h"
#include "llvm/ExecutionEngine/JITSymbol.h"
#include "llvm/ExecutionEngine/Orc/EPCDynamicLibrarySearchGenerator.h"
#include "llvm/ExecutionEngine/Orc/LLJIT.h"
#include "llvm/IR/Module.h"
#include "llvm/Support/Errc.h"
#include "llvm/Support/ErrorHandling.h"
+#include "llvm/Support/Process.h"
#include "llvm/Support/VirtualFileSystem.h"
#include "llvm/Support/raw_ostream.h"
#include "llvm/TargetParser/Host.h"
@@ -303,19 +305,17 @@ IncrementalCompilerBuilder::CreateCpp() {
}
llvm::Expected<std::unique_ptr<CompilerInstance>>
-IncrementalCompilerBuilder::createCuda(bool device) {
+IncrementalCompilerBuilder::createOffload(OffloadType Type, bool device) {
+ const bool HipEnabled = Type == OffloadType::HIP;
std::vector<const char *> Argv;
Argv.reserve(5 + 4 + UserArgs.size());
+ Argv.push_back(HipEnabled ? "-xhip" : "-xcuda");
+ Argv.push_back(device ? "--cuda-device-only" : "--cuda-host-only");
- Argv.push_back("-xcuda");
- if (device)
- Argv.push_back("--cuda-device-only");
- else
- Argv.push_back("--cuda-host-only");
-
- std::string SDKPathArg = "--cuda-path=";
- if (!CudaSDKPath.empty()) {
- SDKPathArg += CudaSDKPath;
+ llvm::StringRef SDKPath = HipEnabled ? RocmSDKPath : CudaSDKPath;
+ std::string SDKPathArg = HipEnabled ? "--rocm-path=" : "--cuda-path=";
+ if (!SDKPath.empty()) {
+ SDKPathArg += SDKPath;
Argv.push_back(SDKPathArg.c_str());
}
@@ -325,6 +325,12 @@ IncrementalCompilerBuilder::createCuda(bool device) {
Argv.push_back(ArchArg.c_str());
}
+ if (OffloadCUID.empty())
+ OffloadCUID = llvm::utohexstr(llvm::sys::Process::GetRandomNumber(),
+ /*LowerCase=*/true);
+ std::string CUIDArg = "-cuid=" + OffloadCUID;
+ Argv.push_back(CUIDArg.c_str());
+
llvm::append_range(Argv, UserArgs);
std::string TT = TargetTriple ? *TargetTriple : llvm::sys::getProcessTriple();
@@ -332,13 +338,13 @@ IncrementalCompilerBuilder::createCuda(bool device) {
}
llvm::Expected<std::unique_ptr<CompilerInstance>>
-IncrementalCompilerBuilder::CreateCudaDevice() {
- return IncrementalCompilerBuilder::createCuda(true);
+IncrementalCompilerBuilder::CreateDevice(OffloadType Type) {
+ return IncrementalCompilerBuilder::createOffload(Type, /*device=*/true);
}
llvm::Expected<std::unique_ptr<CompilerInstance>>
-IncrementalCompilerBuilder::CreateCudaHost() {
- return IncrementalCompilerBuilder::createCuda(false);
+IncrementalCompilerBuilder::CreateHost(OffloadType Type) {
+ return IncrementalCompilerBuilder::createOffload(Type, /*device=*/false);
}
Interpreter::Interpreter(std::unique_ptr<CompilerInstance> Instance,
@@ -473,8 +479,9 @@ llvm::Expected<std::unique_ptr<Interpreter>> Interpreter::create(
}
llvm::Expected<std::unique_ptr<Interpreter>>
-Interpreter::createWithCUDA(std::unique_ptr<CompilerInstance> CI,
- std::unique_ptr<CompilerInstance> DCI) {
+Interpreter::createWithDevice(OffloadType Type,
+ std::unique_ptr<CompilerInstance> CI,
+ std::unique_ptr<CompilerInstance> DCI) {
// avoid writing fat binary to disk using an in-memory virtual file system
llvm::IntrusiveRefCntPtr<llvm::vfs::InMemoryFileSystem> IMVFS =
std::make_unique<llvm::vfs::InMemoryFileSystem>();
@@ -508,14 +515,20 @@ Interpreter::createWithCUDA(std::unique_ptr<CompilerInstance> CI,
Interp->DeviceCI = std::move(DCI);
- auto DeviceParser = std::make_unique<IncrementalCUDADeviceParser>(
- *Interp->DeviceCI, *Interp->getCompilerInstance(),
- Interp->DeviceAct.get(), IMVFS, Err, Interp->PTUs);
+ if (Type == OffloadType::HIP) {
+ // FIXME: HIP device parsing is not supported yet; it should use an
+ // IncrementalHIPDeviceParser once one exists.
+ } else {
+ auto DeviceParser = std::make_unique<IncrementalCUDADeviceParser>(
+ *Interp->DeviceCI, *Interp->getCompilerInstance(),
+ Interp->DeviceAct.get(), IMVFS, Err, Interp->PTUs);
- if (Err)
- return std::move(Err);
+ if (Err)
+ return std::move(Err);
+
+ Interp->DeviceParser = std::move(DeviceParser);
+ }
- Interp->DeviceParser = std::move(DeviceParser);
return std::move(Interp);
}
diff --git a/clang/test/Interpreter/HIP/hip-environment.hip b/clang/test/Interpreter/HIP/hip-environment.hip
new file mode 100644
index 0000000000000..16354ecf44e98
--- /dev/null
+++ b/clang/test/Interpreter/HIP/hip-environment.hip
@@ -0,0 +1,11 @@
+// Check that clang-repl initializes the HIP environment but reports it as
+// unsupported, since HIP execution is not implemented yet. When both -cuda and
+// -hip are passed, -hip wins because it appears later, so the HIP path is taken.
+// An explicit --offload-arch is passed so the test does not rely on GPU
+// auto-detection (--offload-arch=native), which fails on systems that have ROCm
+// installed but no GPU. The test never runs device code, so the arch is
+// arbitrary.
+
+// RUN: not clang-repl -cuda -hip --offload-arch=gfx1100 2>&1 | FileCheck %s
+
+// CHECK: HIP environment is initialized but not supported as of now.
diff --git a/clang/test/Interpreter/HIP/lit.local.cfg b/clang/test/Interpreter/HIP/lit.local.cfg
new file mode 100644
index 0000000000000..70102544ab0fd
--- /dev/null
+++ b/clang/test/Interpreter/HIP/lit.local.cfg
@@ -0,0 +1,2 @@
+if 'host-supports-hip' not in config.available_features:
+ config.unsupported = True
diff --git a/clang/test/lit.cfg.py b/clang/test/lit.cfg.py
index 9b7bd1d329d22..2117bb2594ed3 100644
--- a/clang/test/lit.cfg.py
+++ b/clang/test/lit.cfg.py
@@ -1,5 +1,6 @@
# -*- Python -*-
+import glob
import os
import platform
import re
@@ -222,6 +223,54 @@ def have_host_clang_repl_cuda():
return False
+def _hip_lib_directory():
+ explicit = lit_config.params.get("hip_lib_path")
+ if explicit:
+ candidates = [explicit]
+ else:
+ candidates = []
+ for var in ("ROCM_PATH", "HIP_PATH"):
+ if os.environ.get(var):
+ candidates.append(os.path.join(os.environ[var], "lib"))
+ candidates.append("/opt/rocm/lib")
+ for directory in candidates:
+ if directory and glob.glob(os.path.join(directory, "libamdhip64.so*")):
+ return directory
+ return None
+
+
+def _clang_can_compile_hip(clang, rocm_lib_dir):
+ rocm_root = os.path.dirname(rocm_lib_dir)
+ offload_arch = lit_config.params.get("amdgpu_arch", "gfx906")
+ test_src = b"#include <hip/hip_runtime.h>\n__global__ void k() {}\n"
+ try:
+ proc = subprocess.run(
+ [
+ clang,
+ "-x",
+ "hip",
+ "-fsyntax-only",
+ "-nogpulib",
+ "--offload-arch=" + offload_arch,
+ "--rocm-path=" + rocm_root,
+ "-",
+ ],
+ input=test_src,
+ stdout=subprocess.PIPE,
+ stderr=subprocess.PIPE,
+ )
+ except OSError:
+ return False
+ return proc.returncode == 0
+
+
+def have_host_hip_environment():
+ hip_lib_dir = _hip_lib_directory()
+ if not hip_lib_dir or not config.clang:
+ return False
+ return _clang_can_compile_hip(config.clang, hip_lib_dir)
+
+
skip_clang_repl_checks = lit.util.pythonize_bool(
lit_config.params.get(
"clang_skip_clang_repl_checks",
@@ -234,6 +283,9 @@ def have_host_clang_repl_cuda():
if have_host_clang_repl_cuda():
config.available_features.add('host-supports-cuda')
+
+ if have_host_hip_environment():
+ config.available_features.add("host-supports-hip")
hosttriple = run_clang_repl("--host-jit-triple")
config.substitutions.append(("%host-jit-triple", hosttriple.strip()))
@@ -508,4 +560,4 @@ def user_is_root():
sys.path.append(utilspath)
from update_any_test_checks import utc_lit_plugin
- lit_config.test_updaters.append(utc_lit_plugin)
+ lit_config.test_updaters.append(utc_lit_plugin)
\ No newline at end of file
diff --git a/clang/tools/clang-repl/ClangRepl.cpp b/clang/tools/clang-repl/ClangRepl.cpp
index c9873540a5d66..15eb8ed3034c9 100644
--- a/clang/tools/clang-repl/ClangRepl.cpp
+++ b/clang/tools/clang-repl/ClangRepl.cpp
@@ -52,6 +52,8 @@ LLVM_ATTRIBUTE_USED int __lsan_is_turned_off() { return 1; }
#define DEBUG_TYPE "clang-repl"
+static llvm::cl::opt<bool> HipEnabled("hip", llvm::cl::Hidden);
+static llvm::cl::opt<std::string> RocmPath("rocm-path", llvm::cl::Hidden);
static llvm::cl::opt<bool> CudaEnabled("cuda", llvm::cl::Hidden);
static llvm::cl::opt<std::string> CudaPath("cuda-path", llvm::cl::Hidden);
static llvm::cl::opt<std::string> OffloadArch("offload-arch", llvm::cl::Hidden);
@@ -310,24 +312,34 @@ int main(int argc, const char **argv) {
IEB->SlabAllocateSize = *SizeOrErr;
IEB->UseSharedMemory = UseSharedMemory;
- std::unique_ptr<clang::CompilerInstance> DeviceCI;
- if (CudaEnabled) {
- if (!CudaPath.empty())
- CB.SetCudaSDK(CudaPath);
+ if (HipEnabled && CudaEnabled) {
+ if (HipEnabled.getPosition() > CudaEnabled.getPosition())
+ CudaEnabled = false;
+ else
+ HipEnabled = false;
+ }
- if (OffloadArch.empty()) {
- OffloadArch = "sm_35";
- }
- CB.SetOffloadArch(OffloadArch);
+ bool DeviceEnabled = HipEnabled || CudaEnabled;
+ clang::OffloadType OffloadKind =
+ HipEnabled ? clang::OffloadType::HIP : clang::OffloadType::CUDA;
+ llvm::StringRef DevicePath = HipEnabled ? RocmPath : CudaPath;
+ // For HIP, let the driver auto-detect the GPU via --offload-arch=native.
+ llvm::StringRef DeviceOffloadArch = !OffloadArch.empty()
+ ? llvm::StringRef(OffloadArch)
+ : (HipEnabled ? "native" : "sm_35");
+ std::unique_ptr<clang::CompilerInstance> DeviceCI;
- DeviceCI = ExitOnErr(CB.CreateCudaDevice());
+ if (DeviceEnabled) {
+ CB.SetDeviceSDK(OffloadKind, DevicePath);
+ CB.SetOffloadArch(DeviceOffloadArch);
+ DeviceCI = ExitOnErr(CB.CreateDevice(OffloadKind));
}
// FIXME: Investigate if we could use runToolOnCodeWithArgs from tooling. It
// can replace the boilerplate code for creation of the compiler instance.
std::unique_ptr<clang::CompilerInstance> CI;
- if (CudaEnabled) {
- CI = ExitOnErr(CB.CreateCudaHost());
+ if (DeviceEnabled) {
+ CI = ExitOnErr(CB.CreateHost(OffloadKind));
} else {
CI = ExitOnErr(CB.CreateCpp());
}
@@ -339,20 +351,33 @@ int main(int argc, const char **argv) {
// Load any requested plugins.
CI->LoadRequestedPlugins();
- if (CudaEnabled)
+ if (DeviceEnabled)
DeviceCI->LoadRequestedPlugins();
std::unique_ptr<clang::Interpreter> Interp;
- if (CudaEnabled) {
- Interp = ExitOnErr(
- clang::Interpreter::createWithCUDA(std::move(CI), std::move(DeviceCI)));
+ if (DeviceEnabled) {
+ Interp = ExitOnErr(clang::Interpreter::createWithDevice(
+ OffloadKind, std::move(CI), std::move(DeviceCI)));
- if (CudaPath.empty()) {
- ExitOnErr(Interp->LoadDynamicLibrary("libcudart.so"));
- } else {
- auto CudaRuntimeLibPath = CudaPath + "/lib/libcudart.so";
- ExitOnErr(Interp->LoadDynamicLibrary(CudaRuntimeLibPath.c_str()));
+ if (HipEnabled) {
+ if (RocmPath.empty()) {
+ ExitOnErr(Interp->LoadDynamicLibrary("libamdhip64.so"));
+ } else {
+ auto RocmRuntimeLibPath = RocmPath + "/lib/libamdhip64.so";
+ ExitOnErr(Interp->LoadDynamicLibrary(RocmRuntimeLibPath.c_str()));
+ }
+ llvm::errs()
+ << "HIP environment is initialized but not supported as of now.\n";
+ return EXIT_FAILURE;
+ }
+ if (CudaEnabled) {
+ if (CudaPath.empty()) {
+ ExitOnErr(Interp->LoadDynamicLibrary("libcudart.so"));
+ } else {
+ auto CudaRuntimeLibPath = CudaPath + "/lib/libcudart.so";
+ ExitOnErr(Interp->LoadDynamicLibrary(CudaRuntimeLibPath.c_str()));
+ }
}
} else {
Interp =
More information about the cfe-commits
mailing list