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