[clang] Revert "[clang-repl] Initialized HIP environment for clang-repl (#217582) (PR #228976)
Vassil Vassilev via cfe-commits
cfe-commits at lists.llvm.org
Sun Oct 4 22:23:00 PDT 2026
https://github.com/vgvassilev created https://github.com/llvm/llvm-project/pull/228976
This reverts commit d220d5e8239db918bd68f404dae589385a8029e8.
The change landed without review from the clang-repl code owners and its test fails in some build configurations (see the post-commit discussion on #217582). Revert so the work can go through a proper review, as agreed with the author.
This is not a pure revert: DeviceOffloadTest.cpp, added in #226975 and extended in #226977 on top of the reverted commit, is ported back to the CUDA-specific API (CreateCudaHost, CreateCudaDevice, createWithCUDA). The CUDA fixes from those two commits are kept.
Supersedes #228088.
>From ac2abbf951552a8bea170e16a6c30b37eabd7984 Mon Sep 17 00:00:00 2001
From: Vassil Vassilev <v.g.vassilev at gmail.com>
Date: Mon, 5 Oct 2026 05:16:54 +0000
Subject: [PATCH] Revert "[clang-repl] Initialized HIP environment for
clang-repl (#217582)"
This reverts commit d220d5e8239db918bd68f404dae589385a8029e8.
The change landed without review from the clang-repl code owners and
its test fails in some build configurations (see the post-commit
discussion on #217582). Revert so the work can go through a proper
review, as agreed with the author.
This is not a pure revert: DeviceOffloadTest.cpp, added in #226975 and
extended in #226977 on top of the reverted commit, is ported back to
the CUDA-specific API (CreateCudaHost, CreateCudaDevice,
createWithCUDA). The CUDA fixes from those two commits are kept.
Supersedes #228088.
---
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 ++++++-------------
.../Interpreter/DeviceOffloadTest.cpp | 16 ++---
7 files changed, 58 insertions(+), 186 deletions(-)
delete mode 100644 clang/test/Interpreter/HIP/hip-environment.hip
delete 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 4504b679504e0..c2622b23d5d9c 100644
--- a/clang/include/clang/Interpreter/Interpreter.h
+++ b/clang/include/clang/Interpreter/Interpreter.h
@@ -46,8 +46,6 @@ 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 &);
@@ -67,52 +65,29 @@ class IncrementalCompilerBuilder {
// Offload options
void SetOffloadArch(llvm::StringRef Arch) { OffloadArch = Arch; };
- 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);
- }
+ // CUDA specific
+ void SetCudaSDK(llvm::StringRef path) { CudaSDKPath = path; };
// Hand over the compilation.
void SetDriverCompilationCallback(std::function<DriverCompilationFn> C) {
CompilationCB = C;
}
- 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);
- }
+ llvm::Expected<std::unique_ptr<CompilerInstance>> CreateCudaHost();
+ llvm::Expected<std::unique_ptr<CompilerInstance>> CreateCudaDevice();
private:
llvm::Expected<std::unique_ptr<CompilerInstance>>
create(std::string TT, std::vector<const char *> &ClangArgv);
- llvm::Expected<std::unique_ptr<CompilerInstance>>
- createOffload(OffloadType Type, bool device);
+ llvm::Expected<std::unique_ptr<CompilerInstance>> createCuda(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;
};
@@ -172,8 +147,8 @@ class Interpreter {
create(std::unique_ptr<CompilerInstance> CI,
std::unique_ptr<IncrementalExecutorBuilder> IEB = nullptr);
static llvm::Expected<std::unique_ptr<Interpreter>>
- createWithDevice(OffloadType Type, std::unique_ptr<CompilerInstance> CI,
- std::unique_ptr<CompilerInstance> DCI);
+ createWithCUDA(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 5d5629997e94f..40fbc83db9107 100644
--- a/clang/lib/Interpreter/Interpreter.cpp
+++ b/clang/lib/Interpreter/Interpreter.cpp
@@ -45,14 +45,12 @@
#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"
@@ -306,17 +304,19 @@ IncrementalCompilerBuilder::CreateCpp() {
}
llvm::Expected<std::unique_ptr<CompilerInstance>>
-IncrementalCompilerBuilder::createOffload(OffloadType Type, bool device) {
- const bool HipEnabled = Type == OffloadType::HIP;
+IncrementalCompilerBuilder::createCuda(bool device) {
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");
- llvm::StringRef SDKPath = HipEnabled ? RocmSDKPath : CudaSDKPath;
- std::string SDKPathArg = HipEnabled ? "--rocm-path=" : "--cuda-path=";
- if (!SDKPath.empty()) {
- SDKPathArg += SDKPath;
+ 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;
Argv.push_back(SDKPathArg.c_str());
}
@@ -326,12 +326,6 @@ IncrementalCompilerBuilder::createOffload(OffloadType Type, 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();
@@ -339,13 +333,13 @@ IncrementalCompilerBuilder::createOffload(OffloadType Type, bool device) {
}
llvm::Expected<std::unique_ptr<CompilerInstance>>
-IncrementalCompilerBuilder::CreateDevice(OffloadType Type) {
- return IncrementalCompilerBuilder::createOffload(Type, /*device=*/true);
+IncrementalCompilerBuilder::CreateCudaDevice() {
+ return IncrementalCompilerBuilder::createCuda(true);
}
llvm::Expected<std::unique_ptr<CompilerInstance>>
-IncrementalCompilerBuilder::CreateHost(OffloadType Type) {
- return IncrementalCompilerBuilder::createOffload(Type, /*device=*/false);
+IncrementalCompilerBuilder::CreateCudaHost() {
+ return IncrementalCompilerBuilder::createCuda(false);
}
Interpreter::Interpreter(std::unique_ptr<CompilerInstance> Instance,
@@ -480,9 +474,8 @@ llvm::Expected<std::unique_ptr<Interpreter>> Interpreter::create(
}
llvm::Expected<std::unique_ptr<Interpreter>>
-Interpreter::createWithDevice(OffloadType Type,
- std::unique_ptr<CompilerInstance> CI,
- std::unique_ptr<CompilerInstance> DCI) {
+Interpreter::createWithCUDA(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>();
@@ -519,20 +512,14 @@ Interpreter::createWithDevice(OffloadType Type,
Interp->DeviceCI = std::move(DCI);
- 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);
+ auto DeviceParser = std::make_unique<IncrementalCUDADeviceParser>(
+ *Interp->DeviceCI, *Interp->getCompilerInstance(),
+ Interp->DeviceAct.get(), IMVFS, Err, Interp->PTUs);
- Interp->DeviceParser = std::move(DeviceParser);
- }
+ if (Err)
+ return std::move(Err);
+ 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
deleted file mode 100644
index 16354ecf44e98..0000000000000
--- a/clang/test/Interpreter/HIP/hip-environment.hip
+++ /dev/null
@@ -1,11 +0,0 @@
-// 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
deleted file mode 100644
index 70102544ab0fd..0000000000000
--- a/clang/test/Interpreter/HIP/lit.local.cfg
+++ /dev/null
@@ -1,2 +0,0 @@
-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 ba7d9eba91cd8..af3144f3ddcf9 100644
--- a/clang/test/lit.cfg.py
+++ b/clang/test/lit.cfg.py
@@ -1,6 +1,5 @@
# -*- Python -*-
-import glob
import os
import platform
import re
@@ -223,54 +222,6 @@ 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",
@@ -283,9 +234,6 @@ def have_host_hip_environment():
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.available_features.add("host-jit-triple=" + hosttriple.strip())
config.substitutions.append(("%host-jit-triple", hosttriple.strip()))
@@ -561,4 +509,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)
\ No newline at end of file
+ lit_config.test_updaters.append(utc_lit_plugin)
diff --git a/clang/tools/clang-repl/ClangRepl.cpp b/clang/tools/clang-repl/ClangRepl.cpp
index 15eb8ed3034c9..c9873540a5d66 100644
--- a/clang/tools/clang-repl/ClangRepl.cpp
+++ b/clang/tools/clang-repl/ClangRepl.cpp
@@ -52,8 +52,6 @@ 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);
@@ -312,34 +310,24 @@ int main(int argc, const char **argv) {
IEB->SlabAllocateSize = *SizeOrErr;
IEB->UseSharedMemory = UseSharedMemory;
- if (HipEnabled && CudaEnabled) {
- if (HipEnabled.getPosition() > CudaEnabled.getPosition())
- CudaEnabled = false;
- else
- HipEnabled = false;
- }
-
- 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;
+ if (CudaEnabled) {
+ if (!CudaPath.empty())
+ CB.SetCudaSDK(CudaPath);
+
+ if (OffloadArch.empty()) {
+ OffloadArch = "sm_35";
+ }
+ CB.SetOffloadArch(OffloadArch);
- if (DeviceEnabled) {
- CB.SetDeviceSDK(OffloadKind, DevicePath);
- CB.SetOffloadArch(DeviceOffloadArch);
- DeviceCI = ExitOnErr(CB.CreateDevice(OffloadKind));
+ DeviceCI = ExitOnErr(CB.CreateCudaDevice());
}
// 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 (DeviceEnabled) {
- CI = ExitOnErr(CB.CreateHost(OffloadKind));
+ if (CudaEnabled) {
+ CI = ExitOnErr(CB.CreateCudaHost());
} else {
CI = ExitOnErr(CB.CreateCpp());
}
@@ -351,33 +339,20 @@ int main(int argc, const char **argv) {
// Load any requested plugins.
CI->LoadRequestedPlugins();
- if (DeviceEnabled)
+ if (CudaEnabled)
DeviceCI->LoadRequestedPlugins();
std::unique_ptr<clang::Interpreter> Interp;
- if (DeviceEnabled) {
- Interp = ExitOnErr(clang::Interpreter::createWithDevice(
- OffloadKind, std::move(CI), std::move(DeviceCI)));
+ if (CudaEnabled) {
+ Interp = ExitOnErr(
+ clang::Interpreter::createWithCUDA(std::move(CI), std::move(DeviceCI)));
- 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()));
- }
+ if (CudaPath.empty()) {
+ ExitOnErr(Interp->LoadDynamicLibrary("libcudart.so"));
+ } else {
+ auto CudaRuntimeLibPath = CudaPath + "/lib/libcudart.so";
+ ExitOnErr(Interp->LoadDynamicLibrary(CudaRuntimeLibPath.c_str()));
}
} else {
Interp =
diff --git a/clang/unittests/Interpreter/DeviceOffloadTest.cpp b/clang/unittests/Interpreter/DeviceOffloadTest.cpp
index 8a2385e05a7ff..039c919fe1e1e 100644
--- a/clang/unittests/Interpreter/DeviceOffloadTest.cpp
+++ b/clang/unittests/Interpreter/DeviceOffloadTest.cpp
@@ -62,10 +62,10 @@ TEST_F(DeviceOffloadTest, FirstDeviceModuleVerifies) {
IncrementalCompilerBuilder CB;
CB.SetCompilerArgs(
{"-nocudainc", "-nocudalib", "-fverify-intermediate-code"});
- auto DeviceCI = llvm::cantFail(CB.CreateDevice(OffloadType::CUDA));
- auto HostCI = llvm::cantFail(CB.CreateHost(OffloadType::CUDA));
- auto Interp = llvm::cantFail(Interpreter::createWithDevice(
- OffloadType::CUDA, std::move(HostCI), std::move(DeviceCI)));
+ auto DeviceCI = llvm::cantFail(CB.CreateCudaDevice());
+ auto HostCI = llvm::cantFail(CB.CreateCudaHost());
+ auto Interp = llvm::cantFail(Interpreter::createWithCUDA(
+ std::move(HostCI), std::move(DeviceCI)));
llvm::cantFail(Interp->Parse("__attribute__((device)) void f() {}"));
exit(0);
},
@@ -79,12 +79,12 @@ TEST_F(DeviceOffloadTest, EmptyDeviceModule) {
// Without the runtime headers and libdevice no CUDA toolkit is needed.
IncrementalCompilerBuilder CB;
CB.SetCompilerArgs({"-nocudainc", "-nocudalib"});
- auto DeviceCI = CB.CreateDevice(OffloadType::CUDA);
+ auto DeviceCI = CB.CreateCudaDevice();
ASSERT_THAT_EXPECTED(DeviceCI, llvm::Succeeded());
- auto HostCI = CB.CreateHost(OffloadType::CUDA);
+ auto HostCI = CB.CreateCudaHost();
ASSERT_THAT_EXPECTED(HostCI, llvm::Succeeded());
- auto Interp = Interpreter::createWithDevice(
- OffloadType::CUDA, std::move(*HostCI), std::move(*DeviceCI));
+ auto Interp =
+ Interpreter::createWithCUDA(std::move(*HostCI), std::move(*DeviceCI));
ASSERT_THAT_EXPECTED(Interp, llvm::Succeeded());
// A host-only input leaves the device module without a function. Its PTX
More information about the cfe-commits
mailing list