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 <[email protected]> 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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
