[clang] Connect the HIP device parser into the interpreter (PR #218848)
Aditya Sinha via cfe-commits
cfe-commits at lists.llvm.org
Wed Aug 26 05:14:45 PDT 2026
https://github.com/AdityaSinha149 updated https://github.com/llvm/llvm-project/pull/218848
>From 4ff0da3513f1d9b7c4292f14ca99a9887b082852 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/3] [clang-repl] Hip environment initialized
---
clang/include/clang/Interpreter/Interpreter.h | 25 ++++---
clang/lib/Interpreter/Interpreter.cpp | 69 +++++++++++--------
.../test/Interpreter/HIP/hip-environment.hip | 8 +++
clang/tools/clang-repl/ClangRepl.cpp | 65 +++++++++++------
4 files changed, 108 insertions(+), 59 deletions(-)
create mode 100644 clang/test/Interpreter/HIP/hip-environment.hip
diff --git a/clang/include/clang/Interpreter/Interpreter.h b/clang/include/clang/Interpreter/Interpreter.h
index c2622b23d5d9c..dd9fd9249d1ae 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;
@@ -106,9 +113,9 @@ class Interpreter {
std::unique_ptr<IncrementalExecutor> IncrExecutor;
// An optional parser for CUDA offloading
- std::unique_ptr<IncrementalCUDADeviceParser> DeviceParser;
+ std::unique_ptr<IncrementalCUDADeviceParser> CUDADeviceParser;
- // An optional action for CUDA offloading
+ // An optional action for Device offloading
std::unique_ptr<IncrementalAction> DeviceAct;
/// List containing information about each incrementally parsed piece of code.
@@ -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 ef1ca31538352..1a2a7bd3093d6 100644
--- a/clang/lib/Interpreter/Interpreter.cpp
+++ b/clang/lib/Interpreter/Interpreter.cpp
@@ -302,19 +302,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());
}
@@ -324,6 +321,9 @@ IncrementalCompilerBuilder::createCuda(bool device) {
Argv.push_back(ArchArg.c_str());
}
+ if (device && HipEnabled)
+ Argv.push_back("-O3");
+
llvm::append_range(Argv, UserArgs);
std::string TT = TargetTriple ? *TargetTriple : llvm::sys::getProcessTriple();
@@ -331,13 +331,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,
@@ -403,8 +404,8 @@ Interpreter::Interpreter(std::unique_ptr<CompilerInstance> Instance,
Interpreter::~Interpreter() {
IncrParser.reset();
Act->FinalizeAction();
- if (DeviceParser)
- DeviceParser.reset();
+ if (CUDADeviceParser)
+ CUDADeviceParser.reset();
if (DeviceAct)
DeviceAct->FinalizeAction();
if (IncrExecutor) {
@@ -472,8 +473,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>();
@@ -507,14 +509,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 CUDADeviceParser = 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->CUDADeviceParser = std::move(CUDADeviceParser);
+ }
- Interp->DeviceParser = std::move(DeviceParser);
return std::move(Interp);
}
@@ -554,18 +562,19 @@ llvm::Expected<PartialTranslationUnit &>
Interpreter::Parse(llvm::StringRef Code) {
// If we have a device parser, parse it first. The generated code will be
// included in the host compilation
- if (DeviceParser) {
- llvm::Expected<TranslationUnitDecl *> DeviceTU = DeviceParser->Parse(Code);
+ if (CUDADeviceParser) {
+ llvm::Expected<TranslationUnitDecl *> DeviceTU =
+ CUDADeviceParser->Parse(Code);
if (auto E = DeviceTU.takeError())
return std::move(E);
- DeviceParser->RegisterPTU(*DeviceTU);
+ CUDADeviceParser->RegisterPTU(*DeviceTU);
- llvm::Expected<llvm::StringRef> PTX = DeviceParser->GeneratePTX();
+ llvm::Expected<llvm::StringRef> PTX = CUDADeviceParser->GeneratePTX();
if (!PTX)
return PTX.takeError();
- llvm::Error Err = DeviceParser->GenerateFatbinary();
+ llvm::Error Err = CUDADeviceParser->GenerateFatbinary();
if (Err)
return std::move(Err);
}
@@ -710,4 +719,4 @@ llvm::Error Interpreter::LoadDynamicLibrary(const char *name) {
return EEOrErr->LoadDynamicLibrary(name);
}
-} // end namespace clang
+} // end namespace clang
\ No newline at end of file
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/tools/clang-repl/ClangRepl.cpp b/clang/tools/clang-repl/ClangRepl.cpp
index c9873540a5d66..7b7fc69330409 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 ? "gfx906" : "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,36 @@ 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().changeColor(llvm::raw_ostream::RED, /*Bold=*/true);
+ llvm::errs()
+ << "HIP environment is initialized but not supported as of now.\n";
+ llvm::errs().resetColor();
+ 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 685b2ff442f395cf17404ead2ea07d4846610adb Mon Sep 17 00:00:00 2001
From: AdityaSinha149 <adsinha at amd.com>
Date: Wed, 26 Aug 2026 12:11:34 +0530
Subject: [PATCH 2/3] [clang-repl] Connect the HIP device parser into the
interpreter
---
clang/include/clang/Interpreter/Interpreter.h | 4 ++
clang/lib/Interpreter/Interpreter.cpp | 39 ++++++++++++++++++-
.../HIP/device-function-template.hip | 26 +++++++++++++
.../test/Interpreter/HIP/device-function.hip | 26 +++++++++++++
.../test/Interpreter/HIP/hip-environment.hip | 8 ----
.../test/Interpreter/HIP/host-and-device.hip | 29 ++++++++++++++
clang/test/Interpreter/HIP/lit.local.cfg | 2 +
clang/test/Interpreter/HIP/memory.hip | 25 ++++++++++++
clang/test/Interpreter/HIP/sanity.hip | 13 +++++++
clang/test/lit.cfg.py | 31 +++++++++++++++
10 files changed, 193 insertions(+), 10 deletions(-)
create mode 100644 clang/test/Interpreter/HIP/device-function-template.hip
create mode 100644 clang/test/Interpreter/HIP/device-function.hip
delete mode 100644 clang/test/Interpreter/HIP/hip-environment.hip
create mode 100644 clang/test/Interpreter/HIP/host-and-device.hip
create mode 100644 clang/test/Interpreter/HIP/lit.local.cfg
create mode 100644 clang/test/Interpreter/HIP/memory.hip
create mode 100644 clang/test/Interpreter/HIP/sanity.hip
diff --git a/clang/include/clang/Interpreter/Interpreter.h b/clang/include/clang/Interpreter/Interpreter.h
index dd9fd9249d1ae..bc2a399b8f704 100644
--- a/clang/include/clang/Interpreter/Interpreter.h
+++ b/clang/include/clang/Interpreter/Interpreter.h
@@ -45,6 +45,7 @@ class CXXRecordDecl;
class Decl;
class IncrementalParser;
class IncrementalCUDADeviceParser;
+class IncrementalHIPDeviceParser;
/// Create a pre-configured \c CompilerInstance for incremental processing.
class IncrementalCompilerBuilder {
@@ -115,6 +116,9 @@ class Interpreter {
// An optional parser for CUDA offloading
std::unique_ptr<IncrementalCUDADeviceParser> CUDADeviceParser;
+ // An optional parser for HIP offloading
+ std::unique_ptr<IncrementalHIPDeviceParser> HIPDeviceParser;
+
// An optional action for Device offloading
std::unique_ptr<IncrementalAction> DeviceAct;
diff --git a/clang/lib/Interpreter/Interpreter.cpp b/clang/lib/Interpreter/Interpreter.cpp
index 1a2a7bd3093d6..586b38dfd5c38 100644
--- a/clang/lib/Interpreter/Interpreter.cpp
+++ b/clang/lib/Interpreter/Interpreter.cpp
@@ -406,6 +406,8 @@ Interpreter::~Interpreter() {
Act->FinalizeAction();
if (CUDADeviceParser)
CUDADeviceParser.reset();
+ if (HIPDeviceParser)
+ HIPDeviceParser.reset();
if (DeviceAct)
DeviceAct->FinalizeAction();
if (IncrExecutor) {
@@ -510,8 +512,14 @@ Interpreter::createWithDevice(bool HipEnabled,
Interp->DeviceCI = std::move(DCI);
if (HipEnabled) {
- // FIXME: HIP device parsing is not supported yet; it should use an
- // IncrementalHIPDeviceParser once one exists.
+ auto HIPDeviceParser = std::make_unique<IncrementalHIPDeviceParser>(
+ *Interp->DeviceCI, *Interp->getCompilerInstance(),
+ Interp->DeviceAct.get(), IMVFS, Err, Interp->PTUs);
+
+ if (Err)
+ return std::move(Err);
+
+ Interp->HIPDeviceParser = std::move(HIPDeviceParser);
} else {
auto CUDADeviceParser = std::make_unique<IncrementalCUDADeviceParser>(
*Interp->DeviceCI, *Interp->getCompilerInstance(),
@@ -579,6 +587,33 @@ Interpreter::Parse(llvm::StringRef Code) {
return std::move(Err);
}
+ // HIP device code is compiled into its own code object per chunk. Each PTU is
+ // linked against the device libraries, optimized, and lowered to a
+ // relocatable code object (.hsaco) that is wrapped in a fat binary and
+ // embedded into the host module.
+ if (HIPDeviceParser) {
+ llvm::Expected<TranslationUnitDecl *> DeviceTU =
+ HIPDeviceParser->Parse(Code);
+ if (auto E = DeviceTU.takeError())
+ return std::move(E);
+
+ HIPDeviceParser->RegisterPTU(*DeviceTU);
+
+ if (llvm::Error Err = HIPDeviceParser->LinkDeviceLibs())
+ return std::move(Err);
+
+ if (llvm::Error Err = HIPDeviceParser->optimize())
+ return std::move(Err);
+
+ if (llvm::Expected<llvm::StringRef> HSACO =
+ HIPDeviceParser->GenerateHSACO();
+ !HSACO)
+ return HSACO.takeError();
+
+ if (llvm::Error Err = HIPDeviceParser->GenerateOffloadBundle())
+ return std::move(Err);
+ }
+
// Tell the interpreter sliently ignore unused expressions since value
// printing could cause it.
getCompilerInstance()->getDiagnostics().setSeverity(
diff --git a/clang/test/Interpreter/HIP/device-function-template.hip b/clang/test/Interpreter/HIP/device-function-template.hip
new file mode 100644
index 0000000000000..b93acd8becfb8
--- /dev/null
+++ b/clang/test/Interpreter/HIP/device-function-template.hip
@@ -0,0 +1,26 @@
+// Tests device function templates
+// RUN: cat %s | clang-repl --hip | FileCheck %s
+
+#include <hip/hip_runtime.h>
+
+extern "C" int printf(const char*, ...);
+
+template <typename T> __device__ inline T sum(T a, T b) { return a + b; }
+__global__ void test_kernel(int* value) { *value = sum(40, 2); }
+
+int var;
+int* devptr = nullptr;
+printf("hipMalloc: %d\n", hipMalloc((void **) &devptr, sizeof(int)));
+// CHECK: hipMalloc: 0
+
+test_kernel<<<1,1>>>(devptr);
+printf("HIP Error: %d\n", hipGetLastError());
+// CHECK-NEXT: HIP Error: 0
+
+printf("hipMemcpy: %d\n", hipMemcpy(&var, devptr, sizeof(int), hipMemcpyDeviceToHost));
+// CHECK-NEXT: hipMemcpy: 0
+
+printf("Value: %d\n", var);
+// CHECK-NEXT: Value: 42
+
+%quit
diff --git a/clang/test/Interpreter/HIP/device-function.hip b/clang/test/Interpreter/HIP/device-function.hip
new file mode 100644
index 0000000000000..fc3d159f59579
--- /dev/null
+++ b/clang/test/Interpreter/HIP/device-function.hip
@@ -0,0 +1,26 @@
+// Tests __device__ function calls
+// RUN: cat %s | clang-repl --hip | FileCheck %s
+
+#include <hip/hip_runtime.h>
+
+extern "C" int printf(const char*, ...);
+
+__device__ inline void test_device(int* value) { *value = 42; }
+__global__ void test_kernel(int* value) { test_device(value); }
+
+int var;
+int* devptr = nullptr;
+printf("hipMalloc: %d\n", hipMalloc((void **) &devptr, sizeof(int)));
+// CHECK: hipMalloc: 0
+
+test_kernel<<<1,1>>>(devptr);
+printf("HIP Error: %d\n", hipGetLastError());
+// CHECK-NEXT: HIP Error: 0
+
+printf("hipMemcpy: %d\n", hipMemcpy(&var, devptr, sizeof(int), hipMemcpyDeviceToHost));
+// CHECK-NEXT: hipMemcpy: 0
+
+printf("Value: %d\n", var);
+// CHECK-NEXT: Value: 42
+
+%quit
diff --git a/clang/test/Interpreter/HIP/hip-environment.hip b/clang/test/Interpreter/HIP/hip-environment.hip
deleted file mode 100644
index 350f118bc46db..0000000000000
--- a/clang/test/Interpreter/HIP/hip-environment.hip
+++ /dev/null
@@ -1,8 +0,0 @@
-// 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/host-and-device.hip b/clang/test/Interpreter/HIP/host-and-device.hip
new file mode 100644
index 0000000000000..1996ad543cd9e
--- /dev/null
+++ b/clang/test/Interpreter/HIP/host-and-device.hip
@@ -0,0 +1,29 @@
+// Checks that a function is available in both __host__ and __device__
+// RUN: cat %s | clang-repl --hip | FileCheck %s
+
+#include <hip/hip_runtime.h>
+
+extern "C" int printf(const char*, ...);
+
+__host__ __device__ inline int sum(int a, int b){ return a + b; }
+__global__ void kernel(int * output){ *output = sum(40,2); }
+
+printf("Host sum: %d\n", sum(41,1));
+// CHECK: Host sum: 42
+
+int var = 0;
+int * deviceVar;
+printf("hipMalloc: %d\n", hipMalloc((void **) &deviceVar, sizeof(int)));
+// CHECK-NEXT: hipMalloc: 0
+
+kernel<<<1,1>>>(deviceVar);
+printf("HIP Error: %d\n", hipGetLastError());
+// CHECK-NEXT: HIP Error: 0
+
+printf("hipMemcpy: %d\n", hipMemcpy(&var, deviceVar, sizeof(int), hipMemcpyDeviceToHost));
+// CHECK-NEXT: hipMemcpy: 0
+
+printf("var: %d\n", var);
+// CHECK-NEXT: var: 42
+
+%quit
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/Interpreter/HIP/memory.hip b/clang/test/Interpreter/HIP/memory.hip
new file mode 100644
index 0000000000000..67120eaf2ad10
--- /dev/null
+++ b/clang/test/Interpreter/HIP/memory.hip
@@ -0,0 +1,25 @@
+// Tests hipMemcpy and writes from kernel
+// RUN: cat %s | clang-repl --hip | FileCheck %s
+
+#include <hip/hip_runtime.h>
+
+extern "C" int printf(const char*, ...);
+
+__global__ void test_func(int* value) { *value = 42; }
+
+int var;
+int* devptr = nullptr;
+printf("hipMalloc: %d\n", hipMalloc((void **) &devptr, sizeof(int)));
+// CHECK: hipMalloc: 0
+
+test_func<<<1,1>>>(devptr);
+printf("HIP Error: %d\n", hipGetLastError());
+// CHECK-NEXT: HIP Error: 0
+
+printf("hipMemcpy: %d\n", hipMemcpy(&var, devptr, sizeof(int), hipMemcpyDeviceToHost));
+// CHECK-NEXT: hipMemcpy: 0
+
+printf("Value: %d\n", var);
+// CHECK-NEXT: Value: 42
+
+%quit
diff --git a/clang/test/Interpreter/HIP/sanity.hip b/clang/test/Interpreter/HIP/sanity.hip
new file mode 100644
index 0000000000000..293e5f4b92807
--- /dev/null
+++ b/clang/test/Interpreter/HIP/sanity.hip
@@ -0,0 +1,13 @@
+// RUN: cat %s | clang-repl --hip | FileCheck %s
+
+#include <hip/hip_runtime.h>
+
+extern "C" int printf(const char*, ...);
+
+__global__ void test_func() {}
+
+test_func<<<1,1>>>();
+printf("HIP Error: %d", hipGetLastError());
+// CHECK: HIP Error: 0
+
+%quit
diff --git a/clang/test/lit.cfg.py b/clang/test/lit.cfg.py
index 9b7bd1d329d22..4ec5e94fbf33f 100644
--- a/clang/test/lit.cfg.py
+++ b/clang/test/lit.cfg.py
@@ -221,6 +221,34 @@ 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(
@@ -234,6 +262,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()))
>From 31b9d6d7918b85f0af801c39683bdf8630983cea Mon Sep 17 00:00:00 2001
From: AdityaSinha149 <adsinha at amd.com>
Date: Wed, 26 Aug 2026 12:15:24 +0530
Subject: [PATCH 3/3] mend
---
clang/lib/Interpreter/Interpreter.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/clang/lib/Interpreter/Interpreter.cpp b/clang/lib/Interpreter/Interpreter.cpp
index 586b38dfd5c38..a17b6c7437cb6 100644
--- a/clang/lib/Interpreter/Interpreter.cpp
+++ b/clang/lib/Interpreter/Interpreter.cpp
@@ -754,4 +754,4 @@ llvm::Error Interpreter::LoadDynamicLibrary(const char *name) {
return EEOrErr->LoadDynamicLibrary(name);
}
-} // end namespace clang
\ No newline at end of file
+} // end namespace clang
More information about the cfe-commits
mailing list