[clang] [llvm] [Offload] Add GetErrorName and GetErrorString to LLVMOffloadKernel (PR #212887)
Sophia Herrmann via llvm-commits
llvm-commits at lists.llvm.org
Fri Aug 7 09:55:18 PDT 2026
https://github.com/jellytabby updated https://github.com/llvm/llvm-project/pull/212887
>From 9493ef073072f6b330bf22011d567c7a38308aa0 Mon Sep 17 00:00:00 2001
From: Sophia Herrmann <herrmann15 at llnl.gov>
Date: Wed, 22 Jul 2026 16:57:48 -0700
Subject: [PATCH 1/7] generate LLVMOffloadKernel library
Co-authored-by: Johannes Doerfert <jdoerfert.llvm at gmail.com>
Co-authored-by: Jonas Greifenhain <cadivus at daverkomp.de>
remove old cmake variable
---
offload/CMakeLists.txt | 1 +
offload/languages/CMakeLists.txt | 1 +
offload/languages/kernel/CMakeLists.txt | 35 ++++
offload/languages/kernel/exports | 8 +
.../languages/kernel/include/ExportedAPI.h | 41 ++++
offload/languages/kernel/include/State.h | 106 ++++++++++
offload/languages/kernel/include/Types.h | 28 +++
offload/languages/kernel/src/ExportedAPI.cpp | 99 +++++++++
offload/languages/kernel/src/State.cpp | 198 ++++++++++++++++++
9 files changed, 517 insertions(+)
create mode 100644 offload/languages/CMakeLists.txt
create mode 100644 offload/languages/kernel/CMakeLists.txt
create mode 100644 offload/languages/kernel/exports
create mode 100644 offload/languages/kernel/include/ExportedAPI.h
create mode 100644 offload/languages/kernel/include/State.h
create mode 100644 offload/languages/kernel/include/Types.h
create mode 100644 offload/languages/kernel/src/ExportedAPI.cpp
create mode 100644 offload/languages/kernel/src/State.cpp
diff --git a/offload/CMakeLists.txt b/offload/CMakeLists.txt
index 2885b20f9c1d8..2dd4446979c05 100644
--- a/offload/CMakeLists.txt
+++ b/offload/CMakeLists.txt
@@ -340,6 +340,7 @@ if(BUILD_LIBOMPTARGET)
endif()
add_subdirectory(liboffload)
+add_subdirectory(languages)
# Add tests.
if(OFFLOAD_INCLUDE_TESTS)
diff --git a/offload/languages/CMakeLists.txt b/offload/languages/CMakeLists.txt
new file mode 100644
index 0000000000000..08b2a68082d80
--- /dev/null
+++ b/offload/languages/CMakeLists.txt
@@ -0,0 +1 @@
+add_subdirectory(kernel)
diff --git a/offload/languages/kernel/CMakeLists.txt b/offload/languages/kernel/CMakeLists.txt
new file mode 100644
index 0000000000000..3b0823a44faf1
--- /dev/null
+++ b/offload/languages/kernel/CMakeLists.txt
@@ -0,0 +1,35 @@
+add_llvm_library(
+ LLVMOffloadKernel SHARED
+
+ src/State.cpp
+ src/ExportedAPI.cpp
+
+ LINK_COMPONENTS
+ Support
+ Offload
+ )
+
+if(LIBOMP_HAVE_VERSION_SCRIPT_FLAG)
+ target_link_libraries(LLVMOffloadKernel PRIVATE "-Wl,--version-script=${CMAKE_CURRENT_SOURCE_DIR}/exports")
+endif()
+
+target_include_directories(LLVMOffloadKernel PUBLIC
+ ${CMAKE_CURRENT_SOURCE_DIR}/include
+ ${CMAKE_CURRENT_SOURCE_DIR}/../../liboffload/include
+ ${CMAKE_CURRENT_SOURCE_DIR}/../../liboffload/include/generated
+ ${CMAKE_CURRENT_SOURCE_DIR}/../../include
+ ${CMAKE_CURRENT_SOURCE_DIR}/../../plugins-nextgen/common/include)
+
+target_compile_options(LLVMOffloadKernel PRIVATE ${offload_compile_flags})
+target_link_options(LLVMOffloadKernel PRIVATE ${offload_link_flags})
+
+target_compile_definitions(LLVMOffloadKernel PRIVATE
+ TARGET_NAME="libLLVMOffloadKernel"
+ DEBUG_PREFIX="LLVMOffloadKernel"
+)
+
+set_target_properties(LLVMOffloadKernel PROPERTIES
+ POSITION_INDEPENDENT_CODE ON
+ INSTALL_RPATH "$ORIGIN"
+ BUILD_RPATH "$ORIGIN:${CMAKE_CURRENT_BINARY_DIR}/..")
+install(TARGETS LLVMOffloadKernel LIBRARY COMPONENT offload DESTINATION "${OFFLOAD_INSTALL_LIBDIR}")
diff --git a/offload/languages/kernel/exports b/offload/languages/kernel/exports
new file mode 100644
index 0000000000000..4bc5a4ab9710c
--- /dev/null
+++ b/offload/languages/kernel/exports
@@ -0,0 +1,8 @@
+VERS1.0 {
+ global:
+
+ olK*;
+
+ local:
+ *;
+};
diff --git a/offload/languages/kernel/include/ExportedAPI.h b/offload/languages/kernel/include/ExportedAPI.h
new file mode 100644
index 0000000000000..a4ce168d78fda
--- /dev/null
+++ b/offload/languages/kernel/include/ExportedAPI.h
@@ -0,0 +1,41 @@
+/*===---- ExportedAPI.h - Kernel language runtime - exported api ----------===
+ *
+ * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+ * See https://llvm.org/LICENSE.txt for license information.
+ * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+ *
+ *===-----------------------------------------------------------------------===
+ */
+
+#pragma once
+
+#include "OffloadAPI.h"
+#include "Types.h"
+
+extern "C" {
+ol_device_handle_t olKGetDefaultDevice();
+
+ol_device_handle_t olKGetHostDevice();
+
+int olKGetDeviceCount();
+
+ol_device_handle_t olKGetDevice(int *DeviceNo);
+
+ol_device_handle_t olKSetDefaultDevice(int DeviceNo);
+
+ol_queue_handle_t olKGetDefaultQueue();
+
+CallConfigurationTy *olKGetCallConfiguration();
+
+void olKRegisterKernel(const void *ID, ol_symbol_handle_t Kernel);
+
+void olKUnregisterKernel(const void *ID);
+
+ol_symbol_handle_t olKGetKernel(const void *ID);
+
+void olKRegisterProgram(const void *ID, ol_program_handle_t Program);
+
+ol_program_handle_t olKUnregisterProgram(const void *ID);
+
+ol_program_handle_t olKGetProgram(const void *ID);
+}
diff --git a/offload/languages/kernel/include/State.h b/offload/languages/kernel/include/State.h
new file mode 100644
index 0000000000000..90f323d160a04
--- /dev/null
+++ b/offload/languages/kernel/include/State.h
@@ -0,0 +1,106 @@
+//===------- State.h - Kernel Language (CUDA/HIP) persistent state --------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include "OffloadAPI.h"
+#include "Types.h"
+
+#include "llvm/ADT/ArrayRef.h"
+#include "llvm/ADT/DenseMap.h"
+#include "llvm/ADT/SmallVector.h"
+
+namespace llvm {
+namespace offload {
+
+using KernelIDTy = const void *;
+
+struct ThreadStateTy {
+ ~ThreadStateTy();
+
+ static ThreadStateTy &get();
+
+ static ol_queue_handle_t getDefaultQueue();
+ static ol_device_handle_t getDefaultDevice();
+ static CallConfigurationTy &getCallConfiguration();
+ void setDefaultDevice(ol_device_handle_t Device);
+
+private:
+ void createDefaultQueue(ol_device_handle_t Device);
+
+ ol_device_handle_t DefaultDevice = nullptr;
+ ol_queue_handle_t DefaultQueue = nullptr;
+
+ CallConfigurationTy CC = {};
+
+ ThreadStateTy();
+};
+
+struct StateTy {
+ ~StateTy();
+
+ friend struct ThreadStateTy;
+
+ static StateTy &get();
+ static StateTy *tryGet();
+
+ static ol_device_handle_t getHostDevice() { return get().HostDevice; }
+
+ ArrayRef<ol_device_handle_t> getDevices() const { return Devices; }
+
+ void addDevice(ol_device_handle_t Device) { Devices.push_back(Device); }
+ void setHostDevice(ol_device_handle_t Device) {
+ if (!HostDevice)
+ HostDevice = Device;
+ }
+
+ void addKernel(KernelIDTy KernelID, ol_symbol_handle_t Kernel) {
+ KernelMap[KernelID] = Kernel;
+ }
+
+ void removeKernel(KernelIDTy KernelID) { KernelMap.erase(KernelID); }
+
+ ol_symbol_handle_t getKernel(KernelIDTy KernelID) {
+ return KernelMap[KernelID];
+ }
+
+ void addProgram(const void *Binary, ol_program_handle_t Program) {
+ BinaryRegisterMap[Binary] = Program;
+ }
+
+ ol_program_handle_t removeProgram(const void *Binary) {
+ auto It = BinaryRegisterMap.find(Binary);
+ if (It == BinaryRegisterMap.end())
+ return nullptr;
+ ol_program_handle_t Program = It->second;
+ BinaryRegisterMap.erase(It);
+ return Program;
+ }
+
+ ol_program_handle_t getProgram(const void *Binary) {
+ assert(BinaryRegisterMap.count(Binary));
+ return BinaryRegisterMap[Binary];
+ }
+
+ void destroyRegisteredPrograms();
+
+private:
+ DenseMap<const void *, ol_program_handle_t> BinaryRegisterMap;
+ DenseMap<KernelIDTy, ol_symbol_handle_t> KernelMap;
+ SmallVector<ol_device_handle_t, 8> Devices;
+
+ ol_queue_handle_t DefaultQueue = nullptr;
+ ol_device_handle_t HostDevice = nullptr;
+
+ StateTy();
+};
+
+} // namespace offload
+} // namespace llvm
diff --git a/offload/languages/kernel/include/Types.h b/offload/languages/kernel/include/Types.h
new file mode 100644
index 0000000000000..6c5a0ec19a8f7
--- /dev/null
+++ b/offload/languages/kernel/include/Types.h
@@ -0,0 +1,28 @@
+//===------- Types.h - Kernel Language (CUDA/HIP) api types ---------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include "Types.h"
+#include <cstddef>
+#include <cstdint>
+
+struct uint3 {
+ unsigned x = 0, y = 0, z = 0;
+};
+
+using dim3 = uint3;
+
+struct CallConfigurationTy {
+ dim3 GridSize;
+ dim3 BlockSize;
+ size_t SharedMemory;
+ void *Stream;
+};
diff --git a/offload/languages/kernel/src/ExportedAPI.cpp b/offload/languages/kernel/src/ExportedAPI.cpp
new file mode 100644
index 0000000000000..5c65d3d66b89b
--- /dev/null
+++ b/offload/languages/kernel/src/ExportedAPI.cpp
@@ -0,0 +1,99 @@
+//===------ ExportedAPI.cpp - Kernel Language runtime - exported api ------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+//===----------------------------------------------------------------------===//
+
+#include "ExportedAPI.h"
+
+#include "State.h"
+#include "Types.h"
+
+#include "OffloadAPI.h"
+#include "llvm/ADT/ArrayRef.h"
+
+#include <cstdio>
+#include <stdint.h>
+
+using namespace llvm;
+using namespace offload;
+
+/// Runtime API
+///{
+ol_device_handle_t olKGetDefaultDevice() {
+ ol_device_handle_t DefaultDevice = ThreadStateTy::getDefaultDevice();
+ return DefaultDevice;
+}
+
+ol_device_handle_t olKGetHostDevice() {
+ ol_device_handle_t HostDevice = StateTy::getHostDevice();
+ return HostDevice;
+}
+
+int olKGetDeviceCount() {
+ int DeviceCount = StateTy::get().getDevices().size();
+ return DeviceCount;
+}
+
+ol_device_handle_t olKGetDevice(int *DeviceNo) {
+ ol_device_handle_t DefaultDevice = ThreadStateTy::getDefaultDevice();
+ int DeviceCount = StateTy::get().getDevices().size();
+ ArrayRef<ol_device_handle_t> Devices = StateTy::get().getDevices();
+ for (int i = 0; i < DeviceCount; i++) {
+ if (Devices[i] == DefaultDevice) {
+ *DeviceNo = i;
+ return Devices[i];
+ }
+ }
+ return nullptr;
+}
+
+ol_device_handle_t olKSetDefaultDevice(int DeviceNo) {
+ ArrayRef<ol_device_handle_t> Devices = StateTy::get().getDevices();
+ if (DeviceNo < 0 || DeviceNo >= static_cast<int>(Devices.size()))
+ return nullptr;
+ ol_device_handle_t Device = Devices[DeviceNo];
+ ThreadStateTy::get().setDefaultDevice(Device);
+ return Device;
+}
+
+ol_queue_handle_t olKGetDefaultQueue() {
+ ol_queue_handle_t DefaultQueue = ThreadStateTy::getDefaultQueue();
+ return DefaultQueue;
+}
+
+CallConfigurationTy *olKGetCallConfiguration() {
+ return &ThreadStateTy::getCallConfiguration();
+}
+
+void olKRegisterKernel(const void *ID, ol_symbol_handle_t Kernel) {
+ StateTy::get().addKernel(ID, Kernel);
+}
+
+void olKUnregisterKernel(const void *ID) {
+ if (StateTy *State = StateTy::tryGet())
+ State->removeKernel(ID);
+}
+
+ol_symbol_handle_t olKGetKernel(const void *ID) {
+ return StateTy::get().getKernel(ID);
+}
+
+void olKRegisterProgram(const void *ID, ol_program_handle_t Program) {
+ StateTy::get().addProgram(ID, Program);
+}
+
+ol_program_handle_t olKUnregisterProgram(const void *ID) {
+ if (StateTy *State = StateTy::tryGet())
+ return State->removeProgram(ID);
+ return nullptr;
+}
+
+ol_program_handle_t olKGetProgram(const void *ID) {
+ return StateTy::get().getProgram(ID);
+}
+///}
diff --git a/offload/languages/kernel/src/State.cpp b/offload/languages/kernel/src/State.cpp
new file mode 100644
index 0000000000000..eff3c31257e93
--- /dev/null
+++ b/offload/languages/kernel/src/State.cpp
@@ -0,0 +1,198 @@
+//===------ State.cpp - Kernel Language (CUDA/HIP) persistent state -------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+//===----------------------------------------------------------------------===//
+
+#include "State.h"
+#include "Types.h"
+
+#include "OffloadAPI.h"
+#include "llvm/ADT/SmallPtrSet.h"
+#include "llvm/Support/raw_ostream.h"
+
+#include <atomic>
+#include <cstdio>
+#include <mutex>
+
+#define CHECK_FATAL(Result, ...) \
+ if (Result && Result->Code) { \
+ fprintf(stderr, __VA_ARGS__); \
+ abort(); \
+ }
+
+using namespace llvm;
+using namespace offload;
+
+std::atomic<uint32_t> AnyNonDefaultDevice = 0;
+__attribute__((weak)) uint32_t PerThreadQueue = 0;
+
+thread_local CallConfigurationTy CC = {};
+
+static std::mutex StateLock;
+static std::atomic<StateTy *> StatePtr = nullptr;
+
+static thread_local ThreadStateTy *ThreadState = nullptr;
+
+static std::mutex ThreadStatesLock;
+using ThreadStatesTy = SmallVector<ThreadStateTy *, 64>;
+static ThreadStatesTy *ThreadStatesPtr = nullptr;
+
+static void deleteThreadState() {
+ std::lock_guard<std::mutex> LG(ThreadStatesLock);
+ ThreadStatesTy *ThreadStates = ThreadStatesPtr;
+ ThreadStatesPtr = nullptr;
+ if (!ThreadStates)
+ return;
+
+ for (auto *TS : *ThreadStates)
+ delete (TS);
+ delete ThreadStates;
+ ThreadState = nullptr;
+}
+
+static void deleteState() {
+ StateTy *ST = StatePtr.load();
+ StatePtr.store(nullptr);
+ delete (ST);
+}
+
+static void destroyQueue(ol_queue_handle_t &Queue) {
+ if (!Queue)
+ return;
+
+ olSyncQueue(Queue);
+ olDestroyQueue(Queue);
+ Queue = nullptr;
+}
+
+ThreadStateTy::ThreadStateTy() {
+ if (PerThreadQueue) [[unlikely]]
+ createDefaultQueue(getDefaultDevice());
+ atexit(deleteThreadState);
+}
+ThreadStateTy::~ThreadStateTy() { destroyQueue(DefaultQueue); }
+
+ThreadStateTy &ThreadStateTy::get() {
+ auto *TS = ThreadState;
+ if (!TS) {
+ TS = new ThreadStateTy();
+ ThreadState = TS;
+ std::lock_guard<std::mutex> LG(ThreadStatesLock);
+ if (!ThreadStatesPtr)
+ ThreadStatesPtr = new ThreadStatesTy;
+ ThreadStatesPtr->push_back(TS);
+ }
+ return *TS;
+}
+
+ol_device_handle_t ThreadStateTy::getDefaultDevice() {
+ ol_device_handle_t DD = ThreadStateTy::get().DefaultDevice;
+ if (DD)
+ return DD;
+ for (ol_device_handle_t Device : StateTy::get().getDevices()) {
+ DD = Device;
+ break;
+ }
+ if (AnyNonDefaultDevice.load(std::memory_order_relaxed)) [[unlikely]] {
+ ol_device_handle_t TDD = ThreadStateTy::get().DefaultDevice;
+ if (TDD)
+ DD = TDD;
+ }
+ return DD;
+}
+
+ol_queue_handle_t ThreadStateTy::getDefaultQueue() {
+ if (!PerThreadQueue) [[likely]]
+ return StateTy::get().DefaultQueue;
+ return ThreadStateTy::get().DefaultQueue;
+}
+
+CallConfigurationTy &ThreadStateTy::getCallConfiguration() {
+ return ThreadStateTy::get().CC;
+}
+
+void ThreadStateTy::setDefaultDevice(ol_device_handle_t Device) {
+ DefaultDevice = Device;
+ createDefaultQueue(Device);
+}
+
+void ThreadStateTy::createDefaultQueue(ol_device_handle_t Device) {
+ if (DefaultQueue)
+ olDestroyQueue(DefaultQueue);
+ CHECK_FATAL(olCreateQueue(Device, &DefaultQueue),
+ "Failed to create per-thread default queue");
+}
+
+StateTy &StateTy::get() {
+ StateTy *ST = StatePtr.load();
+ if (!ST) [[unlikely]] {
+ std::lock_guard<std::mutex> LG(StateLock);
+ ST = StatePtr.load();
+ if (!ST) {
+ ST = new StateTy();
+ StatePtr.store(ST);
+ }
+ }
+ return *ST;
+}
+
+StateTy *StateTy::tryGet() { return StatePtr.load(); }
+
+static bool addDevices(ol_device_handle_t Device, void *Payload) {
+ StateTy &State = *reinterpret_cast<StateTy *>(Payload);
+ ol_platform_handle_t Platform;
+ ol_result_t Result;
+
+ Result = olGetDeviceInfo(Device, OL_DEVICE_INFO_PLATFORM, sizeof(Platform),
+ &Platform);
+ if (Result && Result->Code)
+ return true;
+
+ ol_platform_backend_t Backend;
+ Result = olGetPlatformInfo(Platform, OL_PLATFORM_INFO_BACKEND,
+ sizeof(Backend), &Backend);
+ if (Result && Result->Code)
+ return true;
+
+ if (Backend == OL_PLATFORM_BACKEND_HOST)
+ State.setHostDevice(Device);
+ else
+ State.addDevice(Device);
+ return true;
+}
+
+StateTy::StateTy() {
+ CHECK_FATAL(olInit(nullptr), "Failed to initialize the LLVMOffload");
+ CHECK_FATAL(olIterateDevices(addDevices, this), "Failed to identify devices");
+
+ if (!PerThreadQueue) [[likely]]
+ if (!Devices.empty()) [[likely]]
+ CHECK_FATAL(olCreateQueue(Devices.front(), &DefaultQueue),
+ "Failed to create default queue");
+
+ atexit(deleteState);
+}
+
+StateTy::~StateTy() {
+ deleteThreadState();
+ destroyQueue(DefaultQueue);
+ destroyRegisteredPrograms();
+ olShutDown();
+}
+
+void StateTy::destroyRegisteredPrograms() {
+ SmallPtrSet<ol_program_handle_t, 8> Programs;
+ for (auto &It : BinaryRegisterMap)
+ Programs.insert(It.second);
+
+ KernelMap.clear();
+ BinaryRegisterMap.clear();
+
+ for (ol_program_handle_t Program : Programs)
+ olDestroyProgram(Program);
+}
>From cca06271115b71bf9d727b32c778d68b51f64cc3 Mon Sep 17 00:00:00 2001
From: Sophia Herrmann <herrmann15 at llnl.gov>
Date: Fri, 24 Jul 2026 10:20:18 -0700
Subject: [PATCH 2/7] Build one offload language runtime library
Co-authored-by: Johannes Doerfert <jdoerfert.llvm at gmail.com>
Co-authored-by: Jonas Greifenhain <cadivus at daverkomp.de>
---
offload/languages/CMakeLists.txt | 2 +
offload/languages/README.md | 27 ++
offload/languages/cuda/CMakeLists.txt | 1 +
offload/languages/hip/CMakeLists.txt | 1 +
offload/languages/include/cuda/cuda_runtime.h | 24 ++
offload/languages/include/hip/hip_runtime.h | 61 ++++
.../include/kernel/DefineLanguageNames.inc | 51 +++
.../include/kernel/LanguageRuntime.h | 222 ++++++++++++
.../include/kernel/UndefineLanguageNames.inc | 48 +++
offload/languages/kernel/CMakeLists.txt | 59 +++-
offload/languages/kernel/exports | 9 +-
.../languages/kernel/include/ExportedAPI.h | 41 ---
.../kernel/include/LanguageAliases.inc | 70 ++++
.../languages/kernel/include/LanguageLaunch.h | 40 +++
.../kernel/include/LanguageRegistration.h | 64 ++++
offload/languages/kernel/include/State.h | 134 +++++---
offload/languages/kernel/include/Types.h | 10 +-
offload/languages/kernel/src/ExportedAPI.cpp | 99 ------
.../languages/kernel/src/LanguageCommon.cpp | 18 +
.../languages/kernel/src/LanguageLaunch.cpp | 83 +++++
.../kernel/src/LanguageRegistration.cpp | 318 ++++++++++++++++++
.../languages/kernel/src/LanguageRuntime.cpp | 223 ++++++++++++
offload/languages/kernel/src/State.cpp | 138 +++++++-
23 files changed, 1529 insertions(+), 214 deletions(-)
create mode 100644 offload/languages/README.md
create mode 100644 offload/languages/cuda/CMakeLists.txt
create mode 100644 offload/languages/hip/CMakeLists.txt
create mode 100644 offload/languages/include/cuda/cuda_runtime.h
create mode 100644 offload/languages/include/hip/hip_runtime.h
create mode 100644 offload/languages/include/kernel/DefineLanguageNames.inc
create mode 100644 offload/languages/include/kernel/LanguageRuntime.h
create mode 100644 offload/languages/include/kernel/UndefineLanguageNames.inc
delete mode 100644 offload/languages/kernel/include/ExportedAPI.h
create mode 100644 offload/languages/kernel/include/LanguageAliases.inc
create mode 100644 offload/languages/kernel/include/LanguageLaunch.h
create mode 100644 offload/languages/kernel/include/LanguageRegistration.h
delete mode 100644 offload/languages/kernel/src/ExportedAPI.cpp
create mode 100644 offload/languages/kernel/src/LanguageCommon.cpp
create mode 100644 offload/languages/kernel/src/LanguageLaunch.cpp
create mode 100644 offload/languages/kernel/src/LanguageRegistration.cpp
create mode 100644 offload/languages/kernel/src/LanguageRuntime.cpp
diff --git a/offload/languages/CMakeLists.txt b/offload/languages/CMakeLists.txt
index 08b2a68082d80..fbd774e645e21 100644
--- a/offload/languages/CMakeLists.txt
+++ b/offload/languages/CMakeLists.txt
@@ -1 +1,3 @@
add_subdirectory(kernel)
+add_subdirectory(cuda)
+add_subdirectory(hip)
diff --git a/offload/languages/README.md b/offload/languages/README.md
new file mode 100644
index 0000000000000..8361648df1779
--- /dev/null
+++ b/offload/languages/README.md
@@ -0,0 +1,27 @@
+# Offload Language Runtimes
+
+This directory contains the CUDA and HIP runtime layer for LLVM offload. It
+builds one shared library, `LLVMOffloadKernel`.
+
+The installed language headers are the user-facing part of this directory. CUDA
+programs include the CUDA header, HIP programs include the HIP header, and both
+see normal language names such as `cudaMalloc` or `hipMalloc`.
+
+Clang also depends on a small set of runtime entry points for kernel launch and
+device image registration. These symbols are external because generated host
+code calls them, but they are for compiler-generated code rather than for users
+to call directly.
+
+The rest of the runtime is internal. This includes device lookup, queues,
+registered programs, registered kernels, and error conversion.
+
+Some source files are shared by CUDA and HIP. CMake compiles those files once
+with `LANGUAGE=cuda` and once with `LANGUAGE=hip`. The language-name includes
+rename the generic declarations and definitions to the CUDA or HIP spelling for
+that object file.
+
+Source files that do not depend on CUDA or HIP spelling are compiled once and
+shared by both languages.
+
+The `cuda/` and `hip/` directories are kept in case there is ever a need for
+language-specific behavior.
diff --git a/offload/languages/cuda/CMakeLists.txt b/offload/languages/cuda/CMakeLists.txt
new file mode 100644
index 0000000000000..f09e9b46ef038
--- /dev/null
+++ b/offload/languages/cuda/CMakeLists.txt
@@ -0,0 +1 @@
+install(FILES ${CMAKE_CURRENT_SOURCE_DIR}/../include/cuda/cuda_runtime.h DESTINATION ${CMAKE_INSTALL_PREFIX}/include/offload/cuda/)
diff --git a/offload/languages/hip/CMakeLists.txt b/offload/languages/hip/CMakeLists.txt
new file mode 100644
index 0000000000000..ea41dd59adf5c
--- /dev/null
+++ b/offload/languages/hip/CMakeLists.txt
@@ -0,0 +1 @@
+install(FILES ${CMAKE_CURRENT_SOURCE_DIR}/../include/hip/hip_runtime.h DESTINATION ${CMAKE_INSTALL_PREFIX}/include/offload/hip/)
diff --git a/offload/languages/include/cuda/cuda_runtime.h b/offload/languages/include/cuda/cuda_runtime.h
new file mode 100644
index 0000000000000..a140ce1a1aca0
--- /dev/null
+++ b/offload/languages/include/cuda/cuda_runtime.h
@@ -0,0 +1,24 @@
+//===-- cuda_runtime.h - CUDA runtime API declarations --------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_OFFLOAD_LANGUAGES_INCLUDE_CUDA_CUDA_RUNTIME_H
+#define LLVM_OFFLOAD_LANGUAGES_INCLUDE_CUDA_CUDA_RUNTIME_H
+
+#define LANGUAGE cuda
+
+#include "../kernel/DefineLanguageNames.inc"
+
+#include "../kernel/LanguageRuntime.h"
+
+#include "../kernel/UndefineLanguageNames.inc"
+
+#undef LANGUAGE
+
+using cudaDeviceProp = cudaDeviceProp_t;
+
+#endif // LLVM_OFFLOAD_LANGUAGES_INCLUDE_CUDA_CUDA_RUNTIME_H
diff --git a/offload/languages/include/hip/hip_runtime.h b/offload/languages/include/hip/hip_runtime.h
new file mode 100644
index 0000000000000..b56e295c904e3
--- /dev/null
+++ b/offload/languages/include/hip/hip_runtime.h
@@ -0,0 +1,61 @@
+//===-- hip_runtime.h - HIP runtime API declarations ----------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_OFFLOAD_LANGUAGES_INCLUDE_HIP_HIP_RUNTIME_H
+#define LLVM_OFFLOAD_LANGUAGES_INCLUDE_HIP_HIP_RUNTIME_H
+
+#define LANGUAGE hip
+
+#include "../kernel/DefineLanguageNames.inc"
+
+#include "../kernel/LanguageRuntime.h"
+
+#include "../kernel/UndefineLanguageNames.inc"
+
+#undef LANGUAGE
+
+#define hipHostMallocDefault hipHostAllocDefault
+#define hipHostMallocPortable hipHostAllocPortable
+#define hipHostMallocMapped hipHostAllocMapped
+#define hipHostMallocWriteCombined hipHostAllocWriteCombined
+#define hipHostMallocNonCoherent 0x80000000
+
+inline hipError_t hipHostMalloc(void **Ptr, size_t Size, unsigned int Flags) {
+ return hipHostAlloc(Ptr, Size, Flags);
+}
+
+inline hipError_t hipHostFree(void *Ptr) { return ::hipFreeHost(Ptr); }
+
+#ifdef __cplusplus
+template <class T>
+static inline hipError_t hipHostMalloc(T **Ptr, size_t Size,
+ unsigned int Flags) {
+ return ::hipHostMalloc((void **)Ptr, Size, Flags);
+}
+
+template <class T> static inline hipError_t hipHostFree(T *Ptr) {
+ return ::hipHostFree((void *)Ptr);
+}
+#endif
+
+#if defined(__AMDGPU__) || defined(__NVPTX__)
+#define HIP_KERNEL_NAME(...) __VA_ARGS__
+
+extern "C" hipError_t hipLaunchKernel(const char *Kernel, dim3 GridDim,
+ dim3 BlockDim, void **KernelArgs,
+ size_t DynamicSharedMem, void *Stream);
+
+template <typename... AT, typename FT = void (*)(AT...)>
+static inline void hipLaunchKernelGGL(FT Kernel, dim3 GridDim, dim3 BlockDim,
+ size_t DynamicSharedMem, void *Stream,
+ AT... KernelArgs) {
+ Kernel<<<GridDim, BlockDim, DynamicSharedMem, Stream>>>(KernelArgs...);
+}
+#endif
+
+#endif // LLVM_OFFLOAD_LANGUAGES_INCLUDE_HIP_HIP_RUNTIME_H
diff --git a/offload/languages/include/kernel/DefineLanguageNames.inc b/offload/languages/include/kernel/DefineLanguageNames.inc
new file mode 100644
index 0000000000000..7f2b0776a03ff
--- /dev/null
+++ b/offload/languages/include/kernel/DefineLanguageNames.inc
@@ -0,0 +1,51 @@
+//===-- DefineLanguageNames.inc - Kernel language API name definitions ----===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#define COMBINE2(X, Y) X##Y
+#define COMBINE(X, Y) COMBINE2(X, Y)
+
+#define Error_t COMBINE(LANGUAGE, Error_t)
+#define DeviceProp_t COMBINE(LANGUAGE, DeviceProp_t)
+#define Malloc COMBINE(LANGUAGE, Malloc)
+#define Free COMBINE(LANGUAGE, Free)
+#define Memcpy COMBINE(LANGUAGE, Memcpy)
+#define DeviceSynchronize COMBINE(LANGUAGE, DeviceSynchronize)
+#define Success COMBINE(LANGUAGE, Success)
+#define ErrorInvalidValue COMBINE(LANGUAGE, ErrorInvalidValue)
+#define MemcpyKind COMBINE(LANGUAGE, MemcpyKind)
+#define MemcpyHostToHost COMBINE(LANGUAGE, MemcpyHostToHost)
+#define MemcpyHostToDevice COMBINE(LANGUAGE, MemcpyHostToDevice)
+#define MemcpyDeviceToHost COMBINE(LANGUAGE, MemcpyDeviceToHost)
+#define MemcpyDeviceToDevice COMBINE(LANGUAGE, MemcpyDeviceToDevice)
+#define MemcpyDefault COMBINE(LANGUAGE, MemcpyDefault)
+#define GetLastError COMBINE(LANGUAGE, GetLastError)
+#define PeekAtLastError COMBINE(LANGUAGE, PeekAtLastError)
+#define GetErrorName COMBINE(LANGUAGE, GetErrorName)
+#define GetErrorString COMBINE(LANGUAGE, GetErrorString)
+#define GetDevice COMBINE(LANGUAGE, GetDevice)
+#define GetDeviceCount COMBINE(LANGUAGE, GetDeviceCount)
+#define SetDevice COMBINE(LANGUAGE, SetDevice)
+#define HostAlloc COMBINE(LANGUAGE, HostAlloc)
+#define HostAllocDefault COMBINE(LANGUAGE, HostAllocDefault)
+#define HostAllocPortable COMBINE(LANGUAGE, HostAllocPortable)
+#define HostAllocMapped COMBINE(LANGUAGE, HostAllocMapped)
+#define HostAllocWriteCombined COMBINE(LANGUAGE, HostAllocWriteCombined)
+#define MallocHost COMBINE(LANGUAGE, MallocHost)
+#define FreeHost COMBINE(LANGUAGE, FreeHost)
+#define DriverGetVersion COMBINE(LANGUAGE, DriverGetVersion)
+#define GetDeviceProperties COMBINE(LANGUAGE, GetDeviceProperties)
+#define OccupancyMaxPotentialBlockSizeVariableSMem \
+ COMBINE(LANGUAGE, OccupancyMaxPotentialBlockSizeVariableSMem)
+#define Stream_t COMBINE(LANGUAGE, Stream_t)
+#define StreamCreate COMBINE(LANGUAGE, StreamCreate)
+#define StreamCreateWithFlags COMBINE(LANGUAGE, StreamCreateWithFlags)
+#define StreamDestroy COMBINE(LANGUAGE, StreamDestroy)
+#define StreamSynchronize COMBINE(LANGUAGE, StreamSynchronize)
+#define StreamCreateWithFlagsFlags COMBINE(LANGUAGE, StreamCreateWithFlagsFlags)
+#define StreamDefault COMBINE(LANGUAGE, StreamDefault)
+#define StreamNonBlocking COMBINE(LANGUAGE, StreamNonBlocking)
diff --git a/offload/languages/include/kernel/LanguageRuntime.h b/offload/languages/include/kernel/LanguageRuntime.h
new file mode 100644
index 0000000000000..dc00ae4ad5a1c
--- /dev/null
+++ b/offload/languages/include/kernel/LanguageRuntime.h
@@ -0,0 +1,222 @@
+//===-- LanguageRuntime.h - Kernel language runtime API declarations ------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_OFFLOAD_LANGUAGES_INCLUDE_KERNEL_LANGUAGE_RUNTIME_H
+#define LLVM_OFFLOAD_LANGUAGES_INCLUDE_KERNEL_LANGUAGE_RUNTIME_H
+
+#include <cstdio>
+#include <cstdlib>
+#include <stddef.h>
+#include <stdint.h>
+
+enum Error_t : uint32_t {
+ Success = 0,
+ ErrorInvalidValue = 1,
+};
+
+struct DeviceProp_t {
+ char name[256];
+ size_t totalGlobalMem;
+ int warpSize;
+ int multiProcessorCount;
+ int major;
+ int minor;
+ int ECCEnabled;
+ int pciBusID;
+ int pciDeviceID;
+ int pciDomainID;
+ int memoryBusWidth;
+};
+
+enum MemcpyKind {
+ MemcpyHostToHost = 0,
+ MemcpyHostToDevice = 1,
+ MemcpyDeviceToHost = 2,
+ MemcpyDeviceToDevice = 3,
+ MemcpyDefault = 4
+};
+
+enum HostAllocFlags : unsigned int {
+ HostAllocDefault = 0x00,
+ HostAllocPortable = 0x01,
+ HostAllocMapped = 0x02,
+ HostAllocWriteCombined = 0x04,
+};
+
+enum StreamCreateWithFlagsFlags : unsigned int {
+ StreamDefault = 0x00,
+ StreamNonBlocking = 0x01,
+};
+
+typedef struct Stream_st *Stream_t;
+
+/// Malloc, with type template overlay.
+///{
+Error_t Malloc(void **Dev_Ptr, size_t Size);
+
+template <class T> static inline Error_t Malloc(T **dev_Ptr, size_t Size) {
+ return ::Malloc((void **)dev_Ptr, Size);
+}
+
+Error_t HostAlloc(void **Ptr, size_t Size, unsigned int Flags);
+
+template <class T>
+static inline Error_t HostAlloc(T **Ptr, size_t Size, unsigned int Flags) {
+ return ::HostAlloc((void **)Ptr, Size, Flags);
+}
+
+Error_t MallocHost(void **Ptr, size_t Size);
+
+template <class T> static inline Error_t MallocHost(T **Ptr, size_t Size) {
+ return ::MallocHost((void **)Ptr, Size);
+}
+///}
+
+/// Free, no type template necessary.
+Error_t Free(void *Dev_Ptr);
+
+/// Memcpy, with type template overlay.
+///{
+Error_t Memcpy(void *Dst, const void *Src, size_t Size, MemcpyKind Kind);
+
+template <class T>
+static inline Error_t Memcpy(T *Dst, const T *Src, size_t Size,
+ MemcpyKind Kind) {
+ return ::Memcpy((void *)Dst, (const void *)Src, Size, Kind);
+}
+///}
+
+Error_t DeviceSynchronize();
+
+Error_t GetLastError();
+
+Error_t PeekAtLastError();
+
+const char *GetErrorName(Error_t Error);
+
+const char *GetErrorString(Error_t Error);
+
+Error_t GetDevice(int *DeviceNo);
+
+Error_t GetDeviceCount(int *Count);
+
+Error_t SetDevice(int DeviceNo);
+
+Error_t FreeHost(void *Ptr);
+
+Error_t DriverGetVersion(int *Version);
+
+Error_t GetDeviceProperties(DeviceProp_t *DeviceProp, int DeviceNo);
+
+Error_t StreamCreate(Stream_t *stream);
+
+Error_t StreamCreateWithFlags(Stream_t *stream, unsigned int flags);
+
+Error_t StreamDestroy(Stream_t stream);
+
+Error_t StreamSynchronize(Stream_t stream);
+
+template <typename UnaryFunction, class T>
+static inline Error_t OccupancyMaxPotentialBlockSizeVariableSMem(
+ int *minGridSize, int *blockSize, T func,
+ UnaryFunction blockSizeToDynamicSMemSize, int blockSizeLimit = 0) {
+#if defined(__AMDGPU__)
+ // TODO: values taken from AMD Instinct MI250X gfx90a
+ *minGridSize = 220;
+ *blockSize = 1024;
+#elif defined(__NVPTX__)
+ // TODO: values taken from NVIDIA H100 80GB HBM3
+ *minGridSize = 264;
+ *blockSize = 1024;
+#endif
+ return Success;
+}
+
+///
+
+#if defined(__AMDGPU__) || defined(__NVPTX__)
+#include <gpuintrin.h>
+
+/// Define \p FIELD as a property backed by component \p OFFSET of Vec.
+#define __LLVM_OFFLOAD_DEVICE_BUILTIN(FIELD, OFFSET) \
+ __declspec(property(get = __get_##FIELD, \
+ put = __put_##FIELD)) unsigned int FIELD; \
+ __device__ inline __attribute__((always_inline)) T __get_##FIELD(void) \
+ const { \
+ return Vec[OFFSET]; \
+ } \
+ __device__ inline __attribute__((always_inline)) T __put_##FIELD(T V) { \
+ return Vec[OFFSET] = V; \
+ }
+
+/// Common storage for CUDA/HIP vector aliases such as int4 and float3.
+///
+/// Provides array-style indexing and x/y/z/w component properties over Clang
+/// ext_vector_type storage.
+template <class T, int Size> struct BaseVector {
+ using VT = float __attribute__((ext_vector_type(Size)));
+ VT Vec;
+
+ __device__ __host__ BaseVector() = default;
+
+ /// Construct a vector from component values.
+ template <typename... Args>
+ __device__ __host__ BaseVector(Args... args) : BaseVector({args...}) {}
+
+ /// Return component \p Idx.
+ __device__ __host__ T &operator[](int Idx) { return Vec[Idx]; }
+ __device__ __host__ const T &operator[](int Idx) const { return Vec[Idx]; }
+
+ __LLVM_OFFLOAD_DEVICE_BUILTIN(x, 0);
+ __LLVM_OFFLOAD_DEVICE_BUILTIN(y, 1);
+ __LLVM_OFFLOAD_DEVICE_BUILTIN(z, 2);
+ __LLVM_OFFLOAD_DEVICE_BUILTIN(w, 3);
+};
+
+/// Define the vector alias TY##SIZE and its make_TY##SIZE constructor helper.
+#define __VECTOR_DEF_IMPL(TY, SIZE) \
+ using TY##SIZE = BaseVector<TY, SIZE>; \
+ \
+ template <typename... Args> \
+ __device__ __host__ TY##SIZE make_##TY##SIZE(Args... args) { \
+ return TY##SIZE(args...); \
+ }
+
+/// Define the standard 1/2/3/4/8/16 element vector aliases for \p TY.
+#define __VECTOR_DEF(TY) \
+ __VECTOR_DEF_IMPL(TY, 1) \
+ __VECTOR_DEF_IMPL(TY, 2) \
+ __VECTOR_DEF_IMPL(TY, 3) \
+ __VECTOR_DEF_IMPL(TY, 4) \
+ __VECTOR_DEF_IMPL(TY, 8) \
+ __VECTOR_DEF_IMPL(TY, 16)
+
+/// Instantiate CUDA/HIP-style vector types and make_* helpers.
+__VECTOR_DEF(float)
+__VECTOR_DEF(double)
+__VECTOR_DEF(int8_t)
+__VECTOR_DEF(int16_t)
+__VECTOR_DEF(int32_t)
+__VECTOR_DEF(int64_t)
+__VECTOR_DEF(uint8_t)
+__VECTOR_DEF(uint16_t)
+__VECTOR_DEF(uint32_t)
+__VECTOR_DEF(uint64_t)
+__VECTOR_DEF(char)
+__VECTOR_DEF(short)
+__VECTOR_DEF(int)
+__VECTOR_DEF(unsigned)
+__VECTOR_DEF(long)
+
+#undef __VECTOR_DEF_IMPL
+#undef __VECTOR_DEF
+#undef __LLVM_OFFLOAD_DEVICE_BUILTIN
+
+#endif
+
+#endif // LLVM_OFFLOAD_LANGUAGES_INCLUDE_KERNEL_LANGUAGE_RUNTIME_H
diff --git a/offload/languages/include/kernel/UndefineLanguageNames.inc b/offload/languages/include/kernel/UndefineLanguageNames.inc
new file mode 100644
index 0000000000000..ab5f257cdfaea
--- /dev/null
+++ b/offload/languages/include/kernel/UndefineLanguageNames.inc
@@ -0,0 +1,48 @@
+//===-- UndefineLanguageNames.inc - Kernel language API name undefines ----===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#undef Error_t
+#undef DeviceProp_t
+#undef Malloc
+#undef Free
+#undef Memcpy
+#undef DeviceSynchronize
+#undef Success
+#undef ErrorInvalidValue
+#undef MemcpyKind
+#undef MemcpyHostToHost
+#undef MemcpyHostToDevice
+#undef MemcpyDeviceToHost
+#undef MemcpyDeviceToDevice
+#undef MemcpyDefault
+#undef GetLastError
+#undef PeekAtLastError
+#undef GetErrorName
+#undef GetErrorString
+#undef GetDevice
+#undef GetDeviceCount
+#undef SetDevice
+#undef HostAlloc
+#undef HostAllocFlags
+#undef HostAllocDefault
+#undef HostAllocPortable
+#undef HostAllocMapped
+#undef HostAllocWriteCombined
+#undef MallocHost
+#undef FreeHost
+#undef DriverGetVersion
+#undef GetDeviceProperties
+#undef OccupancyMaxPotentialBlockSizeVariableSMem
+#undef Stream_t
+#undef StreamCreate
+#undef StreamCreateWithFlags
+#undef StreamDestroy
+#undef StreamSynchronize
+#undef StreamCreateWithFlagsFlags
+#undef StreamDefault
+#undef StreamNonBlocking
diff --git a/offload/languages/kernel/CMakeLists.txt b/offload/languages/kernel/CMakeLists.txt
index 3b0823a44faf1..81442f9f2c507 100644
--- a/offload/languages/kernel/CMakeLists.txt
+++ b/offload/languages/kernel/CMakeLists.txt
@@ -1,24 +1,56 @@
+set(LLVM_OFFLOAD_KERNEL_INCLUDE_DIRS
+ ${CMAKE_CURRENT_BINARY_DIR}/../include
+ ${CMAKE_CURRENT_SOURCE_DIR}/../include/cuda
+ ${CMAKE_CURRENT_SOURCE_DIR}/../include/hip
+ ${CMAKE_CURRENT_SOURCE_DIR}/../include/kernel
+ ${CMAKE_CURRENT_SOURCE_DIR}/include
+ ${CMAKE_CURRENT_SOURCE_DIR}/../kernel/include
+ ${CMAKE_CURRENT_SOURCE_DIR}/../../liboffload/include
+ ${CMAKE_CURRENT_BINARY_DIR}/../../liboffload/API
+ ${CMAKE_CURRENT_SOURCE_DIR}/../../include
+ ${CMAKE_CURRENT_SOURCE_DIR}/../../plugins-nextgen/common/include)
+
+function(add_llvm_offload_kernel_language_runtime_objects target language)
+ add_library(${target} OBJECT
+ src/LanguageRuntime.cpp
+ )
+ add_dependencies(${target} OffloadAPI)
+ target_include_directories(${target} PRIVATE ${LLVM_OFFLOAD_KERNEL_INCLUDE_DIRS})
+ target_compile_options(${target} PRIVATE ${offload_compile_flags})
+ target_compile_definitions(${target} PRIVATE
+ LANGUAGE=${language}
+ TARGET_NAME="libLLVMOffloadKernel"
+ DEBUG_PREFIX="LLVMOffloadKernel"
+ )
+ set_target_properties(${target} PROPERTIES POSITION_INDEPENDENT_CODE ON)
+endfunction()
+
+add_llvm_offload_kernel_language_runtime_objects(
+ LLVMOffloadKernelCudaRuntimeObjects cuda)
+add_llvm_offload_kernel_language_runtime_objects(
+ LLVMOffloadKernelHipRuntimeObjects hip)
+
add_llvm_library(
LLVMOffloadKernel SHARED
+ $<TARGET_OBJECTS:LLVMOffloadKernelCudaRuntimeObjects>
+ $<TARGET_OBJECTS:LLVMOffloadKernelHipRuntimeObjects>
+ src/LanguageCommon.cpp
+ src/LanguageLaunch.cpp
+ src/LanguageRegistration.cpp
src/State.cpp
- src/ExportedAPI.cpp
LINK_COMPONENTS
Support
Offload
)
-if(LIBOMP_HAVE_VERSION_SCRIPT_FLAG)
+if(LLVM_HAVE_LINK_VERSION_SCRIPT)
target_link_libraries(LLVMOffloadKernel PRIVATE "-Wl,--version-script=${CMAKE_CURRENT_SOURCE_DIR}/exports")
endif()
target_include_directories(LLVMOffloadKernel PUBLIC
- ${CMAKE_CURRENT_SOURCE_DIR}/include
- ${CMAKE_CURRENT_SOURCE_DIR}/../../liboffload/include
- ${CMAKE_CURRENT_SOURCE_DIR}/../../liboffload/include/generated
- ${CMAKE_CURRENT_SOURCE_DIR}/../../include
- ${CMAKE_CURRENT_SOURCE_DIR}/../../plugins-nextgen/common/include)
+ ${LLVM_OFFLOAD_KERNEL_INCLUDE_DIRS})
target_compile_options(LLVMOffloadKernel PRIVATE ${offload_compile_flags})
target_link_options(LLVMOffloadKernel PRIVATE ${offload_link_flags})
@@ -29,7 +61,18 @@ target_compile_definitions(LLVMOffloadKernel PRIVATE
)
set_target_properties(LLVMOffloadKernel PROPERTIES
+ RUNTIME_OUTPUT_DIRECTORY "${LLVM_LIBRARY_OUTPUT_INTDIR}/${OFFLOAD_TARGET_SUBDIR}"
POSITION_INDEPENDENT_CODE ON
INSTALL_RPATH "$ORIGIN"
BUILD_RPATH "$ORIGIN:${CMAKE_CURRENT_BINARY_DIR}/..")
-install(TARGETS LLVMOffloadKernel LIBRARY COMPONENT offload DESTINATION "${OFFLOAD_INSTALL_LIBDIR}")
+install(TARGETS LLVMOffloadKernel
+ COMPONENT offload
+ RUNTIME DESTINATION "${CMAKE_INSTALL_BINDIR}"
+ LIBRARY DESTINATION "${OFFLOAD_INSTALL_LIBDIR}"
+ ARCHIVE DESTINATION "${OFFLOAD_INSTALL_LIBDIR}")
+
+install(FILES
+ ${CMAKE_CURRENT_SOURCE_DIR}/../include/kernel/DefineLanguageNames.inc
+ ${CMAKE_CURRENT_SOURCE_DIR}/../include/kernel/LanguageRuntime.h
+ ${CMAKE_CURRENT_SOURCE_DIR}/../include/kernel/UndefineLanguageNames.inc
+ DESTINATION ${CMAKE_INSTALL_PREFIX}/include/offload/kernel/)
diff --git a/offload/languages/kernel/exports b/offload/languages/kernel/exports
index 4bc5a4ab9710c..3372890ff42c6 100644
--- a/offload/languages/kernel/exports
+++ b/offload/languages/kernel/exports
@@ -1,8 +1,11 @@
VERS1.0 {
global:
-
- olK*;
-
+ *cuda*;
+ *hip*;
+ llvmLaunchKernel*;
+ __llvm*;
+ __tgt_register_lib;
+ __tgt_unregister_lib;
local:
*;
};
diff --git a/offload/languages/kernel/include/ExportedAPI.h b/offload/languages/kernel/include/ExportedAPI.h
deleted file mode 100644
index a4ce168d78fda..0000000000000
--- a/offload/languages/kernel/include/ExportedAPI.h
+++ /dev/null
@@ -1,41 +0,0 @@
-/*===---- ExportedAPI.h - Kernel language runtime - exported api ----------===
- *
- * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
- * See https://llvm.org/LICENSE.txt for license information.
- * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
- *
- *===-----------------------------------------------------------------------===
- */
-
-#pragma once
-
-#include "OffloadAPI.h"
-#include "Types.h"
-
-extern "C" {
-ol_device_handle_t olKGetDefaultDevice();
-
-ol_device_handle_t olKGetHostDevice();
-
-int olKGetDeviceCount();
-
-ol_device_handle_t olKGetDevice(int *DeviceNo);
-
-ol_device_handle_t olKSetDefaultDevice(int DeviceNo);
-
-ol_queue_handle_t olKGetDefaultQueue();
-
-CallConfigurationTy *olKGetCallConfiguration();
-
-void olKRegisterKernel(const void *ID, ol_symbol_handle_t Kernel);
-
-void olKUnregisterKernel(const void *ID);
-
-ol_symbol_handle_t olKGetKernel(const void *ID);
-
-void olKRegisterProgram(const void *ID, ol_program_handle_t Program);
-
-ol_program_handle_t olKUnregisterProgram(const void *ID);
-
-ol_program_handle_t olKGetProgram(const void *ID);
-}
diff --git a/offload/languages/kernel/include/LanguageAliases.inc b/offload/languages/kernel/include/LanguageAliases.inc
new file mode 100644
index 0000000000000..67c188bc16269
--- /dev/null
+++ b/offload/languages/kernel/include/LanguageAliases.inc
@@ -0,0 +1,70 @@
+//===-- LanguageAliases.inc - Language runtime symbol aliases -------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// The LLVM offload kernel runtime implements the shared __llvm* entry points
+// once. These wrappers expose the CUDA/HIP ABI names, such as
+// __cudaRegisterFunction and __hipPopCallConfiguration, as forwarding entry
+// points to the shared implementations.
+//
+//===----------------------------------------------------------------------===//
+
+// This file intentionally has no include guard. It is included once for each
+// LANGUAGE value to emit each language's alias definitions.
+
+#ifndef LANGUAGE
+#error This file should be included, or used, with a LANGUAGE macro set.
+#endif
+
+#define LA_IMPL2(PREFIX, L, NAME) PREFIX##L##NAME
+#define LA_IMPL1(PREFIX, L, NAME) LA_IMPL2(PREFIX, L, NAME)
+#define LANGUAGE_NAME(PREFIX, NAME) LA_IMPL1(PREFIX, LANGUAGE, NAME)
+
+extern "C" void LANGUAGE_NAME(__, RegisterFunction)(
+ const char *Binary, const char *KernelID, char *KernelName,
+ const char *KernelName1, int ThreadLimit, uint3 *Tid, uint3 *Bid,
+ dim3 *BlockDim, dim3 *GridDim, int *WSize) {
+ __llvmRegisterFunction(Binary, KernelID, KernelName, KernelName1, ThreadLimit,
+ Tid, Bid, BlockDim, GridDim, WSize);
+}
+
+extern "C" void LANGUAGE_NAME(__, RegisterVar)(void **Data, char *HostVar,
+ char *DeviceAddress,
+ const char *DeviceName, int Ext,
+ int Size, int Constant,
+ int Global) {
+ __llvmRegisterVar(Data, HostVar, DeviceAddress, DeviceName, Ext, Size,
+ Constant, Global);
+}
+
+extern "C" void LANGUAGE_NAME(__, RegisterManagedVar)(
+ void **Data, char *HostVar, char *DeviceAddress, const char *DeviceName,
+ size_t Size, unsigned Align) {
+ __llvmRegisterManagedVar(Data, HostVar, DeviceAddress, DeviceName, Size,
+ Align);
+}
+
+extern "C" void LANGUAGE_NAME(__, RegisterSurface)(
+ void **Data, const struct surfaceReference *SurfRef, const void **DevPtr,
+ const char *Name, int Dim, int Ext) {
+ __llvmRegisterSurface(Data, SurfRef, DevPtr, Name, Dim, Ext);
+}
+
+extern "C" void LANGUAGE_NAME(__, RegisterTexture)(
+ void **Data, const struct textureReference *TexRef, const void **DevPtr,
+ const char *Name, int Dim, int Norm, int Ext) {
+ __llvmRegisterTexture(Data, TexRef, DevPtr, Name, Dim, Norm, Ext);
+}
+
+extern "C" unsigned LANGUAGE_NAME(__, PopCallConfiguration)(
+ dim3 *GridSize, dim3 *BlockSize, size_t *SharedMemory, void **Stream) {
+ return __llvmPopCallConfiguration(GridSize, BlockSize, SharedMemory, Stream);
+}
+
+#undef LANGUAGE_NAME
+#undef LA_IMPL1
+#undef LA_IMPL2
diff --git a/offload/languages/kernel/include/LanguageLaunch.h b/offload/languages/kernel/include/LanguageLaunch.h
new file mode 100644
index 0000000000000..161aead85bbef
--- /dev/null
+++ b/offload/languages/kernel/include/LanguageLaunch.h
@@ -0,0 +1,40 @@
+//===-- LanguageLaunch.h - Language launch API declarations ---------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_LANGUAGE_LAUNCH_H
+#define LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_LANGUAGE_LAUNCH_H
+
+#include "OffloadAPI.h"
+#include "Types.h"
+
+#include <algorithm> // for std::max
+#include <cstddef>
+#include <cstdint>
+
+extern "C" {
+
+/// Push call configuration for kernel launch
+unsigned __llvmPushCallConfiguration(dim3 __grid_size, dim3 __block_size,
+ size_t __shared_memory, void *__stream);
+
+/// Pop call configuration for kernel launch
+unsigned __llvmPopCallConfiguration(dim3 *__grid_size, dim3 *__block_size,
+ size_t *__shared_memory, void **__stream);
+
+/// Internal kernel launch implementation
+ol_result_t __llvmLaunchKernelImpl(const char *KernelID, dim3 GridDim,
+ dim3 BlockDim, void *KernelArgsPtr,
+ size_t DynamicSharedMem, void *Stream);
+
+/// LLVM-style kernel launch entry point
+unsigned __llvmLaunchKernel(const char *KernelID, dim3 GridDim, dim3 BlockDim,
+ void *KernelArgsPtr, size_t DynamicSharedMem,
+ void *Stream);
+}
+
+#endif // LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_LANGUAGE_LAUNCH_H
diff --git a/offload/languages/kernel/include/LanguageRegistration.h b/offload/languages/kernel/include/LanguageRegistration.h
new file mode 100644
index 0000000000000..248aa74f600bd
--- /dev/null
+++ b/offload/languages/kernel/include/LanguageRegistration.h
@@ -0,0 +1,64 @@
+//===-- LanguageRegistration.h - Language registration API declarations ---===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_LANGUAGE_REGISTRATION_H
+#define LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_LANGUAGE_REGISTRATION_H
+
+#include "OffloadAPI.h"
+#include "Types.h"
+
+#include <cstdint>
+#include <iterator>
+
+#define HIP_FATBIN_MAGIC_STR "__CLANG_OFFLOAD_BUNDLE__"
+constexpr auto HIP_FATBIN_MAGIC_STR_LEN = sizeof(HIP_FATBIN_MAGIC_STR) - 1;
+
+namespace {
+struct FatbinWrapperTy {
+ int Magic;
+ int Version;
+ const char *Data;
+ const char *DataEnd;
+};
+} // namespace
+
+static void readTUFatbin(const char *Binary, const FatbinWrapperTy *FW);
+
+static void readHIPFatbinEntries(const char *Binary, const char *HIPFatbinPtr);
+
+/// Hidden, but exported, Registration API
+///{
+extern "C" {
+
+void __llvmRegisterFunction(const char *Binary, const char *KernelID,
+ char *KernelName, const char *KernelName1, int,
+ uint3 *, uint3 *, dim3 *, dim3 *, int *);
+
+const char *__llvmRegisterFatBinary(const char *Binary);
+
+void __llvmUnregisterFatBinary(void *Handle);
+
+void __llvmRegisterVar(void **, char *, char *, const char *, int, int, int,
+ int);
+
+void __llvmRegisterManagedVar(void **, char *, char *, const char *, size_t,
+ unsigned);
+
+void __llvmRegisterSurface(void **, const struct surfaceReference *,
+ const void **, const char *, int, int);
+
+void __llvmRegisterTexture(void **, const struct textureReference *,
+ const void **, const char *, int, int, int);
+
+struct __tgt_bin_desc;
+void __tgt_register_lib(__tgt_bin_desc *Desc);
+void __tgt_unregister_lib(__tgt_bin_desc *Desc);
+}
+///}
+
+#endif // LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_LANGUAGE_REGISTRATION_H
diff --git a/offload/languages/kernel/include/State.h b/offload/languages/kernel/include/State.h
index 90f323d160a04..12d1abafbf917 100644
--- a/offload/languages/kernel/include/State.h
+++ b/offload/languages/kernel/include/State.h
@@ -1,14 +1,13 @@
-//===------- State.h - Kernel Language (CUDA/HIP) persistent state --------===//
+//===-- State.h - Kernel language persistent state ------------------------===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
-//
-//===----------------------------------------------------------------------===//
-#pragma once
+#ifndef LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_STATE_H
+#define LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_STATE_H
#include "OffloadAPI.h"
#include "Types.h"
@@ -16,23 +15,52 @@
#include "llvm/ADT/ArrayRef.h"
#include "llvm/ADT/DenseMap.h"
#include "llvm/ADT/SmallVector.h"
+#include "llvm/Support/raw_ostream.h"
+
+#define CHECK_FATAL(ResultExpr, ...) \
+ do { \
+ ol_result_t CheckFatalResult = (ResultExpr); \
+ if (CheckFatalResult && CheckFatalResult->Code) { \
+ llvm::errs() << __VA_ARGS__; \
+ if (CheckFatalResult->Details) \
+ llvm::errs() << ": " << CheckFatalResult->Details; \
+ llvm::errs() << '\n'; \
+ abort(); \
+ } \
+ } while (false)
namespace llvm {
namespace offload {
+/// Opaque host-side key used to identify a registered kernel.
+///
+/// This is the address emitted in the offload entry table for the kernel,
+/// not a device number or liboffload handle. The runtime maps it to the
+/// loaded device symbol during registration and uses it again during launch.
using KernelIDTy = const void *;
+/// Per-thread state used by the language runtime entry points.
+///
+/// Tracks the current thread's default device, optional per-thread queue,
+/// and pending kernel launch configuration.
struct ThreadStateTy {
~ThreadStateTy();
- static ThreadStateTy &get();
-
+ /// Return the default queue for the current stream mode.
static ol_queue_handle_t getDefaultQueue();
+
+ /// Return the thread-local default device, or the first discovered device.
static ol_device_handle_t getDefaultDevice();
+
+ /// Return the pending kernel launch configuration for this thread.
static CallConfigurationTy &getCallConfiguration();
- void setDefaultDevice(ol_device_handle_t Device);
+
+ /// Set the thread-local default device to \p Device and recreate its queue.
+ static void setDefaultDevice(ol_device_handle_t Device);
private:
+ static ThreadStateTy &get();
+
void createDefaultQueue(ol_device_handle_t Device);
ol_device_handle_t DefaultDevice = nullptr;
@@ -43,58 +71,78 @@ struct ThreadStateTy {
ThreadStateTy();
};
+/// Process-wide state shared by CUDA and HIP language entry points.
+///
+/// Owns the discovered devices, host device, process default queue, and maps
+/// from registered binaries and kernels to liboffload handles.
struct StateTy {
~StateTy();
friend struct ThreadStateTy;
- static StateTy &get();
- static StateTy *tryGet();
+ /// Return the host device discovered during runtime initialization.
+ static ol_device_handle_t getHostDevice();
- static ol_device_handle_t getHostDevice() { return get().HostDevice; }
+ /// Return the number of non-host devices available to kernel languages.
+ static int getDeviceCount();
- ArrayRef<ol_device_handle_t> getDevices() const { return Devices; }
+ /// Return the thread-local default device and write its number to \p DeviceNo.
+ static ol_device_handle_t getDevice(int *DeviceNo);
- void addDevice(ol_device_handle_t Device) { Devices.push_back(Device); }
- void setHostDevice(ol_device_handle_t Device) {
- if (!HostDevice)
- HostDevice = Device;
- }
+ /// Set the thread-local default device by device number.
+ ///
+ /// \returns the selected device, or nullptr if \p DeviceNo is invalid.
+ static ol_device_handle_t setDefaultDevice(int DeviceNo);
- void addKernel(KernelIDTy KernelID, ol_symbol_handle_t Kernel) {
- KernelMap[KernelID] = Kernel;
- }
+ /// Register \p Kernel for the host-side kernel identifier \p ID.
+ ///
+ /// \p ID is the opaque kernel key emitted by Clang in the offload entry
+ /// table. It is later passed to the launch entry point to recover the
+ /// corresponding liboffload symbol handle.
+ static void registerKernel(const void *ID, ol_symbol_handle_t Kernel);
- void removeKernel(KernelIDTy KernelID) { KernelMap.erase(KernelID); }
+ /// Remove any registered kernel handle for the host-side kernel key \p ID.
+ static void unregisterKernel(const void *ID);
- ol_symbol_handle_t getKernel(KernelIDTy KernelID) {
- return KernelMap[KernelID];
- }
+ /// Return the registered kernel handle for the host-side kernel key \p ID.
+ static ol_symbol_handle_t getKernel(const void *ID);
- void addProgram(const void *Binary, ol_program_handle_t Program) {
- BinaryRegisterMap[Binary] = Program;
- }
+ /// Register \p Program for the binary image identifier \p ID.
+ ///
+ /// \p ID is the device image start address from the offload binary
+ /// descriptor. It keys the loaded program so later function registration
+ /// can look up the program that owns each kernel symbol.
+ static void registerProgram(const void *ID, ol_program_handle_t Program);
- ol_program_handle_t removeProgram(const void *Binary) {
- auto It = BinaryRegisterMap.find(Binary);
- if (It == BinaryRegisterMap.end())
- return nullptr;
- ol_program_handle_t Program = It->second;
- BinaryRegisterMap.erase(It);
- return Program;
- }
+ /// Remove and return the loaded program handle for binary image key \p ID.
+ static ol_program_handle_t unregisterProgram(const void *ID);
- ol_program_handle_t getProgram(const void *Binary) {
- assert(BinaryRegisterMap.count(Binary));
- return BinaryRegisterMap[Binary];
- }
+ /// Return the loaded program handle for binary image key \p ID.
+ static ol_program_handle_t getProgram(const void *ID);
+
+private:
+ static StateTy &get();
+ static StateTy *tryGet();
+ static bool addDevices(ol_device_handle_t Device, void *Payload);
+
+ llvm::ArrayRef<ol_device_handle_t> getDevices() const;
+
+ void addDevice(ol_device_handle_t Device);
+ void setHostDevice(ol_device_handle_t Device);
+
+ void addKernel(KernelIDTy KernelID, ol_symbol_handle_t Kernel);
+ void removeKernel(KernelIDTy KernelID);
+ ol_symbol_handle_t lookupKernel(KernelIDTy KernelID);
+
+ void addProgram(const void *Binary, ol_program_handle_t Program);
+ ol_program_handle_t removeProgram(const void *Binary);
+ ol_program_handle_t lookupProgram(const void *Binary);
void destroyRegisteredPrograms();
-private:
- DenseMap<const void *, ol_program_handle_t> BinaryRegisterMap;
- DenseMap<KernelIDTy, ol_symbol_handle_t> KernelMap;
- SmallVector<ol_device_handle_t, 8> Devices;
+ llvm::DenseMap<const void *, ol_program_handle_t> BinaryRegisterMap;
+ llvm::DenseMap<KernelIDTy, ol_symbol_handle_t> KernelMap;
+ llvm::SmallVector<ol_device_handle_t, 8> Devices;
ol_queue_handle_t DefaultQueue = nullptr;
ol_device_handle_t HostDevice = nullptr;
@@ -104,3 +152,5 @@ struct StateTy {
} // namespace offload
} // namespace llvm
+
+#endif // LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_STATE_H
diff --git a/offload/languages/kernel/include/Types.h b/offload/languages/kernel/include/Types.h
index 6c5a0ec19a8f7..56435daf602fc 100644
--- a/offload/languages/kernel/include/Types.h
+++ b/offload/languages/kernel/include/Types.h
@@ -1,16 +1,14 @@
-//===------- Types.h - Kernel Language (CUDA/HIP) api types ---------------===//
+//===-- Types.h - Kernel language API types -------------------------------===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
-//
-//===----------------------------------------------------------------------===//
-#pragma once
+#ifndef LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_TYPES_H
+#define LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_TYPES_H
-#include "Types.h"
#include <cstddef>
#include <cstdint>
@@ -26,3 +24,5 @@ struct CallConfigurationTy {
size_t SharedMemory;
void *Stream;
};
+
+#endif // LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_TYPES_H
diff --git a/offload/languages/kernel/src/ExportedAPI.cpp b/offload/languages/kernel/src/ExportedAPI.cpp
deleted file mode 100644
index 5c65d3d66b89b..0000000000000
--- a/offload/languages/kernel/src/ExportedAPI.cpp
+++ /dev/null
@@ -1,99 +0,0 @@
-//===------ ExportedAPI.cpp - Kernel Language runtime - exported api ------===//
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-//===----------------------------------------------------------------------===//
-//
-//===----------------------------------------------------------------------===//
-
-#include "ExportedAPI.h"
-
-#include "State.h"
-#include "Types.h"
-
-#include "OffloadAPI.h"
-#include "llvm/ADT/ArrayRef.h"
-
-#include <cstdio>
-#include <stdint.h>
-
-using namespace llvm;
-using namespace offload;
-
-/// Runtime API
-///{
-ol_device_handle_t olKGetDefaultDevice() {
- ol_device_handle_t DefaultDevice = ThreadStateTy::getDefaultDevice();
- return DefaultDevice;
-}
-
-ol_device_handle_t olKGetHostDevice() {
- ol_device_handle_t HostDevice = StateTy::getHostDevice();
- return HostDevice;
-}
-
-int olKGetDeviceCount() {
- int DeviceCount = StateTy::get().getDevices().size();
- return DeviceCount;
-}
-
-ol_device_handle_t olKGetDevice(int *DeviceNo) {
- ol_device_handle_t DefaultDevice = ThreadStateTy::getDefaultDevice();
- int DeviceCount = StateTy::get().getDevices().size();
- ArrayRef<ol_device_handle_t> Devices = StateTy::get().getDevices();
- for (int i = 0; i < DeviceCount; i++) {
- if (Devices[i] == DefaultDevice) {
- *DeviceNo = i;
- return Devices[i];
- }
- }
- return nullptr;
-}
-
-ol_device_handle_t olKSetDefaultDevice(int DeviceNo) {
- ArrayRef<ol_device_handle_t> Devices = StateTy::get().getDevices();
- if (DeviceNo < 0 || DeviceNo >= static_cast<int>(Devices.size()))
- return nullptr;
- ol_device_handle_t Device = Devices[DeviceNo];
- ThreadStateTy::get().setDefaultDevice(Device);
- return Device;
-}
-
-ol_queue_handle_t olKGetDefaultQueue() {
- ol_queue_handle_t DefaultQueue = ThreadStateTy::getDefaultQueue();
- return DefaultQueue;
-}
-
-CallConfigurationTy *olKGetCallConfiguration() {
- return &ThreadStateTy::getCallConfiguration();
-}
-
-void olKRegisterKernel(const void *ID, ol_symbol_handle_t Kernel) {
- StateTy::get().addKernel(ID, Kernel);
-}
-
-void olKUnregisterKernel(const void *ID) {
- if (StateTy *State = StateTy::tryGet())
- State->removeKernel(ID);
-}
-
-ol_symbol_handle_t olKGetKernel(const void *ID) {
- return StateTy::get().getKernel(ID);
-}
-
-void olKRegisterProgram(const void *ID, ol_program_handle_t Program) {
- StateTy::get().addProgram(ID, Program);
-}
-
-ol_program_handle_t olKUnregisterProgram(const void *ID) {
- if (StateTy *State = StateTy::tryGet())
- return State->removeProgram(ID);
- return nullptr;
-}
-
-ol_program_handle_t olKGetProgram(const void *ID) {
- return StateTy::get().getProgram(ID);
-}
-///}
diff --git a/offload/languages/kernel/src/LanguageCommon.cpp b/offload/languages/kernel/src/LanguageCommon.cpp
new file mode 100644
index 0000000000000..f0e00ed9db050
--- /dev/null
+++ b/offload/languages/kernel/src/LanguageCommon.cpp
@@ -0,0 +1,18 @@
+//===-- LanguageCommon.cpp - Shared CUDA/HIP runtime entry points ---------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "LanguageLaunch.h"
+#include "LanguageRegistration.h"
+
+#define LANGUAGE cuda
+#include "LanguageAliases.inc"
+#undef LANGUAGE
+
+#define LANGUAGE hip
+#include "LanguageAliases.inc"
+#undef LANGUAGE
diff --git a/offload/languages/kernel/src/LanguageLaunch.cpp b/offload/languages/kernel/src/LanguageLaunch.cpp
new file mode 100644
index 0000000000000..20779f89f969a
--- /dev/null
+++ b/offload/languages/kernel/src/LanguageLaunch.cpp
@@ -0,0 +1,83 @@
+//===-- LanguageLaunch.cpp - Language launch API --------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "LanguageLaunch.h"
+#include "State.h"
+
+#include <cstdio>
+
+using RuntimeState = llvm::offload::StateTy;
+using ThreadState = llvm::offload::ThreadStateTy;
+
+extern "C" {
+
+/// Push call configuration for kernel launch
+unsigned __llvmPushCallConfiguration(dim3 __grid_size, dim3 __block_size,
+ size_t __shared_memory, void *__stream) {
+ CallConfigurationTy &CC = ThreadState::getCallConfiguration();
+
+ CC.GridSize = __grid_size;
+ CC.BlockSize = __block_size;
+ CC.SharedMemory = __shared_memory;
+ CC.Stream = __stream;
+ return 0;
+}
+
+/// Pop call configuration for kernel launch
+unsigned __llvmPopCallConfiguration(dim3 *__grid_size, dim3 *__block_size,
+ size_t *__shared_memory, void **__stream) {
+ CallConfigurationTy &CC = ThreadState::getCallConfiguration();
+ *__grid_size = CC.GridSize;
+ *__block_size = CC.BlockSize;
+ *__shared_memory = CC.SharedMemory;
+ *__stream = CC.Stream;
+ return 0;
+}
+
+/// Internal kernel launch implementation
+ol_result_t __llvmLaunchKernelImpl(const char *KernelID, dim3 GridDim,
+ dim3 BlockDim, void *KernelArgsPtr,
+ size_t DynamicSharedMem, void *Stream) {
+ ol_device_handle_t Device = ThreadState::getDefaultDevice();
+ ol_symbol_handle_t Kernel = RuntimeState::getKernel(KernelID);
+
+ ol_kernel_launch_size_args_t LaunchSizeArgs;
+ LaunchSizeArgs.Dimensions =
+ 1 + (GridDim.y > 1 || BlockDim.y > 1) + (GridDim.z > 1 || BlockDim.z > 1);
+ LaunchSizeArgs.NumGroups.x = GridDim.x;
+ LaunchSizeArgs.NumGroups.y = std::max(GridDim.y, 1u);
+ LaunchSizeArgs.NumGroups.z = std::max(GridDim.z, 1u);
+ LaunchSizeArgs.GroupSize.x = BlockDim.x;
+ LaunchSizeArgs.GroupSize.y = std::max(BlockDim.y, 1u);
+ LaunchSizeArgs.GroupSize.z = std::max(BlockDim.z, 1u);
+ LaunchSizeArgs.DynSharedMemory = DynamicSharedMem;
+
+ ol_queue_handle_t Queue = Stream ? reinterpret_cast<ol_queue_handle_t>(Stream)
+ : ThreadState::getDefaultQueue();
+
+ struct OffloadKernelArgs {
+ void **Args;
+ size_t NumArgs;
+ size_t *ArgSizes;
+ };
+ OffloadKernelArgs *OKA = static_cast<OffloadKernelArgs *>(KernelArgsPtr);
+
+ return olLaunchKernel(Queue, Device, Kernel, &LaunchSizeArgs,
+ /*Properties=*/nullptr, OKA->NumArgs, OKA->Args,
+ OKA->ArgSizes);
+}
+
+unsigned __llvmLaunchKernel(const char *KernelID, dim3 GridDim, dim3 BlockDim,
+ void *KernelArgsPtr, size_t DynamicSharedMem,
+ void *Stream) {
+ ol_result_t Result = __llvmLaunchKernelImpl(
+ KernelID, GridDim, BlockDim, KernelArgsPtr, DynamicSharedMem, Stream);
+ return Result ? Result->Code : 0;
+}
+
+} // extern "C"
diff --git a/offload/languages/kernel/src/LanguageRegistration.cpp b/offload/languages/kernel/src/LanguageRegistration.cpp
new file mode 100644
index 0000000000000..231eaa691ff0b
--- /dev/null
+++ b/offload/languages/kernel/src/LanguageRegistration.cpp
@@ -0,0 +1,318 @@
+//===-- LanguageRegistration.cpp - Language registration API --------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "LanguageRegistration.h"
+#include "OffloadAPI.h"
+#include "State.h"
+#include "llvm/ADT/StringRef.h"
+#include "llvm/Frontend/Offloading/Utility.h"
+#include "llvm/Support/Error.h"
+#include "llvm/Support/raw_ostream.h"
+#include <cstdio>
+#include <cstring>
+#include <inttypes.h>
+
+using RuntimeState = llvm::offload::StateTy;
+using ThreadState = llvm::offload::ThreadStateTy;
+
+typedef struct __attribute__((__packed__)) {
+ uint32_t Magic;
+ uint16_t Version;
+ uint16_t HeaderSize;
+ uint64_t FatSize;
+} CudaFatbinHeader;
+
+// Inspired by
+// https://github.com/n-eiling/cuda-fatbin-decompression/blob/master/fatbin-decompress.h
+typedef struct __attribute__((__packed__)) {
+ uint16_t Kind;
+ uint16_t Unknown1;
+ uint32_t HeaderSize;
+ uint64_t Size;
+ uint32_t CompressedSize;
+ uint32_t Unknown2;
+ uint16_t Minor;
+ uint16_t Major;
+ uint32_t Arch;
+ uint32_t ObjNameOffset;
+ uint32_t ObjNameLen;
+ uint64_t Flags;
+ uint64_t Zero;
+ uint64_t DecompressedSize;
+} CudaFatbinTextHeader;
+
+// HIP uses this format:
+// https://clang.llvm.org/docs/ClangOffloadBundler.html#bundled-binary-file-layout
+typedef struct __attribute__((__packed__)) {
+ char Magic[24];
+ uint64_t NumBundles;
+} HipFatbinHeader;
+
+typedef struct __attribute__((__packed__)) {
+ uint64_t BundleOffset;
+ uint64_t BundleSize;
+ uint64_t IdLength;
+ char IdString[];
+} HipFatbinBundleEntry;
+
+static void readTUFatbin(const char *Binary, const FatbinWrapperTy *FW) {
+ ol_device_handle_t Device = language_registration::getDefaultDevice();
+
+ const CudaFatbinHeader *Header =
+ reinterpret_cast<const CudaFatbinHeader *>(FW->Data);
+ size_t HeaderSize = static_cast<size_t>(Header->HeaderSize); // Usually 16
+ size_t FatbinSize = static_cast<size_t>(Header->FatSize);
+
+ const void *ProgramData = nullptr;
+ size_t ProgramSize = 0;
+ uint32_t ProgramArch = 0;
+
+ const char *ReadPosition = FW->Data + HeaderSize;
+ while (ReadPosition < (FW->Data + FatbinSize)) {
+ const CudaFatbinTextHeader *TextHeader =
+ reinterpret_cast<const CudaFatbinTextHeader *>(ReadPosition);
+ size_t TextHeaderSize =
+ static_cast<size_t>(TextHeader->HeaderSize); // Usually 64
+ size_t CubinSize = static_cast<size_t>(TextHeader->Size);
+ const void *CubinData =
+ static_cast<const char *>(ReadPosition + TextHeaderSize);
+
+ uint32_t Arch = TextHeader->Arch;
+ bool IsCompatible = false;
+ olIsValidBinary(Device, CubinData, CubinSize, &IsCompatible);
+ if (!IsCompatible) {
+ fprintf(stderr, "Device is not compatible with image.");
+ abort();
+ }
+
+ if (Arch > ProgramArch) {
+ ProgramData = CubinData;
+ ProgramSize = CubinSize;
+ ProgramArch = Arch;
+ }
+
+ ReadPosition += TextHeaderSize + CubinSize;
+ }
+
+ if (ProgramData == nullptr) {
+ fprintf(stderr, "Failed to find compatible binary\n");
+ abort();
+ }
+
+ ol_program_handle_t Program = nullptr;
+
+ ol_result_t Result =
+ olCreateProgram(Device, ProgramData, ProgramSize, &Program);
+
+ if (Result && Result->Code) {
+ fprintf(stderr, "Failed to register device code (%i): %s\n", Result->Code,
+ Result->Details);
+ abort();
+ }
+
+ language_registration::registerProgram(Binary, Program);
+}
+
+static void readHIPFatbinEntries(const char *Binary, const char *HIPFatbinPtr) {
+ ol_device_handle_t Device = language_registration::getDefaultDevice();
+
+ const char *CurrentReadPosition = HIPFatbinPtr;
+
+ const HipFatbinHeader *Header =
+ reinterpret_cast<const HipFatbinHeader *>(CurrentReadPosition);
+ CurrentReadPosition += sizeof(HipFatbinHeader);
+
+ uint64_t NumBundles = Header->NumBundles;
+
+ const void *ProgramData = nullptr;
+ size_t ProgramSize = 0;
+ uint64_t ProgramIdLength = 0;
+ const char *ProgramIdString = nullptr;
+
+ for (uint64_t BundleId = 0; BundleId < NumBundles; ++BundleId) {
+ const HipFatbinBundleEntry *BundleEntry =
+ reinterpret_cast<const HipFatbinBundleEntry *>(CurrentReadPosition);
+
+ uint64_t BundleOffset = BundleEntry->BundleOffset;
+ uint64_t BundleSize = BundleEntry->BundleSize;
+ const char *BundleIdString = BundleEntry->IdString;
+ uint64_t BundleIdLength = BundleEntry->IdLength;
+
+ // Advance by the size of the entry including the ID string
+ CurrentReadPosition += sizeof(HipFatbinBundleEntry) + BundleIdLength;
+
+ if (!BundleSize) {
+ continue;
+ }
+
+ bool IsCompatible = false;
+ olIsValidBinary(Device, HIPFatbinPtr + BundleOffset, BundleSize,
+ &IsCompatible);
+
+ if (!IsCompatible) {
+ fprintf(stderr, "Device is not compatible with image.");
+ abort();
+ }
+
+ llvm::StringRef CurrentBundleId(ProgramIdString, ProgramIdLength);
+ llvm::StringRef NewBundleId(BundleIdString, BundleIdLength);
+ if (NewBundleId.compare(CurrentBundleId) > 0) {
+ ProgramData = HIPFatbinPtr + BundleOffset;
+ ProgramSize = BundleSize;
+ ProgramIdLength = BundleIdLength;
+ ProgramIdString = BundleIdString;
+ }
+ }
+
+ if (ProgramData == nullptr) {
+ fprintf(stderr, "Failed to find compatible binary\n");
+ abort();
+ }
+
+ ol_program_handle_t Program = nullptr;
+ ol_result_t Result =
+ olCreateProgram(Device, ProgramData, ProgramSize, &Program);
+ if (Result && Result->Code) {
+ fprintf(stderr, "Failed to register device code (%i): %s\n", Result->Code,
+ Result->Details);
+ abort();
+ }
+
+ language_registration::registerProgram(Binary, Program);
+}
+
+/// Hidden, but exported, Registration API
+///{
+extern "C" {
+
+void __llvmRegisterFunction(const char *Binary, const char *KernelID,
+ char *KernelName, const char *KernelName1, int,
+ uint3 *, uint3 *, dim3 *, dim3 *, int *) {
+ ol_symbol_handle_t Kernel;
+ ol_program_handle_t Program = RuntimeState::getProgram(Binary);
+ ol_result_t Result = olGetSymbol(
+ Program, KernelName, ol_symbol_kind_t::OL_SYMBOL_KIND_KERNEL, &Kernel);
+ if (Result && Result->Code) {
+ fprintf(stderr, "Failed to register kernel (%i): %s\n", Result->Code,
+ Result->Details);
+ abort();
+ }
+
+ RuntimeState::registerKernel(KernelID, Kernel);
+}
+
+const char *__llvmRegisterFatBinary(const char *Binary) {
+ const auto *FW = reinterpret_cast<const FatbinWrapperTy *>(Binary);
+ if (FW->Magic == 0x466243b1) {
+ readTUFatbin(Binary, FW);
+ } else if (FW->Magic == 0x48495046) {
+ if (!memcmp(FW->Data, HIP_FATBIN_MAGIC_STR, HIP_FATBIN_MAGIC_STR_LEN))
+ readHIPFatbinEntries(Binary, FW->Data);
+ else
+ readTUFatbin(Binary, FW);
+ } else {
+ fprintf(stderr, "Unknown fatbin format");
+ }
+
+ return Binary;
+}
+
+void __llvmUnregisterFatBinary(void *Handle) {
+ if (ol_program_handle_t Program =
+ language_registration::unregisterProgram(Handle))
+ olDestroyProgram(Program);
+}
+
+void __llvmRegisterVar(void **, char *, char *, const char *, int, int, int,
+ int) {
+ llvm::errs() << "RegisterVar is not implemented!" << "\n";
+}
+
+void __llvmRegisterManagedVar(void **, char *, char *, const char *, size_t,
+ unsigned) {
+ llvm::errs() << "RegisterManagedVar is not implemented!" << "\n";
+}
+
+void __llvmRegisterSurface(void **, const struct surfaceReference *,
+ const void **, const char *, int, int) {
+ llvm::errs() << "RegisterSurface is not implemented!" << "\n";
+}
+
+void __llvmRegisterTexture(void **, const struct textureReference *,
+ const void **, const char *, int, int, int) {
+ llvm::errs() << "RegisterTexture is not implemented!" << "\n";
+}
+
+/// This struct is a record of the device image information
+struct __tgt_device_image {
+ void *ImageStart; // Pointer to the target code start
+ void *ImageEnd; // Pointer to the target code end
+ llvm::offloading::EntryTy
+ *EntriesBegin; // Begin of table with all target entries
+ llvm::offloading::EntryTy *EntriesEnd; // End of table (non inclusive)
+};
+
+/// This struct is a record of all the host code that may be offloaded to a
+/// target.
+struct __tgt_bin_desc {
+ int32_t NumDeviceImages; // Number of device types supported
+ __tgt_device_image *DeviceImages; // Array of device images (1 per dev. type)
+ llvm::offloading::EntryTy
+ *HostEntriesBegin; // Begin of table with all host entries
+ llvm::offloading::EntryTy *HostEntriesEnd; // End of table (non inclusive)
+};
+
+void __tgt_register_lib(__tgt_bin_desc *Desc) {
+ // TODO: For each device, lazily.
+ ol_device_handle_t Device = ThreadState::getDefaultDevice();
+
+ for (int32_t I = 0, E = Desc->NumDeviceImages; I < E; ++I) {
+ ol_program_handle_t Program = nullptr;
+
+ __tgt_device_image &DeviceImage = Desc->DeviceImages[I];
+ void *ProgramData = DeviceImage.ImageStart;
+ size_t ProgramSize =
+ (char *)DeviceImage.ImageEnd - (char *)DeviceImage.ImageStart;
+ ol_result_t Result =
+ olCreateProgram(Device, ProgramData, ProgramSize, &Program);
+
+ if (Result && Result->Code) {
+ fprintf(stderr, "Failed to register device code (%i): %s\n", Result->Code,
+ Result->Details);
+ abort();
+ }
+
+ RuntimeState::registerProgram(DeviceImage.ImageStart, Program);
+
+ for (auto *Entry = DeviceImage.EntriesBegin;
+ Entry != DeviceImage.EntriesEnd; ++Entry) {
+ if (!Entry->Size && !Entry->Flags)
+ __llvmRegisterFunction((const char *)DeviceImage.ImageStart,
+ (const char *)Entry->Address, Entry->SymbolName,
+ Entry->SymbolName, 0, nullptr, nullptr, nullptr,
+ nullptr, nullptr);
+ }
+ }
+}
+
+void __tgt_unregister_lib(__tgt_bin_desc *Desc) {
+ for (int32_t I = 0, E = Desc->NumDeviceImages; I < E; ++I) {
+ __tgt_device_image &DeviceImage = Desc->DeviceImages[I];
+ for (auto *Entry = DeviceImage.EntriesBegin;
+ Entry != DeviceImage.EntriesEnd; ++Entry) {
+ if (!Entry->Size && !Entry->Flags)
+ RuntimeState::unregisterKernel((const char *)Entry->Address);
+ }
+
+ if (ol_program_handle_t Program =
+ RuntimeState::unregisterProgram(DeviceImage.ImageStart))
+ olDestroyProgram(Program);
+ }
+}
+}
+///}
diff --git a/offload/languages/kernel/src/LanguageRuntime.cpp b/offload/languages/kernel/src/LanguageRuntime.cpp
new file mode 100644
index 0000000000000..b1b21f57e35f1
--- /dev/null
+++ b/offload/languages/kernel/src/LanguageRuntime.cpp
@@ -0,0 +1,223 @@
+//===-- LanguageRuntime.cpp - Kernel language runtime API implementation --===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LANGUAGE
+#error This file should be included, or used, with a LANGUAGE macro set.
+#endif
+
+// Rename the generic runtime API before declaring or defining language symbols.
+#include "DefineLanguageNames.inc"
+#include "LanguageRuntime.h"
+
+#include "State.h"
+#include "Types.h"
+
+#include "OffloadAPI.h"
+
+#include <cassert>
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+
+#define STR(X) #X
+#define LANGUAGE_STR STR(LANGUAGE)
+
+using RuntimeState = llvm::offload::StateTy;
+using ThreadState = llvm::offload::ThreadStateTy;
+
+static Error_t convertResult(ol_result_t Result) {
+ if (Result == OL_SUCCESS)
+ return Success;
+ switch (Result->Code) {
+ case OL_ERRC_INVALID_VALUE:
+ return ErrorInvalidValue;
+ default:
+ return ErrorInvalidValue;
+ }
+}
+
+Error_t Malloc(void **DevPtr, size_t Size) {
+ ol_device_handle_t Device = ThreadState::getDefaultDevice();
+ ol_result_t Result = olMemAlloc(Device, OL_ALLOC_TYPE_DEVICE, Size, DevPtr);
+ return convertResult(Result);
+}
+
+Error_t Free(void *DevPtr) {
+ ol_result_t Result = olMemFree(DevPtr);
+ return convertResult(Result);
+}
+
+Error_t Memcpy(void *Dst, const void *Src, size_t Size, MemcpyKind Kind) {
+ ol_queue_handle_t Queue = ThreadState::getDefaultQueue();
+
+ ol_result_t Result;
+ switch (Kind) {
+ case MemcpyHostToHost: {
+ ol_device_handle_t Host = RuntimeState::getHostDevice();
+ Result = olMemcpy(nullptr, Dst, Host, const_cast<void *>(Src), Host, Size);
+ break;
+ }
+ case MemcpyHostToDevice: {
+ ol_device_handle_t Device = ThreadState::getDefaultDevice();
+ ol_device_handle_t Host = RuntimeState::getHostDevice();
+ Result = olMemcpy(Queue, Dst, Device, const_cast<void *>(Src), Host, Size);
+ break;
+ }
+ case MemcpyDeviceToHost: {
+ ol_device_handle_t Device = ThreadState::getDefaultDevice();
+ ol_device_handle_t Host = RuntimeState::getHostDevice();
+
+ Result = olMemcpy(Queue, Dst, Host, const_cast<void *>(Src), Device, Size);
+ break;
+ }
+ case MemcpyDeviceToDevice: {
+ ol_device_handle_t Device = ThreadState::getDefaultDevice();
+
+ Result =
+ olMemcpy(Queue, Dst, Device, const_cast<void *>(Src), Device, Size);
+ break;
+ }
+ case MemcpyDefault:
+ fprintf(stderr, LANGUAGE_STR "MemcpyDefault is not implemented yet");
+ abort();
+ };
+
+ Result = olSyncQueue(Queue);
+
+ return convertResult(Result);
+}
+
+Error_t DeviceSynchronize() {
+ // TODO: This is not correct. We likely want to pipe this through to the
+ // plugins.
+ ol_queue_handle_t Queue = ThreadState::getDefaultQueue();
+ ol_result_t Result = olSyncQueue(Queue);
+ return convertResult(Result);
+}
+
+Error_t GetLastError() {
+ // TODO:
+ return Success;
+}
+
+Error_t PeekAtLastError() {
+ // TODO:
+ return Success;
+}
+
+const char *GetErrorName(Error_t Error) {
+ // TODO:
+ return "";
+}
+
+const char *GetErrorString(Error_t Error) {
+ // TODO:
+ return "";
+}
+
+Error_t GetDevice(int *DeviceNo) {
+ ol_device_handle_t Device = RuntimeState::getDevice(DeviceNo);
+ if (!Device)
+ return ErrorInvalidValue;
+ return Success;
+}
+
+Error_t GetDeviceCount(int *Count) {
+ *Count = RuntimeState::getDeviceCount();
+ return Success;
+}
+
+Error_t SetDevice(int DeviceNo) {
+ ol_device_handle_t Device = RuntimeState::setDefaultDevice(DeviceNo);
+ if (!Device)
+ return ErrorInvalidValue;
+ assert(Device == ThreadState::getDefaultDevice() &&
+ "Set Device is not Default Device");
+ return Device ? Success : ErrorInvalidValue;
+}
+
+Error_t HostAlloc(void **Ptr, size_t Size, unsigned int Flags) {
+ ol_device_handle_t Device = ThreadState::getDefaultDevice();
+ ol_result_t Result = olMemAllocHost(Device, Size, Ptr);
+ return convertResult(Result);
+}
+
+Error_t MallocHost(void **Ptr, size_t Size) {
+ return HostAlloc(Ptr, Size, /* HostAllocDefault */ 0);
+}
+
+Error_t FreeHost(void *Ptr) {
+ ol_result_t Result = olMemFree(Ptr);
+ return convertResult(Result);
+}
+
+Error_t DriverGetVersion(int *Version) {
+ // TODO:
+ *Version = 42;
+ return Success;
+}
+
+Error_t GetDeviceProperties(DeviceProp_t *DeviceProp, int DeviceNo) {
+ ol_device_handle_t Device = ThreadState::getDefaultDevice();
+ size_t NameSize = 0;
+ olGetDeviceInfoSize(Device, OL_DEVICE_INFO_NAME, &NameSize);
+ assert(NameSize <= sizeof(DeviceProp->name) &&
+ "Device name is too long for DeviceProp_t");
+ olGetDeviceInfo(Device, OL_DEVICE_INFO_NAME, NameSize, &DeviceProp->name[0]);
+ olGetDeviceInfo(Device, OL_DEVICE_INFO_GLOBAL_MEM_SIZE, sizeof(size_t),
+ &DeviceProp->totalGlobalMem);
+ olGetDeviceInfo(Device, OL_DEVICE_INFO_NUM_COMPUTE_UNITS, sizeof(uint32_t),
+ &DeviceProp->multiProcessorCount);
+ olGetDeviceInfo(Device, OL_DEVICE_INFO_NUM_LANES, sizeof(uint32_t),
+ &DeviceProp->warpSize);
+ return Success;
+}
+
+static Error_t getQueueFromStream(Stream_t Stream, ol_queue_handle_t *Queue) {
+ if (!Stream)
+ return ErrorInvalidValue;
+ *Queue = reinterpret_cast<ol_queue_handle_t>(Stream);
+ return Success;
+}
+
+Error_t StreamCreate(Stream_t *Stream) {
+ ol_queue_handle_t Queue;
+ olCreateQueue(ThreadState::getDefaultDevice(), &Queue);
+ *Stream = reinterpret_cast<Stream_t>(Queue);
+ return Success;
+}
+
+Error_t StreamCreateWithFlags(Stream_t *Stream, unsigned int Flags) {
+ if (Flags == StreamCreateWithFlagsFlags::StreamDefault)
+ // FIXME: [h15] offload streams are non-blocking by default
+ return StreamCreate(Stream);
+ if (Flags == StreamCreateWithFlagsFlags::StreamNonBlocking) {
+ return StreamCreate(Stream);
+ }
+ return ErrorInvalidValue;
+}
+
+Error_t StreamDestroy(Stream_t Stream) {
+ ol_queue_handle_t Queue;
+ Error_t Err = getQueueFromStream(Stream, &Queue);
+ if (Err != Success)
+ return Err;
+ ol_result_t Result = olDestroyQueue(Queue);
+ return convertResult(Result);
+}
+
+Error_t StreamSynchronize(Stream_t Stream) {
+ ol_queue_handle_t Queue;
+ Error_t Err = getQueueFromStream(Stream, &Queue);
+ if (Err != Success)
+ return Err;
+ ol_result_t Result = olSyncQueue(Queue);
+ return convertResult(Result);
+}
+
+#include "UndefineLanguageNames.inc"
diff --git a/offload/languages/kernel/src/State.cpp b/offload/languages/kernel/src/State.cpp
index eff3c31257e93..3dbd1f436ff24 100644
--- a/offload/languages/kernel/src/State.cpp
+++ b/offload/languages/kernel/src/State.cpp
@@ -1,38 +1,41 @@
-//===------ State.cpp - Kernel Language (CUDA/HIP) persistent state -------===//
+//===-- State.cpp - Kernel language persistent state ----------------------===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
-//
-//===----------------------------------------------------------------------===//
#include "State.h"
#include "Types.h"
#include "OffloadAPI.h"
+#include "llvm/ADT/ArrayRef.h"
+#include "llvm/ADT/DenseMap.h"
#include "llvm/ADT/SmallPtrSet.h"
+#include "llvm/ADT/SmallVector.h"
#include "llvm/Support/raw_ostream.h"
#include <atomic>
+#include <cassert>
#include <cstdio>
#include <mutex>
#define CHECK_FATAL(Result, ...) \
if (Result && Result->Code) { \
- fprintf(stderr, __VA_ARGS__); \
+ llvm::errs() << __VA_ARGS__ << '\n'; \
abort(); \
}
using namespace llvm;
using namespace offload;
-std::atomic<uint32_t> AnyNonDefaultDevice = 0;
+// Weak so another runtime object can override the default stream mode.
__attribute__((weak)) uint32_t PerThreadQueue = 0;
thread_local CallConfigurationTy CC = {};
+// Process-wide singleton and thread-state registry.
static std::mutex StateLock;
static std::atomic<StateTy *> StatePtr = nullptr;
@@ -50,7 +53,7 @@ static void deleteThreadState() {
return;
for (auto *TS : *ThreadStates)
- delete (TS);
+ delete TS;
delete ThreadStates;
ThreadState = nullptr;
}
@@ -58,7 +61,7 @@ static void deleteThreadState() {
static void deleteState() {
StateTy *ST = StatePtr.load();
StatePtr.store(nullptr);
- delete (ST);
+ delete ST;
}
static void destroyQueue(ol_queue_handle_t &Queue) {
@@ -70,6 +73,11 @@ static void destroyQueue(ol_queue_handle_t &Queue) {
Queue = nullptr;
}
+namespace llvm {
+namespace offload {
+
+// ThreadStateTy implementation.
+
ThreadStateTy::ThreadStateTy() {
if (PerThreadQueue) [[unlikely]]
createDefaultQueue(getDefaultDevice());
@@ -98,11 +106,6 @@ ol_device_handle_t ThreadStateTy::getDefaultDevice() {
DD = Device;
break;
}
- if (AnyNonDefaultDevice.load(std::memory_order_relaxed)) [[unlikely]] {
- ol_device_handle_t TDD = ThreadStateTy::get().DefaultDevice;
- if (TDD)
- DD = TDD;
- }
return DD;
}
@@ -117,8 +120,9 @@ CallConfigurationTy &ThreadStateTy::getCallConfiguration() {
}
void ThreadStateTy::setDefaultDevice(ol_device_handle_t Device) {
- DefaultDevice = Device;
- createDefaultQueue(Device);
+ ThreadStateTy &State = get();
+ State.DefaultDevice = Device;
+ State.createDefaultQueue(Device);
}
void ThreadStateTy::createDefaultQueue(ol_device_handle_t Device) {
@@ -128,6 +132,8 @@ void ThreadStateTy::createDefaultQueue(ol_device_handle_t Device) {
"Failed to create per-thread default queue");
}
+// StateTy implementation.
+
StateTy &StateTy::get() {
StateTy *ST = StatePtr.load();
if (!ST) [[unlikely]] {
@@ -143,7 +149,103 @@ StateTy &StateTy::get() {
StateTy *StateTy::tryGet() { return StatePtr.load(); }
-static bool addDevices(ol_device_handle_t Device, void *Payload) {
+ol_device_handle_t StateTy::getHostDevice() { return get().HostDevice; }
+
+int StateTy::getDeviceCount() {
+ int DeviceCount = get().getDevices().size();
+ return DeviceCount;
+}
+
+ol_device_handle_t StateTy::getDevice(int *DeviceNo) {
+ ol_device_handle_t DefaultDevice = ThreadStateTy::getDefaultDevice();
+ int DeviceCount = get().getDevices().size();
+ ArrayRef<ol_device_handle_t> Devices = get().getDevices();
+ for (int i = 0; i < DeviceCount; i++) {
+ if (Devices[i] == DefaultDevice) {
+ *DeviceNo = i;
+ return Devices[i];
+ }
+ }
+ return nullptr;
+}
+
+ol_device_handle_t StateTy::setDefaultDevice(int DeviceNo) {
+ ArrayRef<ol_device_handle_t> Devices = get().getDevices();
+ if (DeviceNo < 0 || DeviceNo >= static_cast<int>(Devices.size()))
+ return nullptr;
+ ol_device_handle_t Device = Devices[DeviceNo];
+ ThreadStateTy::setDefaultDevice(Device);
+ return Device;
+}
+
+ArrayRef<ol_device_handle_t> StateTy::getDevices() const { return Devices; }
+
+void StateTy::addDevice(ol_device_handle_t Device) {
+ Devices.push_back(Device);
+}
+
+void StateTy::setHostDevice(ol_device_handle_t Device) {
+ if (!HostDevice)
+ HostDevice = Device;
+}
+
+void StateTy::addKernel(KernelIDTy KernelID, ol_symbol_handle_t Kernel) {
+ KernelMap[KernelID] = Kernel;
+}
+
+void StateTy::removeKernel(KernelIDTy KernelID) { KernelMap.erase(KernelID); }
+
+ol_symbol_handle_t StateTy::lookupKernel(KernelIDTy KernelID) {
+ return KernelMap[KernelID];
+}
+
+void StateTy::registerKernel(const void *ID, ol_symbol_handle_t Kernel) {
+ get().addKernel(ID, Kernel);
+}
+
+void StateTy::unregisterKernel(const void *ID) {
+ if (StateTy *State = tryGet())
+ State->removeKernel(ID);
+}
+
+ol_symbol_handle_t StateTy::getKernel(const void *ID) {
+ return get().lookupKernel(ID);
+}
+
+void StateTy::addProgram(const void *Binary, ol_program_handle_t Program) {
+ BinaryRegisterMap[Binary] = Program;
+}
+
+ol_program_handle_t StateTy::removeProgram(const void *Binary) {
+ auto It = BinaryRegisterMap.find(Binary);
+ if (It == BinaryRegisterMap.end())
+ return nullptr;
+ ol_program_handle_t Program = It->second;
+ BinaryRegisterMap.erase(It);
+ return Program;
+}
+
+ol_program_handle_t StateTy::lookupProgram(const void *Binary) {
+ assert(BinaryRegisterMap.count(Binary) &&
+ "Program not registered for binary");
+ return BinaryRegisterMap[Binary];
+}
+
+void StateTy::registerProgram(const void *ID, ol_program_handle_t Program) {
+ get().addProgram(ID, Program);
+}
+
+ol_program_handle_t StateTy::unregisterProgram(const void *ID) {
+ if (StateTy *State = tryGet())
+ return State->removeProgram(ID);
+ return nullptr;
+}
+
+ol_program_handle_t StateTy::getProgram(const void *ID) {
+ return get().lookupProgram(ID);
+}
+
+bool StateTy::addDevices(ol_device_handle_t Device, void *Payload) {
StateTy &State = *reinterpret_cast<StateTy *>(Payload);
ol_platform_handle_t Platform;
ol_result_t Result;
@@ -168,7 +270,8 @@ static bool addDevices(ol_device_handle_t Device, void *Payload) {
StateTy::StateTy() {
CHECK_FATAL(olInit(nullptr), "Failed to initialize the LLVMOffload");
- CHECK_FATAL(olIterateDevices(addDevices, this), "Failed to identify devices");
+ CHECK_FATAL(olIterateDevices(StateTy::addDevices, this),
+ "Failed to identify devices");
if (!PerThreadQueue) [[likely]]
if (!Devices.empty()) [[likely]]
@@ -196,3 +299,6 @@ void StateTy::destroyRegisteredPrograms() {
for (ol_program_handle_t Program : Programs)
olDestroyProgram(Program);
}
+
+} // namespace offload
+} // namespace llvm
>From 67498fac91e7fe41beed1364d141f4a194ec973e Mon Sep 17 00:00:00 2001
From: Sophia Herrmann <herrmann15 at llnl.gov>
Date: Fri, 24 Jul 2026 10:56:46 -0700
Subject: [PATCH 3/7] remove unimplemented stuff + minor fixes
formatting
remove empty stubs
fix HostToHost bug + SetDevice assertion
add C++ style includes
fix olMemAlloc after rebase
---
.../include/kernel/DefineLanguageNames.inc | 11 ------
.../include/kernel/LanguageRuntime.h | 39 +------------------
.../include/kernel/UndefineLanguageNames.inc | 10 -----
.../languages/kernel/src/LanguageRuntime.cpp | 38 +-----------------
4 files changed, 3 insertions(+), 95 deletions(-)
diff --git a/offload/languages/include/kernel/DefineLanguageNames.inc b/offload/languages/include/kernel/DefineLanguageNames.inc
index 7f2b0776a03ff..cfab6998c28e6 100644
--- a/offload/languages/include/kernel/DefineLanguageNames.inc
+++ b/offload/languages/include/kernel/DefineLanguageNames.inc
@@ -23,10 +23,6 @@
#define MemcpyDeviceToHost COMBINE(LANGUAGE, MemcpyDeviceToHost)
#define MemcpyDeviceToDevice COMBINE(LANGUAGE, MemcpyDeviceToDevice)
#define MemcpyDefault COMBINE(LANGUAGE, MemcpyDefault)
-#define GetLastError COMBINE(LANGUAGE, GetLastError)
-#define PeekAtLastError COMBINE(LANGUAGE, PeekAtLastError)
-#define GetErrorName COMBINE(LANGUAGE, GetErrorName)
-#define GetErrorString COMBINE(LANGUAGE, GetErrorString)
#define GetDevice COMBINE(LANGUAGE, GetDevice)
#define GetDeviceCount COMBINE(LANGUAGE, GetDeviceCount)
#define SetDevice COMBINE(LANGUAGE, SetDevice)
@@ -37,15 +33,8 @@
#define HostAllocWriteCombined COMBINE(LANGUAGE, HostAllocWriteCombined)
#define MallocHost COMBINE(LANGUAGE, MallocHost)
#define FreeHost COMBINE(LANGUAGE, FreeHost)
-#define DriverGetVersion COMBINE(LANGUAGE, DriverGetVersion)
#define GetDeviceProperties COMBINE(LANGUAGE, GetDeviceProperties)
-#define OccupancyMaxPotentialBlockSizeVariableSMem \
- COMBINE(LANGUAGE, OccupancyMaxPotentialBlockSizeVariableSMem)
#define Stream_t COMBINE(LANGUAGE, Stream_t)
#define StreamCreate COMBINE(LANGUAGE, StreamCreate)
-#define StreamCreateWithFlags COMBINE(LANGUAGE, StreamCreateWithFlags)
#define StreamDestroy COMBINE(LANGUAGE, StreamDestroy)
#define StreamSynchronize COMBINE(LANGUAGE, StreamSynchronize)
-#define StreamCreateWithFlagsFlags COMBINE(LANGUAGE, StreamCreateWithFlagsFlags)
-#define StreamDefault COMBINE(LANGUAGE, StreamDefault)
-#define StreamNonBlocking COMBINE(LANGUAGE, StreamNonBlocking)
diff --git a/offload/languages/include/kernel/LanguageRuntime.h b/offload/languages/include/kernel/LanguageRuntime.h
index dc00ae4ad5a1c..b072087ba0e7d 100644
--- a/offload/languages/include/kernel/LanguageRuntime.h
+++ b/offload/languages/include/kernel/LanguageRuntime.h
@@ -9,10 +9,10 @@
#ifndef LLVM_OFFLOAD_LANGUAGES_INCLUDE_KERNEL_LANGUAGE_RUNTIME_H
#define LLVM_OFFLOAD_LANGUAGES_INCLUDE_KERNEL_LANGUAGE_RUNTIME_H
+#include <cstddef>
+#include <cstdint>
#include <cstdio>
#include <cstdlib>
-#include <stddef.h>
-#include <stdint.h>
enum Error_t : uint32_t {
Success = 0,
@@ -48,11 +48,6 @@ enum HostAllocFlags : unsigned int {
HostAllocWriteCombined = 0x04,
};
-enum StreamCreateWithFlagsFlags : unsigned int {
- StreamDefault = 0x00,
- StreamNonBlocking = 0x01,
-};
-
typedef struct Stream_st *Stream_t;
/// Malloc, with type template overlay.
@@ -93,14 +88,6 @@ static inline Error_t Memcpy(T *Dst, const T *Src, size_t Size,
Error_t DeviceSynchronize();
-Error_t GetLastError();
-
-Error_t PeekAtLastError();
-
-const char *GetErrorName(Error_t Error);
-
-const char *GetErrorString(Error_t Error);
-
Error_t GetDevice(int *DeviceNo);
Error_t GetDeviceCount(int *Count);
@@ -109,36 +96,14 @@ Error_t SetDevice(int DeviceNo);
Error_t FreeHost(void *Ptr);
-Error_t DriverGetVersion(int *Version);
-
Error_t GetDeviceProperties(DeviceProp_t *DeviceProp, int DeviceNo);
Error_t StreamCreate(Stream_t *stream);
-Error_t StreamCreateWithFlags(Stream_t *stream, unsigned int flags);
-
Error_t StreamDestroy(Stream_t stream);
Error_t StreamSynchronize(Stream_t stream);
-template <typename UnaryFunction, class T>
-static inline Error_t OccupancyMaxPotentialBlockSizeVariableSMem(
- int *minGridSize, int *blockSize, T func,
- UnaryFunction blockSizeToDynamicSMemSize, int blockSizeLimit = 0) {
-#if defined(__AMDGPU__)
- // TODO: values taken from AMD Instinct MI250X gfx90a
- *minGridSize = 220;
- *blockSize = 1024;
-#elif defined(__NVPTX__)
- // TODO: values taken from NVIDIA H100 80GB HBM3
- *minGridSize = 264;
- *blockSize = 1024;
-#endif
- return Success;
-}
-
-///
-
#if defined(__AMDGPU__) || defined(__NVPTX__)
#include <gpuintrin.h>
diff --git a/offload/languages/include/kernel/UndefineLanguageNames.inc b/offload/languages/include/kernel/UndefineLanguageNames.inc
index ab5f257cdfaea..c1cc75719a18c 100644
--- a/offload/languages/include/kernel/UndefineLanguageNames.inc
+++ b/offload/languages/include/kernel/UndefineLanguageNames.inc
@@ -20,10 +20,6 @@
#undef MemcpyDeviceToHost
#undef MemcpyDeviceToDevice
#undef MemcpyDefault
-#undef GetLastError
-#undef PeekAtLastError
-#undef GetErrorName
-#undef GetErrorString
#undef GetDevice
#undef GetDeviceCount
#undef SetDevice
@@ -35,14 +31,8 @@
#undef HostAllocWriteCombined
#undef MallocHost
#undef FreeHost
-#undef DriverGetVersion
#undef GetDeviceProperties
-#undef OccupancyMaxPotentialBlockSizeVariableSMem
#undef Stream_t
#undef StreamCreate
-#undef StreamCreateWithFlags
#undef StreamDestroy
#undef StreamSynchronize
-#undef StreamCreateWithFlagsFlags
-#undef StreamDefault
-#undef StreamNonBlocking
diff --git a/offload/languages/kernel/src/LanguageRuntime.cpp b/offload/languages/kernel/src/LanguageRuntime.cpp
index b1b21f57e35f1..0859a7955dbe7 100644
--- a/offload/languages/kernel/src/LanguageRuntime.cpp
+++ b/offload/languages/kernel/src/LanguageRuntime.cpp
@@ -100,26 +100,6 @@ Error_t DeviceSynchronize() {
return convertResult(Result);
}
-Error_t GetLastError() {
- // TODO:
- return Success;
-}
-
-Error_t PeekAtLastError() {
- // TODO:
- return Success;
-}
-
-const char *GetErrorName(Error_t Error) {
- // TODO:
- return "";
-}
-
-const char *GetErrorString(Error_t Error) {
- // TODO:
- return "";
-}
-
Error_t GetDevice(int *DeviceNo) {
ol_device_handle_t Device = RuntimeState::getDevice(DeviceNo);
if (!Device)
@@ -138,7 +118,7 @@ Error_t SetDevice(int DeviceNo) {
return ErrorInvalidValue;
assert(Device == ThreadState::getDefaultDevice() &&
"Set Device is not Default Device");
- return Device ? Success : ErrorInvalidValue;
+ return Success;
}
Error_t HostAlloc(void **Ptr, size_t Size, unsigned int Flags) {
@@ -156,12 +136,6 @@ Error_t FreeHost(void *Ptr) {
return convertResult(Result);
}
-Error_t DriverGetVersion(int *Version) {
- // TODO:
- *Version = 42;
- return Success;
-}
-
Error_t GetDeviceProperties(DeviceProp_t *DeviceProp, int DeviceNo) {
ol_device_handle_t Device = ThreadState::getDefaultDevice();
size_t NameSize = 0;
@@ -192,16 +166,6 @@ Error_t StreamCreate(Stream_t *Stream) {
return Success;
}
-Error_t StreamCreateWithFlags(Stream_t *Stream, unsigned int Flags) {
- if (Flags == StreamCreateWithFlagsFlags::StreamDefault)
- // FIXME: [h15] offload streams are non-blocking by default
- return StreamCreate(Stream);
- if (Flags == StreamCreateWithFlagsFlags::StreamNonBlocking) {
- return StreamCreate(Stream);
- }
- return ErrorInvalidValue;
-}
-
Error_t StreamDestroy(Stream_t Stream) {
ol_queue_handle_t Queue;
Error_t Err = getQueueFromStream(Stream, &Queue);
>From c3627d8bbe5b1ad4238de9d373eca660c5bd220b Mon Sep 17 00:00:00 2001
From: Sophia Herrmann <herrmann15 at llnl.gov>
Date: Thu, 30 Jul 2026 14:43:27 -0700
Subject: [PATCH 4/7] remove legacy fatbin registration wrapping
---
.../kernel/include/LanguageAliases.inc | 22 +-
.../kernel/include/LanguageRegistration.h | 20 --
offload/languages/kernel/include/State.h | 9 +-
.../languages/kernel/src/LanguageLaunch.cpp | 24 +--
.../kernel/src/LanguageRegistration.cpp | 195 +-----------------
.../languages/kernel/src/LanguageRuntime.cpp | 2 +
offload/languages/kernel/src/State.cpp | 14 +-
7 files changed, 35 insertions(+), 251 deletions(-)
diff --git a/offload/languages/kernel/include/LanguageAliases.inc b/offload/languages/kernel/include/LanguageAliases.inc
index 67c188bc16269..551f64a5b8fc8 100644
--- a/offload/languages/kernel/include/LanguageAliases.inc
+++ b/offload/languages/kernel/include/LanguageAliases.inc
@@ -24,10 +24,11 @@
#define LA_IMPL1(PREFIX, L, NAME) LA_IMPL2(PREFIX, L, NAME)
#define LANGUAGE_NAME(PREFIX, NAME) LA_IMPL1(PREFIX, LANGUAGE, NAME)
-extern "C" void LANGUAGE_NAME(__, RegisterFunction)(
- const char *Binary, const char *KernelID, char *KernelName,
- const char *KernelName1, int ThreadLimit, uint3 *Tid, uint3 *Bid,
- dim3 *BlockDim, dim3 *GridDim, int *WSize) {
+extern "C" void
+LANGUAGE_NAME(__, RegisterFunction)(const char *Binary, const char *KernelID,
+ char *KernelName, const char *KernelName1,
+ int ThreadLimit, uint3 *Tid, uint3 *Bid,
+ dim3 *BlockDim, dim3 *GridDim, int *WSize) {
__llvmRegisterFunction(Binary, KernelID, KernelName, KernelName1, ThreadLimit,
Tid, Bid, BlockDim, GridDim, WSize);
}
@@ -41,9 +42,11 @@ extern "C" void LANGUAGE_NAME(__, RegisterVar)(void **Data, char *HostVar,
Constant, Global);
}
-extern "C" void LANGUAGE_NAME(__, RegisterManagedVar)(
- void **Data, char *HostVar, char *DeviceAddress, const char *DeviceName,
- size_t Size, unsigned Align) {
+extern "C" void LANGUAGE_NAME(__,
+ RegisterManagedVar)(void **Data, char *HostVar,
+ char *DeviceAddress,
+ const char *DeviceName,
+ size_t Size, unsigned Align) {
__llvmRegisterManagedVar(Data, HostVar, DeviceAddress, DeviceName, Size,
Align);
}
@@ -60,8 +63,9 @@ extern "C" void LANGUAGE_NAME(__, RegisterTexture)(
__llvmRegisterTexture(Data, TexRef, DevPtr, Name, Dim, Norm, Ext);
}
-extern "C" unsigned LANGUAGE_NAME(__, PopCallConfiguration)(
- dim3 *GridSize, dim3 *BlockSize, size_t *SharedMemory, void **Stream) {
+extern "C" unsigned
+LANGUAGE_NAME(__, PopCallConfiguration)(dim3 *GridSize, dim3 *BlockSize,
+ size_t *SharedMemory, void **Stream) {
return __llvmPopCallConfiguration(GridSize, BlockSize, SharedMemory, Stream);
}
diff --git a/offload/languages/kernel/include/LanguageRegistration.h b/offload/languages/kernel/include/LanguageRegistration.h
index 248aa74f600bd..7570ce76de546 100644
--- a/offload/languages/kernel/include/LanguageRegistration.h
+++ b/offload/languages/kernel/include/LanguageRegistration.h
@@ -15,22 +15,6 @@
#include <cstdint>
#include <iterator>
-#define HIP_FATBIN_MAGIC_STR "__CLANG_OFFLOAD_BUNDLE__"
-constexpr auto HIP_FATBIN_MAGIC_STR_LEN = sizeof(HIP_FATBIN_MAGIC_STR) - 1;
-
-namespace {
-struct FatbinWrapperTy {
- int Magic;
- int Version;
- const char *Data;
- const char *DataEnd;
-};
-} // namespace
-
-static void readTUFatbin(const char *Binary, const FatbinWrapperTy *FW);
-
-static void readHIPFatbinEntries(const char *Binary, const char *HIPFatbinPtr);
-
/// Hidden, but exported, Registration API
///{
extern "C" {
@@ -39,10 +23,6 @@ void __llvmRegisterFunction(const char *Binary, const char *KernelID,
char *KernelName, const char *KernelName1, int,
uint3 *, uint3 *, dim3 *, dim3 *, int *);
-const char *__llvmRegisterFatBinary(const char *Binary);
-
-void __llvmUnregisterFatBinary(void *Handle);
-
void __llvmRegisterVar(void **, char *, char *, const char *, int, int, int,
int);
diff --git a/offload/languages/kernel/include/State.h b/offload/languages/kernel/include/State.h
index 12d1abafbf917..3c0eb79811566 100644
--- a/offload/languages/kernel/include/State.h
+++ b/offload/languages/kernel/include/State.h
@@ -34,9 +34,7 @@ namespace offload {
/// Opaque host-side key used to identify a registered kernel.
///
-/// This is the address emitted in the offload entry table for the kernel,
-/// not a device number or liboffload handle. The runtime maps it to the
-/// loaded device symbol during registration and uses it again during launch.
+/// This is the address emitted in the offload entry table for the kernel
using KernelIDTy = const void *;
/// Per-thread state used by the language runtime entry points.
@@ -46,7 +44,7 @@ using KernelIDTy = const void *;
struct ThreadStateTy {
~ThreadStateTy();
- /// Return the default queue for the current stream mode.
+ /// Return the default queue for the current host thread
static ol_queue_handle_t getDefaultQueue();
/// Return the thread-local default device, or the first discovered device.
@@ -86,7 +84,8 @@ struct StateTy {
/// Return the number of non-host devices available to kernel languages.
static int getDeviceCount();
- /// Return the thread-local default device and write its number to \p DeviceNo.
+ /// Return the thread-local default device and write its number to \p
+ /// DeviceNo.
static ol_device_handle_t getDevice(int *DeviceNo);
/// Set the thread-local default device by device number.
diff --git a/offload/languages/kernel/src/LanguageLaunch.cpp b/offload/languages/kernel/src/LanguageLaunch.cpp
index 20779f89f969a..71c2b275beb35 100644
--- a/offload/languages/kernel/src/LanguageLaunch.cpp
+++ b/offload/languages/kernel/src/LanguageLaunch.cpp
@@ -17,25 +17,25 @@ using ThreadState = llvm::offload::ThreadStateTy;
extern "C" {
/// Push call configuration for kernel launch
-unsigned __llvmPushCallConfiguration(dim3 __grid_size, dim3 __block_size,
- size_t __shared_memory, void *__stream) {
+unsigned __llvmPushCallConfiguration(dim3 GridSize, dim3 BlockSize,
+ size_t SharedMemory, void *Stream) {
CallConfigurationTy &CC = ThreadState::getCallConfiguration();
- CC.GridSize = __grid_size;
- CC.BlockSize = __block_size;
- CC.SharedMemory = __shared_memory;
- CC.Stream = __stream;
+ CC.GridSize = GridSize;
+ CC.BlockSize = BlockSize;
+ CC.SharedMemory = SharedMemory;
+ CC.Stream = Stream;
return 0;
}
/// Pop call configuration for kernel launch
-unsigned __llvmPopCallConfiguration(dim3 *__grid_size, dim3 *__block_size,
- size_t *__shared_memory, void **__stream) {
+unsigned __llvmPopCallConfiguration(dim3 *GridSize, dim3 *BlockSize,
+ size_t *SharedMemory, void **Stream) {
CallConfigurationTy &CC = ThreadState::getCallConfiguration();
- *__grid_size = CC.GridSize;
- *__block_size = CC.BlockSize;
- *__shared_memory = CC.SharedMemory;
- *__stream = CC.Stream;
+ *GridSize = CC.GridSize;
+ *BlockSize = CC.BlockSize;
+ *SharedMemory = CC.SharedMemory;
+ *Stream = CC.Stream;
return 0;
}
diff --git a/offload/languages/kernel/src/LanguageRegistration.cpp b/offload/languages/kernel/src/LanguageRegistration.cpp
index 231eaa691ff0b..2219e16ea0365 100644
--- a/offload/languages/kernel/src/LanguageRegistration.cpp
+++ b/offload/languages/kernel/src/LanguageRegistration.cpp
@@ -20,172 +20,6 @@
using RuntimeState = llvm::offload::StateTy;
using ThreadState = llvm::offload::ThreadStateTy;
-typedef struct __attribute__((__packed__)) {
- uint32_t Magic;
- uint16_t Version;
- uint16_t HeaderSize;
- uint64_t FatSize;
-} CudaFatbinHeader;
-
-// Inspired by
-// https://github.com/n-eiling/cuda-fatbin-decompression/blob/master/fatbin-decompress.h
-typedef struct __attribute__((__packed__)) {
- uint16_t Kind;
- uint16_t Unknown1;
- uint32_t HeaderSize;
- uint64_t Size;
- uint32_t CompressedSize;
- uint32_t Unknown2;
- uint16_t Minor;
- uint16_t Major;
- uint32_t Arch;
- uint32_t ObjNameOffset;
- uint32_t ObjNameLen;
- uint64_t Flags;
- uint64_t Zero;
- uint64_t DecompressedSize;
-} CudaFatbinTextHeader;
-
-// HIP uses this format:
-// https://clang.llvm.org/docs/ClangOffloadBundler.html#bundled-binary-file-layout
-typedef struct __attribute__((__packed__)) {
- char Magic[24];
- uint64_t NumBundles;
-} HipFatbinHeader;
-
-typedef struct __attribute__((__packed__)) {
- uint64_t BundleOffset;
- uint64_t BundleSize;
- uint64_t IdLength;
- char IdString[];
-} HipFatbinBundleEntry;
-
-static void readTUFatbin(const char *Binary, const FatbinWrapperTy *FW) {
- ol_device_handle_t Device = language_registration::getDefaultDevice();
-
- const CudaFatbinHeader *Header =
- reinterpret_cast<const CudaFatbinHeader *>(FW->Data);
- size_t HeaderSize = static_cast<size_t>(Header->HeaderSize); // Usually 16
- size_t FatbinSize = static_cast<size_t>(Header->FatSize);
-
- const void *ProgramData = nullptr;
- size_t ProgramSize = 0;
- uint32_t ProgramArch = 0;
-
- const char *ReadPosition = FW->Data + HeaderSize;
- while (ReadPosition < (FW->Data + FatbinSize)) {
- const CudaFatbinTextHeader *TextHeader =
- reinterpret_cast<const CudaFatbinTextHeader *>(ReadPosition);
- size_t TextHeaderSize =
- static_cast<size_t>(TextHeader->HeaderSize); // Usually 64
- size_t CubinSize = static_cast<size_t>(TextHeader->Size);
- const void *CubinData =
- static_cast<const char *>(ReadPosition + TextHeaderSize);
-
- uint32_t Arch = TextHeader->Arch;
- bool IsCompatible = false;
- olIsValidBinary(Device, CubinData, CubinSize, &IsCompatible);
- if (!IsCompatible) {
- fprintf(stderr, "Device is not compatible with image.");
- abort();
- }
-
- if (Arch > ProgramArch) {
- ProgramData = CubinData;
- ProgramSize = CubinSize;
- ProgramArch = Arch;
- }
-
- ReadPosition += TextHeaderSize + CubinSize;
- }
-
- if (ProgramData == nullptr) {
- fprintf(stderr, "Failed to find compatible binary\n");
- abort();
- }
-
- ol_program_handle_t Program = nullptr;
-
- ol_result_t Result =
- olCreateProgram(Device, ProgramData, ProgramSize, &Program);
-
- if (Result && Result->Code) {
- fprintf(stderr, "Failed to register device code (%i): %s\n", Result->Code,
- Result->Details);
- abort();
- }
-
- language_registration::registerProgram(Binary, Program);
-}
-
-static void readHIPFatbinEntries(const char *Binary, const char *HIPFatbinPtr) {
- ol_device_handle_t Device = language_registration::getDefaultDevice();
-
- const char *CurrentReadPosition = HIPFatbinPtr;
-
- const HipFatbinHeader *Header =
- reinterpret_cast<const HipFatbinHeader *>(CurrentReadPosition);
- CurrentReadPosition += sizeof(HipFatbinHeader);
-
- uint64_t NumBundles = Header->NumBundles;
-
- const void *ProgramData = nullptr;
- size_t ProgramSize = 0;
- uint64_t ProgramIdLength = 0;
- const char *ProgramIdString = nullptr;
-
- for (uint64_t BundleId = 0; BundleId < NumBundles; ++BundleId) {
- const HipFatbinBundleEntry *BundleEntry =
- reinterpret_cast<const HipFatbinBundleEntry *>(CurrentReadPosition);
-
- uint64_t BundleOffset = BundleEntry->BundleOffset;
- uint64_t BundleSize = BundleEntry->BundleSize;
- const char *BundleIdString = BundleEntry->IdString;
- uint64_t BundleIdLength = BundleEntry->IdLength;
-
- // Advance by the size of the entry including the ID string
- CurrentReadPosition += sizeof(HipFatbinBundleEntry) + BundleIdLength;
-
- if (!BundleSize) {
- continue;
- }
-
- bool IsCompatible = false;
- olIsValidBinary(Device, HIPFatbinPtr + BundleOffset, BundleSize,
- &IsCompatible);
-
- if (!IsCompatible) {
- fprintf(stderr, "Device is not compatible with image.");
- abort();
- }
-
- llvm::StringRef CurrentBundleId(ProgramIdString, ProgramIdLength);
- llvm::StringRef NewBundleId(BundleIdString, BundleIdLength);
- if (NewBundleId.compare(CurrentBundleId) > 0) {
- ProgramData = HIPFatbinPtr + BundleOffset;
- ProgramSize = BundleSize;
- ProgramIdLength = BundleIdLength;
- ProgramIdString = BundleIdString;
- }
- }
-
- if (ProgramData == nullptr) {
- fprintf(stderr, "Failed to find compatible binary\n");
- abort();
- }
-
- ol_program_handle_t Program = nullptr;
- ol_result_t Result =
- olCreateProgram(Device, ProgramData, ProgramSize, &Program);
- if (Result && Result->Code) {
- fprintf(stderr, "Failed to register device code (%i): %s\n", Result->Code,
- Result->Details);
- abort();
- }
-
- language_registration::registerProgram(Binary, Program);
-}
-
/// Hidden, but exported, Registration API
///{
extern "C" {
@@ -197,37 +31,10 @@ void __llvmRegisterFunction(const char *Binary, const char *KernelID,
ol_program_handle_t Program = RuntimeState::getProgram(Binary);
ol_result_t Result = olGetSymbol(
Program, KernelName, ol_symbol_kind_t::OL_SYMBOL_KIND_KERNEL, &Kernel);
- if (Result && Result->Code) {
- fprintf(stderr, "Failed to register kernel (%i): %s\n", Result->Code,
- Result->Details);
- abort();
- }
-
+ CHECK_FATAL(Result, "Failed to get kernel symbol for " << KernelName);
RuntimeState::registerKernel(KernelID, Kernel);
}
-const char *__llvmRegisterFatBinary(const char *Binary) {
- const auto *FW = reinterpret_cast<const FatbinWrapperTy *>(Binary);
- if (FW->Magic == 0x466243b1) {
- readTUFatbin(Binary, FW);
- } else if (FW->Magic == 0x48495046) {
- if (!memcmp(FW->Data, HIP_FATBIN_MAGIC_STR, HIP_FATBIN_MAGIC_STR_LEN))
- readHIPFatbinEntries(Binary, FW->Data);
- else
- readTUFatbin(Binary, FW);
- } else {
- fprintf(stderr, "Unknown fatbin format");
- }
-
- return Binary;
-}
-
-void __llvmUnregisterFatBinary(void *Handle) {
- if (ol_program_handle_t Program =
- language_registration::unregisterProgram(Handle))
- olDestroyProgram(Program);
-}
-
void __llvmRegisterVar(void **, char *, char *, const char *, int, int, int,
int) {
llvm::errs() << "RegisterVar is not implemented!" << "\n";
diff --git a/offload/languages/kernel/src/LanguageRuntime.cpp b/offload/languages/kernel/src/LanguageRuntime.cpp
index 0859a7955dbe7..058eecf2c97ee 100644
--- a/offload/languages/kernel/src/LanguageRuntime.cpp
+++ b/offload/languages/kernel/src/LanguageRuntime.cpp
@@ -11,8 +11,10 @@
#endif
// Rename the generic runtime API before declaring or defining language symbols.
+// clang-format off
#include "DefineLanguageNames.inc"
#include "LanguageRuntime.h"
+// clang-format on
#include "State.h"
#include "Types.h"
diff --git a/offload/languages/kernel/src/State.cpp b/offload/languages/kernel/src/State.cpp
index 3dbd1f436ff24..9607220ebf3cc 100644
--- a/offload/languages/kernel/src/State.cpp
+++ b/offload/languages/kernel/src/State.cpp
@@ -14,27 +14,18 @@
#include "llvm/ADT/DenseMap.h"
#include "llvm/ADT/SmallPtrSet.h"
#include "llvm/ADT/SmallVector.h"
-#include "llvm/Support/raw_ostream.h"
#include <atomic>
#include <cassert>
#include <cstdio>
#include <mutex>
-#define CHECK_FATAL(Result, ...) \
- if (Result && Result->Code) { \
- llvm::errs() << __VA_ARGS__ << '\n'; \
- abort(); \
- }
-
using namespace llvm;
using namespace offload;
// Weak so another runtime object can override the default stream mode.
__attribute__((weak)) uint32_t PerThreadQueue = 0;
-thread_local CallConfigurationTy CC = {};
-
// Process-wide singleton and thread-state registry.
static std::mutex StateLock;
static std::atomic<StateTy *> StatePtr = nullptr;
@@ -46,6 +37,8 @@ using ThreadStatesTy = SmallVector<ThreadStateTy *, 64>;
static ThreadStatesTy *ThreadStatesPtr = nullptr;
static void deleteThreadState() {
+ // Detach the registry before deletion because deleteThreadState may be called
+ // more than once via atexit and StateTy teardown.
std::lock_guard<std::mutex> LG(ThreadStatesLock);
ThreadStatesTy *ThreadStates = ThreadStatesPtr;
ThreadStatesPtr = nullptr;
@@ -86,10 +79,9 @@ ThreadStateTy::ThreadStateTy() {
ThreadStateTy::~ThreadStateTy() { destroyQueue(DefaultQueue); }
ThreadStateTy &ThreadStateTy::get() {
- auto *TS = ThreadState;
+ auto *&TS = ThreadState;
if (!TS) {
TS = new ThreadStateTy();
- ThreadState = TS;
std::lock_guard<std::mutex> LG(ThreadStatesLock);
if (!ThreadStatesPtr)
ThreadStatesPtr = new ThreadStatesTy;
>From 5590469b3e137ea42a809487223eb464a9bc6cb0 Mon Sep 17 00:00:00 2001
From: Sophia Herrmann <herrmann15 at llnl.gov>
Date: Mon, 27 Jul 2026 15:48:11 -0700
Subject: [PATCH 5/7] decouple front/back end
Co-authored-by: Johannes Doerfert <jdoerfert.llvm at gmail.com>
Co-authored-by: Jonas Greifenhain <cadivus at daverkomp.de>
remove irrelevant artifacts
---
clang/include/clang/Driver/CommonArgs.h | 6 ++
clang/lib/CodeGen/CGCUDANV.cpp | 67 +++++++++-------
clang/lib/Driver/Driver.cpp | 77 +++++++++++--------
clang/lib/Driver/ToolChains/AMDGPU.cpp | 33 +++++++-
clang/lib/Driver/ToolChains/Clang.cpp | 57 ++++++++++----
clang/lib/Driver/ToolChains/CommonArgs.cpp | 19 +++--
clang/lib/Driver/ToolChains/Cuda.cpp | 70 +++++++++++++----
clang/lib/Driver/ToolChains/Gnu.cpp | 1 +
clang/lib/Driver/ToolChains/Linux.cpp | 4 +-
clang/lib/Headers/__clang_gpu_builtin_vars.h | 19 +++++
clang/test/CodeGenCUDA/Inputs/cuda.h | 2 +-
clang/test/CodeGenCUDA/offload_via_llvm.cu | 60 +++++++--------
clang/test/Driver/cuda-via-liboffload.cu | 15 ++--
.../linker-wrapper-image.c | 38 ++++-----
.../ClangLinkerWrapper.cpp | 13 +++-
.../Frontend/Offloading/OffloadWrapper.cpp | 4 +-
16 files changed, 326 insertions(+), 159 deletions(-)
diff --git a/clang/include/clang/Driver/CommonArgs.h b/clang/include/clang/Driver/CommonArgs.h
index 01e358a1d0717..3b53df94fc79f 100644
--- a/clang/include/clang/Driver/CommonArgs.h
+++ b/clang/include/clang/Driver/CommonArgs.h
@@ -148,6 +148,12 @@ void addArchSpecificRPath(const ToolChain &TC, const llvm::opt::ArgList &Args,
void addOpenMPRuntimeLibraryPath(const ToolChain &TC,
const llvm::opt::ArgList &Args,
llvm::opt::ArgStringList &CmdArgs);
+
+bool addLLVMOffloadingRuntime(const Compilation &C,
+ llvm::opt::ArgStringList &CmdArgs,
+ const ToolChain &TC,
+ const llvm::opt::ArgList &Args);
+
/// Returns true, if an OpenMP runtime has been added.
bool addOpenMPRuntime(const Compilation &C, llvm::opt::ArgStringList &CmdArgs,
const ToolChain &TC, const llvm::opt::ArgList &Args,
diff --git a/clang/lib/CodeGen/CGCUDANV.cpp b/clang/lib/CodeGen/CGCUDANV.cpp
index 1e688d29d15a5..29f358c4015ce 100644
--- a/clang/lib/CodeGen/CGCUDANV.cpp
+++ b/clang/lib/CodeGen/CGCUDANV.cpp
@@ -254,9 +254,7 @@ CGNVCUDARuntime::CGNVCUDARuntime(CodeGenModule &CGM)
VoidTy = CGM.VoidTy;
PtrTy = CGM.DefaultPtrTy;
- if (CGM.getLangOpts().OffloadViaLLVM)
- Prefix = "llvm";
- else if (CGM.getLangOpts().HIP)
+ if (CGM.getLangOpts().HIP)
Prefix = "hip";
else
Prefix = "cuda";
@@ -345,41 +343,52 @@ void CGNVCUDARuntime::emitDeviceStub(CodeGenFunction &CGF,
emitDeviceStubBodyLegacy(CGF, Args);
}
-/// Build the input as a sized array of pointers so that it can be launched by
-/// the offloading runtime.
+/// CUDA passes the arguments with a level of indirection. For example, a
+/// (void*, short, void*) is passed as {void **, short *, void **} to the launch
+/// function. For the LLVM/Offload launch we include the number of arguments and
+/// their size. Thus, we pass {{void **, short*, void **}, 3, {sizeof(void*),
+/// sizeof(short), sizeof(void*)}}.
Address CGNVCUDARuntime::prepareKernelArgsLLVMOffload(CodeGenFunction &CGF,
FunctionArgList &Args) {
- SmallVector<llvm::Type *> ArgTypes, KernelLaunchParamsTypes;
- for (auto &Arg : Args)
- ArgTypes.push_back(CGF.ConvertTypeForMem(Arg->getType()));
- llvm::StructType *KernelArgsTy = llvm::StructType::create(ArgTypes);
- llvm::Type *KernelArgsPtrsTy = llvm::ArrayType::get(PtrTy, Args.size());
-
- auto *Int32Ty = CGF.Builder.getInt32Ty();
- KernelLaunchParamsTypes.push_back(Int32Ty);
+ SmallVector<llvm::Type *> KernelLaunchParamsTypes;
+
+ auto *Int64Ty = CGF.Builder.getInt64Ty();
+ KernelLaunchParamsTypes.push_back(PtrTy);
+ KernelLaunchParamsTypes.push_back(Int64Ty);
KernelLaunchParamsTypes.push_back(PtrTy);
llvm::StructType *KernelLaunchParamsTy =
llvm::StructType::create(KernelLaunchParamsTypes);
- Address KernelArgs = CGF.CreateTempAllocaWithoutCast(
- KernelArgsTy, CharUnits::fromQuantity(16), "kernel_args");
- Address KernelArgsPtrs = CGF.CreateTempAllocaWithoutCast(
- KernelArgsPtrsTy, CharUnits::fromQuantity(16), "kernel_args_ptrs");
Address KernelLaunchParams = CGF.CreateTempAllocaWithoutCast(
KernelLaunchParamsTy, CharUnits::fromQuantity(16),
"kernel_launch_params");
+ Address KernelArgs = CGF.CreateTempAlloca(
+ PtrTy, LangAS::Default, CharUnits::fromQuantity(16), "kernel_args",
+ llvm::ConstantInt::get(SizeTy, std::max<size_t>(1, Args.size())));
+ Address KernelArgSizes = CGF.CreateTempAlloca(
+ SizeTy, LangAS::Default, CharUnits::fromQuantity(16), "kernel_arg_sizes",
+ llvm::ConstantInt::get(SizeTy, std::max<size_t>(1, Args.size())));
- CGF.Builder.CreateStore(llvm::ConstantInt::get(Int32Ty, Args.size()),
+ CGF.Builder.CreateStore(KernelArgs.emitRawPointer(CGF),
CGF.Builder.CreateStructGEP(KernelLaunchParams, 0));
- CGF.Builder.CreateStore(KernelArgsPtrs.emitRawPointer(CGF),
+ CGF.Builder.CreateStore(llvm::ConstantInt::get(Int64Ty, Args.size()),
CGF.Builder.CreateStructGEP(KernelLaunchParams, 1));
+ CGF.Builder.CreateStore(KernelArgSizes.emitRawPointer(CGF),
+ CGF.Builder.CreateStructGEP(KernelLaunchParams, 2));
for (unsigned i = 0; i < Args.size(); ++i) {
- auto *ArgVal = CGF.Builder.CreateLoad(CGF.GetAddrOfLocalVar(Args[i]));
- Address ArgAddr = CGF.Builder.CreateStructGEP(KernelArgs, i);
- CGF.Builder.CreateStore(ArgVal, ArgAddr);
- CGF.Builder.CreateStore(ArgAddr.emitRawPointer(CGF),
- CGF.Builder.CreateConstArrayGEP(KernelArgsPtrs, i));
+ llvm::Value *VarPtr = CGF.GetAddrOfLocalVar(Args[i]).emitRawPointer(CGF);
+ llvm::Value *VoidVarPtr = CGF.Builder.CreatePointerCast(VarPtr, PtrTy);
+ CGF.Builder.CreateDefaultAlignedStore(
+ VoidVarPtr, CGF.Builder.CreateConstGEP1_32(
+ PtrTy, KernelArgs.emitRawPointer(CGF), i));
+
+ auto ArgSize = CGM.getDataLayout().getTypeAllocSize(
+ CGM.getTypes().ConvertType(Args[i]->getType()));
+ CGF.Builder.CreateDefaultAlignedStore(
+ llvm::ConstantInt::get(SizeTy, ArgSize),
+ CGF.Builder.CreateConstGEP1_32(SizeTy,
+ KernelArgSizes.emitRawPointer(CGF), i));
}
return KernelLaunchParams;
@@ -408,8 +417,9 @@ Address CGNVCUDARuntime::prepareKernelArgs(CodeGenFunction &CGF,
// array and kernels are launched using cudaLaunchKernel().
void CGNVCUDARuntime::emitDeviceStubBodyNew(CodeGenFunction &CGF,
FunctionArgList &Args) {
+ bool UsesLLVMOffloading = CGF.getLangOpts().OffloadViaLLVM;
// Build the shadow stack entry at the very start of the function.
- Address KernelArgs = CGF.getLangOpts().OffloadViaLLVM
+ Address KernelArgs = UsesLLVMOffloading
? prepareKernelArgsLLVMOffload(CGF, Args)
: prepareKernelArgs(CGF, Args);
@@ -435,7 +445,9 @@ void CGNVCUDARuntime::emitDeviceStubBodyNew(CodeGenFunction &CGF,
else if (CGF.getLangOpts().CUDA)
KernelLaunchAPI = KernelLaunchAPI + "_ptsz";
}
- auto LaunchKernelName = addPrefixToName(KernelLaunchAPI);
+ /// Use __llvmLaunchKernel for LLVMOffload.
+ auto LaunchKernelName = UsesLLVMOffloading ? "__llvm" + KernelLaunchAPI
+ : addPrefixToName(KernelLaunchAPI);
const IdentifierInfo &cudaLaunchKernelII =
CGM.getContext().Idents.get(LaunchKernelName);
FunctionDecl *cudaLaunchKernelFD = nullptr;
@@ -1282,9 +1294,6 @@ void CGNVCUDARuntime::createOffloadingEntries() {
llvm::object::OffloadKind Kind = CGM.getLangOpts().HIP
? llvm::object::OffloadKind::OFK_HIP
: llvm::object::OffloadKind::OFK_Cuda;
- // For now, just spoof this as OpenMP because that's the runtime it uses.
- if (CGM.getLangOpts().OffloadViaLLVM)
- Kind = llvm::object::OffloadKind::OFK_OpenMP;
llvm::Module &M = CGM.getModule();
for (KernelInfo &I : EmittedKernels)
diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp
index 3b5d0c0aad981..a7241bd878dcc 100644
--- a/clang/lib/Driver/Driver.cpp
+++ b/clang/lib/Driver/Driver.cpp
@@ -905,10 +905,15 @@ getSystemOffloadArchs(Compilation &C, Action::OffloadKind Kind) {
if (llvm::ErrorOr<std::string> Executable =
llvm::sys::findProgramByName(Program, {C.getDriver().Dir})) {
llvm::SmallVector<StringRef> Args{*Executable};
- if (Kind == Action::OFK_HIP)
- Args.push_back("--only=amdgpu");
- else if (Kind == Action::OFK_Cuda)
- Args.push_back("--only=nvptx");
+ bool UsesLLVMOffloading =
+ C.getArgs().hasFlag(options::OPT_foffload_via_llvm,
+ options::OPT_fno_offload_via_llvm, false);
+ if (!UsesLLVMOffloading) {
+ if (Kind == Action::OFK_HIP)
+ Args.push_back("--only=amdgpu");
+ else if (Kind == Action::OFK_Cuda)
+ Args.push_back("--only=nvptx");
+ }
auto StdoutOrErr = C.getDriver().executeProgram(Args);
if (!StdoutOrErr) {
@@ -965,15 +970,19 @@ static TripleSet inferOffloadToolchains(Compilation &C,
ID = StringToOffloadArch(
getProcessorFromTargetID(llvm::Triple("amdgcn-amd-amdhsa"), Arch));
- if (Kind == Action::OFK_HIP && !ID.isAMDGPU() && !ID.isSPIRV()) {
- C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch)
- << "HIP" << Arch;
- return {};
- }
- if (Kind == Action::OFK_Cuda && !ID.isNVPTX()) {
- C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch)
- << "CUDA" << Arch;
- return {};
+ bool UsesLLVMOffloading =
+ C.getArgs().hasFlag(options::OPT_foffload_via_llvm,
+ options::OPT_fno_offload_via_llvm, false);
+ if (!UsesLLVMOffloading) {
+ if (Kind == Action::OFK_HIP && !ID.isAMDGPU() && !ID.isSPIRV()) {
+ C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch)
+ << "HIP" << Arch;
+ return {};
+ }
+ if (Kind == Action::OFK_Cuda && !ID.isNVPTX()) {
+ C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch)
+ << "CUDA" << Arch;
+ return {};
}
if (Kind == Action::OFK_OpenMP && (ID.isUnknown() || ID.isUnused())) {
C.getDriver().Diag(clang::diag::err_drv_failed_to_deduce_target_from_arch)
@@ -988,6 +997,8 @@ static TripleSet inferOffloadToolchains(Compilation &C,
llvm::Triple Triple =
OffloadArchToTriple(C.getDefaultToolChain().getTriple(), ID);
+ if (UsesLLVMOffloading)
+ Triple.setEnvironment(llvm::Triple::LLVM);
// Make a new argument that dispatches this argument to the appropriate
// toolchain. This is required when we infer it and create potentially
@@ -1031,32 +1042,30 @@ static TripleSet inferOffloadToolchains(Compilation &C,
void Driver::CreateOffloadingDeviceToolChains(Compilation &C,
InputList &Inputs) {
- bool UseLLVMOffload = C.getInputArgs().hasArg(
- options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
bool IsCuda =
- llvm::any_of(Inputs,
- [](std::pair<types::ID, const llvm::opt::Arg *> &I) {
- return types::isCuda(I.first);
- }) &&
- !UseLLVMOffload;
+ llvm::any_of(Inputs, [](std::pair<types::ID, const llvm::opt::Arg *> &I) {
+ return types::isCuda(I.first);
+ });
bool IsHIP =
(llvm::any_of(Inputs,
[](std::pair<types::ID, const llvm::opt::Arg *> &I) {
return types::isHIP(I.first);
}) ||
C.getInputArgs().hasArg(options::OPT_hip_link) ||
- C.getInputArgs().hasArg(options::OPT_hipstdpar)) &&
- !UseLLVMOffload;
+ C.getInputArgs().hasArg(options::OPT_hipstdpar));
bool IsSYCL = C.getInputArgs().hasFlag(options::OPT_fsycl,
options::OPT_fno_sycl, false);
bool IsOpenMPOffloading =
- UseLLVMOffload ||
(C.getInputArgs().hasFlag(options::OPT_fopenmp, options::OPT_fopenmp_EQ,
options::OPT_fno_openmp, false) &&
(C.getInputArgs().hasArg(options::OPT_offload_targets_EQ) ||
(C.getInputArgs().hasArg(options::OPT_offload_arch_EQ) &&
!(IsCuda || IsHIP))));
+ // We currently don't support any kind of mixed offloading.
+ if (IsOpenMPOffloading)
+ IsCuda = IsHIP = IsSYCL = false;
+
llvm::SmallSet<Action::OffloadKind, 4> Kinds;
const std::pair<bool, Action::OffloadKind> ActiveKinds[] = {
{IsCuda, Action::OFK_Cuda},
@@ -1142,7 +1151,7 @@ void Driver::CreateOffloadingDeviceToolChains(Compilation &C,
C.getDefaultToolChain().getTriple());
// Emit a warning if the detected CUDA version is too new.
- if (Kind == Action::OFK_Cuda) {
+ if (Kind == Action::OFK_Cuda && Target.getOS() == llvm::Triple::CUDA) {
auto &CudaInstallation =
static_cast<const toolchains::CudaToolChain &>(TC).CudaInstallation;
if (CudaInstallation.isValid())
@@ -5067,6 +5076,9 @@ Driver::BuildOffloadingActions(Compilation &C, llvm::opt::DerivedArgList &Args,
getFinalPhase(Args) == phases::Preprocess))
return HostAction;
+ bool UsesLLVMOffloading = Args.hasArg(
+ options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
+
ActionList OffloadActions;
OffloadAction::DeviceDependences DDeps;
@@ -5087,7 +5099,6 @@ Driver::BuildOffloadingActions(Compilation &C, llvm::opt::DerivedArgList &Args,
types::ID InputType = Input.first;
const Arg *InputArg = Input.second;
- // The toolchain can be active for unsupported file types.
if ((Kind == Action::OFK_Cuda && !types::isCuda(InputType)) ||
(Kind == Action::OFK_HIP && !types::isHIP(InputType)))
continue;
@@ -5182,9 +5193,12 @@ Driver::BuildOffloadingActions(Compilation &C, llvm::opt::DerivedArgList &Args,
OffloadAction::DeviceDependences DDep;
DDep.add(*A, *TCAndArch->first, TCAndArch->second, Kind);
- // Compiling CUDA in non-RDC mode uses the PTX output if available.
+ // The legacy CUDA fatbinary path can include PTX alongside the cubin.
+ // The LLVM offload wrapper path feeds these images through a device
+ // linker first, and clang-nvlink-wrapper does not accept PTX as input.
for (Action *Input : A->getInputs())
- if (Kind == Action::OFK_Cuda && A->getType() == types::TY_Object &&
+ if (!UsesLLVMOffloading && Kind == Action::OFK_Cuda &&
+ A->getType() == types::TY_Object &&
!Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc,
false))
DDep.add(*Input, *TCAndArch->first, TCAndArch->second, Kind);
@@ -5212,7 +5226,7 @@ Driver::BuildOffloadingActions(Compilation &C, llvm::opt::DerivedArgList &Args,
return HostAction;
OffloadAction::DeviceDependences DDep;
- if (C.isOffloadingHostKind(Action::OFK_Cuda) &&
+ if (!UsesLLVMOffloading && C.isOffloadingHostKind(Action::OFK_Cuda) &&
(!Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc, false) ||
Args.hasArg(options::OPT_cuda_emit_nvcc_abi))) {
// If we are not in RDC-mode or are targeting the NVCC ABI we just emit the
@@ -5221,7 +5235,7 @@ Driver::BuildOffloadingActions(Compilation &C, llvm::opt::DerivedArgList &Args,
C.MakeAction<LinkJobAction>(OffloadActions, types::TY_CUDA_FATBIN);
DDep.add(*FatbinAction, *C.getSingleOffloadToolChain<Action::OFK_Cuda>(),
/*BA=*/{}, Action::OFK_Cuda);
- } else if (HIPNoRDC && offloadDeviceOnly()) {
+ } else if (!UsesLLVMOffloading && HIPNoRDC && offloadDeviceOnly()) {
// If we are in device-only non-RDC-mode we just emit the final HIP
// fatbinary for each translation unit, linking each input individually.
Action *FatbinAction =
@@ -5229,7 +5243,7 @@ Driver::BuildOffloadingActions(Compilation &C, llvm::opt::DerivedArgList &Args,
DDep.add(*FatbinAction,
*C.getOffloadToolChains<Action::OFK_HIP>().first->second,
/*BA=*/{}, Action::OFK_HIP);
- } else if (HIPNoRDC) {
+ } else if (!UsesLLVMOffloading && HIPNoRDC) {
// Host + device assembly: defer to clang-offload-bundler (see
// BuildActions).
if (HIPAsmBundleDeviceOut &&
@@ -7090,7 +7104,8 @@ const ToolChain &Driver::getOffloadToolChain(
// For AMDHSA offloading (HIP, OpenMP), use the unified AMDGPUToolChain
// This handles both amdgpu-amd-amdhsa and spirv64-amd-amdhsa
// FIXME: This should not key off language or OS.
- if (Kind == Action::OFK_HIP || Kind == Action::OFK_OpenMP)
+ if (Kind == Action::OFK_HIP || Kind == Action::OFK_OpenMP ||
+ Kind == Action::OFK_Cuda)
TC = std::make_unique<toolchains::AMDGPUToolChain>(*this, Target, Args,
HostTC.get(), Kind);
break;
diff --git a/clang/lib/Driver/ToolChains/AMDGPU.cpp b/clang/lib/Driver/ToolChains/AMDGPU.cpp
index 7bce060de0596..5c30417365b94 100644
--- a/clang/lib/Driver/ToolChains/AMDGPU.cpp
+++ b/clang/lib/Driver/ToolChains/AMDGPU.cpp
@@ -515,6 +515,24 @@ void RocmInstallationDetector::AddHIPIncludeArgs(const ArgList &DriverArgs,
!DriverArgs.hasArg(options::OPT_nohipwrapperinc);
bool HasHipStdPar = DriverArgs.hasArg(options::OPT_hipstdpar);
+ if (DriverArgs.hasFlag(options::OPT_foffload_via_llvm,
+ options::OPT_fno_offload_via_llvm, false)) {
+ if (DriverArgs.hasFlag(options::OPT_offload_inc,
+ options::OPT_no_offload_inc, true) &&
+ !DriverArgs.hasArg(options::OPT_nohipwrapperinc) &&
+ !DriverArgs.hasArg(options::OPT_nobuiltininc)) {
+ CC1Args.append({"-include", "__clang_gpu_device_functions.h"});
+
+ SmallString<128> HIPIncludePath(D.ResourceDir);
+ llvm::sys::path::append(HIPIncludePath, "..", "..", "..");
+ llvm::sys::path::append(HIPIncludePath, "include", "offload");
+ CC1Args.push_back("-internal-isystem");
+ CC1Args.push_back(DriverArgs.MakeArgString(HIPIncludePath));
+ CC1Args.append({"-include", "hip/hip_runtime.h"});
+ }
+ return;
+ }
+
if (!DriverArgs.hasArg(options::OPT_nobuiltininc)) {
// HIP header includes standard library wrapper headers under clang
// cuda_wrappers directory. Since these wrapper headers include_next
@@ -699,7 +717,11 @@ AMDGPUToolChain::AMDGPUToolChain(const Driver &D, const llvm::Triple &Triple,
: Generic_ELF(D, Triple, Args),
OptionsDefault(
{{options::OPT_O, "3"}, {options::OPT_cl_std_EQ, "CL1.2"}}),
- HostTC(HostTC_), UseHIPLinker(Kind == Action::OFK_HIP),
+ HostTC(HostTC_),
+ UseHIPLinker(Kind == Action::OFK_HIP ||
+ (Kind == Action::OFK_Cuda &&
+ Args.hasFlag(options::OPT_foffload_via_llvm,
+ options::OPT_fno_offload_via_llvm, false))),
ShouldLinkDeviceLibs(ShouldLinkDeviceLibs) {
loadMultilibsFromYAML(Args, D);
@@ -709,8 +731,10 @@ AMDGPUToolChain::AMDGPUToolChain(const Driver &D, const llvm::Triple &Triple,
// each tool invocation.
checkAMDGPUCodeObjectVersion(D, Args);
+ bool UsesLLVMOffloading = Args.hasFlag(
+ options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
if (Triple.getOS() == llvm::Triple::AMDHSA &&
- Triple.getEnvironment() != llvm::Triple::LLVM)
+ Triple.getEnvironment() != llvm::Triple::LLVM && !UsesLLVMOffloading)
RocmInstallation->detectDeviceLibrary();
if (HostTC)
@@ -889,7 +913,10 @@ bool AMDGPUToolChain::isWave64(const llvm::opt::ArgList &DriverArgs,
void AMDGPUToolChain::addClangTargetOptions(
const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args,
BoundArch BA, Action::OffloadKind DeviceOffloadingKind) const {
- if (DeviceOffloadingKind == Action::OFK_HIP) {
+ bool UsesLLVMOffloading = DriverArgs.hasFlag(
+ options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
+ if (DeviceOffloadingKind == Action::OFK_HIP ||
+ (DeviceOffloadingKind == Action::OFK_Cuda && UsesLLVMOffloading)) {
CC1Args.append({"-fcuda-is-device", "-fno-threadsafe-statics"});
if (!DriverArgs.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc,
diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
index 85fe99dbf8b69..09f6c09c340e8 100644
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -52,6 +52,7 @@
#include "llvm/Support/Path.h"
#include "llvm/Support/Process.h"
#include "llvm/Support/YAMLParser.h"
+#include "llvm/Support/raw_ostream.h"
#include "llvm/TargetParser/AArch64TargetParser.h"
#include "llvm/TargetParser/ARMTargetParserCommon.h"
#include "llvm/TargetParser/Host.h"
@@ -919,9 +920,12 @@ void Clang::AddPreprocessingOptions(Compilation &C, const JobAction &JA,
// before we -I or -include anything else, because we must pick up the
// CUDA/HIP/SYCL headers from the particular CUDA/ROCm/SYCL installation,
// rather than from e.g. /usr/local/include.
- if (JA.isOffloading(Action::OFK_Cuda))
+ bool UsesLLVMOffloading = Args.hasFlag(
+ options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
+ if (JA.isOffloading(Action::OFK_Cuda) && !UsesLLVMOffloading) {
getToolChain().AddCudaIncludeArgs(Args, CmdArgs);
- if (JA.isOffloading(Action::OFK_HIP))
+ }
+ if (JA.isOffloading(Action::OFK_HIP) && !UsesLLVMOffloading)
getToolChain().AddHIPIncludeArgs(Args, CmdArgs);
if (JA.isOffloading(Action::OFK_SYCL))
getToolChain().addSYCLIncludeArgs(Args, CmdArgs);
@@ -946,17 +950,35 @@ void Clang::AddPreprocessingOptions(Compilation &C, const JobAction &JA,
CmdArgs.push_back("-include");
CmdArgs.push_back("__clang_openmp_device_functions.h");
}
+ bool isCudaInput = llvm::any_of(
+ Inputs, [](const InputInfo &I) { return types::isCuda(I.getType()); });
+ if (UsesLLVMOffloading && JA.isOffloading(Action::OFK_Cuda) &&
+ Args.hasFlag(options::OPT_offload_inc, options::OPT_no_offload_inc,
+ true) &&
+ !Args.hasArg(options::OPT_nobuiltininc) && isCudaInput) {
+ CmdArgs.append({"-include", "__clang_gpu_device_functions.h"});
+
+ SmallString<128> OffloadCudaInclude(D.Dir);
+ llvm::sys::path::append(OffloadCudaInclude, "..", "include", "offload",
+ "cuda");
+ CmdArgs.append({"-internal-isystem", Args.MakeArgString(OffloadCudaInclude),
+ "-include"});
+ CmdArgs.push_back("cuda_runtime.h");
+ }
+ bool isHIPInput = llvm::any_of(
+ Inputs, [](const InputInfo &I) { return types::isHIP(I.getType()); });
+ if (UsesLLVMOffloading && JA.isOffloading(Action::OFK_HIP) &&
+ Args.hasFlag(options::OPT_offload_inc, options::OPT_no_offload_inc,
+ true) &&
+ !Args.hasArg(options::OPT_nohipwrapperinc) &&
+ !Args.hasArg(options::OPT_nobuiltininc) && isHIPInput) {
+ CmdArgs.append({"-include", "__clang_gpu_device_functions.h"});
- if (Args.hasArg(options::OPT_foffload_via_llvm)) {
- // Add llvm_wrappers/* to our system include path. This lets us wrap
- // standard library headers and other headers.
- SmallString<128> P(D.ResourceDir);
- llvm::sys::path::append(P, "include", "llvm_offload_wrappers");
- CmdArgs.append({"-internal-isystem", Args.MakeArgString(P), "-include"});
- if (JA.isDeviceOffloading(Action::OFK_OpenMP))
- CmdArgs.push_back("__llvm_offload_device.h");
- else
- CmdArgs.push_back("__llvm_offload_host.h");
+ SmallString<128> OffloadHIPInclude(D.Dir);
+ llvm::sys::path::append(OffloadHIPInclude, "..", "include", "offload");
+ CmdArgs.append({"-internal-isystem", Args.MakeArgString(OffloadHIPInclude),
+ "-include"});
+ CmdArgs.push_back("hip/hip_runtime.h");
}
// Add -i* options, and automatically translate to
@@ -1138,7 +1160,7 @@ void Clang::AddPreprocessingOptions(Compilation &C, const JobAction &JA,
Args.hasFlag(options::OPT_offload_inc, options::OPT_no_offload_inc,
true) &&
!Args.hasArg(options::OPT_nobuiltininc) &&
- (C.getActiveOffloadKinds() == Action::OFK_OpenMP)) {
+ JA.isDeviceOffloading(Action::OFK_OpenMP)) {
// TODO: CUDA / HIP include their own headers for some common functions
// implemented here. We'll need to clean those up so they do not conflict.
SmallString<128> P(D.ResourceDir);
@@ -5161,6 +5183,8 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
bool IsSYCLDevice = JA.isDeviceOffloading(Action::OFK_SYCL);
bool IsOpenMPDevice = JA.isDeviceOffloading(Action::OFK_OpenMP);
bool IsExtractAPI = isa<ExtractAPIJobAction>(JA);
+ bool UsesLLVMOffloading = Args.hasFlag(
+ options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
bool IsDeviceOffloadAction = !(JA.isDeviceOffloading(Action::OFK_None) ||
JA.isDeviceOffloading(Action::OFK_Host));
bool IsHostOffloadingAction =
@@ -5279,7 +5303,7 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
}
}
- if (IsCuda && !IsCudaDevice) {
+ if (IsCuda && !IsCudaDevice && !UsesLLVMOffloading) {
// We need to figure out which CUDA version we're compiling for, as that
// determines how we load and launch GPU kernels.
auto *CTC = static_cast<const toolchains::CudaToolChain *>(
@@ -8286,12 +8310,13 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
// Host-side offloading compilation receives all device-side outputs. Include
// them in the host compilation depending on the target. If the host inputs
// are not empty we use the new-driver scheme, otherwise use the old scheme.
- if ((IsCuda || IsHIP) && CudaDeviceInput) {
+ if ((IsCuda || IsHIP) && !UsesLLVMOffloading && CudaDeviceInput) {
CmdArgs.push_back("-fcuda-include-gpubinary");
CmdArgs.push_back(CudaDeviceInput->getFilename());
} else if (!HostOffloadingInputs.empty()) {
if ((IsCuda || IsHIP) &&
- (!IsRDCMode || Args.hasArg(options::OPT_cuda_emit_nvcc_abi))) {
+ (!IsRDCMode || Args.hasArg(options::OPT_cuda_emit_nvcc_abi))
+ && !UsesLLVMOffloading) {
assert(HostOffloadingInputs.size() == 1 && "Only one input expected");
CmdArgs.push_back("-fcuda-include-gpubinary");
CmdArgs.push_back(HostOffloadingInputs.front().getFilename());
diff --git a/clang/lib/Driver/ToolChains/CommonArgs.cpp b/clang/lib/Driver/ToolChains/CommonArgs.cpp
index 883296e43111b..a76f4aa6ae853 100644
--- a/clang/lib/Driver/ToolChains/CommonArgs.cpp
+++ b/clang/lib/Driver/ToolChains/CommonArgs.cpp
@@ -1491,18 +1491,25 @@ void tools::addArchSpecificRPath(const ToolChain &TC, const ArgList &Args,
}
}
+bool tools::addLLVMOffloadingRuntime(const Compilation &C,
+ ArgStringList &CmdArgs,
+ const ToolChain &TC, const ArgList &Args) {
+
+ if (!Args.hasFlag(options::OPT_foffload_via_llvm,
+ options::OPT_fno_offload_via_llvm, false))
+ return false;
+
+ CmdArgs.push_back("-lLLVMOffloadKernel");
+ return true;
+}
+
bool tools::addOpenMPRuntime(const Compilation &C, ArgStringList &CmdArgs,
const ToolChain &TC, const ArgList &Args,
bool ForceStaticHostRuntime, bool IsOffloadingHost,
bool GompNeedsRT) {
if (!Args.hasFlag(options::OPT_fopenmp, options::OPT_fopenmp_EQ,
- options::OPT_fno_openmp, false)) {
- // We need libomptarget (liboffload) if it's the choosen offloading runtime.
- if (Args.hasFlag(options::OPT_foffload_via_llvm,
- options::OPT_fno_offload_via_llvm, false))
- CmdArgs.push_back("-lomptarget");
+ options::OPT_fno_openmp, false))
return false;
- }
Driver::OpenMPRuntimeKind RTKind = TC.getDriver().getOpenMPRuntime(Args);
diff --git a/clang/lib/Driver/ToolChains/Cuda.cpp b/clang/lib/Driver/ToolChains/Cuda.cpp
index 84dcf180e30a9..d64044fdf2b1f 100644
--- a/clang/lib/Driver/ToolChains/Cuda.cpp
+++ b/clang/lib/Driver/ToolChains/Cuda.cpp
@@ -303,6 +303,23 @@ CudaInstallationDetector::CudaInstallationDetector(
void CudaInstallationDetector::AddCudaIncludeArgs(
const ArgList &DriverArgs, ArgStringList &CC1Args) const {
+ if (DriverArgs.hasFlag(options::OPT_foffload_via_llvm,
+ options::OPT_fno_offload_via_llvm, false)) {
+ if (DriverArgs.hasFlag(options::OPT_offload_inc,
+ options::OPT_no_offload_inc, true) &&
+ !DriverArgs.hasArg(options::OPT_nobuiltininc)) {
+ CC1Args.append({"-include", "__clang_gpu_device_functions.h"});
+
+ SmallString<128> CudaIncludePath(D.ResourceDir);
+ llvm::sys::path::append(CudaIncludePath, "..", "..", "..");
+ llvm::sys::path::append(CudaIncludePath, "include", "offload", "cuda");
+ CC1Args.push_back("-internal-isystem");
+ CC1Args.push_back(DriverArgs.MakeArgString(CudaIncludePath));
+ CC1Args.append({"-include", "cuda_runtime.h"});
+ }
+ return;
+ }
+
if (!DriverArgs.hasArg(options::OPT_nobuiltininc)) {
// Add cuda_wrappers/* to our system include path. This lets us wrap
// standard library headers.
@@ -398,7 +415,9 @@ void NVPTX::Assembler::ConstructJob(Compilation &C, const JobAction &JA,
const char *LinkingOutput) const {
const auto &TC =
static_cast<const toolchains::NVPTXToolChain &>(getToolChain());
- assert(TC.getTriple().isNVPTX() && "Wrong platform");
+
+ bool UsesLLVMOffloading = Args.hasFlag(
+ options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
BoundArch GPUArch;
// If this is a CUDA action we need to extract the device architecture
@@ -421,7 +440,7 @@ void NVPTX::Assembler::ConstructJob(Compilation &C, const JobAction &JA,
"Device action expected to have an architecture.");
// Check that our installation's ptxas supports gpu_arch.
- if (!Args.hasArg(options::OPT_no_cuda_version_check)) {
+ if (!UsesLLVMOffloading && !Args.hasArg(options::OPT_no_cuda_version_check)) {
TC.CudaInstallation.CheckCudaVersionSupportsArch(GPUArch.Arch);
}
@@ -494,7 +513,8 @@ void NVPTX::Assembler::ConstructJob(Compilation &C, const JobAction &JA,
/*Default=*/true);
else if (JA.isOffloading(Action::OFK_Cuda))
// In CUDA we generate relocatable code by default.
- Relocatable = Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc,
+ Relocatable = UsesLLVMOffloading ||
+ Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc,
/*Default=*/false);
else
// Otherwise, we are compiling directly and should create linkable output.
@@ -543,7 +563,9 @@ void NVPTX::FatBinary::ConstructJob(Compilation &C, const JobAction &JA,
const char *LinkingOutput) const {
const auto &TC =
static_cast<const toolchains::CudaToolChain &>(getToolChain());
- assert(TC.getTriple().isNVPTX() && "Wrong platform");
+ bool UsesLLVMOffloading = Args.hasFlag(
+ options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
+ assert((UsesLLVMOffloading || TC.getTriple().isNVPTX()) && "Wrong platform");
ArgStringList CmdArgs;
if (TC.CudaInstallation.version() <= CudaVersion::CUDA_100)
@@ -591,7 +613,9 @@ void NVPTX::Linker::ConstructJob(Compilation &C, const JobAction &JA,
static_cast<const toolchains::NVPTXToolChain &>(getToolChain());
ArgStringList CmdArgs;
- assert(TC.getTriple().isNVPTX() && "Wrong platform");
+ bool UsesLLVMOffloading = Args.hasFlag(
+ options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
+ assert((UsesLLVMOffloading || TC.getTriple().isNVPTX()) && "Wrong platform");
assert((Output.isFilename() || Output.isNothing()) && "Invalid output.");
if (Output.isFilename()) {
@@ -897,9 +921,12 @@ void CudaToolChain::addClangTargetOptions(
BoundArch BA, Action::OffloadKind DeviceOffloadingKind) const {
HostTC.addClangTargetOptions(DriverArgs, CC1Args, BA, DeviceOffloadingKind);
+ bool UsesLLVMOffloading = DriverArgs.hasFlag(
+ options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
+
StringRef GpuArch = DriverArgs.getLastArgValue(options::OPT_march_EQ);
assert((DeviceOffloadingKind == Action::OFK_OpenMP ||
- DeviceOffloadingKind == Action::OFK_Cuda) &&
+ DeviceOffloadingKind == Action::OFK_Cuda || UsesLLVMOffloading) &&
"Only OpenMP or CUDA offloading kinds are supported for NVIDIA GPUs.");
CC1Args.append({"-fcuda-is-device", "-mllvm",
@@ -918,6 +945,9 @@ void CudaToolChain::addClangTargetOptions(
DriverArgs.hasArg(options::OPT_S))
return;
+ if (UsesLLVMOffloading)
+ return;
+
std::string LibDeviceFile = CudaInstallation.getLibDeviceFile(GpuArch);
if (LibDeviceFile.empty()) {
getDriver().Diag(diag::err_drv_no_cuda_libdevice) << GpuArch;
@@ -927,13 +957,6 @@ void CudaToolChain::addClangTargetOptions(
CC1Args.push_back("-mlink-builtin-bitcode");
CC1Args.push_back(DriverArgs.MakeArgString(LibDeviceFile));
- // For now, we don't use any Offload/OpenMP device runtime when we offload
- // CUDA via LLVM/Offload. We should split the Offload/OpenMP device runtime
- // and include the "generic" (or CUDA-specific) parts.
- if (DriverArgs.hasFlag(options::OPT_foffload_via_llvm,
- options::OPT_fno_offload_via_llvm, false))
- return;
-
clang::CudaVersion CudaInstallationVersion = CudaInstallation.version();
if (CudaInstallationVersion >= CudaVersion::UNKNOWN)
@@ -974,6 +997,23 @@ llvm::DenormalMode CudaToolChain::getDefaultDenormalModeForType(
void CudaToolChain::AddCudaIncludeArgs(const ArgList &DriverArgs,
ArgStringList &CC1Args) const {
+ if (DriverArgs.hasFlag(options::OPT_foffload_via_llvm,
+ options::OPT_fno_offload_via_llvm, false)) {
+ if (DriverArgs.hasFlag(options::OPT_offload_inc,
+ options::OPT_no_offload_inc, true) &&
+ !DriverArgs.hasArg(options::OPT_nobuiltininc)) {
+ CC1Args.append({"-include", "__clang_gpu_device_functions.h"});
+
+ SmallString<128> CudaIncludePath(getDriver().ResourceDir);
+ llvm::sys::path::append(CudaIncludePath, "..", "..", "..");
+ llvm::sys::path::append(CudaIncludePath, "include", "offload", "cuda");
+ CC1Args.push_back("-internal-isystem");
+ CC1Args.push_back(DriverArgs.MakeArgString(CudaIncludePath));
+ CC1Args.append({"-include", "cuda_runtime.h"});
+ }
+ return;
+ }
+
// Check our CUDA version if we're going to include the CUDA headers.
if (DriverArgs.hasFlag(options::OPT_offload_inc, options::OPT_no_offload_inc,
true) &&
@@ -1046,6 +1086,10 @@ CudaToolChain::GetCXXStdlibType(const ArgList &Args) const {
void CudaToolChain::AddClangSystemIncludeArgs(const ArgList &DriverArgs,
ArgStringList &CC1Args) const {
+ if (DriverArgs.hasFlag(options::OPT_foffload_via_llvm,
+ options::OPT_fno_offload_via_llvm, false))
+ return;
+
HostTC.AddClangSystemIncludeArgs(DriverArgs, CC1Args);
if (DriverArgs.hasFlag(options::OPT_offload_inc, options::OPT_no_offload_inc,
diff --git a/clang/lib/Driver/ToolChains/Gnu.cpp b/clang/lib/Driver/ToolChains/Gnu.cpp
index 24076d8814322..72affac131701 100644
--- a/clang/lib/Driver/ToolChains/Gnu.cpp
+++ b/clang/lib/Driver/ToolChains/Gnu.cpp
@@ -510,6 +510,7 @@ void tools::gnutools::Linker::ConstructJob(Compilation &C, const JobAction &JA,
// FIXME: Does this really make sense for all GNU toolchains?
WantPthread = true;
+ addLLVMOffloadingRuntime(C, CmdArgs, ToolChain, Args);
AddRunTimeLibs(ToolChain, D, CmdArgs, Args);
// LLVM support for atomics on 32-bit SPARC V8+ is incomplete, so
diff --git a/clang/lib/Driver/ToolChains/Linux.cpp b/clang/lib/Driver/ToolChains/Linux.cpp
index 1ab385a9ea001..03017f20e391d 100644
--- a/clang/lib/Driver/ToolChains/Linux.cpp
+++ b/clang/lib/Driver/ToolChains/Linux.cpp
@@ -881,7 +881,9 @@ void Linux::addOffloadRTLibs(unsigned ActiveKinds, const ArgList &Args,
if (!Args.hasFlag(options::OPT_offloadlib, options::OPT_no_offloadlib,
true) ||
Args.hasArg(options::OPT_nostdlib) ||
- Args.hasArg(options::OPT_no_hip_rt) || Args.hasArg(options::OPT_r))
+ Args.hasArg(options::OPT_no_hip_rt) || Args.hasArg(options::OPT_r) ||
+ Args.hasFlag(options::OPT_foffload_via_llvm,
+ options::OPT_fno_offload_via_llvm, false))
return;
llvm::SmallVector<std::pair<StringRef, StringRef>> Libraries;
diff --git a/clang/lib/Headers/__clang_gpu_builtin_vars.h b/clang/lib/Headers/__clang_gpu_builtin_vars.h
index b80248dcd2be3..f3137be0a3181 100644
--- a/clang/lib/Headers/__clang_gpu_builtin_vars.h
+++ b/clang/lib/Headers/__clang_gpu_builtin_vars.h
@@ -6,6 +6,8 @@
//
//===-----------------------------------------------------------------------===
+#include <stddef.h>
+
#ifndef __CLANG_GPU_BUILTIN_VARS_H__
#define __CLANG_GPU_BUILTIN_VARS_H__
@@ -20,6 +22,23 @@ static inline __attribute__((device)) const struct {
}
} warpSize{};
+extern "C" {
+
+typedef struct dim3 {
+ dim3() {}
+ dim3(unsigned x) : x(x) {}
+ unsigned x = 0, y = 0, z = 0;
+} dim3;
+
+// TODO: For some reason the CUDA device compilation requires this declaration
+// to be present on the device while it is only used on the host.
+unsigned __llvmPushCallConfiguration(dim3 gridDim, dim3 blockDim,
+ size_t sharedMem = 0, void *stream = 0);
+unsigned __llvmLaunchKernel(const void *func, dim3 gridDim, dim3 blockDim,
+ void **args, size_t sharedMem = 0,
+ void *stream = 0);
+}
+
// Make sure nobody can create instances of the coordinate types, take their
// address, copy, or assign them.
#pragma push_macro("__GPU_DISALLOW_BUILTINVAR_ACCESS")
diff --git a/clang/test/CodeGenCUDA/Inputs/cuda.h b/clang/test/CodeGenCUDA/Inputs/cuda.h
index 421fa4dd7dbae..83bb7b7bdbb7f 100644
--- a/clang/test/CodeGenCUDA/Inputs/cuda.h
+++ b/clang/test/CodeGenCUDA/Inputs/cuda.h
@@ -55,7 +55,7 @@ extern "C" hipError_t hipLaunchKernel_spt(const void *func, dim3 gridDim,
#elif __OFFLOAD_VIA_LLVM__
extern "C" unsigned __llvmPushCallConfiguration(dim3 gridDim, dim3 blockDim,
size_t sharedMem = 0, void *stream = 0);
-extern "C" unsigned llvmLaunchKernel(const void *func, dim3 gridDim, dim3 blockDim,
+extern "C" unsigned __llvmLaunchKernel(const void *func, dim3 gridDim, dim3 blockDim,
void **args, size_t sharedMem = 0, void *stream = 0);
#else
typedef struct cudaStream *cudaStream_t;
diff --git a/clang/test/CodeGenCUDA/offload_via_llvm.cu b/clang/test/CodeGenCUDA/offload_via_llvm.cu
index b13a64c81b775..c58e793d11826 100644
--- a/clang/test/CodeGenCUDA/offload_via_llvm.cu
+++ b/clang/test/CodeGenCUDA/offload_via_llvm.cu
@@ -14,9 +14,7 @@
// HST-NEXT: [[DOTADDR1:%.*]] = alloca i16, align 2
// HST-NEXT: [[DOTADDR2:%.*]] = alloca ptr, align 4
// HST-NEXT: [[DOTADDR3:%.*]] = alloca ptr, align 4
-// HST-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[TMP0]], align 16
-// HST-NEXT: [[KERNEL_ARGS_PTRS:%.*]] = alloca [4 x ptr], align 16
-// HST-NEXT: [[KERNEL_LAUNCH_PARAMS:%.*]] = alloca [[TMP1]], align 16
+// HST-NEXT: [[KERNEL_LAUNCH_PARAMS:%.*]] = alloca [[TMP0]], align 16
// HST-NEXT: [[GRID_DIM:%.*]] = alloca [[STRUCT_DIM3:%.*]], align 8
// HST-NEXT: [[BLOCK_DIM:%.*]] = alloca [[STRUCT_DIM3]], align 8
// HST-NEXT: [[SHMEM_SIZE:%.*]] = alloca i32, align 4
@@ -25,34 +23,34 @@
// HST-NEXT: store i16 [[TMP1]], ptr [[DOTADDR1]], align 2
// HST-NEXT: store ptr [[TMP2]], ptr [[DOTADDR2]], align 4
// HST-NEXT: store ptr [[TMP3]], ptr [[DOTADDR3]], align 4
-// HST-NEXT: [[TMP4:%.*]] = getelementptr inbounds nuw [[TMP1]], ptr [[KERNEL_LAUNCH_PARAMS]], i32 0, i32 0
-// HST-NEXT: store i32 4, ptr [[TMP4]], align 16
-// HST-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[TMP1]], ptr [[KERNEL_LAUNCH_PARAMS]], i32 0, i32 1
-// HST-NEXT: store ptr [[KERNEL_ARGS_PTRS]], ptr [[TMP5]], align 4
-// HST-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTADDR]], align 4
-// HST-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_ARGS]], i32 0, i32 0
-// HST-NEXT: store i32 [[TMP6]], ptr [[TMP7]], align 16
-// HST-NEXT: [[TMP8:%.*]] = getelementptr inbounds [4 x ptr], ptr [[KERNEL_ARGS_PTRS]], i32 0, i32 0
-// HST-NEXT: store ptr [[TMP7]], ptr [[TMP8]], align 16
-// HST-NEXT: [[TMP9:%.*]] = load i16, ptr [[DOTADDR1]], align 2
-// HST-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_ARGS]], i32 0, i32 1
-// HST-NEXT: store i16 [[TMP9]], ptr [[TMP10]], align 4
-// HST-NEXT: [[TMP11:%.*]] = getelementptr inbounds [4 x ptr], ptr [[KERNEL_ARGS_PTRS]], i32 0, i32 1
-// HST-NEXT: store ptr [[TMP10]], ptr [[TMP11]], align 4
-// HST-NEXT: [[TMP12:%.*]] = load ptr, ptr [[DOTADDR2]], align 4
-// HST-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_ARGS]], i32 0, i32 2
-// HST-NEXT: store ptr [[TMP12]], ptr [[TMP13]], align 8
-// HST-NEXT: [[TMP14:%.*]] = getelementptr inbounds [4 x ptr], ptr [[KERNEL_ARGS_PTRS]], i32 0, i32 2
-// HST-NEXT: store ptr [[TMP13]], ptr [[TMP14]], align 8
-// HST-NEXT: [[TMP15:%.*]] = load ptr, ptr [[DOTADDR3]], align 4
-// HST-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_ARGS]], i32 0, i32 3
-// HST-NEXT: store ptr [[TMP15]], ptr [[TMP16]], align 4
-// HST-NEXT: [[TMP17:%.*]] = getelementptr inbounds [4 x ptr], ptr [[KERNEL_ARGS_PTRS]], i32 0, i32 3
-// HST-NEXT: store ptr [[TMP16]], ptr [[TMP17]], align 4
-// HST-NEXT: [[TMP18:%.*]] = call i32 @__llvmPopCallConfiguration(ptr [[GRID_DIM]], ptr [[BLOCK_DIM]], ptr [[SHMEM_SIZE]], ptr [[STREAM]])
-// HST-NEXT: [[TMP19:%.*]] = load i32, ptr [[SHMEM_SIZE]], align 4
-// HST-NEXT: [[TMP20:%.*]] = load ptr, ptr [[STREAM]], align 4
-// HST-NEXT: [[CALL:%.*]] = call noundef i32 @llvmLaunchKernel(ptr noundef @_Z18__device_stub__fooisPvS_, ptr noundef byval([[STRUCT_DIM3]]) align 4 [[GRID_DIM]], ptr noundef byval([[STRUCT_DIM3]]) align 4 [[BLOCK_DIM]], ptr noundef [[KERNEL_LAUNCH_PARAMS]], i32 noundef [[TMP19]], ptr noundef [[TMP20]]) #[[ATTR3:[0-9]+]]
+// HST-NEXT: [[KERNEL_ARGS:%.*]] = alloca ptr, i32 4, align 16
+// HST-NEXT: [[KERNEL_ARG_SIZES:%.*]] = alloca i32, i32 4, align 16
+// HST-NEXT: [[TMP4:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_LAUNCH_PARAMS]], i32 0, i32 0
+// HST-NEXT: store ptr [[KERNEL_ARGS]], ptr [[TMP4]], align 16
+// HST-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_LAUNCH_PARAMS]], i32 0, i32 1
+// HST-NEXT: store i64 4, ptr [[TMP5]], align 8
+// HST-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_LAUNCH_PARAMS]], i32 0, i32 2
+// HST-NEXT: store ptr [[KERNEL_ARG_SIZES]], ptr [[TMP6]], align 16
+// HST-NEXT: [[TMP7:%.*]] = getelementptr ptr, ptr [[KERNEL_ARGS]], i32 0
+// HST-NEXT: store ptr [[DOTADDR]], ptr [[TMP7]], align 4
+// HST-NEXT: [[TMP8:%.*]] = getelementptr i32, ptr [[KERNEL_ARG_SIZES]], i32 0
+// HST-NEXT: store i32 4, ptr [[TMP8]], align 4
+// HST-NEXT: [[TMP9:%.*]] = getelementptr ptr, ptr [[KERNEL_ARGS]], i32 1
+// HST-NEXT: store ptr [[DOTADDR1]], ptr [[TMP9]], align 4
+// HST-NEXT: [[TMP10:%.*]] = getelementptr i32, ptr [[KERNEL_ARG_SIZES]], i32 1
+// HST-NEXT: store i32 2, ptr [[TMP10]], align 4
+// HST-NEXT: [[TMP11:%.*]] = getelementptr ptr, ptr [[KERNEL_ARGS]], i32 2
+// HST-NEXT: store ptr [[DOTADDR2]], ptr [[TMP11]], align 4
+// HST-NEXT: [[TMP12:%.*]] = getelementptr i32, ptr [[KERNEL_ARG_SIZES]], i32 2
+// HST-NEXT: store i32 4, ptr [[TMP12]], align 4
+// HST-NEXT: [[TMP13:%.*]] = getelementptr ptr, ptr [[KERNEL_ARGS]], i32 3
+// HST-NEXT: store ptr [[DOTADDR3]], ptr [[TMP13]], align 4
+// HST-NEXT: [[TMP14:%.*]] = getelementptr i32, ptr [[KERNEL_ARG_SIZES]], i32 3
+// HST-NEXT: store i32 4, ptr [[TMP14]], align 4
+// HST-NEXT: [[TMP15:%.*]] = call i32 @__cudaPopCallConfiguration(ptr [[GRID_DIM]], ptr [[BLOCK_DIM]], ptr [[SHMEM_SIZE]], ptr [[STREAM]])
+// HST-NEXT: [[TMP16:%.*]] = load i32, ptr [[SHMEM_SIZE]], align 4
+// HST-NEXT: [[TMP17:%.*]] = load ptr, ptr [[STREAM]], align 4
+// HST-NEXT: [[CALL:%.*]] = call noundef i32 @__llvmLaunchKernel(ptr noundef @_Z18__device_stub__fooisPvS_, ptr noundef byval([[STRUCT_DIM3]]) align 4 [[GRID_DIM]], ptr noundef byval([[STRUCT_DIM3]]) align 4 [[BLOCK_DIM]], ptr noundef [[KERNEL_LAUNCH_PARAMS]], i32 noundef [[TMP16]], ptr noundef [[TMP17]]) #[[ATTR3:[0-9]+]]
// HST-NEXT: br label %[[SETUP_END:.*]]
// HST: [[SETUP_END]]:
// HST-NEXT: ret void
diff --git a/clang/test/Driver/cuda-via-liboffload.cu b/clang/test/Driver/cuda-via-liboffload.cu
index 68dc963e906b2..d30e529f0ce12 100644
--- a/clang/test/Driver/cuda-via-liboffload.cu
+++ b/clang/test/Driver/cuda-via-liboffload.cu
@@ -2,21 +2,20 @@
// RUN: --offload-arch=sm_35 --offload-arch=sm_70 %s 2>&1 \
// RUN: | FileCheck -check-prefix BINDINGS %s
-// BINDINGS: "x86_64-unknown-linux-gnu" - "clang", inputs: ["[[INPUT:.+]]"], output: "[[HOST_BC:.+]]"
-// BINDINGS-NEXT: "nvptx64-nvidia-cuda" - "clang", inputs: ["[[INPUT]]", "[[HOST_BC]]"], output: "[[PTX_SM_35:.+]]"
-// BINDINGS-NEXT: "nvptx64-nvidia-cuda" - "NVPTX::Assembler", inputs: ["[[PTX_SM_35]]"], output: "[[CUBIN_SM_35:.+]]"
-// BINDINGS-NEXT: "nvptx64-nvidia-cuda" - "clang", inputs: ["[[INPUT]]", "[[HOST_BC]]"], output: "[[PTX_SM_70:.+]]"
-// BINDINGS-NEXT: "nvptx64-nvidia-cuda" - "NVPTX::Assembler", inputs: ["[[PTX_SM_70:.+]]"], output: "[[CUBIN_SM_70:.+]]"
+// BINDINGS: "nvptx64-nvidia-cuda-llvm" - "clang", inputs: ["[[INPUT:.+]]"], output: "[[PTX_SM_35:.+]]"
+// BINDINGS-NEXT: "nvptx64-nvidia-cuda-llvm" - "NVPTX::Assembler", inputs: ["[[PTX_SM_35]]"], output: "[[CUBIN_SM_35:.+]]"
+// BINDINGS-NEXT: "nvptx64-nvidia-cuda-llvm" - "clang", inputs: ["[[INPUT]]"], output: "[[PTX_SM_70:.+]]"
+// BINDINGS-NEXT: "nvptx64-nvidia-cuda-llvm" - "NVPTX::Assembler", inputs: ["[[PTX_SM_70:.+]]"], output: "[[CUBIN_SM_70:.+]]"
// BINDINGS-NEXT: "x86_64-unknown-linux-gnu" - "Offload::Packager", inputs: ["[[CUBIN_SM_35]]", "[[CUBIN_SM_70]]"], output: "[[BINARY:.+]]"
-// BINDINGS-NEXT: "x86_64-unknown-linux-gnu" - "clang", inputs: ["[[HOST_BC]]", "[[BINARY]]"], output: "[[HOST_OBJ:.+]]"
+// BINDINGS-NEXT: "x86_64-unknown-linux-gnu" - "clang", inputs: ["[[INPUT]]", "[[BINARY]]"], output: "[[HOST_OBJ:.+]]"
// BINDINGS-NEXT: "x86_64-unknown-linux-gnu" - "Offload::Linker", inputs: ["[[HOST_OBJ]]"], output: "a.out"
// RUN: %clang -### -target x86_64-linux-gnu -foffload-via-llvm -ccc-print-bindings \
// RUN: --offload-arch=sm_35 --offload-arch=sm_70 %s 2>&1 \
// RUN: | FileCheck -check-prefix BINDINGS-DEVICE %s
-// BINDINGS-DEVICE: # "nvptx64-nvidia-cuda" - "clang", inputs: ["[[INPUT:.+]]"], output: "[[PTX:.+]]"
-// BINDINGS-DEVICE: # "nvptx64-nvidia-cuda" - "NVPTX::Assembler", inputs: ["[[PTX]]"], output: "[[CUBIN:.+]]"
+// BINDINGS-DEVICE: # "nvptx64-nvidia-cuda-llvm" - "clang", inputs: ["[[INPUT:.+]]"], output: "[[PTX:.+]]"
+// BINDINGS-DEVICE: # "nvptx64-nvidia-cuda-llvm" - "NVPTX::Assembler", inputs: ["[[PTX]]"], output: "[[CUBIN:.+]]"
// RUN: %clang -### -target x86_64-linux-gnu -ccc-print-bindings --offload-link -foffload-via-llvm %s 2>&1 | FileCheck -check-prefix DEVICE-LINK %s
diff --git a/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-image.c b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-image.c
index d05e9d54a108a..083d3340f6f81 100644
--- a/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-image.c
+++ b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-image.c
@@ -94,7 +94,7 @@
// CUDA-NEXT: br i1 %1, label %while.entry, label %while.end
//
// CUDA: while.entry:
-// CUDA-NEXT: %entry1 = phi ptr [ @__start_llvm_offload_entries, %entry ], [ %16, %if.end ]
+// CUDA-NEXT: %entry1 = phi ptr [ @__start_llvm_offload_entries, %entry ], [ %17, %if.end ]
// CUDA-NEXT: %2 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i32 0, i32 4
// CUDA-NEXT: %addr = load ptr, ptr %2, align 8
// CUDA-NEXT: %3 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i32 0, i32 8
@@ -117,15 +117,16 @@
// CUDA-NEXT: %constant = lshr i32 %11, 4
// CUDA-NEXT: %12 = and i32 %flags, 32
// CUDA-NEXT: %normalized = lshr i32 %12, 5
-// CUDA-NEXT: %13 = icmp eq i16 %kind, 2
-// CUDA-NEXT: br i1 %13, label %if.kind, label %if.end
+// CUDA-NEXT: %13 = and i16 %kind, 2
+// CUDA-NEXT: %14 = icmp ne i16 %13, 0
+// CUDA-NEXT: br i1 %14, label %if.kind, label %if.end
//
// CUDA: if.kind:
-// CUDA-NEXT: %14 = icmp eq i64 %size, 0
-// CUDA-NEXT: br i1 %14, label %if.then, label %if.else
+// CUDA-NEXT: %15 = icmp eq i64 %size, 0
+// CUDA-NEXT: br i1 %15, label %if.then, label %if.else
//
// CUDA: if.then:
-// CUDA-NEXT: %15 = call i32 @__cudaRegisterFunction(ptr %0, ptr %addr, ptr %name, ptr %name, i32 -1, ptr null, ptr null, ptr null, ptr null, ptr null)
+// CUDA-NEXT: %16 = call i32 @__cudaRegisterFunction(ptr %0, ptr %addr, ptr %name, ptr %name, i32 -1, ptr null, ptr null, ptr null, ptr null, ptr null)
// CUDA-NEXT: br label %if.end
//
// CUDA: if.else:
@@ -151,9 +152,9 @@
// CUDA-NEXT: br label %if.end
//
// CUDA: if.end:
-// CUDA-NEXT: %16 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i64 1
-// CUDA-NEXT: %17 = icmp eq ptr %16, @__stop_llvm_offload_entries
-// CUDA-NEXT: br i1 %17, label %while.end, label %while.entry
+// CUDA-NEXT: %17 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i64 1
+// CUDA-NEXT: %18 = icmp eq ptr %17, @__stop_llvm_offload_entries
+// CUDA-NEXT: br i1 %18, label %while.end, label %while.entry
//
// CUDA: while.end:
// CUDA-NEXT: ret void
@@ -236,7 +237,7 @@
// HIP-NEXT: br i1 %1, label %while.entry, label %while.end
//
// HIP: while.entry:
-// HIP-NEXT: %entry1 = phi ptr [ @{{.*offload_entries.*}}, %entry ], [ %16, %if.end ]
+// HIP-NEXT: %entry1 = phi ptr [ @{{.*offload_entries.*}}, %entry ], [ %17, %if.end ]
// HIP-NEXT: %2 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i32 0, i32 4
// HIP-NEXT: %addr = load ptr, ptr %2, align 8
// HIP-NEXT: %3 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i32 0, i32 8
@@ -259,15 +260,16 @@
// HIP-NEXT: %constant = lshr i32 %11, 4
// HIP-NEXT: %12 = and i32 %flags, 32
// HIP-NEXT: %normalized = lshr i32 %12, 5
-// HIP-NEXT: %13 = icmp eq i16 %kind, 4
-// HIP-NEXT: br i1 %13, label %if.kind, label %if.end
+// HIP-NEXT: %13 = and i16 %kind, 4
+// HIP-NEXT: %14 = icmp ne i16 %13, 0
+// HIP-NEXT: br i1 %14, label %if.kind, label %if.end
//
// HIP: if.kind:
-// HIP-NEXT: %14 = icmp eq i64 %size, 0
-// HIP-NEXT: br i1 %14, label %if.then, label %if.else
+// HIP-NEXT: %15 = icmp eq i64 %size, 0
+// HIP-NEXT: br i1 %15, label %if.then, label %if.else
//
// HIP: if.then:
-// HIP-NEXT: %15 = call i32 @__hipRegisterFunction(ptr %0, ptr %addr, ptr %name, ptr %name, i32 -1, ptr null, ptr null, ptr null, ptr null, ptr null)
+// HIP-NEXT: %16 = call i32 @__hipRegisterFunction(ptr %0, ptr %addr, ptr %name, ptr %name, i32 -1, ptr null, ptr null, ptr null, ptr null, ptr null)
// HIP-NEXT: br label %if.end
//
// HIP: if.else:
@@ -295,9 +297,9 @@
// HIP-NEXT: br label %if.end
//
// HIP: if.end:
-// HIP-NEXT: %16 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i64 1
-// HIP-NEXT: %17 = icmp eq ptr %16, @{{.*offload_entries.*}}
-// HIP-NEXT: br i1 %17, label %while.end, label %while.entry
+// HIP-NEXT: %17 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i64 1
+// HIP-NEXT: %18 = icmp eq ptr %17, @{{.*offload_entries.*}}
+// HIP-NEXT: br i1 %18, label %while.end, label %while.entry
//
// HIP: while.end:
// HIP-NEXT: ret void
diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
index f2a58774e99af..21aee2121f255 100644
--- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
+++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
@@ -139,6 +139,13 @@ static bool CanonicalPrefixes = true;
using OffloadingImage = OffloadBinary::OffloadingImage;
+static bool usesLLVMOffloadWrapper(ArrayRef<OffloadingImage> Images) {
+ return llvm::any_of(Images, [](const OffloadingImage &Image) {
+ return Triple(Image.StringData.lookup("triple")).getEnvironment() ==
+ Triple::LLVM;
+ });
+}
+
namespace llvm {
// Provide DenseMapInfo so that OffloadKind can be used in a DenseMap.
template <> struct DenseMapInfo<OffloadKind> {
@@ -977,6 +984,9 @@ Expected<SmallVector<std::unique_ptr<MemoryBuffer>>>
bundleLinkedOutput(ArrayRef<OffloadingImage> Images, const ArgList &Args,
OffloadKind Kind) {
llvm::TimeTraceScope TimeScope("Bundle linked output");
+ if (usesLLVMOffloadWrapper(Images))
+ return bundleOpenMP(Images);
+
switch (Kind) {
case OFK_OpenMP:
return (Verbose && SaveTemps) ? bundleOpenMPVerbose(Images)
@@ -1220,7 +1230,8 @@ linkAndWrapDeviceFiles(ArrayRef<SmallVector<OffloadFile>> LinkerInputFiles,
continue;
}
- auto OutputOrErr = wrapDeviceImages(*BundledImagesOrErr, Args, Kind);
+ OffloadKind WrapperKind = usesLLVMOffloadWrapper(Input) ? OFK_OpenMP : Kind;
+ auto OutputOrErr = wrapDeviceImages(*BundledImagesOrErr, Args, WrapperKind);
if (!OutputOrErr)
return OutputOrErr.takeError();
WrappedOutput.push_back(*OutputOrErr);
diff --git a/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp b/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp
index ded603a1e00e3..037b81a7c42fb 100644
--- a/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp
+++ b/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp
@@ -485,10 +485,12 @@ Function *createRegisterGlobalsFunction(Module &M, bool IsHIP,
llvm::offloading::OffloadGlobalNormalized));
auto *Normalized = Builder.CreateLShr(
NormalizedBit, ConstantInt::get(Type::getInt32Ty(C), 5), "normalized");
- auto *KindCond = Builder.CreateICmpEQ(
+ auto *KindAnd = Builder.CreateAnd(
Kind, ConstantInt::get(Type::getInt16Ty(C),
IsHIP ? object::OffloadKind::OFK_HIP
: object::OffloadKind::OFK_Cuda));
+ auto *KindCond =
+ Builder.CreateICmpNE(KindAnd, ConstantInt::get(Type::getInt16Ty(C), 0));
Builder.CreateCondBr(KindCond, IfKindBB, IfEndBB);
Builder.SetInsertPoint(IfKindBB);
auto *FnCond = Builder.CreateICmpEQ(
>From 4a00037fbe2dfda13ffb3cadba33149f4cd57e0b Mon Sep 17 00:00:00 2001
From: Sophia Herrmann <herrmann15 at llnl.gov>
Date: Tue, 28 Jul 2026 17:49:31 -0700
Subject: [PATCH 6/7] add unittests
---
clang/lib/Driver/Driver.cpp | 19 +++----
clang/lib/Driver/ToolChains/Clang.cpp | 4 +-
offload/test/lit.cfg | 10 +++-
offload/test/offloading/CUDA/basic_launch.cu | 25 +++++----
.../CUDA/basic_launch_blocks_and_threads.cu | 22 ++++----
.../offloading/CUDA/basic_launch_multi_arg.cu | 36 +++++++------
offload/test/offloading/CUDA/device_api.cu | 45 ++++++++++++++++
.../test/offloading/CUDA/device_properties.cu | 40 +++++++++++++++
offload/test/offloading/CUDA/host_alloc.cu | 40 +++++++++++++++
offload/test/offloading/CUDA/launch_tu.cu | 25 +++++----
offload/test/offloading/CUDA/memcpy_kinds.cu | 51 +++++++++++++++++++
offload/test/offloading/CUDA/stream_api.cu | 46 +++++++++++++++++
offload/test/offloading/CUDA/syncthreads.cu | 40 +++++++++++++++
.../offloading/CUDA/thread_and_block_id.cu | 44 ++++++++++++++++
offload/test/offloading/HIP/basic_launch.hip | 30 +++++++++++
.../HIP/basic_launch_blocks_and_threads.hip | 31 +++++++++++
.../offloading/HIP/basic_launch_multi_arg.hip | 39 ++++++++++++++
offload/test/offloading/HIP/device_api.hip | 45 ++++++++++++++++
.../test/offloading/HIP/device_properties.hip | 40 +++++++++++++++
offload/test/offloading/HIP/host_alloc.hip | 40 +++++++++++++++
offload/test/offloading/HIP/kernel_tu.hip.inc | 1 +
offload/test/offloading/HIP/launch_tu.hip | 30 +++++++++++
offload/test/offloading/HIP/memcpy_kinds.hip | 51 +++++++++++++++++++
offload/test/offloading/HIP/stream_api.hip | 46 +++++++++++++++++
offload/test/offloading/HIP/syncthreads.hip | 40 +++++++++++++++
.../offloading/HIP/thread_and_block_id.hip | 44 ++++++++++++++++
26 files changed, 815 insertions(+), 69 deletions(-)
create mode 100644 offload/test/offloading/CUDA/device_api.cu
create mode 100644 offload/test/offloading/CUDA/device_properties.cu
create mode 100644 offload/test/offloading/CUDA/host_alloc.cu
create mode 100644 offload/test/offloading/CUDA/memcpy_kinds.cu
create mode 100644 offload/test/offloading/CUDA/stream_api.cu
create mode 100644 offload/test/offloading/CUDA/syncthreads.cu
create mode 100644 offload/test/offloading/CUDA/thread_and_block_id.cu
create mode 100644 offload/test/offloading/HIP/basic_launch.hip
create mode 100644 offload/test/offloading/HIP/basic_launch_blocks_and_threads.hip
create mode 100644 offload/test/offloading/HIP/basic_launch_multi_arg.hip
create mode 100644 offload/test/offloading/HIP/device_api.hip
create mode 100644 offload/test/offloading/HIP/device_properties.hip
create mode 100644 offload/test/offloading/HIP/host_alloc.hip
create mode 100644 offload/test/offloading/HIP/kernel_tu.hip.inc
create mode 100644 offload/test/offloading/HIP/launch_tu.hip
create mode 100644 offload/test/offloading/HIP/memcpy_kinds.hip
create mode 100644 offload/test/offloading/HIP/stream_api.hip
create mode 100644 offload/test/offloading/HIP/syncthreads.hip
create mode 100644 offload/test/offloading/HIP/thread_and_block_id.hip
diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp
index a7241bd878dcc..c05724fd646d4 100644
--- a/clang/lib/Driver/Driver.cpp
+++ b/clang/lib/Driver/Driver.cpp
@@ -974,15 +974,16 @@ static TripleSet inferOffloadToolchains(Compilation &C,
C.getArgs().hasFlag(options::OPT_foffload_via_llvm,
options::OPT_fno_offload_via_llvm, false);
if (!UsesLLVMOffloading) {
- if (Kind == Action::OFK_HIP && !ID.isAMDGPU() && !ID.isSPIRV()) {
- C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch)
- << "HIP" << Arch;
- return {};
- }
- if (Kind == Action::OFK_Cuda && !ID.isNVPTX()) {
- C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch)
- << "CUDA" << Arch;
- return {};
+ if (Kind == Action::OFK_HIP && !ID.isAMDGPU() && !ID.isSPIRV()) {
+ C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch)
+ << "HIP" << Arch;
+ return {};
+ }
+ if (Kind == Action::OFK_Cuda && !ID.isNVPTX()) {
+ C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch)
+ << "CUDA" << Arch;
+ return {};
+ }
}
if (Kind == Action::OFK_OpenMP && (ID.isUnknown() || ID.isUnused())) {
C.getDriver().Diag(clang::diag::err_drv_failed_to_deduce_target_from_arch)
diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
index 09f6c09c340e8..f9f6cb07ccc37 100644
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -8315,8 +8315,8 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
CmdArgs.push_back(CudaDeviceInput->getFilename());
} else if (!HostOffloadingInputs.empty()) {
if ((IsCuda || IsHIP) &&
- (!IsRDCMode || Args.hasArg(options::OPT_cuda_emit_nvcc_abi))
- && !UsesLLVMOffloading) {
+ (!IsRDCMode || Args.hasArg(options::OPT_cuda_emit_nvcc_abi)) &&
+ !UsesLLVMOffloading) {
assert(HostOffloadingInputs.size() == 1 && "Only one input expected");
CmdArgs.push_back("-fcuda-include-gpubinary");
CmdArgs.push_back(HostOffloadingInputs.front().getFilename());
diff --git a/offload/test/lit.cfg b/offload/test/lit.cfg
index ace2b1ea8a749..2981cf5d10fea 100644
--- a/offload/test/lit.cfg
+++ b/offload/test/lit.cfg
@@ -83,7 +83,7 @@ def remove_suffix_if_present(name):
config.name = 'libomptarget :: ' + config.libomptarget_current_target
# suffixes: A list of file extensions to treat as test files.
-config.suffixes = ['.c', '.cpp', '.cc', '.f90', '.cu', '.td']
+config.suffixes = ['.c', '.cpp', '.cc', '.f90', '.cu', '.hip', '.td']
# excludes: A list of directories to exclude from the testuites.
config.excludes = ['Inputs', 'unit']
@@ -91,6 +91,11 @@ config.excludes = ['Inputs', 'unit']
# test_source_root: The root path where tests are located.
config.test_source_root = os.path.dirname(__file__)
+# language includes
+config.test_language_includes = os.path.join(config.test_source_root, "../languages/include")
+config.test_language_cuda_includes = os.path.join(config.test_language_includes, "cuda")
+config.test_language_hip_includes = os.path.join(config.test_language_includes, "hip")
+
# test_exec_root: The root object directory where output is placed
config.test_exec_root = config.libomptarget_obj_root
@@ -100,6 +105,9 @@ config.test_format = lit.formats.ShTest()
# compiler flags
config.test_flags = " -I " + config.test_source_root + \
" -I " + config.omp_header_directory + \
+ " -I " + config.test_language_includes + \
+ " -I " + config.test_language_cuda_includes + \
+ " -I " + config.test_language_hip_includes + \
" -L " + config.library_dir + \
" -L " + config.llvm_library_intdir + \
" -L " + config.llvm_lib_directory
diff --git a/offload/test/offloading/CUDA/basic_launch.cu b/offload/test/offloading/CUDA/basic_launch.cu
index e017241bb9a74..5ecfc3e9d5601 100644
--- a/offload/test/offloading/CUDA/basic_launch.cu
+++ b/offload/test/offloading/CUDA/basic_launch.cu
@@ -7,25 +7,24 @@
// UNSUPPORTED: aarch64-unknown-linux-gnu
// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
// UNSUPPORTED: intelgpu
#include <stdio.h>
-extern "C" {
-void *llvm_omp_target_alloc_shared(size_t Size, int DeviceNum);
-void llvm_omp_target_free_shared(void *DevicePtr, int DeviceNum);
-}
-
__global__ void square(int *A) { *A = 42; }
int main(int argc, char **argv) {
- int DevNo = 0;
- int *Ptr = reinterpret_cast<int *>(llvm_omp_target_alloc_shared(4, DevNo));
- *Ptr = 7;
- printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
- // CHECK: Ptr [[Ptr:0x.*]], *Ptr: 7
+ int *Ptr;
+ cudaMalloc(&Ptr, 4);
+ printf("Ptr %p\n", Ptr);
+ // CHECK: Ptr [[Ptr:0x.*]]
square<<<1, 1>>>(Ptr);
- printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
- // CHECK: Ptr [[Ptr]], *Ptr: 42
- llvm_omp_target_free_shared(Ptr, DevNo);
+ int I = 0;
+ cudaDeviceSynchronize();
+ cudaMemcpy(&I, Ptr, sizeof(int), cudaMemcpyDeviceToHost);
+ printf("I: %i\n", I);
+ // CHECK: I: 42
}
diff --git a/offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu b/offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu
index a428e25d82359..0bfc9e231ffde 100644
--- a/offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu
+++ b/offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu
@@ -7,27 +7,25 @@
// UNSUPPORTED: aarch64-unknown-linux-gnu
// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
// UNSUPPORTED: intelgpu
#include <stdio.h>
-extern "C" {
-void *llvm_omp_target_alloc_shared(size_t Size, int DeviceNum);
-void llvm_omp_target_free_shared(void *DevicePtr, int DeviceNum);
-}
-
__global__ void square(int *A) {
__scoped_atomic_fetch_add(A, 1, __ATOMIC_SEQ_CST, __MEMORY_SCOPE_DEVICE);
}
int main(int argc, char **argv) {
int DevNo = 0;
- int *Ptr = reinterpret_cast<int *>(llvm_omp_target_alloc_shared(4, DevNo));
- *Ptr = 0;
- printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
- // CHECK: Ptr [[Ptr:0x.*]], *Ptr: 0
+ int *Ptr, I;
+ cudaMalloc(&Ptr, 4);
+ printf("Ptr %p\n", Ptr);
+ // CHECK: Ptr [[Ptr:0x.*]]
square<<<7, 6>>>(Ptr);
- printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
- // CHECK: Ptr [[Ptr]], *Ptr: 42
- llvm_omp_target_free_shared(Ptr, DevNo);
+ cudaMemcpy(&I, Ptr, sizeof(int), cudaMemcpyDeviceToHost);
+ printf("I: %i\n", I);
+ // CHECK: I: 42
}
diff --git a/offload/test/offloading/CUDA/basic_launch_multi_arg.cu b/offload/test/offloading/CUDA/basic_launch_multi_arg.cu
index db2a1e48371b0..505a9f9379c08 100644
--- a/offload/test/offloading/CUDA/basic_launch_multi_arg.cu
+++ b/offload/test/offloading/CUDA/basic_launch_multi_arg.cu
@@ -6,15 +6,13 @@
// clang-format on
// REQUIRES: gpu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
// UNSUPPORTED: intelgpu
#include <stdio.h>
-extern "C" {
-void *llvm_omp_target_alloc_shared(size_t Size, int DeviceNum);
-void llvm_omp_target_free_shared(void *DevicePtr, int DeviceNum);
-}
-
__global__ void square(int *Dst, short Q, int *Src, short P) {
*Dst = (Src[0] + Src[1]) * (Q + P);
Src[0] = Q;
@@ -23,19 +21,19 @@ __global__ void square(int *Dst, short Q, int *Src, short P) {
int main(int argc, char **argv) {
int DevNo = 0;
- int *Ptr = reinterpret_cast<int *>(llvm_omp_target_alloc_shared(4, DevNo));
- int *Src = reinterpret_cast<int *>(llvm_omp_target_alloc_shared(8, DevNo));
- *Ptr = 7;
- Src[0] = -2;
- Src[1] = 8;
- printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
- // CHECK: Ptr [[Ptr:0x.*]], *Ptr: 7
- printf("Src: %i : %i\n", Src[0], Src[1]);
- // CHECK: Src: -2 : 8
+ int *Src, *Ptr;
+ cudaMalloc(&Ptr, 4);
+ cudaMalloc(&Src, 8);
+
+ int I = 7;
+ int HostSrc[2] = {-2, 8};
+ cudaMemcpy(Ptr, &I, sizeof(int), cudaMemcpyHostToDevice);
+ cudaMemcpy(Src, &HostSrc[0], 2 * sizeof(int), cudaMemcpyHostToDevice);
square<<<1, 1>>>(Ptr, 3, Src, 4);
- printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
- // CHECK: Ptr [[Ptr]], *Ptr: 42
- printf("Src: %i : %i\n", Src[0], Src[1]);
- // CHECK: Src: 3 : 4
- llvm_omp_target_free_shared(Ptr, DevNo);
+ cudaMemcpy(&I, Ptr, sizeof(int), cudaMemcpyDeviceToHost);
+ cudaMemcpy(&HostSrc[0], Src, 2 * sizeof(int), cudaMemcpyDeviceToHost);
+ printf("I: %i\n", I);
+ // CHECK: I: 42
+ printf("Src: %i, %i\n", HostSrc[0], HostSrc[1]);
+ // CHECK: Src: 3, 4
}
diff --git a/offload/test/offloading/CUDA/device_api.cu b/offload/test/offloading/CUDA/device_api.cu
new file mode 100644
index 0000000000000..184167f2d17e4
--- /dev/null
+++ b/offload/test/offloading/CUDA/device_api.cu
@@ -0,0 +1,45 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+int main(int argc, char **argv) {
+ int Count = 0;
+ if (cudaGetDeviceCount(&Count) != cudaSuccess)
+ return 1;
+
+ printf("device count: %d\n", Count);
+ // CHECK: device count: {{[1-9][0-9]*}}
+
+ int Device = -1;
+ if (cudaGetDevice(&Device) != cudaSuccess)
+ return 1;
+
+ printf("device: %d\n", Device);
+ // CHECK: device: {{[0-9]+}}
+
+ if (cudaSetDevice(Device) != cudaSuccess)
+ return 1;
+
+ int After = -1;
+ if (cudaGetDevice(&After) != cudaSuccess)
+ return 1;
+
+ printf("device after set: %d\n", After);
+ // CHECK: device after set: {{[0-9]+}}
+
+ cudaError_t Err = cudaSetDevice(-1);
+ printf("set invalid device: %u\n", Err);
+ // CHECK: set invalid device: 1
+}
diff --git a/offload/test/offloading/CUDA/device_properties.cu b/offload/test/offloading/CUDA/device_properties.cu
new file mode 100644
index 0000000000000..8f625f6ccabe1
--- /dev/null
+++ b/offload/test/offloading/CUDA/device_properties.cu
@@ -0,0 +1,40 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+int main(int argc, char **argv) {
+ cudaDeviceProp Prop = {};
+ cudaError_t Err = cudaGetDeviceProperties(&Prop, 0);
+ if (Err != cudaSuccess) {
+ printf("cudaGetDeviceProperties failed: %u\n", Err);
+ return 1;
+ }
+
+ printf("Device name: %s\n", Prop.name);
+ // CHECK: Device name:
+ printf("Total global memory: %zu\n", Prop.totalGlobalMem);
+ // CHECK: Total global memory:
+ printf("Multiprocessors: %i\n", Prop.multiProcessorCount);
+ // CHECK: Multiprocessors:
+ printf("Warp size: %i\n", Prop.warpSize);
+ // CHECK: Warp size:
+
+ if (!Prop.name[0] || !Prop.totalGlobalMem || !Prop.multiProcessorCount ||
+ !Prop.warpSize)
+ return 1;
+
+ printf("Device properties are populated.\n");
+ // CHECK: Device properties are populated.
+}
diff --git a/offload/test/offloading/CUDA/host_alloc.cu b/offload/test/offloading/CUDA/host_alloc.cu
new file mode 100644
index 0000000000000..f23eac7604582
--- /dev/null
+++ b/offload/test/offloading/CUDA/host_alloc.cu
@@ -0,0 +1,40 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+int main(int argc, char **argv) {
+ int *HostAllocPtr = nullptr;
+ if (cudaHostAlloc(&HostAllocPtr, sizeof(int), cudaHostAllocDefault) !=
+ cudaSuccess)
+ return 1;
+
+ *HostAllocPtr = 17;
+ printf("cudaHostAlloc value: %d\n", *HostAllocPtr);
+ // CHECK: cudaHostAlloc value: 17
+
+ if (cudaFreeHost(HostAllocPtr) != cudaSuccess)
+ return 1;
+
+ int *MallocHostPtr = nullptr;
+ if (cudaMallocHost(&MallocHostPtr, sizeof(int)) != cudaSuccess)
+ return 1;
+
+ *MallocHostPtr = 23;
+ printf("cudaMallocHost value: %d\n", *MallocHostPtr);
+ // CHECK: cudaMallocHost value: 23
+
+ if (cudaFreeHost(MallocHostPtr) != cudaSuccess)
+ return 1;
+}
diff --git a/offload/test/offloading/CUDA/launch_tu.cu b/offload/test/offloading/CUDA/launch_tu.cu
index a46472b514a6c..8b92194ba435e 100644
--- a/offload/test/offloading/CUDA/launch_tu.cu
+++ b/offload/test/offloading/CUDA/launch_tu.cu
@@ -1,31 +1,30 @@
// clang-format off
// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t.launch_tu.o -c
// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda %S/kernel_tu.cu.inc -o %t.kernel_tu.o -c
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %t.launch_tu.o %t.kernel_tu.o -o %t
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native --offload-link %t.launch_tu.o %t.kernel_tu.o -o %t
// RUN: %t | %fcheck-generic
// clang-format on
// UNSUPPORTED: aarch64-unknown-linux-gnu
// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
// UNSUPPORTED: intelgpu
#include <stdio.h>
-extern "C" {
-void *llvm_omp_target_alloc_shared(size_t Size, int DeviceNum);
-void llvm_omp_target_free_shared(void *DevicePtr, int DeviceNum);
-}
-
extern __global__ void square(int *A);
int main(int argc, char **argv) {
int DevNo = 0;
- int *Ptr = reinterpret_cast<int *>(llvm_omp_target_alloc_shared(4, DevNo));
- *Ptr = 7;
- printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
- // CHECK: Ptr [[Ptr:0x.*]], *Ptr: 7
+ int *Ptr;
+ cudaMalloc(&Ptr, 4);
+ printf("Ptr %p\n", Ptr);
+ // CHECK: Ptr [[Ptr:0x.*]]
square<<<1, 1>>>(Ptr);
- printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
- // CHECK: Ptr [[Ptr]], *Ptr: 42
- llvm_omp_target_free_shared(Ptr, DevNo);
+ int I;
+ cudaMemcpy(&I, Ptr, sizeof(int), cudaMemcpyDeviceToHost);
+ printf("I: %i\n", I);
+ // CHECK: I: 42
}
diff --git a/offload/test/offloading/CUDA/memcpy_kinds.cu b/offload/test/offloading/CUDA/memcpy_kinds.cu
new file mode 100644
index 0000000000000..a4288ee51ee3b
--- /dev/null
+++ b/offload/test/offloading/CUDA/memcpy_kinds.cu
@@ -0,0 +1,51 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+int main(int argc, char **argv) {
+ int HostSrc = 11;
+ int HostDst = 0;
+ if (cudaMemcpy(&HostDst, &HostSrc, sizeof(int), cudaMemcpyHostToHost) !=
+ cudaSuccess)
+ return 1;
+
+ printf("host to host: %d\n", HostDst);
+ // CHECK: host to host: 11
+
+ int *DevSrc = nullptr;
+ int *DevDst = nullptr;
+ int Result = 0;
+ if (cudaMalloc(&DevSrc, sizeof(int)) != cudaSuccess)
+ return 1;
+ if (cudaMalloc(&DevDst, sizeof(int)) != cudaSuccess)
+ return 1;
+
+ HostSrc = 42;
+ if (cudaMemcpy(DevSrc, &HostSrc, sizeof(int), cudaMemcpyHostToDevice) !=
+ cudaSuccess)
+ return 1;
+ if (cudaMemcpy(DevDst, DevSrc, sizeof(int), cudaMemcpyDeviceToDevice) !=
+ cudaSuccess)
+ return 1;
+ if (cudaMemcpy(&Result, DevDst, sizeof(int), cudaMemcpyDeviceToHost) !=
+ cudaSuccess)
+ return 1;
+
+ printf("device to device: %d\n", Result);
+ // CHECK: device to device: 42
+
+ cudaFree(DevSrc);
+ cudaFree(DevDst);
+}
diff --git a/offload/test/offloading/CUDA/stream_api.cu b/offload/test/offloading/CUDA/stream_api.cu
new file mode 100644
index 0000000000000..7202751f8207e
--- /dev/null
+++ b/offload/test/offloading/CUDA/stream_api.cu
@@ -0,0 +1,46 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+__global__ void setValue(int *Out) { *Out = 42; }
+
+int main(int argc, char **argv) {
+ cudaStream_t Stream = nullptr;
+ if (cudaStreamCreate(&Stream) != cudaSuccess)
+ return 1;
+
+ printf("stream created: %d\n", Stream != nullptr);
+ // CHECK: stream created: 1
+
+ int *DevPtr = nullptr;
+ int Result = 0;
+ if (cudaMalloc(&DevPtr, sizeof(int)) != cudaSuccess)
+ return 1;
+
+ setValue<<<1, 1, 0, Stream>>>(DevPtr);
+
+ if (cudaStreamSynchronize(Stream) != cudaSuccess)
+ return 1;
+ if (cudaMemcpy(&Result, DevPtr, sizeof(int), cudaMemcpyDeviceToHost) !=
+ cudaSuccess)
+ return 1;
+
+ printf("stream result: %d\n", Result);
+ // CHECK: stream result: 42
+
+ if (cudaStreamDestroy(Stream) != cudaSuccess)
+ return 1;
+ cudaFree(DevPtr);
+}
diff --git a/offload/test/offloading/CUDA/syncthreads.cu b/offload/test/offloading/CUDA/syncthreads.cu
new file mode 100644
index 0000000000000..0c6048c32f824
--- /dev/null
+++ b/offload/test/offloading/CUDA/syncthreads.cu
@@ -0,0 +1,40 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+__global__ void reduceBlock(int *Out) {
+ __shared__ int Scratch[64];
+ int Tid = threadIdx.x;
+ Scratch[Tid] = Tid;
+ __syncthreads();
+
+ if (Tid == 0) {
+ int Sum = 0;
+ for (int I = 0; I < 64; ++I)
+ Sum += Scratch[I];
+ Out[0] = Sum;
+ }
+}
+
+int main(int argc, char **argv) {
+ int *DevPtr;
+ int Result = 0;
+ cudaMalloc(&DevPtr, sizeof(int));
+ reduceBlock<<<1, 64>>>(DevPtr);
+ cudaMemcpy(&Result, DevPtr, sizeof(int), cudaMemcpyDeviceToHost);
+
+ printf("sum: %i\n", Result);
+ // CHECK: sum: 2016
+}
diff --git a/offload/test/offloading/CUDA/thread_and_block_id.cu b/offload/test/offloading/CUDA/thread_and_block_id.cu
new file mode 100644
index 0000000000000..30b87659d2eed
--- /dev/null
+++ b/offload/test/offloading/CUDA/thread_and_block_id.cu
@@ -0,0 +1,44 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: aarch64-unknown-linux-gnu-LTO
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu-LTO
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+
+#include <stdio.h>
+#include <stdlib.h>
+
+__global__ void fill(int *A) {
+ int tid = threadIdx.x + blockDim.x * blockIdx.x;
+ A[tid] = 42;
+}
+
+int main(int argc, char **argv) {
+ int NThreads = 128;
+ int NBlocks = 512;
+ int Size = sizeof(int) * NThreads * NBlocks;
+ int *Ptr = (int *)calloc(1, Size);
+ int *DevPtr;
+ cudaMalloc(&DevPtr, Size);
+ cudaMemcpy(DevPtr, Ptr, Size, cudaMemcpyHostToDevice);
+ printf("DevPtr %p\n", DevPtr);
+ // CHECK: DevPtr [[DevPtr:0x.*]]
+ fill<<<NBlocks, NThreads>>>(DevPtr);
+ cudaMemcpy(Ptr, DevPtr, Size, cudaMemcpyDeviceToHost);
+
+ for (int I = 0; I < NBlocks * NThreads; ++I) {
+ if (Ptr[I] == 42)
+ continue;
+ printf("Error at %i: %i vs %i\n", I, Ptr[I], 42);
+ return 1;
+ }
+ return 0;
+}
diff --git a/offload/test/offloading/HIP/basic_launch.hip b/offload/test/offloading/HIP/basic_launch.hip
new file mode 100644
index 0000000000000..bd2f2a6078671
--- /dev/null
+++ b/offload/test/offloading/HIP/basic_launch.hip
@@ -0,0 +1,30 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+__global__ void square(int *A) { *A = 42; }
+
+int main(int argc, char **argv) {
+ int *Ptr;
+ hipMalloc(&Ptr, 4);
+ printf("Ptr %p\n", Ptr);
+ // CHECK: Ptr [[Ptr:0x.*]]
+ square<<<1, 1>>>(Ptr);
+ int I = 0;
+ hipDeviceSynchronize();
+ hipMemcpy(&I, Ptr, sizeof(int), hipMemcpyDeviceToHost);
+ printf("I: %i\n", I);
+ // CHECK: I: 42
+}
diff --git a/offload/test/offloading/HIP/basic_launch_blocks_and_threads.hip b/offload/test/offloading/HIP/basic_launch_blocks_and_threads.hip
new file mode 100644
index 0000000000000..344b98b1636f1
--- /dev/null
+++ b/offload/test/offloading/HIP/basic_launch_blocks_and_threads.hip
@@ -0,0 +1,31 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+__global__ void square(int *A) {
+ __scoped_atomic_fetch_add(A, 1, __ATOMIC_SEQ_CST, __MEMORY_SCOPE_DEVICE);
+}
+
+int main(int argc, char **argv) {
+ int DevNo = 0;
+ int *Ptr, I;
+ hipMalloc(&Ptr, 4);
+ printf("Ptr %p\n", Ptr);
+ // CHECK: Ptr [[Ptr:0x.*]]
+ square<<<7, 6>>>(Ptr);
+ hipMemcpy(&I, Ptr, sizeof(int), hipMemcpyDeviceToHost);
+ printf("I: %i\n", I);
+ // CHECK: I: 42
+}
diff --git a/offload/test/offloading/HIP/basic_launch_multi_arg.hip b/offload/test/offloading/HIP/basic_launch_multi_arg.hip
new file mode 100644
index 0000000000000..6e599d6704598
--- /dev/null
+++ b/offload/test/offloading/HIP/basic_launch_multi_arg.hip
@@ -0,0 +1,39 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// REQUIRES: gpu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+__global__ void square(int *Dst, short Q, int *Src, short P) {
+ *Dst = (Src[0] + Src[1]) * (Q + P);
+ Src[0] = Q;
+ Src[1] = P;
+}
+
+int main(int argc, char **argv) {
+ int DevNo = 0;
+ int *Src, *Ptr;
+ hipMalloc(&Ptr, 4);
+ hipMalloc(&Src, 8);
+
+ int I = 7;
+ int HostSrc[2] = {-2,8};
+ hipMemcpy(Ptr, &I, sizeof(int), hipMemcpyHostToDevice);
+ hipMemcpy(Src, &HostSrc[0], 2*sizeof(int), hipMemcpyHostToDevice);
+ square<<<1, 1>>>(Ptr, 3, Src, 4);
+ hipMemcpy(&I, Ptr, sizeof(int), hipMemcpyDeviceToHost);
+ hipMemcpy(&HostSrc[0], Src, 2 * sizeof(int), hipMemcpyDeviceToHost);
+ printf("I: %i\n", I);
+ // CHECK: I: 42
+ printf("Src: %i, %i\n", HostSrc[0], HostSrc[1]);
+ // CHECK: Src: 3, 4
+}
diff --git a/offload/test/offloading/HIP/device_api.hip b/offload/test/offloading/HIP/device_api.hip
new file mode 100644
index 0000000000000..031e3703e66c1
--- /dev/null
+++ b/offload/test/offloading/HIP/device_api.hip
@@ -0,0 +1,45 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+int main(int argc, char **argv) {
+ int Count = 0;
+ if (hipGetDeviceCount(&Count) != hipSuccess)
+ return 1;
+
+ printf("device count: %d\n", Count);
+ // CHECK: device count: {{[1-9][0-9]*}}
+
+ int Device = -1;
+ if (hipGetDevice(&Device) != hipSuccess)
+ return 1;
+
+ printf("device: %d\n", Device);
+ // CHECK: device: {{[0-9]+}}
+
+ if (hipSetDevice(Device) != hipSuccess)
+ return 1;
+
+ int After = -1;
+ if (hipGetDevice(&After) != hipSuccess)
+ return 1;
+
+ printf("device after set: %d\n", After);
+ // CHECK: device after set: {{[0-9]+}}
+
+ hipError_t Err = hipSetDevice(-1);
+ printf("set invalid device: %u\n", Err);
+ // CHECK: set invalid device: 1
+}
diff --git a/offload/test/offloading/HIP/device_properties.hip b/offload/test/offloading/HIP/device_properties.hip
new file mode 100644
index 0000000000000..1a9b9a70f8ea9
--- /dev/null
+++ b/offload/test/offloading/HIP/device_properties.hip
@@ -0,0 +1,40 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+int main(int argc, char **argv) {
+ hipDeviceProp_t Prop = {};
+ hipError_t Err = hipGetDeviceProperties(&Prop, 0);
+ if (Err != hipSuccess) {
+ printf("hipGetDeviceProperties failed: %u\n", Err);
+ return 1;
+ }
+
+ printf("Device name: %s\n", Prop.name);
+ // CHECK: Device name:
+ printf("Total global memory: %zu\n", Prop.totalGlobalMem);
+ // CHECK: Total global memory:
+ printf("Multiprocessors: %i\n", Prop.multiProcessorCount);
+ // CHECK: Multiprocessors:
+ printf("Warp size: %i\n", Prop.warpSize);
+ // CHECK: Warp size:
+
+ if (!Prop.name[0] || !Prop.totalGlobalMem || !Prop.multiProcessorCount ||
+ !Prop.warpSize)
+ return 1;
+
+ printf("Device properties are populated.\n");
+ // CHECK: Device properties are populated.
+}
diff --git a/offload/test/offloading/HIP/host_alloc.hip b/offload/test/offloading/HIP/host_alloc.hip
new file mode 100644
index 0000000000000..8b067b39f2f81
--- /dev/null
+++ b/offload/test/offloading/HIP/host_alloc.hip
@@ -0,0 +1,40 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+int main(int argc, char **argv) {
+ int *HostAllocPtr = nullptr;
+ if (hipHostAlloc(&HostAllocPtr, sizeof(int), hipHostAllocDefault) !=
+ hipSuccess)
+ return 1;
+
+ *HostAllocPtr = 17;
+ printf("hipHostAlloc value: %d\n", *HostAllocPtr);
+ // CHECK: hipHostAlloc value: 17
+
+ if (hipFreeHost(HostAllocPtr) != hipSuccess)
+ return 1;
+
+ int *MallocHostPtr = nullptr;
+ if (hipMallocHost(&MallocHostPtr, sizeof(int)) != hipSuccess)
+ return 1;
+
+ *MallocHostPtr = 23;
+ printf("hipMallocHost value: %d\n", *MallocHostPtr);
+ // CHECK: hipMallocHost value: 23
+
+ if (hipFreeHost(MallocHostPtr) != hipSuccess)
+ return 1;
+}
diff --git a/offload/test/offloading/HIP/kernel_tu.hip.inc b/offload/test/offloading/HIP/kernel_tu.hip.inc
new file mode 100644
index 0000000000000..d7d28a109dfc5
--- /dev/null
+++ b/offload/test/offloading/HIP/kernel_tu.hip.inc
@@ -0,0 +1 @@
+__global__ void square(int *A) { *A = 42; }
diff --git a/offload/test/offloading/HIP/launch_tu.hip b/offload/test/offloading/HIP/launch_tu.hip
new file mode 100644
index 0000000000000..03073029ca211
--- /dev/null
+++ b/offload/test/offloading/HIP/launch_tu.hip
@@ -0,0 +1,30 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t.launch_tu.o -c
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip %S/kernel_tu.hip.inc -o %t.kernel_tu.o -c
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native --offload-link %t.launch_tu.o %t.kernel_tu.o -o %t
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+extern __global__ void square(int *A);
+
+int main(int argc, char **argv) {
+ int DevNo = 0;
+ int *Ptr;
+ hipMalloc(&Ptr, 4);
+ printf("Ptr %p\n", Ptr);
+ // CHECK: Ptr [[Ptr:0x.*]]
+ square<<<1, 1>>>(Ptr);
+ int I;
+ hipMemcpy(&I, Ptr, sizeof(int), hipMemcpyDeviceToHost);
+ printf("I: %i\n", I);
+ // CHECK: I: 42
+}
diff --git a/offload/test/offloading/HIP/memcpy_kinds.hip b/offload/test/offloading/HIP/memcpy_kinds.hip
new file mode 100644
index 0000000000000..6755a55aa0794
--- /dev/null
+++ b/offload/test/offloading/HIP/memcpy_kinds.hip
@@ -0,0 +1,51 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+int main(int argc, char **argv) {
+ int HostSrc = 11;
+ int HostDst = 0;
+ if (hipMemcpy(&HostDst, &HostSrc, sizeof(int), hipMemcpyHostToHost) !=
+ hipSuccess)
+ return 1;
+
+ printf("host to host: %d\n", HostDst);
+ // CHECK: host to host: 11
+
+ int *DevSrc = nullptr;
+ int *DevDst = nullptr;
+ int Result = 0;
+ if (hipMalloc(&DevSrc, sizeof(int)) != hipSuccess)
+ return 1;
+ if (hipMalloc(&DevDst, sizeof(int)) != hipSuccess)
+ return 1;
+
+ HostSrc = 42;
+ if (hipMemcpy(DevSrc, &HostSrc, sizeof(int), hipMemcpyHostToDevice) !=
+ hipSuccess)
+ return 1;
+ if (hipMemcpy(DevDst, DevSrc, sizeof(int), hipMemcpyDeviceToDevice) !=
+ hipSuccess)
+ return 1;
+ if (hipMemcpy(&Result, DevDst, sizeof(int), hipMemcpyDeviceToHost) !=
+ hipSuccess)
+ return 1;
+
+ printf("device to device: %d\n", Result);
+ // CHECK: device to device: 42
+
+ hipFree(DevSrc);
+ hipFree(DevDst);
+}
diff --git a/offload/test/offloading/HIP/stream_api.hip b/offload/test/offloading/HIP/stream_api.hip
new file mode 100644
index 0000000000000..c0e2699822814
--- /dev/null
+++ b/offload/test/offloading/HIP/stream_api.hip
@@ -0,0 +1,46 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+__global__ void setValue(int *Out) { *Out = 42; }
+
+int main(int argc, char **argv) {
+ hipStream_t Stream = nullptr;
+ if (hipStreamCreate(&Stream) != hipSuccess)
+ return 1;
+
+ printf("stream created: %d\n", Stream != nullptr);
+ // CHECK: stream created: 1
+
+ int *DevPtr = nullptr;
+ int Result = 0;
+ if (hipMalloc(&DevPtr, sizeof(int)) != hipSuccess)
+ return 1;
+
+ setValue<<<1, 1, 0, Stream>>>(DevPtr);
+
+ if (hipStreamSynchronize(Stream) != hipSuccess)
+ return 1;
+ if (hipMemcpy(&Result, DevPtr, sizeof(int), hipMemcpyDeviceToHost) !=
+ hipSuccess)
+ return 1;
+
+ printf("stream result: %d\n", Result);
+ // CHECK: stream result: 42
+
+ if (hipStreamDestroy(Stream) != hipSuccess)
+ return 1;
+ hipFree(DevPtr);
+}
diff --git a/offload/test/offloading/HIP/syncthreads.hip b/offload/test/offloading/HIP/syncthreads.hip
new file mode 100644
index 0000000000000..5962ab5468b86
--- /dev/null
+++ b/offload/test/offloading/HIP/syncthreads.hip
@@ -0,0 +1,40 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+__global__ void reduceBlock(int *Out) {
+ __shared__ int Scratch[64];
+ int Tid = threadIdx.x;
+ Scratch[Tid] = Tid;
+ __syncthreads();
+
+ if (Tid == 0) {
+ int Sum = 0;
+ for (int I = 0; I < 64; ++I)
+ Sum += Scratch[I];
+ Out[0] = Sum;
+ }
+}
+
+int main(int argc, char **argv) {
+ int *DevPtr;
+ int Result = 0;
+ hipMalloc(&DevPtr, sizeof(int));
+ reduceBlock<<<1, 64>>>(DevPtr);
+ hipMemcpy(&Result, DevPtr, sizeof(int), hipMemcpyDeviceToHost);
+
+ printf("sum: %i\n", Result);
+ // CHECK: sum: 2016
+}
diff --git a/offload/test/offloading/HIP/thread_and_block_id.hip b/offload/test/offloading/HIP/thread_and_block_id.hip
new file mode 100644
index 0000000000000..c9c33c55d72fb
--- /dev/null
+++ b/offload/test/offloading/HIP/thread_and_block_id.hip
@@ -0,0 +1,44 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: aarch64-unknown-linux-gnu-LTO
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu-LTO
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+
+#include <stdio.h>
+#include <stdlib.h>
+
+__global__ void fill(int *A) {
+ int tid = threadIdx.x + blockDim.x * blockIdx.x;
+ A[tid] = 42;
+}
+
+int main(int argc, char **argv) {
+ int NThreads = 128;
+ int NBlocks = 512;
+ int Size = sizeof(int) * NThreads * NBlocks;
+ int *Ptr = (int*)calloc(1, Size);
+ int *DevPtr;
+ hipMalloc(&DevPtr, Size);
+ hipMemcpy(DevPtr, Ptr, Size, hipMemcpyHostToDevice);
+ printf("DevPtr %p\n", DevPtr);
+ // CHECK: DevPtr [[DevPtr:0x.*]]
+ fill<<<NBlocks, NThreads>>>(DevPtr);
+ hipMemcpy(Ptr, DevPtr, Size, hipMemcpyDeviceToHost);
+
+ for (int I = 0; I < NBlocks * NThreads; ++I) {
+ if (Ptr[I] == 42)
+ continue;
+ printf("Error at %i: %i vs %i\n", I, Ptr[I], 42);
+ return 1;
+ }
+ return 0;
+}
>From 770c550c494b834fc19d3283c600f01c4e314505 Mon Sep 17 00:00:00 2001
From: Sophia Herrmann <herrmann15 at llnl.gov>
Date: Wed, 29 Jul 2026 15:12:50 -0700
Subject: [PATCH 7/7] add getErrorName + getErrorString + Test
---
.../include/kernel/DefineLanguageNames.inc | 4 ++
.../languages/include/kernel/LanguageErrors.h | 25 ++++++++
.../include/kernel/LanguageRuntime.h | 5 +-
.../include/kernel/UndefineLanguageNames.inc | 4 ++
offload/languages/kernel/CMakeLists.txt | 2 +
.../languages/kernel/include/LanguageUtils.h | 43 +++++++++++++
.../languages/kernel/src/LanguageErrors.cpp | 51 ++++++++++++++++
.../languages/kernel/src/LanguageRuntime.cpp | 30 +++------
offload/test/offloading/CUDA/device_api.cu | 2 +-
offload/test/offloading/CUDA/error_kinds.cu | 61 +++++++++++++++++++
offload/test/offloading/HIP/device_api.hip | 2 +-
offload/test/offloading/HIP/error_kinds.hip | 61 +++++++++++++++++++
12 files changed, 261 insertions(+), 29 deletions(-)
create mode 100644 offload/languages/include/kernel/LanguageErrors.h
create mode 100644 offload/languages/kernel/include/LanguageUtils.h
create mode 100644 offload/languages/kernel/src/LanguageErrors.cpp
create mode 100644 offload/test/offloading/CUDA/error_kinds.cu
create mode 100644 offload/test/offloading/HIP/error_kinds.hip
diff --git a/offload/languages/include/kernel/DefineLanguageNames.inc b/offload/languages/include/kernel/DefineLanguageNames.inc
index cfab6998c28e6..989d119ae6ed7 100644
--- a/offload/languages/include/kernel/DefineLanguageNames.inc
+++ b/offload/languages/include/kernel/DefineLanguageNames.inc
@@ -17,6 +17,10 @@
#define DeviceSynchronize COMBINE(LANGUAGE, DeviceSynchronize)
#define Success COMBINE(LANGUAGE, Success)
#define ErrorInvalidValue COMBINE(LANGUAGE, ErrorInvalidValue)
+#define ErrorInvalidDevice COMBINE(LANGUAGE, ErrorInvalidDevice)
+#define ErrorUnknown COMBINE(LANGUAGE, ErrorUnknown)
+#define GetErrorName COMBINE(LANGUAGE, GetErrorName)
+#define GetErrorString COMBINE(LANGUAGE, GetErrorString)
#define MemcpyKind COMBINE(LANGUAGE, MemcpyKind)
#define MemcpyHostToHost COMBINE(LANGUAGE, MemcpyHostToHost)
#define MemcpyHostToDevice COMBINE(LANGUAGE, MemcpyHostToDevice)
diff --git a/offload/languages/include/kernel/LanguageErrors.h b/offload/languages/include/kernel/LanguageErrors.h
new file mode 100644
index 0000000000000..0c2a83cf5ef5b
--- /dev/null
+++ b/offload/languages/include/kernel/LanguageErrors.h
@@ -0,0 +1,25 @@
+//===-- LanguageErrors.h - Kernel language error API declarations ---------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_OFFLOAD_LANGUAGES_INCLUDE_KERNEL_LANGUAGE_ERRORS_H
+#define LLVM_OFFLOAD_LANGUAGES_INCLUDE_KERNEL_LANGUAGE_ERRORS_H
+
+#include <cstdint>
+
+enum Error_t : uint32_t {
+ Success = 0,
+ ErrorInvalidValue = 1,
+ ErrorInvalidDevice = 2,
+ ErrorUnknown = 3,
+};
+
+const char *GetErrorName(Error_t Error);
+
+const char *GetErrorString(Error_t Error);
+
+#endif // LLVM_OFFLOAD_LANGUAGES_INCLUDE_KERNEL_LANGUAGE_ERRORS_H
diff --git a/offload/languages/include/kernel/LanguageRuntime.h b/offload/languages/include/kernel/LanguageRuntime.h
index b072087ba0e7d..0b7f6e4778bfb 100644
--- a/offload/languages/include/kernel/LanguageRuntime.h
+++ b/offload/languages/include/kernel/LanguageRuntime.h
@@ -14,10 +14,7 @@
#include <cstdio>
#include <cstdlib>
-enum Error_t : uint32_t {
- Success = 0,
- ErrorInvalidValue = 1,
-};
+#include "LanguageErrors.h"
struct DeviceProp_t {
char name[256];
diff --git a/offload/languages/include/kernel/UndefineLanguageNames.inc b/offload/languages/include/kernel/UndefineLanguageNames.inc
index c1cc75719a18c..b2a9f2fd846ae 100644
--- a/offload/languages/include/kernel/UndefineLanguageNames.inc
+++ b/offload/languages/include/kernel/UndefineLanguageNames.inc
@@ -14,6 +14,10 @@
#undef DeviceSynchronize
#undef Success
#undef ErrorInvalidValue
+#undef ErrorInvalidDevice
+#undef ErrorUnknown
+#undef GetErrorName
+#undef GetErrorString
#undef MemcpyKind
#undef MemcpyHostToHost
#undef MemcpyHostToDevice
diff --git a/offload/languages/kernel/CMakeLists.txt b/offload/languages/kernel/CMakeLists.txt
index 81442f9f2c507..2a344bd1e211c 100644
--- a/offload/languages/kernel/CMakeLists.txt
+++ b/offload/languages/kernel/CMakeLists.txt
@@ -13,6 +13,7 @@ set(LLVM_OFFLOAD_KERNEL_INCLUDE_DIRS
function(add_llvm_offload_kernel_language_runtime_objects target language)
add_library(${target} OBJECT
src/LanguageRuntime.cpp
+ src/LanguageErrors.cpp
)
add_dependencies(${target} OffloadAPI)
target_include_directories(${target} PRIVATE ${LLVM_OFFLOAD_KERNEL_INCLUDE_DIRS})
@@ -73,6 +74,7 @@ install(TARGETS LLVMOffloadKernel
install(FILES
${CMAKE_CURRENT_SOURCE_DIR}/../include/kernel/DefineLanguageNames.inc
+ ${CMAKE_CURRENT_SOURCE_DIR}/../include/kernel/LanguageErrors.h
${CMAKE_CURRENT_SOURCE_DIR}/../include/kernel/LanguageRuntime.h
${CMAKE_CURRENT_SOURCE_DIR}/../include/kernel/UndefineLanguageNames.inc
DESTINATION ${CMAKE_INSTALL_PREFIX}/include/offload/kernel/)
diff --git a/offload/languages/kernel/include/LanguageUtils.h b/offload/languages/kernel/include/LanguageUtils.h
new file mode 100644
index 0000000000000..3b937dce16356
--- /dev/null
+++ b/offload/languages/kernel/include/LanguageUtils.h
@@ -0,0 +1,43 @@
+//===-- LanguageUtils.h - Kernel Language utility functions ---------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_LANGUAGE_UTILS_H
+#define LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_LANGUAGE_UTILS_H
+
+#ifndef LANGUAGE
+#error This file should be included, or used, with a LANGUAGE macro set.
+#endif
+
+#include "OffloadAPI.h"
+
+/// Convert an ol_result_t to the active language's Error_t.
+static inline Error_t convertResult(ol_result_t Result) {
+ if (Result == OL_SUCCESS)
+ return Success;
+ switch (Result->Code) {
+ case OL_ERRC_INVALID_VALUE:
+ case OL_ERRC_INVALID_ARGUMENT:
+ case OL_ERRC_INVALID_NULL_POINTER:
+ return ErrorInvalidValue;
+ case OL_ERRC_INVALID_DEVICE:
+ return ErrorInvalidDevice;
+ default:
+ return ErrorUnknown;
+ }
+}
+
+/// Convert a Stream_t to an ol_queue_handle_t.
+static inline Error_t getQueueFromStream(Stream_t Stream,
+ ol_queue_handle_t *Queue) {
+ if (!Stream)
+ return ErrorInvalidValue;
+ *Queue = reinterpret_cast<ol_queue_handle_t>(Stream);
+ return Success;
+}
+
+#endif // LLVM_OFFLOAD_LANGUAGES_KERNEL_INCLUDE_LANGUAGE_UTILS_H
diff --git a/offload/languages/kernel/src/LanguageErrors.cpp b/offload/languages/kernel/src/LanguageErrors.cpp
new file mode 100644
index 0000000000000..f725e5f029d28
--- /dev/null
+++ b/offload/languages/kernel/src/LanguageErrors.cpp
@@ -0,0 +1,51 @@
+//===-- LanguageErrors.cpp - Kernel language error API implementation -----===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LANGUAGE
+#error This file should be included, or used, with a LANGUAGE macro set.
+#endif
+
+// Rename the generic error API before declaring or defining language symbols.
+// clang-format off
+#include "DefineLanguageNames.inc"
+#include "LanguageErrors.h"
+// clang-format on
+
+const char *GetErrorName(Error_t Error) {
+ switch (Error) {
+#define LLVM_OFFLOAD_STRINGIFY_IMPL(NAME) #NAME
+#define LLVM_OFFLOAD_STRINGIFY(NAME) LLVM_OFFLOAD_STRINGIFY_IMPL(NAME)
+#define LLVM_OFFLOAD_ERR_STR(NAME) \
+ case NAME: \
+ return LLVM_OFFLOAD_STRINGIFY(NAME);
+ LLVM_OFFLOAD_ERR_STR(Success)
+ LLVM_OFFLOAD_ERR_STR(ErrorInvalidValue)
+ LLVM_OFFLOAD_ERR_STR(ErrorInvalidDevice)
+#undef LLVM_OFFLOAD_ERR_STR
+#undef LLVM_OFFLOAD_STRINGIFY
+#undef LLVM_OFFLOAD_STRINGIFY_IMPL
+ default:
+ return "Unrecognized error";
+ };
+}
+
+const char *GetErrorString(Error_t Error) {
+ switch (Error) {
+ case Success:
+ return "No error";
+ case ErrorInvalidValue:
+ return "Invalid argument value";
+ case ErrorInvalidDevice:
+ return "Invalid device number";
+ case ErrorUnknown:
+ return "Unknown error";
+ }
+ return "Unrecognized error";
+}
+
+#include "UndefineLanguageNames.inc"
diff --git a/offload/languages/kernel/src/LanguageRuntime.cpp b/offload/languages/kernel/src/LanguageRuntime.cpp
index 058eecf2c97ee..57d786448dabb 100644
--- a/offload/languages/kernel/src/LanguageRuntime.cpp
+++ b/offload/languages/kernel/src/LanguageRuntime.cpp
@@ -16,6 +16,7 @@
#include "LanguageRuntime.h"
// clang-format on
+#include "LanguageUtils.h"
#include "State.h"
#include "Types.h"
@@ -32,17 +33,6 @@
using RuntimeState = llvm::offload::StateTy;
using ThreadState = llvm::offload::ThreadStateTy;
-static Error_t convertResult(ol_result_t Result) {
- if (Result == OL_SUCCESS)
- return Success;
- switch (Result->Code) {
- case OL_ERRC_INVALID_VALUE:
- return ErrorInvalidValue;
- default:
- return ErrorInvalidValue;
- }
-}
-
Error_t Malloc(void **DevPtr, size_t Size) {
ol_device_handle_t Device = ThreadState::getDefaultDevice();
ol_result_t Result = olMemAlloc(Device, OL_ALLOC_TYPE_DEVICE, Size, DevPtr);
@@ -105,7 +95,7 @@ Error_t DeviceSynchronize() {
Error_t GetDevice(int *DeviceNo) {
ol_device_handle_t Device = RuntimeState::getDevice(DeviceNo);
if (!Device)
- return ErrorInvalidValue;
+ return ErrorInvalidDevice;
return Success;
}
@@ -117,7 +107,7 @@ Error_t GetDeviceCount(int *Count) {
Error_t SetDevice(int DeviceNo) {
ol_device_handle_t Device = RuntimeState::setDefaultDevice(DeviceNo);
if (!Device)
- return ErrorInvalidValue;
+ return ErrorInvalidDevice;
assert(Device == ThreadState::getDefaultDevice() &&
"Set Device is not Default Device");
return Success;
@@ -154,18 +144,12 @@ Error_t GetDeviceProperties(DeviceProp_t *DeviceProp, int DeviceNo) {
return Success;
}
-static Error_t getQueueFromStream(Stream_t Stream, ol_queue_handle_t *Queue) {
- if (!Stream)
- return ErrorInvalidValue;
- *Queue = reinterpret_cast<ol_queue_handle_t>(Stream);
- return Success;
-}
-
Error_t StreamCreate(Stream_t *Stream) {
ol_queue_handle_t Queue;
- olCreateQueue(ThreadState::getDefaultDevice(), &Queue);
- *Stream = reinterpret_cast<Stream_t>(Queue);
- return Success;
+ ol_result_t Result = olCreateQueue(ThreadState::getDefaultDevice(), &Queue);
+ if (Result == OL_SUCCESS)
+ *Stream = reinterpret_cast<Stream_t>(Queue);
+ return convertResult(Result);
}
Error_t StreamDestroy(Stream_t Stream) {
diff --git a/offload/test/offloading/CUDA/device_api.cu b/offload/test/offloading/CUDA/device_api.cu
index 184167f2d17e4..af2b046eee397 100644
--- a/offload/test/offloading/CUDA/device_api.cu
+++ b/offload/test/offloading/CUDA/device_api.cu
@@ -41,5 +41,5 @@ int main(int argc, char **argv) {
cudaError_t Err = cudaSetDevice(-1);
printf("set invalid device: %u\n", Err);
- // CHECK: set invalid device: 1
+ // CHECK: set invalid device: 2
}
diff --git a/offload/test/offloading/CUDA/error_kinds.cu b/offload/test/offloading/CUDA/error_kinds.cu
new file mode 100644
index 0000000000000..5fa3035fe52f0
--- /dev/null
+++ b/offload/test/offloading/CUDA/error_kinds.cu
@@ -0,0 +1,61 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+static void print_error(const char *Label, cudaError_t Error) {
+ printf("%s value: %u\n", Label, static_cast<unsigned>(Error));
+ printf("%s name: %s\n", Label, cudaGetErrorName(Error));
+ printf("%s string: %s\n", Label, cudaGetErrorString(Error));
+}
+
+int main() {
+ print_error("success", cudaSuccess);
+ // CHECK: success value: 0
+ // CHECK: success name: cudaSuccess
+ // CHECK: success string: No error
+
+ print_error("invalid value", cudaErrorInvalidValue);
+ // CHECK: invalid value value: 1
+ // CHECK: invalid value name: cudaErrorInvalidValue
+ // CHECK: invalid value string: Invalid argument value
+
+ print_error("invalid device", cudaErrorInvalidDevice);
+ // CHECK: invalid device value: 2
+ // CHECK: invalid device name: cudaErrorInvalidDevice
+ // CHECK: invalid device string: Invalid device number
+
+ print_error("unknown", cudaErrorUnknown);
+ // CHECK: unknown value: 3
+ // CHECK: unknown name: Unrecognized error
+ // CHECK: unknown string: Unknown error
+
+ cudaError_t Unrecognized = static_cast<cudaError_t>(999);
+ print_error("unrecognized", Unrecognized);
+ // CHECK: unrecognized value: 999
+ // CHECK: unrecognized name: Unrecognized error
+ // CHECK: unrecognized string: Unrecognized error
+
+ print_error("set invalid device", cudaSetDevice(-1));
+ // CHECK: set invalid device value: 2
+ // CHECK: set invalid device name: cudaErrorInvalidDevice
+ // CHECK: set invalid device string: Invalid device number
+
+ print_error("null stream destroy", cudaStreamDestroy(nullptr));
+ // CHECK: null stream destroy value: 1
+ // CHECK: null stream destroy name: cudaErrorInvalidValue
+ // CHECK: null stream destroy string: Invalid argument value
+
+ return 0;
+}
diff --git a/offload/test/offloading/HIP/device_api.hip b/offload/test/offloading/HIP/device_api.hip
index 031e3703e66c1..5fb66e6e45eeb 100644
--- a/offload/test/offloading/HIP/device_api.hip
+++ b/offload/test/offloading/HIP/device_api.hip
@@ -41,5 +41,5 @@ int main(int argc, char **argv) {
hipError_t Err = hipSetDevice(-1);
printf("set invalid device: %u\n", Err);
- // CHECK: set invalid device: 1
+ // CHECK: set invalid device: 2
}
diff --git a/offload/test/offloading/HIP/error_kinds.hip b/offload/test/offloading/HIP/error_kinds.hip
new file mode 100644
index 0000000000000..af760a4cd4399
--- /dev/null
+++ b/offload/test/offloading/HIP/error_kinds.hip
@@ -0,0 +1,61 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+static void print_error(const char *Label, hipError_t Error) {
+ printf("%s value: %u\n", Label, static_cast<unsigned>(Error));
+ printf("%s name: %s\n", Label, hipGetErrorName(Error));
+ printf("%s string: %s\n", Label, hipGetErrorString(Error));
+}
+
+int main() {
+ print_error("success", hipSuccess);
+ // CHECK: success value: 0
+ // CHECK: success name: hipSuccess
+ // CHECK: success string: No error
+
+ print_error("invalid value", hipErrorInvalidValue);
+ // CHECK: invalid value value: 1
+ // CHECK: invalid value name: hipErrorInvalidValue
+ // CHECK: invalid value string: Invalid argument value
+
+ print_error("invalid device", hipErrorInvalidDevice);
+ // CHECK: invalid device value: 2
+ // CHECK: invalid device name: hipErrorInvalidDevice
+ // CHECK: invalid device string: Invalid device number
+
+ print_error("unknown", hipErrorUnknown);
+ // CHECK: unknown value: 3
+ // CHECK: unknown name: Unrecognized error
+ // CHECK: unknown string: Unknown error
+
+ hipError_t Unrecognized = static_cast<hipError_t>(999);
+ print_error("unrecognized", Unrecognized);
+ // CHECK: unrecognized value: 999
+ // CHECK: unrecognized name: Unrecognized error
+ // CHECK: unrecognized string: Unrecognized error
+
+ print_error("set invalid device", hipSetDevice(-1));
+ // CHECK: set invalid device value: 2
+ // CHECK: set invalid device name: hipErrorInvalidDevice
+ // CHECK: set invalid device string: Invalid device number
+
+ print_error("null stream destroy", hipStreamDestroy(nullptr));
+ // CHECK: null stream destroy value: 1
+ // CHECK: null stream destroy name: hipErrorInvalidValue
+ // CHECK: null stream destroy string: Invalid argument value
+
+ return 0;
+}
More information about the llvm-commits
mailing list