llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang Author: Aditya Sinha (AdityaSinha149) <details> <summary>Changes</summary> This PR initializes the HIP environment for clang-repl. Since an IncrementalHipDeviceParser does not exist yet, a message is printed after the environment setup indicating that the HIP environment was set up but is not yet supported. It also adds a test file that runs clang-repl with the -hip option and checks that this message is printed. Once the parser is built, the current test file will be removed and replaced with tests that verify the complete flow. This PR has a prerequisite, #<!-- -->217228, which skips emitting the `__hip_cuid_` global under incremental extensions. Without it, clang-repl's HIP mode would emit the same symbol in every incremental module and fail at JIT link with a duplicate-symbol error. Assisted by Claude Opus 4.8 --- Full diff: https://github.com/llvm/llvm-project/pull/217582.diff 6 Files Affected: - (modified) clang/include/clang/Interpreter/Interpreter.h (+14-7) - (modified) clang/lib/Interpreter/Interpreter.cpp (+27-22) - (added) clang/test/Interpreter/HIP/hip-environment.hip (+8) - (added) clang/test/Interpreter/HIP/lit.local.cfg (+2) - (modified) clang/test/lit.cfg.py (+37-1) - (modified) clang/tools/clang-repl/ClangRepl.cpp (+42-20) ``````````diff 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 = `````````` </details> https://github.com/llvm/llvm-project/pull/217582 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
