https://github.com/AdityaSinha149 updated https://github.com/llvm/llvm-project/pull/218848
>From 0d99803d44048692ea4c501443aba4a299e0e212 Mon Sep 17 00:00:00 2001 From: AdityaSinha149 <[email protected]> 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 <[email protected]> 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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
