[llvm] [Offload] generate LLVMOffloadKernel library (PR #211694)
Sophia Herrmann via llvm-commits
llvm-commits at lists.llvm.org
Fri Jul 31 10:56:18 PDT 2026
https://github.com/jellytabby updated https://github.com/llvm/llvm-project/pull/211694
>From a5a4cf58bd175860457b112a440d88da29ffeeb4 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 01/10] generate LLVMOffloadKernel library
Co-authored-by: Johannes Doerfert <jdoerfert.llvm at gmail.com>
Co-authored-by: Jonas Greifenhain <cadivus at daverkomp.de>
---
offload/CMakeLists.txt | 10 +
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, 526 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..f72ffb10a4ad6 100644
--- a/offload/CMakeLists.txt
+++ b/offload/CMakeLists.txt
@@ -315,6 +315,15 @@ endif()
pythonize_bool(LIBOMPTARGET_OMPT_SUPPORT)
+if(${LLVM_LIBC_GPU_BUILD})
+ set(LIBOMPTARGET_HAS_LIBC TRUE)
+else()
+ set(LIBOMPTARGET_HAS_LIBC FALSE)
+endif()
+set(LIBOMPTARGET_GPU_LIBC_SUPPORT ${LIBOMPTARGET_HAS_LIBC} CACHE BOOL
+ "Libomptarget support for the GPU libc")
+pythonize_bool(LIBOMPTARGET_GPU_LIBC_SUPPORT)
+
set(LIBOMPTARGET_INCLUDE_DIR ${CMAKE_CURRENT_SOURCE_DIR}/include)
set(LIBOMPTARGET_BINARY_INCLUDE_DIR ${CMAKE_CURRENT_BINARY_DIR}/include)
message(STATUS "OpenMP tools dir in libomptarget: ${LIBOMP_OMP_TOOLS_INCLUDE_DIR}")
@@ -340,6 +349,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 c7c6534be022fcf46570e86e30cb9b46bd7fb724 Mon Sep 17 00:00:00 2001
From: Sophia Herrmann <herrmann15 at llnl.gov>
Date: Thu, 23 Jul 2026 08:53:04 -0700
Subject: [PATCH 02/10] remove old cmake variable
---
offload/CMakeLists.txt | 9 ---------
1 file changed, 9 deletions(-)
diff --git a/offload/CMakeLists.txt b/offload/CMakeLists.txt
index f72ffb10a4ad6..2dd4446979c05 100644
--- a/offload/CMakeLists.txt
+++ b/offload/CMakeLists.txt
@@ -315,15 +315,6 @@ endif()
pythonize_bool(LIBOMPTARGET_OMPT_SUPPORT)
-if(${LLVM_LIBC_GPU_BUILD})
- set(LIBOMPTARGET_HAS_LIBC TRUE)
-else()
- set(LIBOMPTARGET_HAS_LIBC FALSE)
-endif()
-set(LIBOMPTARGET_GPU_LIBC_SUPPORT ${LIBOMPTARGET_HAS_LIBC} CACHE BOOL
- "Libomptarget support for the GPU libc")
-pythonize_bool(LIBOMPTARGET_GPU_LIBC_SUPPORT)
-
set(LIBOMPTARGET_INCLUDE_DIR ${CMAKE_CURRENT_SOURCE_DIR}/include)
set(LIBOMPTARGET_BINARY_INCLUDE_DIR ${CMAKE_CURRENT_BINARY_DIR}/include)
message(STATUS "OpenMP tools dir in libomptarget: ${LIBOMP_OMP_TOOLS_INCLUDE_DIR}")
>From 8ae9993fd2c38a3a20c8748262deb70fdd0ce66f 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 03/10] 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/cuda/CMakeLists.txt | 1 +
offload/languages/cuda/src/cuda_runtime.cpp | 20 ++
offload/languages/hip/CMakeLists.txt | 1 +
offload/languages/hip/src/hip_runtime.cpp | 24 ++
offload/languages/include/cuda/cuda_runtime.h | 22 ++
offload/languages/include/hip/hip_runtime.h | 59 ++++
.../include/kernel/DefineLanguageNames.inc | 53 +++
.../include/kernel/LanguageRuntime.h | 216 ++++++++++++
.../include/kernel/UndefineLanguageNames.inc | 50 +++
offload/languages/kernel/CMakeLists.txt | 27 +-
offload/languages/kernel/exports | 9 +-
.../languages/kernel/include/ExportedAPI.h | 41 ---
.../kernel/include/LanguageAliases.h | 42 +++
.../languages/kernel/include/LanguageLaunch.h | 50 +++
.../kernel/include/LanguageRegistration.h | 61 ++++
.../languages/kernel/include/Registration.h | 17 +
offload/languages/kernel/include/RuntimeAPI.h | 47 +++
.../languages/kernel/src/LanguageCommon.cpp | 18 +
.../languages/kernel/src/LanguageLaunch.cpp | 96 ++++++
.../kernel/src/LanguageRegistration.cpp | 318 ++++++++++++++++++
.../languages/kernel/src/LanguageRuntime.cpp | 220 ++++++++++++
.../src/{ExportedAPI.cpp => RuntimeAPI.cpp} | 42 +--
23 files changed, 1368 insertions(+), 68 deletions(-)
create mode 100644 offload/languages/cuda/CMakeLists.txt
create mode 100644 offload/languages/cuda/src/cuda_runtime.cpp
create mode 100644 offload/languages/hip/CMakeLists.txt
create mode 100644 offload/languages/hip/src/hip_runtime.cpp
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.h
create mode 100644 offload/languages/kernel/include/LanguageLaunch.h
create mode 100644 offload/languages/kernel/include/LanguageRegistration.h
create mode 100644 offload/languages/kernel/include/Registration.h
create mode 100644 offload/languages/kernel/include/RuntimeAPI.h
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
rename offload/languages/kernel/src/{ExportedAPI.cpp => RuntimeAPI.cpp} (69%)
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/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/cuda/src/cuda_runtime.cpp b/offload/languages/cuda/src/cuda_runtime.cpp
new file mode 100644
index 0000000000000..00d11c76d668b
--- /dev/null
+++ b/offload/languages/cuda/src/cuda_runtime.cpp
@@ -0,0 +1,20 @@
+/*===---- cuda_runtime.cpp - CUDA runtime api implementations --------------===
+ *
+ * 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 "cuda_runtime.h"
+
+#include "OffloadAPI.h"
+
+#define LANGUAGE cuda
+
+#include "../../kernel/src/LanguageRuntime.cpp"
+
+extern "C" {
+void __cudaRegisterFatBinaryEnd(void *) {}
+}
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/hip/src/hip_runtime.cpp b/offload/languages/hip/src/hip_runtime.cpp
new file mode 100644
index 0000000000000..082f5552bcf15
--- /dev/null
+++ b/offload/languages/hip/src/hip_runtime.cpp
@@ -0,0 +1,24 @@
+/*===---- hip_runtime.cpp - HIP runtime api implementations ----------------===
+ *
+ * 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 "hip_runtime.h"
+
+#include "LanguageLaunch.h"
+#include "OffloadAPI.h"
+
+#define LANGUAGE hip
+
+#include "../../kernel/src/LanguageRuntime.cpp"
+
+extern "C" hipError_t hipLaunchKernel(const char *KernelID, dim3 GridDim,
+ dim3 BlockDim, void **KernelArgsPtr,
+ size_t DynamicSharedMem, void *Stream) {
+ return convertResult(__llvmLaunchKernelImpl(
+ KernelID, GridDim, BlockDim, KernelArgsPtr, DynamicSharedMem, Stream));
+}
diff --git a/offload/languages/include/cuda/cuda_runtime.h b/offload/languages/include/cuda/cuda_runtime.h
new file mode 100644
index 0000000000000..4594784b0b8a1
--- /dev/null
+++ b/offload/languages/include/cuda/cuda_runtime.h
@@ -0,0 +1,22 @@
+/*===---- 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
+ *
+ *===-----------------------------------------------------------------------===
+ */
+
+#pragma once
+
+#define LANGUAGE cuda
+
+#include "../kernel/DefineLanguageNames.inc"
+
+#include "../kernel/LanguageRuntime.h"
+
+#include "../kernel/UndefineLanguageNames.inc"
+
+#undef LANGUAGE
+
+using cudaDeviceProp = cudaDeviceProp_t;
diff --git a/offload/languages/include/hip/hip_runtime.h b/offload/languages/include/hip/hip_runtime.h
new file mode 100644
index 0000000000000..3492fa08e8461
--- /dev/null
+++ b/offload/languages/include/hip/hip_runtime.h
@@ -0,0 +1,59 @@
+/*===---- 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
+ *
+ *===-----------------------------------------------------------------------===
+ */
+
+#pragma once
+
+#define LANGUAGE hip
+
+#include "../kernel/DefineLanguageNames.inc"
+
+#include "../kernel/LanguageRuntime.h"
+
+#include "../kernel/UndefineLanguageNames.inc"
+
+#undef LANGUAGE
+
+enum hipHostMallocFlag_t : unsigned int {
+ hipHostMallocDefault = hipHostAllocDefault,
+ hipHostMallocPortable = hipHostAllocPortable,
+ hipHostMallocMapped = hipHostAllocMapped,
+ hipHostMallocWriteCombined = hipHostAllocWriteCombined,
+ 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); }
+
+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);
+}
+
+#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
diff --git a/offload/languages/include/kernel/DefineLanguageNames.inc b/offload/languages/include/kernel/DefineLanguageNames.inc
new file mode 100644
index 0000000000000..3d68405896fe5
--- /dev/null
+++ b/offload/languages/include/kernel/DefineLanguageNames.inc
@@ -0,0 +1,53 @@
+//===-- LanguageNames.inc - Kernel Language runtime API names -------------===//
+//
+// 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..6bdae2329f536
--- /dev/null
+++ b/offload/languages/include/kernel/LanguageRuntime.h
@@ -0,0 +1,216 @@
+/*===---- language_runtime.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
+ *
+ *===-----------------------------------------------------------------------===
+ */
+
+#pragma once
+
+#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);
+}
+///}
+
+/// DeviceSynchronize.
+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 __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; \
+ }
+
+template <class T, int Size> struct BaseVector {
+ using VT = float __attribute__((ext_vector_type(Size)));
+ VT Vec;
+
+ __device__ __host__ BaseVector() = default;
+ // __device__ __host__ BaseVector(std::initializer_list<T> List) {
+ // auto It = List.begin();
+ // for (int I = 0, E = List.size(); I < E; ++I, ++It)
+ // Vec[I] = *It;
+ // }
+
+ template <typename... Args>
+ __device__ __host__ BaseVector(Args... args) : BaseVector({args...}) {}
+
+ __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 __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 __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)
+
+__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
diff --git a/offload/languages/include/kernel/UndefineLanguageNames.inc b/offload/languages/include/kernel/UndefineLanguageNames.inc
new file mode 100644
index 0000000000000..08155f689f722
--- /dev/null
+++ b/offload/languages/include/kernel/UndefineLanguageNames.inc
@@ -0,0 +1,50 @@
+//===-- LanguageNames.inc - Kernel Language runtime API names -------------===//
+//
+// 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..093297184e8e4 100644
--- a/offload/languages/kernel/CMakeLists.txt
+++ b/offload/languages/kernel/CMakeLists.txt
@@ -1,22 +1,30 @@
add_llvm_library(
LLVMOffloadKernel SHARED
+ ../cuda/src/cuda_runtime.cpp
+ ../hip/src/hip_runtime.cpp
+ src/LanguageCommon.cpp
src/State.cpp
- src/ExportedAPI.cpp
+ src/RuntimeAPI.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_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_SOURCE_DIR}/../../liboffload/include/generated
+ ${CMAKE_CURRENT_BINARY_DIR}/../../liboffload/API
${CMAKE_CURRENT_SOURCE_DIR}/../../include
${CMAKE_CURRENT_SOURCE_DIR}/../../plugins-nextgen/common/include)
@@ -29,7 +37,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.h b/offload/languages/kernel/include/LanguageAliases.h
new file mode 100644
index 0000000000000..280099841315d
--- /dev/null
+++ b/offload/languages/kernel/include/LanguageAliases.h
@@ -0,0 +1,42 @@
+//===------- Aliases.h --- Helpers to make 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
+//
+//===----------------------------------------------------------------------===//
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LANGUAGE
+#error This file should be included, or used, with a LANGUAGE macro set.
+#endif
+
+#define MA_IMPL2(PREFIX, L, RTY, NAME, ...) \
+ extern "C" [[gnu::alias("__llvm" #NAME)]] RTY PREFIX##L##NAME(__VA_ARGS__);
+
+#define MA_IMPL1(PREFIX, L, RTY, NAME, ...) \
+ MA_IMPL2(PREFIX, L, RTY, NAME, __VA_ARGS__)
+
+#define MAKE_ALIAS(PREFIX, RTY, NAME, ...) \
+ MA_IMPL1(PREFIX, LANGUAGE, RTY, NAME, __VA_ARGS__)
+
+MAKE_ALIAS(__, void, RegisterFunction, const char *, const char *, char *,
+ const char *, int, uint3 *, uint3 *, dim3 *, dim3 *, int *)
+MAKE_ALIAS(__, const char *, RegisterFatBinary, const char *)
+MAKE_ALIAS(__, void, UnregisterFatBinary, void *)
+MAKE_ALIAS(__, void, RegisterVar, void **, char *, char *, const char *, int,
+ int, int, int)
+MAKE_ALIAS(__, void, RegisterManagedVar, void **, char *, char *, const char *,
+ size_t, unsigned)
+MAKE_ALIAS(__, void, RegisterSurface, void **, const struct surfaceReference *,
+ const void **, const char *, int, int)
+MAKE_ALIAS(__, void, RegisterTexture, void **, const struct textureReference *,
+ const void **, const char *, int, int, int)
+
+MAKE_ALIAS(__, unsigned, PushCallConfiguration, dim3, dim3, size_t, void *)
+MAKE_ALIAS(__, unsigned, PopCallConfiguration, dim3 *, dim3 *, size_t *, void *)
+
+#undef MAKE_ALIAS
+#undef MA_IMPL1
+#undef MA_IMPL2
diff --git a/offload/languages/kernel/include/LanguageLaunch.h b/offload/languages/kernel/include/LanguageLaunch.h
new file mode 100644
index 0000000000000..4be6f2499eb04
--- /dev/null
+++ b/offload/languages/kernel/include/LanguageLaunch.h
@@ -0,0 +1,50 @@
+//===------ LanguageLaunch.h - Header for LanguageLaunch.cpp ------------===//
+//
+// 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_LANGUAGE_LAUNCH_H
+#define LLVM_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 points
+unsigned __llvmLaunchKernel(const char *KernelID, dim3 GridDim, dim3 BlockDim,
+ void *KernelArgsPtr, size_t DynamicSharedMem,
+ void *Stream);
+
+unsigned __llvmLaunchKernel_spt(const char *KernelID, dim3 GridDim,
+ dim3 BlockDim, void *KernelArgsPtr,
+ size_t DynamicSharedMem, void *Stream);
+
+unsigned __llvmLaunchKernel_ptsz(const char *KernelID, dim3 GridDim,
+ dim3 BlockDim, void *KernelArgsPtr,
+ size_t DynamicSharedMem, void *Stream);
+}
+
+#endif // LLVM_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..f871c1072c49c
--- /dev/null
+++ b/offload/languages/kernel/include/LanguageRegistration.h
@@ -0,0 +1,61 @@
+//===---- LanguageRegistration.h - Language (CUDA/HIP) 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 "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);
+}
+///}
diff --git a/offload/languages/kernel/include/Registration.h b/offload/languages/kernel/include/Registration.h
new file mode 100644
index 0000000000000..70f6a641f538d
--- /dev/null
+++ b/offload/languages/kernel/include/Registration.h
@@ -0,0 +1,17 @@
+//===-- Registration.h - Kernel Language (CUDA/HIP) registration handling -===//
+//
+// 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
+
+namespace llvm {
+namespace offload {
+void readHIPFatbin(const char *Binary, const char *HIPFatbinPtr);
+} // namespace offload
+} // namespace llvm
diff --git a/offload/languages/kernel/include/RuntimeAPI.h b/offload/languages/kernel/include/RuntimeAPI.h
new file mode 100644
index 0000000000000..2aff8c27150aa
--- /dev/null
+++ b/offload/languages/kernel/include/RuntimeAPI.h
@@ -0,0 +1,47 @@
+/*===---- RuntimeAPI.h - Kernel language runtime internals ----------------===
+ *
+ * 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"
+
+namespace llvm {
+namespace offload {
+namespace kernel {
+
+ol_device_handle_t getDefaultDevice();
+
+ol_device_handle_t getHostDevice();
+
+int getDeviceCount();
+
+ol_device_handle_t getDevice(int *DeviceNo);
+
+ol_device_handle_t setDefaultDevice(int DeviceNo);
+
+ol_queue_handle_t getDefaultQueue();
+
+CallConfigurationTy *getCallConfiguration();
+
+void registerKernel(const void *ID, ol_symbol_handle_t Kernel);
+
+void unregisterKernel(const void *ID);
+
+ol_symbol_handle_t getKernel(const void *ID);
+
+void registerProgram(const void *ID, ol_program_handle_t Program);
+
+ol_program_handle_t unregisterProgram(const void *ID);
+
+ol_program_handle_t getProgram(const void *ID);
+
+} // namespace kernel
+} // namespace offload
+} // namespace llvm
diff --git a/offload/languages/kernel/src/LanguageCommon.cpp b/offload/languages/kernel/src/LanguageCommon.cpp
new file mode 100644
index 0000000000000..b0dc712bb5b41
--- /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 "LanguageRegistration.cpp"
+#include "LanguageLaunch.cpp"
+
+#define LANGUAGE cuda
+#include "LanguageAliases.h"
+#undef LANGUAGE
+
+#define LANGUAGE hip
+#include "LanguageAliases.h"
+#undef LANGUAGE
diff --git a/offload/languages/kernel/src/LanguageLaunch.cpp b/offload/languages/kernel/src/LanguageLaunch.cpp
new file mode 100644
index 0000000000000..9d7fb26d4768a
--- /dev/null
+++ b/offload/languages/kernel/src/LanguageLaunch.cpp
@@ -0,0 +1,96 @@
+//===------ LanguageLaunch.cpp - Language (CUDA/HIP) 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 "RuntimeAPI.h"
+
+#include <cstdio>
+
+namespace language_launch = llvm::offload::kernel;
+
+extern "C" {
+
+/// Push call configuration for kernel launch
+unsigned __llvmPushCallConfiguration(dim3 __grid_size, dim3 __block_size,
+ size_t __shared_memory, void *__stream) {
+ CallConfigurationTy &CC = *language_launch::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 = *language_launch::getCallConfiguration();
+ *__grid_size = CC.GridSize;
+ *__block_size = CC.BlockSize;
+ *__shared_memory = CC.SharedMemory;
+ *((void **)__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 = language_launch::getDefaultDevice();
+ ol_symbol_handle_t Kernel = language_launch::getKernel(KernelID);
+
+ ol_dimensions_t GridDimensions, BlockDimensions;
+ ol_kernel_launch_size_args_t LaunchSizeArgs;
+ LaunchSizeArgs.Dimensions =
+ 1 + !!(GridDim.y * BlockDim.y > 1) + !!(GridDim.z * BlockDim.z > 1);
+ GridDimensions.x = GridDim.x;
+ GridDimensions.y = std::max(GridDim.y, 1u);
+ GridDimensions.z = std::max(GridDim.z, 1u);
+ LaunchSizeArgs.NumGroups = GridDimensions;
+ BlockDimensions.x = BlockDim.x;
+ BlockDimensions.y = std::max(BlockDim.y, 1u);
+ BlockDimensions.z = std::max(BlockDim.z, 1u);
+ LaunchSizeArgs.GroupSize = BlockDimensions;
+ LaunchSizeArgs.DynSharedMemory = DynamicSharedMem;
+
+ ol_queue_handle_t Queue = Stream ? reinterpret_cast<ol_queue_handle_t>(Stream)
+ : language_launch::getDefaultQueue();
+
+ ol_kernel_launch_prop_t Properties = {.type = OL_KERNEL_LAUNCH_PROP_TYPE_NONE,
+ .data = nullptr};
+
+ struct OffloadKernelArgs {
+ void **Args;
+ size_t NumArgs;
+ size_t *ArgSizes;
+ };
+ OffloadKernelArgs *OKA = reinterpret_cast<OffloadKernelArgs *>(KernelArgsPtr);
+
+ ol_result_t Result;
+ Result = olLaunchKernel(Queue, Device, Kernel, &LaunchSizeArgs, &Properties,
+ OKA->NumArgs, OKA->Args, OKA->ArgSizes);
+ return Result;
+}
+
+#define LLVM_STYLE_LAUNCH(SUFFIX, PER_THREAD_STREAM) \
+ unsigned __llvmLaunchKernel##SUFFIX(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; \
+ }
+
+LLVM_STYLE_LAUNCH(, false);
+LLVM_STYLE_LAUNCH(_spt, true);
+LLVM_STYLE_LAUNCH(_ptsz, true);
+
+} // extern "C"
diff --git a/offload/languages/kernel/src/LanguageRegistration.cpp b/offload/languages/kernel/src/LanguageRegistration.cpp
new file mode 100644
index 0000000000000..5db739bbe3b51
--- /dev/null
+++ b/offload/languages/kernel/src/LanguageRegistration.cpp
@@ -0,0 +1,318 @@
+//===---- LanguageRegistration.h - Language (CUDA/HIP) 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 "RuntimeAPI.h"
+#include "llvm/ADT/StringRef.h"
+#include "llvm/Frontend/Offloading/Utility.h"
+#include "llvm/Support/Error.h"
+#include <cstdio>
+#include <cstring>
+#include <inttypes.h>
+
+namespace language_registration = llvm::offload::kernel;
+
+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 = language_registration::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();
+ }
+
+ language_registration::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) {
+ fprintf(stderr, "RegisterVar is not implemented!");
+}
+
+void __llvmRegisterManagedVar(void **, char *, char *, const char *, size_t,
+ unsigned) {
+ fprintf(stderr, "RegisterManagedVar is not implemented!");
+}
+
+void __llvmRegisterSurface(void **, const struct surfaceReference *,
+ const void **, const char *, int, int) {
+ fprintf(stderr, "RegisterSurface is not implemented!");
+}
+
+void __llvmRegisterTexture(void **, const struct textureReference *,
+ const void **, const char *, int, int, int) {
+ fprintf(stderr, "RegisterTexture is not implemented!");
+}
+
+/// 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 = language_registration::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();
+ }
+
+ language_registration::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)
+ language_registration::unregisterKernel((const char *)Entry->Address);
+ }
+
+ if (ol_program_handle_t Program =
+ language_registration::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..3852e666b191d
--- /dev/null
+++ b/offload/languages/kernel/src/LanguageRuntime.cpp
@@ -0,0 +1,220 @@
+//===-- 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
+//
+//===----------------------------------------------------------------------===//
+//
+//===----------------------------------------------------------------------===//
+
+#include "LanguageRuntime.h"
+#include <cassert>
+
+#ifndef LANGUAGE
+#error This file should be included, or used, with a LANGUAGE macro set.
+#endif
+
+#include "RuntimeAPI.h"
+#include "Types.h"
+
+#include "OffloadAPI.h"
+
+#include "DefineLanguageNames.inc"
+
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+
+#define STR(X) #X
+#define LANGUAGE_STR STR(LANGUAGE)
+
+namespace language_runtime = llvm::offload::kernel;
+
+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 = language_runtime::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 = language_runtime::getDefaultQueue();
+
+ ol_result_t Result;
+ switch (Kind) {
+ case MemcpyHostToHost: {
+ ol_device_handle_t Host = language_runtime::getHostDevice();
+ Result = olMemcpy(Queue, Dst, Host, const_cast<void *>(Src), Host, Size);
+ break;
+ }
+ case MemcpyHostToDevice: {
+ ol_device_handle_t Device = language_runtime::getDefaultDevice();
+ ol_device_handle_t Host = language_runtime::getHostDevice();
+ Result = olMemcpy(Queue, Dst, Device, const_cast<void *>(Src), Host, Size);
+ break;
+ }
+ case MemcpyDeviceToHost: {
+ ol_device_handle_t Device = language_runtime::getDefaultDevice();
+ ol_device_handle_t Host = language_runtime::getHostDevice();
+
+ Result = olMemcpy(Queue, Dst, Host, const_cast<void *>(Src), Device, Size);
+ break;
+ }
+ case MemcpyDeviceToDevice: {
+ ol_device_handle_t Device = language_runtime::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 = language_runtime::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 = language_runtime::getDevice(DeviceNo);
+ if (!Device)
+ return ErrorInvalidValue;
+ return Success;
+}
+
+Error_t GetDeviceCount(int *Count) {
+ *Count = language_runtime::getDeviceCount();
+ return Success;
+}
+
+Error_t SetDevice(int DeviceNo) {
+ ol_device_handle_t Device = language_runtime::setDefaultDevice(DeviceNo);
+ assert(Device == language_runtime::getDefaultDevice() &&
+ "Set Device is not Default Device");
+ return Device ? Success : ErrorInvalidValue;
+}
+
+Error_t HostAlloc(void **Ptr, size_t Size, unsigned int Flags) {
+ // TODO:
+ ol_device_handle_t Device = language_runtime::getDefaultDevice();
+ ol_result_t Result = olMemAlloc(Device, OL_ALLOC_TYPE_HOST, 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) {
+ // TODO: [h15] add remaining pci/mem fields
+ ol_device_handle_t Device = language_runtime::getDefaultDevice();
+ size_t nameSize = 0;
+ olGetDeviceInfoSize(Device, OL_DEVICE_INFO_NAME, &nameSize);
+ 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(language_runtime::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);
+}
diff --git a/offload/languages/kernel/src/ExportedAPI.cpp b/offload/languages/kernel/src/RuntimeAPI.cpp
similarity index 69%
rename from offload/languages/kernel/src/ExportedAPI.cpp
rename to offload/languages/kernel/src/RuntimeAPI.cpp
index 5c65d3d66b89b..8fecfc3d2b388 100644
--- a/offload/languages/kernel/src/ExportedAPI.cpp
+++ b/offload/languages/kernel/src/RuntimeAPI.cpp
@@ -1,4 +1,4 @@
-//===------ ExportedAPI.cpp - Kernel Language runtime - exported api ------===//
+//===------ RuntimeAPI.cpp - Kernel language runtime internals ------------===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
@@ -8,7 +8,7 @@
//
//===----------------------------------------------------------------------===//
-#include "ExportedAPI.h"
+#include "RuntimeAPI.h"
#include "State.h"
#include "Types.h"
@@ -19,27 +19,26 @@
#include <cstdio>
#include <stdint.h>
-using namespace llvm;
-using namespace offload;
+namespace llvm {
+namespace offload {
+namespace kernel {
-/// Runtime API
-///{
-ol_device_handle_t olKGetDefaultDevice() {
+ol_device_handle_t getDefaultDevice() {
ol_device_handle_t DefaultDevice = ThreadStateTy::getDefaultDevice();
return DefaultDevice;
}
-ol_device_handle_t olKGetHostDevice() {
+ol_device_handle_t getHostDevice() {
ol_device_handle_t HostDevice = StateTy::getHostDevice();
return HostDevice;
}
-int olKGetDeviceCount() {
+int getDeviceCount() {
int DeviceCount = StateTy::get().getDevices().size();
return DeviceCount;
}
-ol_device_handle_t olKGetDevice(int *DeviceNo) {
+ol_device_handle_t getDevice(int *DeviceNo) {
ol_device_handle_t DefaultDevice = ThreadStateTy::getDefaultDevice();
int DeviceCount = StateTy::get().getDevices().size();
ArrayRef<ol_device_handle_t> Devices = StateTy::get().getDevices();
@@ -52,7 +51,7 @@ ol_device_handle_t olKGetDevice(int *DeviceNo) {
return nullptr;
}
-ol_device_handle_t olKSetDefaultDevice(int DeviceNo) {
+ol_device_handle_t setDefaultDevice(int DeviceNo) {
ArrayRef<ol_device_handle_t> Devices = StateTy::get().getDevices();
if (DeviceNo < 0 || DeviceNo >= static_cast<int>(Devices.size()))
return nullptr;
@@ -61,39 +60,42 @@ ol_device_handle_t olKSetDefaultDevice(int DeviceNo) {
return Device;
}
-ol_queue_handle_t olKGetDefaultQueue() {
+ol_queue_handle_t getDefaultQueue() {
ol_queue_handle_t DefaultQueue = ThreadStateTy::getDefaultQueue();
return DefaultQueue;
}
-CallConfigurationTy *olKGetCallConfiguration() {
+CallConfigurationTy *getCallConfiguration() {
return &ThreadStateTy::getCallConfiguration();
}
-void olKRegisterKernel(const void *ID, ol_symbol_handle_t Kernel) {
+void registerKernel(const void *ID, ol_symbol_handle_t Kernel) {
StateTy::get().addKernel(ID, Kernel);
}
-void olKUnregisterKernel(const void *ID) {
+void unregisterKernel(const void *ID) {
if (StateTy *State = StateTy::tryGet())
State->removeKernel(ID);
}
-ol_symbol_handle_t olKGetKernel(const void *ID) {
+ol_symbol_handle_t getKernel(const void *ID) {
return StateTy::get().getKernel(ID);
}
-void olKRegisterProgram(const void *ID, ol_program_handle_t Program) {
+void registerProgram(const void *ID, ol_program_handle_t Program) {
StateTy::get().addProgram(ID, Program);
}
-ol_program_handle_t olKUnregisterProgram(const void *ID) {
+ol_program_handle_t unregisterProgram(const void *ID) {
if (StateTy *State = StateTy::tryGet())
return State->removeProgram(ID);
return nullptr;
}
-ol_program_handle_t olKGetProgram(const void *ID) {
+ol_program_handle_t getProgram(const void *ID) {
return StateTy::get().getProgram(ID);
}
-///}
+
+} // namespace kernel
+} // namespace offload
+} // namespace llvm
>From a030c34ba1e6f988623c5c9ed4d04ed1e252252e 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 04/10] Modularize
---
.../include/kernel/DefineLanguageNames.inc | 6 -----
.../include/kernel/LanguageRuntime.h | 25 -------------------
.../include/kernel/UndefineLanguageNames.inc | 5 ----
.../languages/kernel/src/LanguageRuntime.cpp | 11 --------
4 files changed, 47 deletions(-)
diff --git a/offload/languages/include/kernel/DefineLanguageNames.inc b/offload/languages/include/kernel/DefineLanguageNames.inc
index 3d68405896fe5..869b17101b0fd 100644
--- a/offload/languages/include/kernel/DefineLanguageNames.inc
+++ b/offload/languages/include/kernel/DefineLanguageNames.inc
@@ -41,13 +41,7 @@
#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 6bdae2329f536..24b2c2aeb965e 100644
--- a/offload/languages/include/kernel/LanguageRuntime.h
+++ b/offload/languages/include/kernel/LanguageRuntime.h
@@ -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.
@@ -116,30 +111,10 @@ 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 08155f689f722..cb295b5120516 100644
--- a/offload/languages/include/kernel/UndefineLanguageNames.inc
+++ b/offload/languages/include/kernel/UndefineLanguageNames.inc
@@ -39,12 +39,7 @@
#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 3852e666b191d..3c502513cd1e3 100644
--- a/offload/languages/kernel/src/LanguageRuntime.cpp
+++ b/offload/languages/kernel/src/LanguageRuntime.cpp
@@ -163,7 +163,6 @@ Error_t DriverGetVersion(int *Version) {
}
Error_t GetDeviceProperties(DeviceProp_t *DeviceProp, int DeviceNo) {
- // TODO: [h15] add remaining pci/mem fields
ol_device_handle_t Device = language_runtime::getDefaultDevice();
size_t nameSize = 0;
olGetDeviceInfoSize(Device, OL_DEVICE_INFO_NAME, &nameSize);
@@ -191,16 +190,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 285844dcddc0f4c856337f236e46db9d0f251e20 Mon Sep 17 00:00:00 2001
From: Sophia Herrmann <herrmann15 at llnl.gov>
Date: Fri, 24 Jul 2026 11:33:47 -0700
Subject: [PATCH 05/10] formatting
---
offload/languages/kernel/src/LanguageCommon.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/offload/languages/kernel/src/LanguageCommon.cpp b/offload/languages/kernel/src/LanguageCommon.cpp
index b0dc712bb5b41..9b376e979d8f6 100644
--- a/offload/languages/kernel/src/LanguageCommon.cpp
+++ b/offload/languages/kernel/src/LanguageCommon.cpp
@@ -6,8 +6,8 @@
//
//===----------------------------------------------------------------------===//
-#include "LanguageRegistration.cpp"
#include "LanguageLaunch.cpp"
+#include "LanguageRegistration.cpp"
#define LANGUAGE cuda
#include "LanguageAliases.h"
>From 26c2a5219cca88e02201fb2678fe25f6d1fd5668 Mon Sep 17 00:00:00 2001
From: Sophia Herrmann <herrmann15 at llnl.gov>
Date: Tue, 28 Jul 2026 17:50:40 -0700
Subject: [PATCH 06/10] remove empty stubs
---
.../include/kernel/DefineLanguageNames.inc | 5 ----
.../include/kernel/LanguageRuntime.h | 10 -------
.../include/kernel/UndefineLanguageNames.inc | 5 ----
.../languages/kernel/src/LanguageRuntime.cpp | 26 -------------------
4 files changed, 46 deletions(-)
diff --git a/offload/languages/include/kernel/DefineLanguageNames.inc b/offload/languages/include/kernel/DefineLanguageNames.inc
index 869b17101b0fd..a07ab0ad7e522 100644
--- a/offload/languages/include/kernel/DefineLanguageNames.inc
+++ b/offload/languages/include/kernel/DefineLanguageNames.inc
@@ -25,10 +25,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)
@@ -39,7 +35,6 @@
#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 Stream_t COMBINE(LANGUAGE, Stream_t)
#define StreamCreate COMBINE(LANGUAGE, StreamCreate)
diff --git a/offload/languages/include/kernel/LanguageRuntime.h b/offload/languages/include/kernel/LanguageRuntime.h
index 24b2c2aeb965e..8e4f793d10423 100644
--- a/offload/languages/include/kernel/LanguageRuntime.h
+++ b/offload/languages/include/kernel/LanguageRuntime.h
@@ -89,14 +89,6 @@ static inline Error_t Memcpy(T *Dst, const T *Src, size_t Size,
/// DeviceSynchronize.
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);
@@ -105,8 +97,6 @@ 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);
diff --git a/offload/languages/include/kernel/UndefineLanguageNames.inc b/offload/languages/include/kernel/UndefineLanguageNames.inc
index cb295b5120516..4e6b433bbd8b5 100644
--- a/offload/languages/include/kernel/UndefineLanguageNames.inc
+++ b/offload/languages/include/kernel/UndefineLanguageNames.inc
@@ -22,10 +22,6 @@
#undef MemcpyDeviceToHost
#undef MemcpyDeviceToDevice
#undef MemcpyDefault
-#undef GetLastError
-#undef PeekAtLastError
-#undef GetErrorName
-#undef GetErrorString
#undef GetDevice
#undef GetDeviceCount
#undef SetDevice
@@ -37,7 +33,6 @@
#undef HostAllocWriteCombined
#undef MallocHost
#undef FreeHost
-#undef DriverGetVersion
#undef GetDeviceProperties
#undef Stream_t
#undef StreamCreate
diff --git a/offload/languages/kernel/src/LanguageRuntime.cpp b/offload/languages/kernel/src/LanguageRuntime.cpp
index 3c502513cd1e3..0d177e8718e3b 100644
--- a/offload/languages/kernel/src/LanguageRuntime.cpp
+++ b/offload/languages/kernel/src/LanguageRuntime.cpp
@@ -101,26 +101,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 = language_runtime::getDevice(DeviceNo);
if (!Device)
@@ -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 = language_runtime::getDefaultDevice();
size_t nameSize = 0;
>From f1cb98b385e5989be1fe1e6daf6749d1827959d3 Mon Sep 17 00:00:00 2001
From: Sophia Herrmann <herrmann15 at llnl.gov>
Date: Tue, 28 Jul 2026 17:51:09 -0700
Subject: [PATCH 07/10] fix HostToHost bug + SetDevice assertion
---
offload/languages/kernel/src/LanguageRuntime.cpp | 6 ++++--
1 file changed, 4 insertions(+), 2 deletions(-)
diff --git a/offload/languages/kernel/src/LanguageRuntime.cpp b/offload/languages/kernel/src/LanguageRuntime.cpp
index 0d177e8718e3b..461487669435a 100644
--- a/offload/languages/kernel/src/LanguageRuntime.cpp
+++ b/offload/languages/kernel/src/LanguageRuntime.cpp
@@ -60,7 +60,7 @@ Error_t Memcpy(void *Dst, const void *Src, size_t Size, MemcpyKind Kind) {
switch (Kind) {
case MemcpyHostToHost: {
ol_device_handle_t Host = language_runtime::getHostDevice();
- Result = olMemcpy(Queue, Dst, Host, const_cast<void *>(Src), Host, Size);
+ Result = olMemcpy(nullptr, Dst, Host, const_cast<void *>(Src), Host, Size);
break;
}
case MemcpyHostToDevice: {
@@ -115,9 +115,11 @@ Error_t GetDeviceCount(int *Count) {
Error_t SetDevice(int DeviceNo) {
ol_device_handle_t Device = language_runtime::setDefaultDevice(DeviceNo);
+ if (!Device)
+ return ErrorInvalidValue;
assert(Device == language_runtime::getDefaultDevice() &&
"Set Device is not Default Device");
- return Device ? Success : ErrorInvalidValue;
+ return Success;
}
Error_t HostAlloc(void **Ptr, size_t Size, unsigned int Flags) {
>From eb7f6c2fd97841152f846a4aa66bbc2def9a46d4 Mon Sep 17 00:00:00 2001
From: Sophia Herrmann <herrmann15 at llnl.gov>
Date: Thu, 30 Jul 2026 14:42:45 -0700
Subject: [PATCH 08/10] add C++ style includes
---
offload/languages/include/kernel/LanguageRuntime.h | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/offload/languages/include/kernel/LanguageRuntime.h b/offload/languages/include/kernel/LanguageRuntime.h
index 8e4f793d10423..8b114e0d3d054 100644
--- a/offload/languages/include/kernel/LanguageRuntime.h
+++ b/offload/languages/include/kernel/LanguageRuntime.h
@@ -9,10 +9,10 @@
#pragma once
+#include <cstddef>
+#include <cstdint>
#include <cstdio>
#include <cstdlib>
-#include <stddef.h>
-#include <stdint.h>
enum Error_t : uint32_t {
Success = 0,
>From 4c19a034461e4efa2b7263aad6d38848f48c2806 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 09/10] remove legacy fatbin registration wrapping
---
offload/languages/cuda/src/cuda_runtime.cpp | 4 -
.../kernel/include/LanguageAliases.h | 2 -
.../kernel/include/LanguageRegistration.h | 20 --
.../kernel/src/LanguageRegistration.cpp | 188 ------------------
4 files changed, 214 deletions(-)
diff --git a/offload/languages/cuda/src/cuda_runtime.cpp b/offload/languages/cuda/src/cuda_runtime.cpp
index 00d11c76d668b..536fc8e38bb72 100644
--- a/offload/languages/cuda/src/cuda_runtime.cpp
+++ b/offload/languages/cuda/src/cuda_runtime.cpp
@@ -14,7 +14,3 @@
#define LANGUAGE cuda
#include "../../kernel/src/LanguageRuntime.cpp"
-
-extern "C" {
-void __cudaRegisterFatBinaryEnd(void *) {}
-}
diff --git a/offload/languages/kernel/include/LanguageAliases.h b/offload/languages/kernel/include/LanguageAliases.h
index 280099841315d..897914d8d0cf1 100644
--- a/offload/languages/kernel/include/LanguageAliases.h
+++ b/offload/languages/kernel/include/LanguageAliases.h
@@ -23,8 +23,6 @@
MAKE_ALIAS(__, void, RegisterFunction, const char *, const char *, char *,
const char *, int, uint3 *, uint3 *, dim3 *, dim3 *, int *)
-MAKE_ALIAS(__, const char *, RegisterFatBinary, const char *)
-MAKE_ALIAS(__, void, UnregisterFatBinary, void *)
MAKE_ALIAS(__, void, RegisterVar, void **, char *, char *, const char *, int,
int, int, int)
MAKE_ALIAS(__, void, RegisterManagedVar, void **, char *, char *, const char *,
diff --git a/offload/languages/kernel/include/LanguageRegistration.h b/offload/languages/kernel/include/LanguageRegistration.h
index f871c1072c49c..651dbf2ef586a 100644
--- a/offload/languages/kernel/include/LanguageRegistration.h
+++ b/offload/languages/kernel/include/LanguageRegistration.h
@@ -14,22 +14,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" {
@@ -38,10 +22,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/src/LanguageRegistration.cpp b/offload/languages/kernel/src/LanguageRegistration.cpp
index 5db739bbe3b51..152cad14fc6f3 100644
--- a/offload/languages/kernel/src/LanguageRegistration.cpp
+++ b/offload/languages/kernel/src/LanguageRegistration.cpp
@@ -20,172 +20,6 @@
namespace language_registration = llvm::offload::kernel;
-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" {
@@ -206,28 +40,6 @@ void __llvmRegisterFunction(const char *Binary, const char *KernelID,
language_registration::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) {
fprintf(stderr, "RegisterVar is not implemented!");
>From 5ed1b8e11b74e518aa8a72315c77fb9676bd19ed Mon Sep 17 00:00:00 2001
From: Sophia Herrmann <herrmann15 at llnl.gov>
Date: Thu, 30 Jul 2026 14:44:31 -0700
Subject: [PATCH 10/10] fix olMemAlloc after rebase
---
offload/languages/kernel/src/LanguageRuntime.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/offload/languages/kernel/src/LanguageRuntime.cpp b/offload/languages/kernel/src/LanguageRuntime.cpp
index 461487669435a..0f80c31719176 100644
--- a/offload/languages/kernel/src/LanguageRuntime.cpp
+++ b/offload/languages/kernel/src/LanguageRuntime.cpp
@@ -125,7 +125,7 @@ Error_t SetDevice(int DeviceNo) {
Error_t HostAlloc(void **Ptr, size_t Size, unsigned int Flags) {
// TODO:
ol_device_handle_t Device = language_runtime::getDefaultDevice();
- ol_result_t Result = olMemAlloc(Device, OL_ALLOC_TYPE_HOST, Size, Ptr);
+ ol_result_t Result = olMemAllocHost(Device, Size, Ptr);
return convertResult(Result);
}
More information about the llvm-commits
mailing list