[clang] [clang-repl] Initialized HIP environment for clang-repl (PR #217582)
Aditya Sinha via cfe-commits
cfe-commits at lists.llvm.org
Mon Aug 31 01:22:05 PDT 2026
https://github.com/AdityaSinha149 updated https://github.com/llvm/llvm-project/pull/217582
>From 365f118dd935b3da664011ec43357c1cc6e78c77 Mon Sep 17 00:00:00 2001
From: AdityaSinha149 <adsinha at amd.com>
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 <adsinha at amd.com>
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 <adsinha at amd.com>
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 <adsinha at amd.com>
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()) {
More information about the cfe-commits
mailing list