[clang] Connect the HIP device parser into the interpreter (PR #218848)
Aditya Sinha via cfe-commits
cfe-commits at lists.llvm.org
Sat Sep 26 00:28:14 PDT 2026
https://github.com/AdityaSinha149 updated https://github.com/llvm/llvm-project/pull/218848
>From 0d99803d44048692ea4c501443aba4a299e0e212 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/2] [clang-repl] Hip environment initialized
---
clang/include/clang/Interpreter/Interpreter.h | 2 ++
1 file changed, 2 insertions(+)
diff --git a/clang/include/clang/Interpreter/Interpreter.h b/clang/include/clang/Interpreter/Interpreter.h
index b45c61a199f35..750583bf7ef56 100644
--- a/clang/include/clang/Interpreter/Interpreter.h
+++ b/clang/include/clang/Interpreter/Interpreter.h
@@ -48,6 +48,8 @@ class IncrementalDeviceParser;
enum class OffloadType { CUDA, HIP };
+enum class OffloadType { CUDA, HIP };
+
/// Create a pre-configured \c CompilerInstance for incremental processing.
class IncrementalCompilerBuilder {
using DriverCompilationFn = llvm::Error(const driver::Compilation &);
>From 0cb68137532f78a2fed52294cac976d65b57c3c1 Mon Sep 17 00:00:00 2001
From: AdityaSinha149 <adsinha at amd.com>
Date: Tue, 8 Sep 2026 15:59:58 +0530
Subject: [PATCH 2/2] [clang-repl] Connection between Hip Environment and Hip
Parser
---
clang/include/clang/Interpreter/Interpreter.h | 12 ++++---
clang/lib/Interpreter/Interpreter.cpp | 32 +++++++++++++++++--
.../HIP/device-function-template.hip | 26 +++++++++++++++
.../test/Interpreter/HIP/device-function.hip | 26 +++++++++++++++
.../test/Interpreter/HIP/host-and-device.hip | 29 +++++++++++++++++
clang/test/Interpreter/HIP/memory.hip | 25 +++++++++++++++
clang/test/Interpreter/HIP/sanity.hip | 13 ++++++++
7 files changed, 156 insertions(+), 7 deletions(-)
create mode 100644 clang/test/Interpreter/HIP/device-function-template.hip
create mode 100644 clang/test/Interpreter/HIP/device-function.hip
create mode 100644 clang/test/Interpreter/HIP/host-and-device.hip
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 750583bf7ef56..fef8e8abd279c 100644
--- a/clang/include/clang/Interpreter/Interpreter.h
+++ b/clang/include/clang/Interpreter/Interpreter.h
@@ -44,9 +44,8 @@ class CompilerInstance;
class CXXRecordDecl;
class Decl;
class IncrementalParser;
-class IncrementalDeviceParser;
-
-enum class OffloadType { CUDA, HIP };
+class IncrementalCUDADeviceParser;
+class IncrementalHIPDeviceParser;
enum class OffloadType { CUDA, HIP };
@@ -133,9 +132,12 @@ class Interpreter {
std::unique_ptr<IncrementalExecutor> IncrExecutor;
// An optional parser for CUDA offloading
- std::unique_ptr<IncrementalDeviceParser> DeviceParser;
+ std::unique_ptr<IncrementalCUDADeviceParser> DeviceParser;
+
+ // An optional parser for HIP offloading
+ std::unique_ptr<IncrementalHIPDeviceParser> HIPDeviceParser;
- // 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.
diff --git a/clang/lib/Interpreter/Interpreter.cpp b/clang/lib/Interpreter/Interpreter.cpp
index 8aafbc867ddaf..8e9e322f09007 100644
--- a/clang/lib/Interpreter/Interpreter.cpp
+++ b/clang/lib/Interpreter/Interpreter.cpp
@@ -413,6 +413,8 @@ Interpreter::~Interpreter() {
Act->FinalizeAction();
if (DeviceParser)
DeviceParser.reset();
+ if (HIPDeviceParser)
+ HIPDeviceParser.reset();
if (DeviceAct)
DeviceAct->FinalizeAction();
if (IncrExecutor) {
@@ -517,8 +519,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.
+ 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 DeviceParser = std::make_unique<IncrementalCUDADeviceParser>(
*Interp->DeviceCI, *Interp->getCompilerInstance(),
@@ -580,6 +588,26 @@ Interpreter::Parse(llvm::StringRef Code) {
return std::move(Err);
}
+ // If we have a HIP device parser, parse and lower the device code first so
+ // that the generated offload bundle is available to the host compilation.
+ 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->optimize())
+ return std::move(Err);
+
+ llvm::Expected<llvm::StringRef> HSACO = HIPDeviceParser->GenerateHSACO();
+ if (!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/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/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
More information about the cfe-commits
mailing list