https://github.com/AdityaSinha149 updated https://github.com/llvm/llvm-project/pull/217582
>From 365f118dd935b3da664011ec43357c1cc6e78c77 Mon Sep 17 00:00:00 2001 From: AdityaSinha149 <[email protected]> Date: Wed, 19 Aug 2026 22:51:56 +0530 Subject: [PATCH 1/4] [clang-repl] Hip environment initialized --- clang/include/clang/Interpreter/Interpreter.h | 21 ++++--- clang/lib/Interpreter/Interpreter.cpp | 49 ++++++++------- .../test/Interpreter/HIP/hip-environment.hip | 8 +++ clang/test/Interpreter/HIP/lit.local.cfg | 2 + clang/test/lit.cfg.py | 38 +++++++++++- clang/tools/clang-repl/ClangRepl.cpp | 62 +++++++++++++------ 6 files changed, 130 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..757e74e3bafb8 100644 --- a/clang/include/clang/Interpreter/Interpreter.h +++ b/clang/include/clang/Interpreter/Interpreter.h @@ -65,27 +65,34 @@ class IncrementalCompilerBuilder { // Offload options void SetOffloadArch(llvm::StringRef Arch) { OffloadArch = Arch; }; - // CUDA specific - void SetCudaSDK(llvm::StringRef path) { CudaSDKPath = path; }; + void SetDeviceSDK(llvm::StringRef Path, bool HipEnabled) { + if (HipEnabled) + RocmSDKPath = Path; + else + CudaSDKPath = 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(bool HipEnabled); + llvm::Expected<std::unique_ptr<CompilerInstance>> + CreateDevice(bool HipEnabled); 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(bool HipEnabled, bool device); std::vector<const char *> UserArgs; std::optional<std::string> TargetTriple; llvm::StringRef OffloadArch; + llvm::StringRef RocmSDKPath; llvm::StringRef CudaSDKPath; std::optional<std::function<DriverCompilationFn>> CompilationCB; @@ -147,8 +154,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(bool HipEnabled, 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..25f664bbdf423 100644 --- a/clang/lib/Interpreter/Interpreter.cpp +++ b/clang/lib/Interpreter/Interpreter.cpp @@ -303,19 +303,16 @@ IncrementalCompilerBuilder::CreateCpp() { } llvm::Expected<std::unique_ptr<CompilerInstance>> -IncrementalCompilerBuilder::createCuda(bool device) { +IncrementalCompilerBuilder::createOffload(bool HipEnabled, 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"); - 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()); } @@ -332,13 +329,14 @@ IncrementalCompilerBuilder::createCuda(bool device) { } llvm::Expected<std::unique_ptr<CompilerInstance>> -IncrementalCompilerBuilder::CreateCudaDevice() { - return IncrementalCompilerBuilder::createCuda(true); +IncrementalCompilerBuilder::CreateDevice(bool HipEnabled) { + return IncrementalCompilerBuilder::createOffload(HipEnabled, /*device=*/true); } llvm::Expected<std::unique_ptr<CompilerInstance>> -IncrementalCompilerBuilder::CreateCudaHost() { - return IncrementalCompilerBuilder::createCuda(false); +IncrementalCompilerBuilder::CreateHost(bool HipEnabled) { + return IncrementalCompilerBuilder::createOffload(HipEnabled, + /*device=*/false); } Interpreter::Interpreter(std::unique_ptr<CompilerInstance> Instance, @@ -473,8 +471,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(bool HipEnabled, + 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 +507,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 (HipEnabled) { + // 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..350f118bc46db --- /dev/null +++ b/clang/test/Interpreter/HIP/hip-environment.hip @@ -0,0 +1,8 @@ +// Check that clang-repl initializes the HIP environment. HIP execution is not +// supported yet, so this only verifies that the environment is set up and that +// clang-repl reports it as unsupported. When both -cuda and -hip are passed, +// -hip wins (it appears later), so the HIP path is taken. + +// RUN: not clang-repl -cuda -hip 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..9827c80589237 100644 --- a/clang/test/lit.cfg.py +++ b/clang/test/lit.cfg.py @@ -222,6 +222,39 @@ def have_host_clang_repl_cuda(): return False +def have_host_clang_repl_hip(): + clang_repl_exe = lit.util.which("clang-repl", config.clang_tools_dir) + + if not clang_repl_exe: + return False + + testcode = b"\n".join( + [ + b"#include <hip/hip_runtime.h>", + b"__global__ void test_func() {}", + b"test_func<<<1,1>>>();", + b'extern "C" int puts(const char *s);', + b'puts(hipGetLastError() ? "failure" : "success");', + b"%quit", + ] + ) + try: + clang_repl_cmd = subprocess.run( + [clang_repl_exe, "--hip"], + stdout=subprocess.PIPE, + stderr=subprocess.PIPE, + input=testcode, + ) + except OSError: + return False + + if clang_repl_cmd.returncode == 0: + if clang_repl_cmd.stdout.find(b"success") != -1: + return True + + return False + + skip_clang_repl_checks = lit.util.pythonize_bool( lit_config.params.get( "clang_skip_clang_repl_checks", @@ -234,6 +267,9 @@ def have_host_clang_repl_cuda(): if have_host_clang_repl_cuda(): config.available_features.add('host-supports-cuda') + + if have_host_clang_repl_hip(): + config.available_features.add("host-supports-hip") hosttriple = run_clang_repl("--host-jit-triple") config.substitutions.append(("%host-jit-triple", hosttriple.strip())) @@ -508,4 +544,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..1367c39acda3c 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,31 @@ 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; + llvm::StringRef DevicePath = HipEnabled ? RocmPath : CudaPath; + llvm::StringRef DeviceOffloadArch = !OffloadArch.empty() + ? llvm::StringRef(OffloadArch) + : (HipEnabled ? "gfx90a" : "sm_35"); + std::unique_ptr<clang::CompilerInstance> DeviceCI; - DeviceCI = ExitOnErr(CB.CreateCudaDevice()); + if (DeviceEnabled) { + CB.SetDeviceSDK(DevicePath, HipEnabled); + CB.SetOffloadArch(DeviceOffloadArch); + DeviceCI = ExitOnErr(CB.CreateDevice(HipEnabled)); } // 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(HipEnabled)); } else { CI = ExitOnErr(CB.CreateCpp()); } @@ -339,20 +348,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( + HipEnabled, 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 = >From 7c771bfe63a247933f86c6f7062686557e8bb943 Mon Sep 17 00:00:00 2001 From: AdityaSinha149 <[email protected]> Date: Mon, 31 Aug 2026 13:15:15 +0530 Subject: [PATCH 2/4] only external hip env check --- clang/test/lit.cfg.py | 64 +++++++++++++++++++++++++++---------------- 1 file changed, 40 insertions(+), 24 deletions(-) diff --git a/clang/test/lit.cfg.py b/clang/test/lit.cfg.py index 9827c80589237..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,37 +223,52 @@ def have_host_clang_repl_cuda(): return False -def have_host_clang_repl_hip(): - clang_repl_exe = lit.util.which("clang-repl", config.clang_tools_dir) - - if not clang_repl_exe: - return False - - testcode = b"\n".join( - [ - b"#include <hip/hip_runtime.h>", - b"__global__ void test_func() {}", - b"test_func<<<1,1>>>();", - b'extern "C" int puts(const char *s);', - b'puts(hipGetLastError() ? "failure" : "success");', - b"%quit", - ] - ) +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: - clang_repl_cmd = subprocess.run( - [clang_repl_exe, "--hip"], + 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, - input=testcode, ) except OSError: return False + return proc.returncode == 0 - if clang_repl_cmd.returncode == 0: - if clang_repl_cmd.stdout.find(b"success") != -1: - return True - return False +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( @@ -268,7 +284,7 @@ def have_host_clang_repl_hip(): if have_host_clang_repl_cuda(): config.available_features.add('host-supports-cuda') - if have_host_clang_repl_hip(): + 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())) >From ce10fbefb2bed1d5e17b3d45184f868b3351efcb Mon Sep 17 00:00:00 2001 From: AdityaSinha149 <[email protected]> Date: Mon, 31 Aug 2026 13:39:40 +0530 Subject: [PATCH 3/4] made common cuid for both host and device --- clang/include/clang/Interpreter/Interpreter.h | 2 ++ clang/lib/Interpreter/Interpreter.cpp | 8 ++++++++ 2 files changed, 10 insertions(+) diff --git a/clang/include/clang/Interpreter/Interpreter.h b/clang/include/clang/Interpreter/Interpreter.h index 757e74e3bafb8..e682e8f1c2f00 100644 --- a/clang/include/clang/Interpreter/Interpreter.h +++ b/clang/include/clang/Interpreter/Interpreter.h @@ -95,6 +95,8 @@ class IncrementalCompilerBuilder { llvm::StringRef RocmSDKPath; llvm::StringRef CudaSDKPath; + std::string OffloadCUID; + std::optional<std::function<DriverCompilationFn>> CompilationCB; }; diff --git a/clang/lib/Interpreter/Interpreter.cpp b/clang/lib/Interpreter/Interpreter.cpp index 25f664bbdf423..de43a6d12a900 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" @@ -322,6 +324,12 @@ IncrementalCompilerBuilder::createOffload(bool HipEnabled, 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(); >From c5cedf0cc8334936315f61e87e7550d2e4cfddf8 Mon Sep 17 00:00:00 2001 From: AdityaSinha149 <[email protected]> Date: Mon, 31 Aug 2026 13:51:30 +0530 Subject: [PATCH 4/4] made offload-kind enum --- clang/include/clang/Interpreter/Interpreter.h | 27 ++++++++++++++----- clang/lib/Interpreter/Interpreter.cpp | 16 +++++------ clang/tools/clang-repl/ClangRepl.cpp | 10 ++++--- 3 files changed, 35 insertions(+), 18 deletions(-) diff --git a/clang/include/clang/Interpreter/Interpreter.h b/clang/include/clang/Interpreter/Interpreter.h index e682e8f1c2f00..454c3f691eaa7 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,28 +67,41 @@ class IncrementalCompilerBuilder { // Offload options void SetOffloadArch(llvm::StringRef Arch) { OffloadArch = Arch; }; - void SetDeviceSDK(llvm::StringRef Path, bool HipEnabled) { - if (HipEnabled) + 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>> CreateHost(bool HipEnabled); + llvm::Expected<std::unique_ptr<CompilerInstance>> CreateHost(OffloadType Type); llvm::Expected<std::unique_ptr<CompilerInstance>> - CreateDevice(bool HipEnabled); + 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>> - createOffload(bool HipEnabled, bool device); + createOffload(OffloadType Type, bool device); std::vector<const char *> UserArgs; std::optional<std::string> TargetTriple; @@ -156,7 +171,7 @@ class Interpreter { create(std::unique_ptr<CompilerInstance> CI, std::unique_ptr<IncrementalExecutorBuilder> IEB = nullptr); static llvm::Expected<std::unique_ptr<Interpreter>> - createWithDevice(bool HipEnabled, std::unique_ptr<CompilerInstance> CI, + createWithDevice(OffloadType Type, std::unique_ptr<CompilerInstance> CI, std::unique_ptr<CompilerInstance> DCI); const ASTContext &getASTContext() const; diff --git a/clang/lib/Interpreter/Interpreter.cpp b/clang/lib/Interpreter/Interpreter.cpp index de43a6d12a900..684e256f20830 100644 --- a/clang/lib/Interpreter/Interpreter.cpp +++ b/clang/lib/Interpreter/Interpreter.cpp @@ -305,7 +305,8 @@ IncrementalCompilerBuilder::CreateCpp() { } llvm::Expected<std::unique_ptr<CompilerInstance>> -IncrementalCompilerBuilder::createOffload(bool HipEnabled, 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"); @@ -337,14 +338,13 @@ IncrementalCompilerBuilder::createOffload(bool HipEnabled, bool device) { } llvm::Expected<std::unique_ptr<CompilerInstance>> -IncrementalCompilerBuilder::CreateDevice(bool HipEnabled) { - return IncrementalCompilerBuilder::createOffload(HipEnabled, /*device=*/true); +IncrementalCompilerBuilder::CreateDevice(OffloadType Type) { + return IncrementalCompilerBuilder::createOffload(Type, /*device=*/true); } llvm::Expected<std::unique_ptr<CompilerInstance>> -IncrementalCompilerBuilder::CreateHost(bool HipEnabled) { - return IncrementalCompilerBuilder::createOffload(HipEnabled, - /*device=*/false); +IncrementalCompilerBuilder::CreateHost(OffloadType Type) { + return IncrementalCompilerBuilder::createOffload(Type, /*device=*/false); } Interpreter::Interpreter(std::unique_ptr<CompilerInstance> Instance, @@ -479,7 +479,7 @@ llvm::Expected<std::unique_ptr<Interpreter>> Interpreter::create( } llvm::Expected<std::unique_ptr<Interpreter>> -Interpreter::createWithDevice(bool HipEnabled, +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 @@ -515,7 +515,7 @@ Interpreter::createWithDevice(bool HipEnabled, Interp->DeviceCI = std::move(DCI); - if (HipEnabled) { + if (Type == OffloadType::HIP) { // FIXME: HIP device parsing is not supported yet; it should use an // IncrementalHIPDeviceParser once one exists. } else { diff --git a/clang/tools/clang-repl/ClangRepl.cpp b/clang/tools/clang-repl/ClangRepl.cpp index 1367c39acda3c..eae6ff951df36 100644 --- a/clang/tools/clang-repl/ClangRepl.cpp +++ b/clang/tools/clang-repl/ClangRepl.cpp @@ -320,6 +320,8 @@ int main(int argc, const char **argv) { } bool DeviceEnabled = HipEnabled || CudaEnabled; + clang::OffloadType OffloadKind = + HipEnabled ? clang::OffloadType::HIP : clang::OffloadType::CUDA; llvm::StringRef DevicePath = HipEnabled ? RocmPath : CudaPath; llvm::StringRef DeviceOffloadArch = !OffloadArch.empty() ? llvm::StringRef(OffloadArch) @@ -327,16 +329,16 @@ int main(int argc, const char **argv) { std::unique_ptr<clang::CompilerInstance> DeviceCI; if (DeviceEnabled) { - CB.SetDeviceSDK(DevicePath, HipEnabled); + CB.SetDeviceSDK(OffloadKind, DevicePath); CB.SetOffloadArch(DeviceOffloadArch); - DeviceCI = ExitOnErr(CB.CreateDevice(HipEnabled)); + 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 (DeviceEnabled) { - CI = ExitOnErr(CB.CreateHost(HipEnabled)); + CI = ExitOnErr(CB.CreateHost(OffloadKind)); } else { CI = ExitOnErr(CB.CreateCpp()); } @@ -355,7 +357,7 @@ int main(int argc, const char **argv) { if (DeviceEnabled) { Interp = ExitOnErr(clang::Interpreter::createWithDevice( - HipEnabled, std::move(CI), std::move(DeviceCI))); + OffloadKind, std::move(CI), std::move(DeviceCI))); if (HipEnabled) { if (RocmPath.empty()) { _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
