[clang] [llvm] [LLVMOffload] Get LastError for bad Kernels (PR #213390)
Sophia Herrmann via cfe-commits
cfe-commits at lists.llvm.org
Fri Jul 31 17:46:59 PDT 2026
https://github.com/jellytabby created https://github.com/llvm/llvm-project/pull/213390
This PR moves Kernel launches into the language specific layer in order to retrieve errors generated by them. It also adds extra guards to the internal kernel launch. It also adds tests for #213389.
This PR depends on #213389 and its predecessors but because I do not have commit access I cannot stack the PR. For review only consider the LAST ONE commit.
>From a46fd13176b19764576a6a20b4c7914d247c7b26 Mon Sep 17 00:00:00 2001
From: Sophia Herrmann <herrmann15 at llnl.gov>
Date: Wed, 22 Jul 2026 16:57:48 -0700
Subject: [PATCH 1/9] generate LLVMOffloadKernel library
Co-authored-by: Johannes Doerfert <jdoerfert.llvm at gmail.com>
Co-authored-by: Jonas Greifenhain <cadivus at daverkomp.de>
remove old cmake variable
---
offload/CMakeLists.txt | 1 +
offload/languages/CMakeLists.txt | 1 +
offload/languages/kernel/CMakeLists.txt | 35 ++++
offload/languages/kernel/exports | 8 +
.../languages/kernel/include/ExportedAPI.h | 41 ++++
offload/languages/kernel/include/State.h | 106 ++++++++++
offload/languages/kernel/include/Types.h | 28 +++
offload/languages/kernel/src/ExportedAPI.cpp | 99 +++++++++
offload/languages/kernel/src/State.cpp | 198 ++++++++++++++++++
9 files changed, 517 insertions(+)
create mode 100644 offload/languages/CMakeLists.txt
create mode 100644 offload/languages/kernel/CMakeLists.txt
create mode 100644 offload/languages/kernel/exports
create mode 100644 offload/languages/kernel/include/ExportedAPI.h
create mode 100644 offload/languages/kernel/include/State.h
create mode 100644 offload/languages/kernel/include/Types.h
create mode 100644 offload/languages/kernel/src/ExportedAPI.cpp
create mode 100644 offload/languages/kernel/src/State.cpp
diff --git a/offload/CMakeLists.txt b/offload/CMakeLists.txt
index 2885b20f9c1d8..2dd4446979c05 100644
--- a/offload/CMakeLists.txt
+++ b/offload/CMakeLists.txt
@@ -340,6 +340,7 @@ if(BUILD_LIBOMPTARGET)
endif()
add_subdirectory(liboffload)
+add_subdirectory(languages)
# Add tests.
if(OFFLOAD_INCLUDE_TESTS)
diff --git a/offload/languages/CMakeLists.txt b/offload/languages/CMakeLists.txt
new file mode 100644
index 0000000000000..08b2a68082d80
--- /dev/null
+++ b/offload/languages/CMakeLists.txt
@@ -0,0 +1 @@
+add_subdirectory(kernel)
diff --git a/offload/languages/kernel/CMakeLists.txt b/offload/languages/kernel/CMakeLists.txt
new file mode 100644
index 0000000000000..3b0823a44faf1
--- /dev/null
+++ b/offload/languages/kernel/CMakeLists.txt
@@ -0,0 +1,35 @@
+add_llvm_library(
+ LLVMOffloadKernel SHARED
+
+ src/State.cpp
+ src/ExportedAPI.cpp
+
+ LINK_COMPONENTS
+ Support
+ Offload
+ )
+
+if(LIBOMP_HAVE_VERSION_SCRIPT_FLAG)
+ target_link_libraries(LLVMOffloadKernel PRIVATE "-Wl,--version-script=${CMAKE_CURRENT_SOURCE_DIR}/exports")
+endif()
+
+target_include_directories(LLVMOffloadKernel PUBLIC
+ ${CMAKE_CURRENT_SOURCE_DIR}/include
+ ${CMAKE_CURRENT_SOURCE_DIR}/../../liboffload/include
+ ${CMAKE_CURRENT_SOURCE_DIR}/../../liboffload/include/generated
+ ${CMAKE_CURRENT_SOURCE_DIR}/../../include
+ ${CMAKE_CURRENT_SOURCE_DIR}/../../plugins-nextgen/common/include)
+
+target_compile_options(LLVMOffloadKernel PRIVATE ${offload_compile_flags})
+target_link_options(LLVMOffloadKernel PRIVATE ${offload_link_flags})
+
+target_compile_definitions(LLVMOffloadKernel PRIVATE
+ TARGET_NAME="libLLVMOffloadKernel"
+ DEBUG_PREFIX="LLVMOffloadKernel"
+)
+
+set_target_properties(LLVMOffloadKernel PROPERTIES
+ POSITION_INDEPENDENT_CODE ON
+ INSTALL_RPATH "$ORIGIN"
+ BUILD_RPATH "$ORIGIN:${CMAKE_CURRENT_BINARY_DIR}/..")
+install(TARGETS LLVMOffloadKernel LIBRARY COMPONENT offload DESTINATION "${OFFLOAD_INSTALL_LIBDIR}")
diff --git a/offload/languages/kernel/exports b/offload/languages/kernel/exports
new file mode 100644
index 0000000000000..4bc5a4ab9710c
--- /dev/null
+++ b/offload/languages/kernel/exports
@@ -0,0 +1,8 @@
+VERS1.0 {
+ global:
+
+ olK*;
+
+ local:
+ *;
+};
diff --git a/offload/languages/kernel/include/ExportedAPI.h b/offload/languages/kernel/include/ExportedAPI.h
new file mode 100644
index 0000000000000..a4ce168d78fda
--- /dev/null
+++ b/offload/languages/kernel/include/ExportedAPI.h
@@ -0,0 +1,41 @@
+/*===---- ExportedAPI.h - Kernel language runtime - exported api ----------===
+ *
+ * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+ * See https://llvm.org/LICENSE.txt for license information.
+ * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+ *
+ *===-----------------------------------------------------------------------===
+ */
+
+#pragma once
+
+#include "OffloadAPI.h"
+#include "Types.h"
+
+extern "C" {
+ol_device_handle_t olKGetDefaultDevice();
+
+ol_device_handle_t olKGetHostDevice();
+
+int olKGetDeviceCount();
+
+ol_device_handle_t olKGetDevice(int *DeviceNo);
+
+ol_device_handle_t olKSetDefaultDevice(int DeviceNo);
+
+ol_queue_handle_t olKGetDefaultQueue();
+
+CallConfigurationTy *olKGetCallConfiguration();
+
+void olKRegisterKernel(const void *ID, ol_symbol_handle_t Kernel);
+
+void olKUnregisterKernel(const void *ID);
+
+ol_symbol_handle_t olKGetKernel(const void *ID);
+
+void olKRegisterProgram(const void *ID, ol_program_handle_t Program);
+
+ol_program_handle_t olKUnregisterProgram(const void *ID);
+
+ol_program_handle_t olKGetProgram(const void *ID);
+}
diff --git a/offload/languages/kernel/include/State.h b/offload/languages/kernel/include/State.h
new file mode 100644
index 0000000000000..90f323d160a04
--- /dev/null
+++ b/offload/languages/kernel/include/State.h
@@ -0,0 +1,106 @@
+//===------- State.h - Kernel Language (CUDA/HIP) persistent state --------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include "OffloadAPI.h"
+#include "Types.h"
+
+#include "llvm/ADT/ArrayRef.h"
+#include "llvm/ADT/DenseMap.h"
+#include "llvm/ADT/SmallVector.h"
+
+namespace llvm {
+namespace offload {
+
+using KernelIDTy = const void *;
+
+struct ThreadStateTy {
+ ~ThreadStateTy();
+
+ static ThreadStateTy &get();
+
+ static ol_queue_handle_t getDefaultQueue();
+ static ol_device_handle_t getDefaultDevice();
+ static CallConfigurationTy &getCallConfiguration();
+ void setDefaultDevice(ol_device_handle_t Device);
+
+private:
+ void createDefaultQueue(ol_device_handle_t Device);
+
+ ol_device_handle_t DefaultDevice = nullptr;
+ ol_queue_handle_t DefaultQueue = nullptr;
+
+ CallConfigurationTy CC = {};
+
+ ThreadStateTy();
+};
+
+struct StateTy {
+ ~StateTy();
+
+ friend struct ThreadStateTy;
+
+ static StateTy &get();
+ static StateTy *tryGet();
+
+ static ol_device_handle_t getHostDevice() { return get().HostDevice; }
+
+ ArrayRef<ol_device_handle_t> getDevices() const { return Devices; }
+
+ void addDevice(ol_device_handle_t Device) { Devices.push_back(Device); }
+ void setHostDevice(ol_device_handle_t Device) {
+ if (!HostDevice)
+ HostDevice = Device;
+ }
+
+ void addKernel(KernelIDTy KernelID, ol_symbol_handle_t Kernel) {
+ KernelMap[KernelID] = Kernel;
+ }
+
+ void removeKernel(KernelIDTy KernelID) { KernelMap.erase(KernelID); }
+
+ ol_symbol_handle_t getKernel(KernelIDTy KernelID) {
+ return KernelMap[KernelID];
+ }
+
+ void addProgram(const void *Binary, ol_program_handle_t Program) {
+ BinaryRegisterMap[Binary] = Program;
+ }
+
+ ol_program_handle_t removeProgram(const void *Binary) {
+ auto It = BinaryRegisterMap.find(Binary);
+ if (It == BinaryRegisterMap.end())
+ return nullptr;
+ ol_program_handle_t Program = It->second;
+ BinaryRegisterMap.erase(It);
+ return Program;
+ }
+
+ ol_program_handle_t getProgram(const void *Binary) {
+ assert(BinaryRegisterMap.count(Binary));
+ return BinaryRegisterMap[Binary];
+ }
+
+ void destroyRegisteredPrograms();
+
+private:
+ DenseMap<const void *, ol_program_handle_t> BinaryRegisterMap;
+ DenseMap<KernelIDTy, ol_symbol_handle_t> KernelMap;
+ SmallVector<ol_device_handle_t, 8> Devices;
+
+ ol_queue_handle_t DefaultQueue = nullptr;
+ ol_device_handle_t HostDevice = nullptr;
+
+ StateTy();
+};
+
+} // namespace offload
+} // namespace llvm
diff --git a/offload/languages/kernel/include/Types.h b/offload/languages/kernel/include/Types.h
new file mode 100644
index 0000000000000..6c5a0ec19a8f7
--- /dev/null
+++ b/offload/languages/kernel/include/Types.h
@@ -0,0 +1,28 @@
+//===------- Types.h - Kernel Language (CUDA/HIP) api types ---------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include "Types.h"
+#include <cstddef>
+#include <cstdint>
+
+struct uint3 {
+ unsigned x = 0, y = 0, z = 0;
+};
+
+using dim3 = uint3;
+
+struct CallConfigurationTy {
+ dim3 GridSize;
+ dim3 BlockSize;
+ size_t SharedMemory;
+ void *Stream;
+};
diff --git a/offload/languages/kernel/src/ExportedAPI.cpp b/offload/languages/kernel/src/ExportedAPI.cpp
new file mode 100644
index 0000000000000..5c65d3d66b89b
--- /dev/null
+++ b/offload/languages/kernel/src/ExportedAPI.cpp
@@ -0,0 +1,99 @@
+//===------ ExportedAPI.cpp - Kernel Language runtime - exported api ------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+//===----------------------------------------------------------------------===//
+
+#include "ExportedAPI.h"
+
+#include "State.h"
+#include "Types.h"
+
+#include "OffloadAPI.h"
+#include "llvm/ADT/ArrayRef.h"
+
+#include <cstdio>
+#include <stdint.h>
+
+using namespace llvm;
+using namespace offload;
+
+/// Runtime API
+///{
+ol_device_handle_t olKGetDefaultDevice() {
+ ol_device_handle_t DefaultDevice = ThreadStateTy::getDefaultDevice();
+ return DefaultDevice;
+}
+
+ol_device_handle_t olKGetHostDevice() {
+ ol_device_handle_t HostDevice = StateTy::getHostDevice();
+ return HostDevice;
+}
+
+int olKGetDeviceCount() {
+ int DeviceCount = StateTy::get().getDevices().size();
+ return DeviceCount;
+}
+
+ol_device_handle_t olKGetDevice(int *DeviceNo) {
+ ol_device_handle_t DefaultDevice = ThreadStateTy::getDefaultDevice();
+ int DeviceCount = StateTy::get().getDevices().size();
+ ArrayRef<ol_device_handle_t> Devices = StateTy::get().getDevices();
+ for (int i = 0; i < DeviceCount; i++) {
+ if (Devices[i] == DefaultDevice) {
+ *DeviceNo = i;
+ return Devices[i];
+ }
+ }
+ return nullptr;
+}
+
+ol_device_handle_t olKSetDefaultDevice(int DeviceNo) {
+ ArrayRef<ol_device_handle_t> Devices = StateTy::get().getDevices();
+ if (DeviceNo < 0 || DeviceNo >= static_cast<int>(Devices.size()))
+ return nullptr;
+ ol_device_handle_t Device = Devices[DeviceNo];
+ ThreadStateTy::get().setDefaultDevice(Device);
+ return Device;
+}
+
+ol_queue_handle_t olKGetDefaultQueue() {
+ ol_queue_handle_t DefaultQueue = ThreadStateTy::getDefaultQueue();
+ return DefaultQueue;
+}
+
+CallConfigurationTy *olKGetCallConfiguration() {
+ return &ThreadStateTy::getCallConfiguration();
+}
+
+void olKRegisterKernel(const void *ID, ol_symbol_handle_t Kernel) {
+ StateTy::get().addKernel(ID, Kernel);
+}
+
+void olKUnregisterKernel(const void *ID) {
+ if (StateTy *State = StateTy::tryGet())
+ State->removeKernel(ID);
+}
+
+ol_symbol_handle_t olKGetKernel(const void *ID) {
+ return StateTy::get().getKernel(ID);
+}
+
+void olKRegisterProgram(const void *ID, ol_program_handle_t Program) {
+ StateTy::get().addProgram(ID, Program);
+}
+
+ol_program_handle_t olKUnregisterProgram(const void *ID) {
+ if (StateTy *State = StateTy::tryGet())
+ return State->removeProgram(ID);
+ return nullptr;
+}
+
+ol_program_handle_t olKGetProgram(const void *ID) {
+ return StateTy::get().getProgram(ID);
+}
+///}
diff --git a/offload/languages/kernel/src/State.cpp b/offload/languages/kernel/src/State.cpp
new file mode 100644
index 0000000000000..eff3c31257e93
--- /dev/null
+++ b/offload/languages/kernel/src/State.cpp
@@ -0,0 +1,198 @@
+//===------ State.cpp - Kernel Language (CUDA/HIP) persistent state -------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+//===----------------------------------------------------------------------===//
+
+#include "State.h"
+#include "Types.h"
+
+#include "OffloadAPI.h"
+#include "llvm/ADT/SmallPtrSet.h"
+#include "llvm/Support/raw_ostream.h"
+
+#include <atomic>
+#include <cstdio>
+#include <mutex>
+
+#define CHECK_FATAL(Result, ...) \
+ if (Result && Result->Code) { \
+ fprintf(stderr, __VA_ARGS__); \
+ abort(); \
+ }
+
+using namespace llvm;
+using namespace offload;
+
+std::atomic<uint32_t> AnyNonDefaultDevice = 0;
+__attribute__((weak)) uint32_t PerThreadQueue = 0;
+
+thread_local CallConfigurationTy CC = {};
+
+static std::mutex StateLock;
+static std::atomic<StateTy *> StatePtr = nullptr;
+
+static thread_local ThreadStateTy *ThreadState = nullptr;
+
+static std::mutex ThreadStatesLock;
+using ThreadStatesTy = SmallVector<ThreadStateTy *, 64>;
+static ThreadStatesTy *ThreadStatesPtr = nullptr;
+
+static void deleteThreadState() {
+ std::lock_guard<std::mutex> LG(ThreadStatesLock);
+ ThreadStatesTy *ThreadStates = ThreadStatesPtr;
+ ThreadStatesPtr = nullptr;
+ if (!ThreadStates)
+ return;
+
+ for (auto *TS : *ThreadStates)
+ delete (TS);
+ delete ThreadStates;
+ ThreadState = nullptr;
+}
+
+static void deleteState() {
+ StateTy *ST = StatePtr.load();
+ StatePtr.store(nullptr);
+ delete (ST);
+}
+
+static void destroyQueue(ol_queue_handle_t &Queue) {
+ if (!Queue)
+ return;
+
+ olSyncQueue(Queue);
+ olDestroyQueue(Queue);
+ Queue = nullptr;
+}
+
+ThreadStateTy::ThreadStateTy() {
+ if (PerThreadQueue) [[unlikely]]
+ createDefaultQueue(getDefaultDevice());
+ atexit(deleteThreadState);
+}
+ThreadStateTy::~ThreadStateTy() { destroyQueue(DefaultQueue); }
+
+ThreadStateTy &ThreadStateTy::get() {
+ auto *TS = ThreadState;
+ if (!TS) {
+ TS = new ThreadStateTy();
+ ThreadState = TS;
+ std::lock_guard<std::mutex> LG(ThreadStatesLock);
+ if (!ThreadStatesPtr)
+ ThreadStatesPtr = new ThreadStatesTy;
+ ThreadStatesPtr->push_back(TS);
+ }
+ return *TS;
+}
+
+ol_device_handle_t ThreadStateTy::getDefaultDevice() {
+ ol_device_handle_t DD = ThreadStateTy::get().DefaultDevice;
+ if (DD)
+ return DD;
+ for (ol_device_handle_t Device : StateTy::get().getDevices()) {
+ DD = Device;
+ break;
+ }
+ if (AnyNonDefaultDevice.load(std::memory_order_relaxed)) [[unlikely]] {
+ ol_device_handle_t TDD = ThreadStateTy::get().DefaultDevice;
+ if (TDD)
+ DD = TDD;
+ }
+ return DD;
+}
+
+ol_queue_handle_t ThreadStateTy::getDefaultQueue() {
+ if (!PerThreadQueue) [[likely]]
+ return StateTy::get().DefaultQueue;
+ return ThreadStateTy::get().DefaultQueue;
+}
+
+CallConfigurationTy &ThreadStateTy::getCallConfiguration() {
+ return ThreadStateTy::get().CC;
+}
+
+void ThreadStateTy::setDefaultDevice(ol_device_handle_t Device) {
+ DefaultDevice = Device;
+ createDefaultQueue(Device);
+}
+
+void ThreadStateTy::createDefaultQueue(ol_device_handle_t Device) {
+ if (DefaultQueue)
+ olDestroyQueue(DefaultQueue);
+ CHECK_FATAL(olCreateQueue(Device, &DefaultQueue),
+ "Failed to create per-thread default queue");
+}
+
+StateTy &StateTy::get() {
+ StateTy *ST = StatePtr.load();
+ if (!ST) [[unlikely]] {
+ std::lock_guard<std::mutex> LG(StateLock);
+ ST = StatePtr.load();
+ if (!ST) {
+ ST = new StateTy();
+ StatePtr.store(ST);
+ }
+ }
+ return *ST;
+}
+
+StateTy *StateTy::tryGet() { return StatePtr.load(); }
+
+static bool addDevices(ol_device_handle_t Device, void *Payload) {
+ StateTy &State = *reinterpret_cast<StateTy *>(Payload);
+ ol_platform_handle_t Platform;
+ ol_result_t Result;
+
+ Result = olGetDeviceInfo(Device, OL_DEVICE_INFO_PLATFORM, sizeof(Platform),
+ &Platform);
+ if (Result && Result->Code)
+ return true;
+
+ ol_platform_backend_t Backend;
+ Result = olGetPlatformInfo(Platform, OL_PLATFORM_INFO_BACKEND,
+ sizeof(Backend), &Backend);
+ if (Result && Result->Code)
+ return true;
+
+ if (Backend == OL_PLATFORM_BACKEND_HOST)
+ State.setHostDevice(Device);
+ else
+ State.addDevice(Device);
+ return true;
+}
+
+StateTy::StateTy() {
+ CHECK_FATAL(olInit(nullptr), "Failed to initialize the LLVMOffload");
+ CHECK_FATAL(olIterateDevices(addDevices, this), "Failed to identify devices");
+
+ if (!PerThreadQueue) [[likely]]
+ if (!Devices.empty()) [[likely]]
+ CHECK_FATAL(olCreateQueue(Devices.front(), &DefaultQueue),
+ "Failed to create default queue");
+
+ atexit(deleteState);
+}
+
+StateTy::~StateTy() {
+ deleteThreadState();
+ destroyQueue(DefaultQueue);
+ destroyRegisteredPrograms();
+ olShutDown();
+}
+
+void StateTy::destroyRegisteredPrograms() {
+ SmallPtrSet<ol_program_handle_t, 8> Programs;
+ for (auto &It : BinaryRegisterMap)
+ Programs.insert(It.second);
+
+ KernelMap.clear();
+ BinaryRegisterMap.clear();
+
+ for (ol_program_handle_t Program : Programs)
+ olDestroyProgram(Program);
+}
>From 0e25bbc4eb9e4544e1566b62d07d971ce53a06b7 Mon Sep 17 00:00:00 2001
From: Sophia Herrmann <herrmann15 at llnl.gov>
Date: Fri, 24 Jul 2026 10:20:18 -0700
Subject: [PATCH 2/9] 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 c6b7e7396b7456e7233c0576e51a61e5da354159 Mon Sep 17 00:00:00 2001
From: Sophia Herrmann <herrmann15 at llnl.gov>
Date: Fri, 24 Jul 2026 10:56:46 -0700
Subject: [PATCH 3/9] remove unimplemented stuff + minor fixes
formatting
remove empty stubs
fix HostToHost bug + SetDevice assertion
add C++ style includes
fix olMemAlloc after rebase
---
.../include/kernel/DefineLanguageNames.inc | 11 -----
.../include/kernel/LanguageRuntime.h | 39 +---------------
.../include/kernel/UndefineLanguageNames.inc | 10 -----
.../languages/kernel/src/LanguageCommon.cpp | 2 +-
.../languages/kernel/src/LanguageRuntime.cpp | 45 +++----------------
5 files changed, 8 insertions(+), 99 deletions(-)
diff --git a/offload/languages/include/kernel/DefineLanguageNames.inc b/offload/languages/include/kernel/DefineLanguageNames.inc
index 3d68405896fe5..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,15 +35,8 @@
#define HostAllocWriteCombined COMBINE(LANGUAGE, HostAllocWriteCombined)
#define MallocHost COMBINE(LANGUAGE, MallocHost)
#define FreeHost COMBINE(LANGUAGE, FreeHost)
-#define DriverGetVersion COMBINE(LANGUAGE, DriverGetVersion)
#define GetDeviceProperties COMBINE(LANGUAGE, GetDeviceProperties)
-#define OccupancyMaxPotentialBlockSizeVariableSMem \
- COMBINE(LANGUAGE, OccupancyMaxPotentialBlockSizeVariableSMem)
#define Stream_t COMBINE(LANGUAGE, Stream_t)
#define StreamCreate COMBINE(LANGUAGE, StreamCreate)
-#define StreamCreateWithFlags COMBINE(LANGUAGE, StreamCreateWithFlags)
#define StreamDestroy COMBINE(LANGUAGE, StreamDestroy)
#define StreamSynchronize COMBINE(LANGUAGE, StreamSynchronize)
-#define StreamCreateWithFlagsFlags COMBINE(LANGUAGE, StreamCreateWithFlagsFlags)
-#define StreamDefault COMBINE(LANGUAGE, StreamDefault)
-#define StreamNonBlocking COMBINE(LANGUAGE, StreamNonBlocking)
diff --git a/offload/languages/include/kernel/LanguageRuntime.h b/offload/languages/include/kernel/LanguageRuntime.h
index 6bdae2329f536..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,
@@ -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.
@@ -94,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);
@@ -110,36 +97,14 @@ Error_t SetDevice(int DeviceNo);
Error_t FreeHost(void *Ptr);
-Error_t DriverGetVersion(int *Version);
-
Error_t GetDeviceProperties(DeviceProp_t *DeviceProp, int DeviceNo);
Error_t StreamCreate(Stream_t *stream);
-Error_t StreamCreateWithFlags(Stream_t *stream, unsigned int flags);
-
Error_t StreamDestroy(Stream_t stream);
Error_t StreamSynchronize(Stream_t stream);
-template <typename UnaryFunction, class T>
-static inline Error_t OccupancyMaxPotentialBlockSizeVariableSMem(
- int *minGridSize, int *blockSize, T func,
- UnaryFunction blockSizeToDynamicSMemSize, int blockSizeLimit = 0) {
-#if defined(__AMDGPU__)
- // TODO: values taken from AMD Instinct MI250X gfx90a
- *minGridSize = 220;
- *blockSize = 1024;
-#elif defined(__NVPTX__)
- // TODO: values taken from NVIDIA H100 80GB HBM3
- *minGridSize = 264;
- *blockSize = 1024;
-#endif
- return Success;
-}
-
-///
-
#if defined(__AMDGPU__) || defined(__NVPTX__)
#include <gpuintrin.h>
diff --git a/offload/languages/include/kernel/UndefineLanguageNames.inc b/offload/languages/include/kernel/UndefineLanguageNames.inc
index 08155f689f722..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,14 +33,8 @@
#undef HostAllocWriteCombined
#undef MallocHost
#undef FreeHost
-#undef DriverGetVersion
#undef GetDeviceProperties
-#undef OccupancyMaxPotentialBlockSizeVariableSMem
#undef Stream_t
#undef StreamCreate
-#undef StreamCreateWithFlags
#undef StreamDestroy
#undef StreamSynchronize
-#undef StreamCreateWithFlagsFlags
-#undef StreamDefault
-#undef StreamNonBlocking
diff --git a/offload/languages/kernel/src/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"
diff --git a/offload/languages/kernel/src/LanguageRuntime.cpp b/offload/languages/kernel/src/LanguageRuntime.cpp
index 3852e666b191d..0f80c31719176 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: {
@@ -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)
@@ -135,15 +115,17 @@ 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) {
// 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);
}
@@ -156,14 +138,7 @@ 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) {
- // 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 +166,6 @@ Error_t StreamCreate(Stream_t *Stream) {
return Success;
}
-Error_t StreamCreateWithFlags(Stream_t *Stream, unsigned int Flags) {
- if (Flags == StreamCreateWithFlagsFlags::StreamDefault)
- // FIXME: [h15] offload streams are non-blocking by default
- return StreamCreate(Stream);
- if (Flags == StreamCreateWithFlagsFlags::StreamNonBlocking) {
- return StreamCreate(Stream);
- }
- return ErrorInvalidValue;
-}
-
Error_t StreamDestroy(Stream_t Stream) {
ol_queue_handle_t Queue;
Error_t Err = getQueueFromStream(Stream, &Queue);
>From 2f800de53717a7c4eed7e7d9462cab2cef7aaf7a Mon Sep 17 00:00:00 2001
From: Sophia Herrmann <herrmann15 at llnl.gov>
Date: Thu, 30 Jul 2026 14:43:27 -0700
Subject: [PATCH 4/9] remove legacy fatbin registration wrapping
---
offload/languages/cuda/src/cuda_runtime.cpp | 4 -
offload/languages/hip/src/hip_runtime.cpp | 7 -
.../kernel/include/LanguageAliases.h | 2 -
.../kernel/include/LanguageRegistration.h | 20 --
.../kernel/src/LanguageRegistration.cpp | 188 ------------------
5 files changed, 221 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/hip/src/hip_runtime.cpp b/offload/languages/hip/src/hip_runtime.cpp
index 082f5552bcf15..881eca894a364 100644
--- a/offload/languages/hip/src/hip_runtime.cpp
+++ b/offload/languages/hip/src/hip_runtime.cpp
@@ -15,10 +15,3 @@
#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/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 915de926ae6a5b74e10370c43a59b25a03d8708b Mon Sep 17 00:00:00 2001
From: Sophia Herrmann <herrmann15 at llnl.gov>
Date: Mon, 27 Jul 2026 15:48:11 -0700
Subject: [PATCH 5/9] decouple front/back end
Co-authored-by: Johannes Doerfert <jdoerfert.llvm at gmail.com>
Co-authored-by: Jonas Greifenhain <cadivus at daverkomp.de>
remove irrelevant artifacts
---
clang/include/clang/Driver/CommonArgs.h | 6 ++
clang/lib/CodeGen/CGCUDANV.cpp | 67 +++++++++-------
clang/lib/Driver/Driver.cpp | 78 +++++++++++--------
clang/lib/Driver/ToolChains/AMDGPU.cpp | 33 +++++++-
clang/lib/Driver/ToolChains/Clang.cpp | 56 +++++++++----
clang/lib/Driver/ToolChains/CommonArgs.cpp | 19 +++--
clang/lib/Driver/ToolChains/Cuda.cpp | 70 +++++++++++++----
clang/lib/Driver/ToolChains/Gnu.cpp | 1 +
clang/lib/Driver/ToolChains/Linux.cpp | 4 +-
clang/lib/Headers/__clang_gpu_builtin_vars.h | 19 +++++
clang/test/CodeGenCUDA/Inputs/cuda.h | 2 +-
clang/test/CodeGenCUDA/offload_via_llvm.cu | 60 +++++++-------
clang/test/Driver/cuda-via-liboffload.cu | 15 ++--
.../linker-wrapper-image.c | 38 ++++-----
.../ClangLinkerWrapper.cpp | 13 +++-
.../Frontend/Offloading/OffloadWrapper.cpp | 4 +-
16 files changed, 326 insertions(+), 159 deletions(-)
diff --git a/clang/include/clang/Driver/CommonArgs.h b/clang/include/clang/Driver/CommonArgs.h
index 8c861df793311..ad1912247e01f 100644
--- a/clang/include/clang/Driver/CommonArgs.h
+++ b/clang/include/clang/Driver/CommonArgs.h
@@ -144,6 +144,12 @@ void addArchSpecificRPath(const ToolChain &TC, const llvm::opt::ArgList &Args,
void addOpenMPRuntimeLibraryPath(const ToolChain &TC,
const llvm::opt::ArgList &Args,
llvm::opt::ArgStringList &CmdArgs);
+
+bool addLLVMOffloadingRuntime(const Compilation &C,
+ llvm::opt::ArgStringList &CmdArgs,
+ const ToolChain &TC,
+ const llvm::opt::ArgList &Args);
+
/// Returns true, if an OpenMP runtime has been added.
bool addOpenMPRuntime(const Compilation &C, llvm::opt::ArgStringList &CmdArgs,
const ToolChain &TC, const llvm::opt::ArgList &Args,
diff --git a/clang/lib/CodeGen/CGCUDANV.cpp b/clang/lib/CodeGen/CGCUDANV.cpp
index 416ed935c1b30..0ea3ed36fae83 100644
--- a/clang/lib/CodeGen/CGCUDANV.cpp
+++ b/clang/lib/CodeGen/CGCUDANV.cpp
@@ -254,9 +254,7 @@ CGNVCUDARuntime::CGNVCUDARuntime(CodeGenModule &CGM)
VoidTy = CGM.VoidTy;
PtrTy = CGM.DefaultPtrTy;
- if (CGM.getLangOpts().OffloadViaLLVM)
- Prefix = "llvm";
- else if (CGM.getLangOpts().HIP)
+ if (CGM.getLangOpts().HIP)
Prefix = "hip";
else
Prefix = "cuda";
@@ -345,41 +343,52 @@ void CGNVCUDARuntime::emitDeviceStub(CodeGenFunction &CGF,
emitDeviceStubBodyLegacy(CGF, Args);
}
-/// Build the input as a sized array of pointers so that it can be launched by
-/// the offloading runtime.
+/// CUDA passes the arguments with a level of indirection. For example, a
+/// (void*, short, void*) is passed as {void **, short *, void **} to the launch
+/// function. For the LLVM/Offload launch we include the number of arguments and
+/// their size. Thus, we pass {{void **, short*, void **}, 3, {sizeof(void*),
+/// sizeof(short), sizeof(void*)}}.
Address CGNVCUDARuntime::prepareKernelArgsLLVMOffload(CodeGenFunction &CGF,
FunctionArgList &Args) {
- SmallVector<llvm::Type *> ArgTypes, KernelLaunchParamsTypes;
- for (auto &Arg : Args)
- ArgTypes.push_back(CGF.ConvertTypeForMem(Arg->getType()));
- llvm::StructType *KernelArgsTy = llvm::StructType::create(ArgTypes);
- llvm::Type *KernelArgsPtrsTy = llvm::ArrayType::get(PtrTy, Args.size());
-
- auto *Int32Ty = CGF.Builder.getInt32Ty();
- KernelLaunchParamsTypes.push_back(Int32Ty);
+ SmallVector<llvm::Type *> KernelLaunchParamsTypes;
+
+ auto *Int64Ty = CGF.Builder.getInt64Ty();
+ KernelLaunchParamsTypes.push_back(PtrTy);
+ KernelLaunchParamsTypes.push_back(Int64Ty);
KernelLaunchParamsTypes.push_back(PtrTy);
llvm::StructType *KernelLaunchParamsTy =
llvm::StructType::create(KernelLaunchParamsTypes);
- Address KernelArgs = CGF.CreateTempAllocaWithoutCast(
- KernelArgsTy, CharUnits::fromQuantity(16), "kernel_args");
- Address KernelArgsPtrs = CGF.CreateTempAllocaWithoutCast(
- KernelArgsPtrsTy, CharUnits::fromQuantity(16), "kernel_args_ptrs");
Address KernelLaunchParams = CGF.CreateTempAllocaWithoutCast(
KernelLaunchParamsTy, CharUnits::fromQuantity(16),
"kernel_launch_params");
+ Address KernelArgs = CGF.CreateTempAlloca(
+ PtrTy, LangAS::Default, CharUnits::fromQuantity(16), "kernel_args",
+ llvm::ConstantInt::get(SizeTy, std::max<size_t>(1, Args.size())));
+ Address KernelArgSizes = CGF.CreateTempAlloca(
+ SizeTy, LangAS::Default, CharUnits::fromQuantity(16), "kernel_arg_sizes",
+ llvm::ConstantInt::get(SizeTy, std::max<size_t>(1, Args.size())));
- CGF.Builder.CreateStore(llvm::ConstantInt::get(Int32Ty, Args.size()),
+ CGF.Builder.CreateStore(KernelArgs.emitRawPointer(CGF),
CGF.Builder.CreateStructGEP(KernelLaunchParams, 0));
- CGF.Builder.CreateStore(KernelArgsPtrs.emitRawPointer(CGF),
+ CGF.Builder.CreateStore(llvm::ConstantInt::get(Int64Ty, Args.size()),
CGF.Builder.CreateStructGEP(KernelLaunchParams, 1));
+ CGF.Builder.CreateStore(KernelArgSizes.emitRawPointer(CGF),
+ CGF.Builder.CreateStructGEP(KernelLaunchParams, 2));
for (unsigned i = 0; i < Args.size(); ++i) {
- auto *ArgVal = CGF.Builder.CreateLoad(CGF.GetAddrOfLocalVar(Args[i]));
- Address ArgAddr = CGF.Builder.CreateStructGEP(KernelArgs, i);
- CGF.Builder.CreateStore(ArgVal, ArgAddr);
- CGF.Builder.CreateStore(ArgAddr.emitRawPointer(CGF),
- CGF.Builder.CreateConstArrayGEP(KernelArgsPtrs, i));
+ llvm::Value *VarPtr = CGF.GetAddrOfLocalVar(Args[i]).emitRawPointer(CGF);
+ llvm::Value *VoidVarPtr = CGF.Builder.CreatePointerCast(VarPtr, PtrTy);
+ CGF.Builder.CreateDefaultAlignedStore(
+ VoidVarPtr, CGF.Builder.CreateConstGEP1_32(
+ PtrTy, KernelArgs.emitRawPointer(CGF), i));
+
+ auto ArgSize = CGM.getDataLayout().getTypeAllocSize(
+ CGM.getTypes().ConvertType(Args[i]->getType()));
+ CGF.Builder.CreateDefaultAlignedStore(
+ llvm::ConstantInt::get(SizeTy, ArgSize),
+ CGF.Builder.CreateConstGEP1_32(SizeTy,
+ KernelArgSizes.emitRawPointer(CGF), i));
}
return KernelLaunchParams;
@@ -408,8 +417,9 @@ Address CGNVCUDARuntime::prepareKernelArgs(CodeGenFunction &CGF,
// array and kernels are launched using cudaLaunchKernel().
void CGNVCUDARuntime::emitDeviceStubBodyNew(CodeGenFunction &CGF,
FunctionArgList &Args) {
+ bool UsesLLVMOffloading = CGF.getLangOpts().OffloadViaLLVM;
// Build the shadow stack entry at the very start of the function.
- Address KernelArgs = CGF.getLangOpts().OffloadViaLLVM
+ Address KernelArgs = UsesLLVMOffloading
? prepareKernelArgsLLVMOffload(CGF, Args)
: prepareKernelArgs(CGF, Args);
@@ -435,7 +445,9 @@ void CGNVCUDARuntime::emitDeviceStubBodyNew(CodeGenFunction &CGF,
else if (CGF.getLangOpts().CUDA)
KernelLaunchAPI = KernelLaunchAPI + "_ptsz";
}
- auto LaunchKernelName = addPrefixToName(KernelLaunchAPI);
+ /// Use __llvmLaunchKernel for LLVMOffload.
+ auto LaunchKernelName = UsesLLVMOffloading ? "__llvm" + KernelLaunchAPI
+ : addPrefixToName(KernelLaunchAPI);
const IdentifierInfo &cudaLaunchKernelII =
CGM.getContext().Idents.get(LaunchKernelName);
FunctionDecl *cudaLaunchKernelFD = nullptr;
@@ -1282,9 +1294,6 @@ void CGNVCUDARuntime::createOffloadingEntries() {
llvm::object::OffloadKind Kind = CGM.getLangOpts().HIP
? llvm::object::OffloadKind::OFK_HIP
: llvm::object::OffloadKind::OFK_Cuda;
- // For now, just spoof this as OpenMP because that's the runtime it uses.
- if (CGM.getLangOpts().OffloadViaLLVM)
- Kind = llvm::object::OffloadKind::OFK_OpenMP;
llvm::Module &M = CGM.getModule();
for (KernelInfo &I : EmittedKernels)
diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp
index 38795f7c2ae7a..f27557c09fca3 100644
--- a/clang/lib/Driver/Driver.cpp
+++ b/clang/lib/Driver/Driver.cpp
@@ -905,10 +905,15 @@ getSystemOffloadArchs(Compilation &C, Action::OffloadKind Kind) {
if (llvm::ErrorOr<std::string> Executable =
llvm::sys::findProgramByName(Program, {C.getDriver().Dir})) {
llvm::SmallVector<StringRef> Args{*Executable};
- if (Kind == Action::OFK_HIP)
- Args.push_back("--only=amdgpu");
- else if (Kind == Action::OFK_Cuda)
- Args.push_back("--only=nvptx");
+ bool UsesLLVMOffloading =
+ C.getArgs().hasFlag(options::OPT_foffload_via_llvm,
+ options::OPT_fno_offload_via_llvm, false);
+ if (!UsesLLVMOffloading) {
+ if (Kind == Action::OFK_HIP)
+ Args.push_back("--only=amdgpu");
+ else if (Kind == Action::OFK_Cuda)
+ Args.push_back("--only=nvptx");
+ }
auto StdoutOrErr = C.getDriver().executeProgram(Args);
if (!StdoutOrErr) {
@@ -965,15 +970,20 @@ static TripleSet inferOffloadToolchains(Compilation &C,
ID = StringToOffloadArch(
getProcessorFromTargetID(llvm::Triple("amdgcn-amd-amdhsa"), Arch));
- if (Kind == Action::OFK_HIP && !IsAMDOffloadArch(ID)) {
- C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch)
- << "HIP" << Arch;
- return {};
- }
- if (Kind == Action::OFK_Cuda && !IsNVIDIAOffloadArch(ID)) {
- C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch)
- << "CUDA" << Arch;
- return {};
+ bool UsesLLVMOffloading =
+ C.getArgs().hasFlag(options::OPT_foffload_via_llvm,
+ options::OPT_fno_offload_via_llvm, false);
+ if (!UsesLLVMOffloading) {
+ if (Kind == Action::OFK_HIP && !IsAMDOffloadArch(ID)) {
+ C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch)
+ << "HIP" << Arch;
+ return {};
+ }
+ if (Kind == Action::OFK_Cuda && !IsNVIDIAOffloadArch(ID)) {
+ C.getDriver().Diag(clang::diag::err_drv_offload_bad_gpu_arch)
+ << "CUDA" << Arch;
+ return {};
+ }
}
if (Kind == Action::OFK_OpenMP &&
(ID == OffloadArch::Unknown || ID == OffloadArch::Unused)) {
@@ -989,6 +999,8 @@ static TripleSet inferOffloadToolchains(Compilation &C,
llvm::Triple Triple =
OffloadArchToTriple(C.getDefaultToolChain().getTriple(), ID);
+ if (UsesLLVMOffloading)
+ Triple.setEnvironment(llvm::Triple::LLVM);
// Make a new argument that dispatches this argument to the appropriate
// toolchain. This is required when we infer it and create potentially
@@ -1032,32 +1044,30 @@ static TripleSet inferOffloadToolchains(Compilation &C,
void Driver::CreateOffloadingDeviceToolChains(Compilation &C,
InputList &Inputs) {
- bool UseLLVMOffload = C.getInputArgs().hasArg(
- options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
bool IsCuda =
- llvm::any_of(Inputs,
- [](std::pair<types::ID, const llvm::opt::Arg *> &I) {
- return types::isCuda(I.first);
- }) &&
- !UseLLVMOffload;
+ llvm::any_of(Inputs, [](std::pair<types::ID, const llvm::opt::Arg *> &I) {
+ return types::isCuda(I.first);
+ });
bool IsHIP =
(llvm::any_of(Inputs,
[](std::pair<types::ID, const llvm::opt::Arg *> &I) {
return types::isHIP(I.first);
}) ||
C.getInputArgs().hasArg(options::OPT_hip_link) ||
- C.getInputArgs().hasArg(options::OPT_hipstdpar)) &&
- !UseLLVMOffload;
+ C.getInputArgs().hasArg(options::OPT_hipstdpar));
bool IsSYCL = C.getInputArgs().hasFlag(options::OPT_fsycl,
options::OPT_fno_sycl, false);
bool IsOpenMPOffloading =
- UseLLVMOffload ||
(C.getInputArgs().hasFlag(options::OPT_fopenmp, options::OPT_fopenmp_EQ,
options::OPT_fno_openmp, false) &&
(C.getInputArgs().hasArg(options::OPT_offload_targets_EQ) ||
(C.getInputArgs().hasArg(options::OPT_offload_arch_EQ) &&
!(IsCuda || IsHIP))));
+ // We currently don't support any kind of mixed offloading.
+ if (IsOpenMPOffloading)
+ IsCuda = IsHIP = IsSYCL = false;
+
llvm::SmallSet<Action::OffloadKind, 4> Kinds;
const std::pair<bool, Action::OffloadKind> ActiveKinds[] = {
{IsCuda, Action::OFK_Cuda},
@@ -1143,7 +1153,7 @@ void Driver::CreateOffloadingDeviceToolChains(Compilation &C,
C.getDefaultToolChain().getTriple());
// Emit a warning if the detected CUDA version is too new.
- if (Kind == Action::OFK_Cuda) {
+ if (Kind == Action::OFK_Cuda && Target.getOS() == llvm::Triple::CUDA) {
auto &CudaInstallation =
static_cast<const toolchains::CudaToolChain &>(TC).CudaInstallation;
if (CudaInstallation.isValid())
@@ -5069,6 +5079,9 @@ Driver::BuildOffloadingActions(Compilation &C, llvm::opt::DerivedArgList &Args,
getFinalPhase(Args) == phases::Preprocess))
return HostAction;
+ bool UsesLLVMOffloading = Args.hasArg(
+ options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
+
ActionList OffloadActions;
OffloadAction::DeviceDependences DDeps;
@@ -5089,7 +5102,6 @@ Driver::BuildOffloadingActions(Compilation &C, llvm::opt::DerivedArgList &Args,
types::ID InputType = Input.first;
const Arg *InputArg = Input.second;
- // The toolchain can be active for unsupported file types.
if ((Kind == Action::OFK_Cuda && !types::isCuda(InputType)) ||
(Kind == Action::OFK_HIP && !types::isHIP(InputType)))
continue;
@@ -5184,9 +5196,12 @@ Driver::BuildOffloadingActions(Compilation &C, llvm::opt::DerivedArgList &Args,
OffloadAction::DeviceDependences DDep;
DDep.add(*A, *TCAndArch->first, TCAndArch->second, Kind);
- // Compiling CUDA in non-RDC mode uses the PTX output if available.
+ // The legacy CUDA fatbinary path can include PTX alongside the cubin.
+ // The LLVM offload wrapper path feeds these images through a device
+ // linker first, and clang-nvlink-wrapper does not accept PTX as input.
for (Action *Input : A->getInputs())
- if (Kind == Action::OFK_Cuda && A->getType() == types::TY_Object &&
+ if (!UsesLLVMOffloading && Kind == Action::OFK_Cuda &&
+ A->getType() == types::TY_Object &&
!Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc,
false))
DDep.add(*Input, *TCAndArch->first, TCAndArch->second, Kind);
@@ -5214,7 +5229,7 @@ Driver::BuildOffloadingActions(Compilation &C, llvm::opt::DerivedArgList &Args,
return HostAction;
OffloadAction::DeviceDependences DDep;
- if (C.isOffloadingHostKind(Action::OFK_Cuda) &&
+ if (!UsesLLVMOffloading && C.isOffloadingHostKind(Action::OFK_Cuda) &&
!Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc, false)) {
// If we are not in RDC-mode we just emit the final CUDA fatbinary for
// each translation unit without requiring any linking.
@@ -5222,7 +5237,7 @@ Driver::BuildOffloadingActions(Compilation &C, llvm::opt::DerivedArgList &Args,
C.MakeAction<LinkJobAction>(OffloadActions, types::TY_CUDA_FATBIN);
DDep.add(*FatbinAction, *C.getSingleOffloadToolChain<Action::OFK_Cuda>(),
/*BA=*/{}, Action::OFK_Cuda);
- } else if (HIPNoRDC && offloadDeviceOnly()) {
+ } else if (!UsesLLVMOffloading && HIPNoRDC && offloadDeviceOnly()) {
// If we are in device-only non-RDC-mode we just emit the final HIP
// fatbinary for each translation unit, linking each input individually.
Action *FatbinAction =
@@ -5230,7 +5245,7 @@ Driver::BuildOffloadingActions(Compilation &C, llvm::opt::DerivedArgList &Args,
DDep.add(*FatbinAction,
*C.getOffloadToolChains<Action::OFK_HIP>().first->second,
/*BA=*/{}, Action::OFK_HIP);
- } else if (HIPNoRDC) {
+ } else if (!UsesLLVMOffloading && HIPNoRDC) {
// Host + device assembly: defer to clang-offload-bundler (see
// BuildActions).
if (HIPAsmBundleDeviceOut &&
@@ -7091,7 +7106,8 @@ const ToolChain &Driver::getOffloadToolChain(
// For AMDHSA offloading (HIP, OpenMP), use the unified AMDGPUToolChain
// This handles both amdgpu-amd-amdhsa and spirv64-amd-amdhsa
// FIXME: This should not key off language or OS.
- if (Kind == Action::OFK_HIP || Kind == Action::OFK_OpenMP)
+ if (Kind == Action::OFK_HIP || Kind == Action::OFK_OpenMP ||
+ Kind == Action::OFK_Cuda)
TC = std::make_unique<toolchains::AMDGPUToolChain>(*this, Target, Args,
HostTC.get(), Kind);
break;
diff --git a/clang/lib/Driver/ToolChains/AMDGPU.cpp b/clang/lib/Driver/ToolChains/AMDGPU.cpp
index 7bce060de0596..5c30417365b94 100644
--- a/clang/lib/Driver/ToolChains/AMDGPU.cpp
+++ b/clang/lib/Driver/ToolChains/AMDGPU.cpp
@@ -515,6 +515,24 @@ void RocmInstallationDetector::AddHIPIncludeArgs(const ArgList &DriverArgs,
!DriverArgs.hasArg(options::OPT_nohipwrapperinc);
bool HasHipStdPar = DriverArgs.hasArg(options::OPT_hipstdpar);
+ if (DriverArgs.hasFlag(options::OPT_foffload_via_llvm,
+ options::OPT_fno_offload_via_llvm, false)) {
+ if (DriverArgs.hasFlag(options::OPT_offload_inc,
+ options::OPT_no_offload_inc, true) &&
+ !DriverArgs.hasArg(options::OPT_nohipwrapperinc) &&
+ !DriverArgs.hasArg(options::OPT_nobuiltininc)) {
+ CC1Args.append({"-include", "__clang_gpu_device_functions.h"});
+
+ SmallString<128> HIPIncludePath(D.ResourceDir);
+ llvm::sys::path::append(HIPIncludePath, "..", "..", "..");
+ llvm::sys::path::append(HIPIncludePath, "include", "offload");
+ CC1Args.push_back("-internal-isystem");
+ CC1Args.push_back(DriverArgs.MakeArgString(HIPIncludePath));
+ CC1Args.append({"-include", "hip/hip_runtime.h"});
+ }
+ return;
+ }
+
if (!DriverArgs.hasArg(options::OPT_nobuiltininc)) {
// HIP header includes standard library wrapper headers under clang
// cuda_wrappers directory. Since these wrapper headers include_next
@@ -699,7 +717,11 @@ AMDGPUToolChain::AMDGPUToolChain(const Driver &D, const llvm::Triple &Triple,
: Generic_ELF(D, Triple, Args),
OptionsDefault(
{{options::OPT_O, "3"}, {options::OPT_cl_std_EQ, "CL1.2"}}),
- HostTC(HostTC_), UseHIPLinker(Kind == Action::OFK_HIP),
+ HostTC(HostTC_),
+ UseHIPLinker(Kind == Action::OFK_HIP ||
+ (Kind == Action::OFK_Cuda &&
+ Args.hasFlag(options::OPT_foffload_via_llvm,
+ options::OPT_fno_offload_via_llvm, false))),
ShouldLinkDeviceLibs(ShouldLinkDeviceLibs) {
loadMultilibsFromYAML(Args, D);
@@ -709,8 +731,10 @@ AMDGPUToolChain::AMDGPUToolChain(const Driver &D, const llvm::Triple &Triple,
// each tool invocation.
checkAMDGPUCodeObjectVersion(D, Args);
+ bool UsesLLVMOffloading = Args.hasFlag(
+ options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
if (Triple.getOS() == llvm::Triple::AMDHSA &&
- Triple.getEnvironment() != llvm::Triple::LLVM)
+ Triple.getEnvironment() != llvm::Triple::LLVM && !UsesLLVMOffloading)
RocmInstallation->detectDeviceLibrary();
if (HostTC)
@@ -889,7 +913,10 @@ bool AMDGPUToolChain::isWave64(const llvm::opt::ArgList &DriverArgs,
void AMDGPUToolChain::addClangTargetOptions(
const llvm::opt::ArgList &DriverArgs, llvm::opt::ArgStringList &CC1Args,
BoundArch BA, Action::OffloadKind DeviceOffloadingKind) const {
- if (DeviceOffloadingKind == Action::OFK_HIP) {
+ bool UsesLLVMOffloading = DriverArgs.hasFlag(
+ options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
+ if (DeviceOffloadingKind == Action::OFK_HIP ||
+ (DeviceOffloadingKind == Action::OFK_Cuda && UsesLLVMOffloading)) {
CC1Args.append({"-fcuda-is-device", "-fno-threadsafe-statics"});
if (!DriverArgs.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc,
diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
index 94f9a26aac39f..035bc3f5b4273 100644
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -53,6 +53,7 @@
#include "llvm/Support/Path.h"
#include "llvm/Support/Process.h"
#include "llvm/Support/YAMLParser.h"
+#include "llvm/Support/raw_ostream.h"
#include "llvm/TargetParser/AArch64TargetParser.h"
#include "llvm/TargetParser/ARMTargetParserCommon.h"
#include "llvm/TargetParser/Host.h"
@@ -952,9 +953,12 @@ void Clang::AddPreprocessingOptions(Compilation &C, const JobAction &JA,
// before we -I or -include anything else, because we must pick up the
// CUDA/HIP/SYCL headers from the particular CUDA/ROCm/SYCL installation,
// rather than from e.g. /usr/local/include.
- if (JA.isOffloading(Action::OFK_Cuda))
+ bool UsesLLVMOffloading = Args.hasFlag(
+ options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
+ if (JA.isOffloading(Action::OFK_Cuda) && !UsesLLVMOffloading) {
getToolChain().AddCudaIncludeArgs(Args, CmdArgs);
- if (JA.isOffloading(Action::OFK_HIP))
+ }
+ if (JA.isOffloading(Action::OFK_HIP) && !UsesLLVMOffloading)
getToolChain().AddHIPIncludeArgs(Args, CmdArgs);
if (JA.isOffloading(Action::OFK_SYCL))
getToolChain().addSYCLIncludeArgs(Args, CmdArgs);
@@ -979,17 +983,35 @@ void Clang::AddPreprocessingOptions(Compilation &C, const JobAction &JA,
CmdArgs.push_back("-include");
CmdArgs.push_back("__clang_openmp_device_functions.h");
}
+ bool isCudaInput = llvm::any_of(
+ Inputs, [](const InputInfo &I) { return types::isCuda(I.getType()); });
+ if (UsesLLVMOffloading && JA.isOffloading(Action::OFK_Cuda) &&
+ Args.hasFlag(options::OPT_offload_inc, options::OPT_no_offload_inc,
+ true) &&
+ !Args.hasArg(options::OPT_nobuiltininc) && isCudaInput) {
+ CmdArgs.append({"-include", "__clang_gpu_device_functions.h"});
+
+ SmallString<128> OffloadCudaInclude(D.Dir);
+ llvm::sys::path::append(OffloadCudaInclude, "..", "include", "offload",
+ "cuda");
+ CmdArgs.append({"-internal-isystem", Args.MakeArgString(OffloadCudaInclude),
+ "-include"});
+ CmdArgs.push_back("cuda_runtime.h");
+ }
+ bool isHIPInput = llvm::any_of(
+ Inputs, [](const InputInfo &I) { return types::isHIP(I.getType()); });
+ if (UsesLLVMOffloading && JA.isOffloading(Action::OFK_HIP) &&
+ Args.hasFlag(options::OPT_offload_inc, options::OPT_no_offload_inc,
+ true) &&
+ !Args.hasArg(options::OPT_nohipwrapperinc) &&
+ !Args.hasArg(options::OPT_nobuiltininc) && isHIPInput) {
+ CmdArgs.append({"-include", "__clang_gpu_device_functions.h"});
- if (Args.hasArg(options::OPT_foffload_via_llvm)) {
- // Add llvm_wrappers/* to our system include path. This lets us wrap
- // standard library headers and other headers.
- SmallString<128> P(D.ResourceDir);
- llvm::sys::path::append(P, "include", "llvm_offload_wrappers");
- CmdArgs.append({"-internal-isystem", Args.MakeArgString(P), "-include"});
- if (JA.isDeviceOffloading(Action::OFK_OpenMP))
- CmdArgs.push_back("__llvm_offload_device.h");
- else
- CmdArgs.push_back("__llvm_offload_host.h");
+ SmallString<128> OffloadHIPInclude(D.Dir);
+ llvm::sys::path::append(OffloadHIPInclude, "..", "include", "offload");
+ CmdArgs.append({"-internal-isystem", Args.MakeArgString(OffloadHIPInclude),
+ "-include"});
+ CmdArgs.push_back("hip/hip_runtime.h");
}
// Add -i* options, and automatically translate to
@@ -1171,7 +1193,7 @@ void Clang::AddPreprocessingOptions(Compilation &C, const JobAction &JA,
Args.hasFlag(options::OPT_offload_inc, options::OPT_no_offload_inc,
true) &&
!Args.hasArg(options::OPT_nobuiltininc) &&
- (C.getActiveOffloadKinds() == Action::OFK_OpenMP)) {
+ JA.isDeviceOffloading(Action::OFK_OpenMP)) {
// TODO: CUDA / HIP include their own headers for some common functions
// implemented here. We'll need to clean those up so they do not conflict.
SmallString<128> P(D.ResourceDir);
@@ -5194,6 +5216,8 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
bool IsSYCLDevice = JA.isDeviceOffloading(Action::OFK_SYCL);
bool IsOpenMPDevice = JA.isDeviceOffloading(Action::OFK_OpenMP);
bool IsExtractAPI = isa<ExtractAPIJobAction>(JA);
+ bool UsesLLVMOffloading = Args.hasFlag(
+ options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
bool IsDeviceOffloadAction = !(JA.isDeviceOffloading(Action::OFK_None) ||
JA.isDeviceOffloading(Action::OFK_Host));
bool IsHostOffloadingAction =
@@ -5312,7 +5336,7 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
}
}
- if (IsCuda && !IsCudaDevice) {
+ if (IsCuda && !IsCudaDevice && !UsesLLVMOffloading) {
// We need to figure out which CUDA version we're compiling for, as that
// determines how we load and launch GPU kernels.
auto *CTC = static_cast<const toolchains::CudaToolChain *>(
@@ -8313,11 +8337,11 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
// Host-side offloading compilation receives all device-side outputs. Include
// them in the host compilation depending on the target. If the host inputs
// are not empty we use the new-driver scheme, otherwise use the old scheme.
- if ((IsCuda || IsHIP) && CudaDeviceInput) {
+ if ((IsCuda || IsHIP) && !UsesLLVMOffloading && CudaDeviceInput) {
CmdArgs.push_back("-fcuda-include-gpubinary");
CmdArgs.push_back(CudaDeviceInput->getFilename());
} else if (!HostOffloadingInputs.empty()) {
- if ((IsCuda || IsHIP) && !IsRDCMode) {
+ if ((IsCuda || IsHIP) && !UsesLLVMOffloading && !IsRDCMode) {
assert(HostOffloadingInputs.size() == 1 && "Only one input expected");
CmdArgs.push_back("-fcuda-include-gpubinary");
CmdArgs.push_back(HostOffloadingInputs.front().getFilename());
diff --git a/clang/lib/Driver/ToolChains/CommonArgs.cpp b/clang/lib/Driver/ToolChains/CommonArgs.cpp
index a2b99dffc383e..37a4392f2941f 100644
--- a/clang/lib/Driver/ToolChains/CommonArgs.cpp
+++ b/clang/lib/Driver/ToolChains/CommonArgs.cpp
@@ -1459,18 +1459,25 @@ void tools::addArchSpecificRPath(const ToolChain &TC, const ArgList &Args,
}
}
+bool tools::addLLVMOffloadingRuntime(const Compilation &C,
+ ArgStringList &CmdArgs,
+ const ToolChain &TC, const ArgList &Args) {
+
+ if (!Args.hasFlag(options::OPT_foffload_via_llvm,
+ options::OPT_fno_offload_via_llvm, false))
+ return false;
+
+ CmdArgs.push_back("-lLLVMOffloadKernel");
+ return true;
+}
+
bool tools::addOpenMPRuntime(const Compilation &C, ArgStringList &CmdArgs,
const ToolChain &TC, const ArgList &Args,
bool ForceStaticHostRuntime, bool IsOffloadingHost,
bool GompNeedsRT) {
if (!Args.hasFlag(options::OPT_fopenmp, options::OPT_fopenmp_EQ,
- options::OPT_fno_openmp, false)) {
- // We need libomptarget (liboffload) if it's the choosen offloading runtime.
- if (Args.hasFlag(options::OPT_foffload_via_llvm,
- options::OPT_fno_offload_via_llvm, false))
- CmdArgs.push_back("-lomptarget");
+ options::OPT_fno_openmp, false))
return false;
- }
Driver::OpenMPRuntimeKind RTKind = TC.getDriver().getOpenMPRuntime(Args);
diff --git a/clang/lib/Driver/ToolChains/Cuda.cpp b/clang/lib/Driver/ToolChains/Cuda.cpp
index 5040e451ed9e5..f79109f71efb3 100644
--- a/clang/lib/Driver/ToolChains/Cuda.cpp
+++ b/clang/lib/Driver/ToolChains/Cuda.cpp
@@ -301,6 +301,23 @@ CudaInstallationDetector::CudaInstallationDetector(
void CudaInstallationDetector::AddCudaIncludeArgs(
const ArgList &DriverArgs, ArgStringList &CC1Args) const {
+ if (DriverArgs.hasFlag(options::OPT_foffload_via_llvm,
+ options::OPT_fno_offload_via_llvm, false)) {
+ if (DriverArgs.hasFlag(options::OPT_offload_inc,
+ options::OPT_no_offload_inc, true) &&
+ !DriverArgs.hasArg(options::OPT_nobuiltininc)) {
+ CC1Args.append({"-include", "__clang_gpu_device_functions.h"});
+
+ SmallString<128> CudaIncludePath(D.ResourceDir);
+ llvm::sys::path::append(CudaIncludePath, "..", "..", "..");
+ llvm::sys::path::append(CudaIncludePath, "include", "offload", "cuda");
+ CC1Args.push_back("-internal-isystem");
+ CC1Args.push_back(DriverArgs.MakeArgString(CudaIncludePath));
+ CC1Args.append({"-include", "cuda_runtime.h"});
+ }
+ return;
+ }
+
if (!DriverArgs.hasArg(options::OPT_nobuiltininc)) {
// Add cuda_wrappers/* to our system include path. This lets us wrap
// standard library headers.
@@ -395,7 +412,9 @@ void NVPTX::Assembler::ConstructJob(Compilation &C, const JobAction &JA,
const char *LinkingOutput) const {
const auto &TC =
static_cast<const toolchains::NVPTXToolChain &>(getToolChain());
- assert(TC.getTriple().isNVPTX() && "Wrong platform");
+
+ bool UsesLLVMOffloading = Args.hasFlag(
+ options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
BoundArch GPUArch;
// If this is a CUDA action we need to extract the device architecture
@@ -418,7 +437,7 @@ void NVPTX::Assembler::ConstructJob(Compilation &C, const JobAction &JA,
"Device action expected to have an architecture.");
// Check that our installation's ptxas supports gpu_arch.
- if (!Args.hasArg(options::OPT_no_cuda_version_check)) {
+ if (!UsesLLVMOffloading && !Args.hasArg(options::OPT_no_cuda_version_check)) {
TC.CudaInstallation.CheckCudaVersionSupportsArch(GPUArch.Arch);
}
@@ -491,7 +510,8 @@ void NVPTX::Assembler::ConstructJob(Compilation &C, const JobAction &JA,
/*Default=*/true);
else if (JA.isOffloading(Action::OFK_Cuda))
// In CUDA we generate relocatable code by default.
- Relocatable = Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc,
+ Relocatable = UsesLLVMOffloading ||
+ Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc,
/*Default=*/false);
else
// Otherwise, we are compiling directly and should create linkable output.
@@ -540,7 +560,9 @@ void NVPTX::FatBinary::ConstructJob(Compilation &C, const JobAction &JA,
const char *LinkingOutput) const {
const auto &TC =
static_cast<const toolchains::CudaToolChain &>(getToolChain());
- assert(TC.getTriple().isNVPTX() && "Wrong platform");
+ bool UsesLLVMOffloading = Args.hasFlag(
+ options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
+ assert((UsesLLVMOffloading || TC.getTriple().isNVPTX()) && "Wrong platform");
ArgStringList CmdArgs;
if (TC.CudaInstallation.version() <= CudaVersion::CUDA_100)
@@ -588,7 +610,9 @@ void NVPTX::Linker::ConstructJob(Compilation &C, const JobAction &JA,
static_cast<const toolchains::NVPTXToolChain &>(getToolChain());
ArgStringList CmdArgs;
- assert(TC.getTriple().isNVPTX() && "Wrong platform");
+ bool UsesLLVMOffloading = Args.hasFlag(
+ options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
+ assert((UsesLLVMOffloading || TC.getTriple().isNVPTX()) && "Wrong platform");
assert((Output.isFilename() || Output.isNothing()) && "Invalid output.");
if (Output.isFilename()) {
@@ -886,9 +910,12 @@ void CudaToolChain::addClangTargetOptions(
BoundArch BA, Action::OffloadKind DeviceOffloadingKind) const {
HostTC.addClangTargetOptions(DriverArgs, CC1Args, BA, DeviceOffloadingKind);
+ bool UsesLLVMOffloading = DriverArgs.hasFlag(
+ options::OPT_foffload_via_llvm, options::OPT_fno_offload_via_llvm, false);
+
StringRef GpuArch = DriverArgs.getLastArgValue(options::OPT_march_EQ);
assert((DeviceOffloadingKind == Action::OFK_OpenMP ||
- DeviceOffloadingKind == Action::OFK_Cuda) &&
+ DeviceOffloadingKind == Action::OFK_Cuda || UsesLLVMOffloading) &&
"Only OpenMP or CUDA offloading kinds are supported for NVIDIA GPUs.");
CC1Args.append({"-fcuda-is-device", "-mllvm",
@@ -907,6 +934,9 @@ void CudaToolChain::addClangTargetOptions(
DriverArgs.hasArg(options::OPT_S))
return;
+ if (UsesLLVMOffloading)
+ return;
+
std::string LibDeviceFile = CudaInstallation.getLibDeviceFile(GpuArch);
if (LibDeviceFile.empty()) {
getDriver().Diag(diag::err_drv_no_cuda_libdevice) << GpuArch;
@@ -916,13 +946,6 @@ void CudaToolChain::addClangTargetOptions(
CC1Args.push_back("-mlink-builtin-bitcode");
CC1Args.push_back(DriverArgs.MakeArgString(LibDeviceFile));
- // For now, we don't use any Offload/OpenMP device runtime when we offload
- // CUDA via LLVM/Offload. We should split the Offload/OpenMP device runtime
- // and include the "generic" (or CUDA-specific) parts.
- if (DriverArgs.hasFlag(options::OPT_foffload_via_llvm,
- options::OPT_fno_offload_via_llvm, false))
- return;
-
clang::CudaVersion CudaInstallationVersion = CudaInstallation.version();
if (CudaInstallationVersion >= CudaVersion::UNKNOWN)
@@ -963,6 +986,23 @@ llvm::DenormalMode CudaToolChain::getDefaultDenormalModeForType(
void CudaToolChain::AddCudaIncludeArgs(const ArgList &DriverArgs,
ArgStringList &CC1Args) const {
+ if (DriverArgs.hasFlag(options::OPT_foffload_via_llvm,
+ options::OPT_fno_offload_via_llvm, false)) {
+ if (DriverArgs.hasFlag(options::OPT_offload_inc,
+ options::OPT_no_offload_inc, true) &&
+ !DriverArgs.hasArg(options::OPT_nobuiltininc)) {
+ CC1Args.append({"-include", "__clang_gpu_device_functions.h"});
+
+ SmallString<128> CudaIncludePath(getDriver().ResourceDir);
+ llvm::sys::path::append(CudaIncludePath, "..", "..", "..");
+ llvm::sys::path::append(CudaIncludePath, "include", "offload", "cuda");
+ CC1Args.push_back("-internal-isystem");
+ CC1Args.push_back(DriverArgs.MakeArgString(CudaIncludePath));
+ CC1Args.append({"-include", "cuda_runtime.h"});
+ }
+ return;
+ }
+
// Check our CUDA version if we're going to include the CUDA headers.
if (DriverArgs.hasFlag(options::OPT_offload_inc, options::OPT_no_offload_inc,
true) &&
@@ -1035,6 +1075,10 @@ CudaToolChain::GetCXXStdlibType(const ArgList &Args) const {
void CudaToolChain::AddClangSystemIncludeArgs(const ArgList &DriverArgs,
ArgStringList &CC1Args) const {
+ if (DriverArgs.hasFlag(options::OPT_foffload_via_llvm,
+ options::OPT_fno_offload_via_llvm, false))
+ return;
+
HostTC.AddClangSystemIncludeArgs(DriverArgs, CC1Args);
if (DriverArgs.hasFlag(options::OPT_offload_inc, options::OPT_no_offload_inc,
diff --git a/clang/lib/Driver/ToolChains/Gnu.cpp b/clang/lib/Driver/ToolChains/Gnu.cpp
index 24076d8814322..72affac131701 100644
--- a/clang/lib/Driver/ToolChains/Gnu.cpp
+++ b/clang/lib/Driver/ToolChains/Gnu.cpp
@@ -510,6 +510,7 @@ void tools::gnutools::Linker::ConstructJob(Compilation &C, const JobAction &JA,
// FIXME: Does this really make sense for all GNU toolchains?
WantPthread = true;
+ addLLVMOffloadingRuntime(C, CmdArgs, ToolChain, Args);
AddRunTimeLibs(ToolChain, D, CmdArgs, Args);
// LLVM support for atomics on 32-bit SPARC V8+ is incomplete, so
diff --git a/clang/lib/Driver/ToolChains/Linux.cpp b/clang/lib/Driver/ToolChains/Linux.cpp
index 1ab385a9ea001..03017f20e391d 100644
--- a/clang/lib/Driver/ToolChains/Linux.cpp
+++ b/clang/lib/Driver/ToolChains/Linux.cpp
@@ -881,7 +881,9 @@ void Linux::addOffloadRTLibs(unsigned ActiveKinds, const ArgList &Args,
if (!Args.hasFlag(options::OPT_offloadlib, options::OPT_no_offloadlib,
true) ||
Args.hasArg(options::OPT_nostdlib) ||
- Args.hasArg(options::OPT_no_hip_rt) || Args.hasArg(options::OPT_r))
+ Args.hasArg(options::OPT_no_hip_rt) || Args.hasArg(options::OPT_r) ||
+ Args.hasFlag(options::OPT_foffload_via_llvm,
+ options::OPT_fno_offload_via_llvm, false))
return;
llvm::SmallVector<std::pair<StringRef, StringRef>> Libraries;
diff --git a/clang/lib/Headers/__clang_gpu_builtin_vars.h b/clang/lib/Headers/__clang_gpu_builtin_vars.h
index b80248dcd2be3..f3137be0a3181 100644
--- a/clang/lib/Headers/__clang_gpu_builtin_vars.h
+++ b/clang/lib/Headers/__clang_gpu_builtin_vars.h
@@ -6,6 +6,8 @@
//
//===-----------------------------------------------------------------------===
+#include <stddef.h>
+
#ifndef __CLANG_GPU_BUILTIN_VARS_H__
#define __CLANG_GPU_BUILTIN_VARS_H__
@@ -20,6 +22,23 @@ static inline __attribute__((device)) const struct {
}
} warpSize{};
+extern "C" {
+
+typedef struct dim3 {
+ dim3() {}
+ dim3(unsigned x) : x(x) {}
+ unsigned x = 0, y = 0, z = 0;
+} dim3;
+
+// TODO: For some reason the CUDA device compilation requires this declaration
+// to be present on the device while it is only used on the host.
+unsigned __llvmPushCallConfiguration(dim3 gridDim, dim3 blockDim,
+ size_t sharedMem = 0, void *stream = 0);
+unsigned __llvmLaunchKernel(const void *func, dim3 gridDim, dim3 blockDim,
+ void **args, size_t sharedMem = 0,
+ void *stream = 0);
+}
+
// Make sure nobody can create instances of the coordinate types, take their
// address, copy, or assign them.
#pragma push_macro("__GPU_DISALLOW_BUILTINVAR_ACCESS")
diff --git a/clang/test/CodeGenCUDA/Inputs/cuda.h b/clang/test/CodeGenCUDA/Inputs/cuda.h
index 421fa4dd7dbae..83bb7b7bdbb7f 100644
--- a/clang/test/CodeGenCUDA/Inputs/cuda.h
+++ b/clang/test/CodeGenCUDA/Inputs/cuda.h
@@ -55,7 +55,7 @@ extern "C" hipError_t hipLaunchKernel_spt(const void *func, dim3 gridDim,
#elif __OFFLOAD_VIA_LLVM__
extern "C" unsigned __llvmPushCallConfiguration(dim3 gridDim, dim3 blockDim,
size_t sharedMem = 0, void *stream = 0);
-extern "C" unsigned llvmLaunchKernel(const void *func, dim3 gridDim, dim3 blockDim,
+extern "C" unsigned __llvmLaunchKernel(const void *func, dim3 gridDim, dim3 blockDim,
void **args, size_t sharedMem = 0, void *stream = 0);
#else
typedef struct cudaStream *cudaStream_t;
diff --git a/clang/test/CodeGenCUDA/offload_via_llvm.cu b/clang/test/CodeGenCUDA/offload_via_llvm.cu
index b13a64c81b775..c58e793d11826 100644
--- a/clang/test/CodeGenCUDA/offload_via_llvm.cu
+++ b/clang/test/CodeGenCUDA/offload_via_llvm.cu
@@ -14,9 +14,7 @@
// HST-NEXT: [[DOTADDR1:%.*]] = alloca i16, align 2
// HST-NEXT: [[DOTADDR2:%.*]] = alloca ptr, align 4
// HST-NEXT: [[DOTADDR3:%.*]] = alloca ptr, align 4
-// HST-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[TMP0]], align 16
-// HST-NEXT: [[KERNEL_ARGS_PTRS:%.*]] = alloca [4 x ptr], align 16
-// HST-NEXT: [[KERNEL_LAUNCH_PARAMS:%.*]] = alloca [[TMP1]], align 16
+// HST-NEXT: [[KERNEL_LAUNCH_PARAMS:%.*]] = alloca [[TMP0]], align 16
// HST-NEXT: [[GRID_DIM:%.*]] = alloca [[STRUCT_DIM3:%.*]], align 8
// HST-NEXT: [[BLOCK_DIM:%.*]] = alloca [[STRUCT_DIM3]], align 8
// HST-NEXT: [[SHMEM_SIZE:%.*]] = alloca i32, align 4
@@ -25,34 +23,34 @@
// HST-NEXT: store i16 [[TMP1]], ptr [[DOTADDR1]], align 2
// HST-NEXT: store ptr [[TMP2]], ptr [[DOTADDR2]], align 4
// HST-NEXT: store ptr [[TMP3]], ptr [[DOTADDR3]], align 4
-// HST-NEXT: [[TMP4:%.*]] = getelementptr inbounds nuw [[TMP1]], ptr [[KERNEL_LAUNCH_PARAMS]], i32 0, i32 0
-// HST-NEXT: store i32 4, ptr [[TMP4]], align 16
-// HST-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[TMP1]], ptr [[KERNEL_LAUNCH_PARAMS]], i32 0, i32 1
-// HST-NEXT: store ptr [[KERNEL_ARGS_PTRS]], ptr [[TMP5]], align 4
-// HST-NEXT: [[TMP6:%.*]] = load i32, ptr [[DOTADDR]], align 4
-// HST-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_ARGS]], i32 0, i32 0
-// HST-NEXT: store i32 [[TMP6]], ptr [[TMP7]], align 16
-// HST-NEXT: [[TMP8:%.*]] = getelementptr inbounds [4 x ptr], ptr [[KERNEL_ARGS_PTRS]], i32 0, i32 0
-// HST-NEXT: store ptr [[TMP7]], ptr [[TMP8]], align 16
-// HST-NEXT: [[TMP9:%.*]] = load i16, ptr [[DOTADDR1]], align 2
-// HST-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_ARGS]], i32 0, i32 1
-// HST-NEXT: store i16 [[TMP9]], ptr [[TMP10]], align 4
-// HST-NEXT: [[TMP11:%.*]] = getelementptr inbounds [4 x ptr], ptr [[KERNEL_ARGS_PTRS]], i32 0, i32 1
-// HST-NEXT: store ptr [[TMP10]], ptr [[TMP11]], align 4
-// HST-NEXT: [[TMP12:%.*]] = load ptr, ptr [[DOTADDR2]], align 4
-// HST-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_ARGS]], i32 0, i32 2
-// HST-NEXT: store ptr [[TMP12]], ptr [[TMP13]], align 8
-// HST-NEXT: [[TMP14:%.*]] = getelementptr inbounds [4 x ptr], ptr [[KERNEL_ARGS_PTRS]], i32 0, i32 2
-// HST-NEXT: store ptr [[TMP13]], ptr [[TMP14]], align 8
-// HST-NEXT: [[TMP15:%.*]] = load ptr, ptr [[DOTADDR3]], align 4
-// HST-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_ARGS]], i32 0, i32 3
-// HST-NEXT: store ptr [[TMP15]], ptr [[TMP16]], align 4
-// HST-NEXT: [[TMP17:%.*]] = getelementptr inbounds [4 x ptr], ptr [[KERNEL_ARGS_PTRS]], i32 0, i32 3
-// HST-NEXT: store ptr [[TMP16]], ptr [[TMP17]], align 4
-// HST-NEXT: [[TMP18:%.*]] = call i32 @__llvmPopCallConfiguration(ptr [[GRID_DIM]], ptr [[BLOCK_DIM]], ptr [[SHMEM_SIZE]], ptr [[STREAM]])
-// HST-NEXT: [[TMP19:%.*]] = load i32, ptr [[SHMEM_SIZE]], align 4
-// HST-NEXT: [[TMP20:%.*]] = load ptr, ptr [[STREAM]], align 4
-// HST-NEXT: [[CALL:%.*]] = call noundef i32 @llvmLaunchKernel(ptr noundef @_Z18__device_stub__fooisPvS_, ptr noundef byval([[STRUCT_DIM3]]) align 4 [[GRID_DIM]], ptr noundef byval([[STRUCT_DIM3]]) align 4 [[BLOCK_DIM]], ptr noundef [[KERNEL_LAUNCH_PARAMS]], i32 noundef [[TMP19]], ptr noundef [[TMP20]]) #[[ATTR3:[0-9]+]]
+// HST-NEXT: [[KERNEL_ARGS:%.*]] = alloca ptr, i32 4, align 16
+// HST-NEXT: [[KERNEL_ARG_SIZES:%.*]] = alloca i32, i32 4, align 16
+// HST-NEXT: [[TMP4:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_LAUNCH_PARAMS]], i32 0, i32 0
+// HST-NEXT: store ptr [[KERNEL_ARGS]], ptr [[TMP4]], align 16
+// HST-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_LAUNCH_PARAMS]], i32 0, i32 1
+// HST-NEXT: store i64 4, ptr [[TMP5]], align 8
+// HST-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[TMP0]], ptr [[KERNEL_LAUNCH_PARAMS]], i32 0, i32 2
+// HST-NEXT: store ptr [[KERNEL_ARG_SIZES]], ptr [[TMP6]], align 16
+// HST-NEXT: [[TMP7:%.*]] = getelementptr ptr, ptr [[KERNEL_ARGS]], i32 0
+// HST-NEXT: store ptr [[DOTADDR]], ptr [[TMP7]], align 4
+// HST-NEXT: [[TMP8:%.*]] = getelementptr i32, ptr [[KERNEL_ARG_SIZES]], i32 0
+// HST-NEXT: store i32 4, ptr [[TMP8]], align 4
+// HST-NEXT: [[TMP9:%.*]] = getelementptr ptr, ptr [[KERNEL_ARGS]], i32 1
+// HST-NEXT: store ptr [[DOTADDR1]], ptr [[TMP9]], align 4
+// HST-NEXT: [[TMP10:%.*]] = getelementptr i32, ptr [[KERNEL_ARG_SIZES]], i32 1
+// HST-NEXT: store i32 2, ptr [[TMP10]], align 4
+// HST-NEXT: [[TMP11:%.*]] = getelementptr ptr, ptr [[KERNEL_ARGS]], i32 2
+// HST-NEXT: store ptr [[DOTADDR2]], ptr [[TMP11]], align 4
+// HST-NEXT: [[TMP12:%.*]] = getelementptr i32, ptr [[KERNEL_ARG_SIZES]], i32 2
+// HST-NEXT: store i32 4, ptr [[TMP12]], align 4
+// HST-NEXT: [[TMP13:%.*]] = getelementptr ptr, ptr [[KERNEL_ARGS]], i32 3
+// HST-NEXT: store ptr [[DOTADDR3]], ptr [[TMP13]], align 4
+// HST-NEXT: [[TMP14:%.*]] = getelementptr i32, ptr [[KERNEL_ARG_SIZES]], i32 3
+// HST-NEXT: store i32 4, ptr [[TMP14]], align 4
+// HST-NEXT: [[TMP15:%.*]] = call i32 @__cudaPopCallConfiguration(ptr [[GRID_DIM]], ptr [[BLOCK_DIM]], ptr [[SHMEM_SIZE]], ptr [[STREAM]])
+// HST-NEXT: [[TMP16:%.*]] = load i32, ptr [[SHMEM_SIZE]], align 4
+// HST-NEXT: [[TMP17:%.*]] = load ptr, ptr [[STREAM]], align 4
+// HST-NEXT: [[CALL:%.*]] = call noundef i32 @__llvmLaunchKernel(ptr noundef @_Z18__device_stub__fooisPvS_, ptr noundef byval([[STRUCT_DIM3]]) align 4 [[GRID_DIM]], ptr noundef byval([[STRUCT_DIM3]]) align 4 [[BLOCK_DIM]], ptr noundef [[KERNEL_LAUNCH_PARAMS]], i32 noundef [[TMP16]], ptr noundef [[TMP17]]) #[[ATTR3:[0-9]+]]
// HST-NEXT: br label %[[SETUP_END:.*]]
// HST: [[SETUP_END]]:
// HST-NEXT: ret void
diff --git a/clang/test/Driver/cuda-via-liboffload.cu b/clang/test/Driver/cuda-via-liboffload.cu
index 68dc963e906b2..d30e529f0ce12 100644
--- a/clang/test/Driver/cuda-via-liboffload.cu
+++ b/clang/test/Driver/cuda-via-liboffload.cu
@@ -2,21 +2,20 @@
// RUN: --offload-arch=sm_35 --offload-arch=sm_70 %s 2>&1 \
// RUN: | FileCheck -check-prefix BINDINGS %s
-// BINDINGS: "x86_64-unknown-linux-gnu" - "clang", inputs: ["[[INPUT:.+]]"], output: "[[HOST_BC:.+]]"
-// BINDINGS-NEXT: "nvptx64-nvidia-cuda" - "clang", inputs: ["[[INPUT]]", "[[HOST_BC]]"], output: "[[PTX_SM_35:.+]]"
-// BINDINGS-NEXT: "nvptx64-nvidia-cuda" - "NVPTX::Assembler", inputs: ["[[PTX_SM_35]]"], output: "[[CUBIN_SM_35:.+]]"
-// BINDINGS-NEXT: "nvptx64-nvidia-cuda" - "clang", inputs: ["[[INPUT]]", "[[HOST_BC]]"], output: "[[PTX_SM_70:.+]]"
-// BINDINGS-NEXT: "nvptx64-nvidia-cuda" - "NVPTX::Assembler", inputs: ["[[PTX_SM_70:.+]]"], output: "[[CUBIN_SM_70:.+]]"
+// BINDINGS: "nvptx64-nvidia-cuda-llvm" - "clang", inputs: ["[[INPUT:.+]]"], output: "[[PTX_SM_35:.+]]"
+// BINDINGS-NEXT: "nvptx64-nvidia-cuda-llvm" - "NVPTX::Assembler", inputs: ["[[PTX_SM_35]]"], output: "[[CUBIN_SM_35:.+]]"
+// BINDINGS-NEXT: "nvptx64-nvidia-cuda-llvm" - "clang", inputs: ["[[INPUT]]"], output: "[[PTX_SM_70:.+]]"
+// BINDINGS-NEXT: "nvptx64-nvidia-cuda-llvm" - "NVPTX::Assembler", inputs: ["[[PTX_SM_70:.+]]"], output: "[[CUBIN_SM_70:.+]]"
// BINDINGS-NEXT: "x86_64-unknown-linux-gnu" - "Offload::Packager", inputs: ["[[CUBIN_SM_35]]", "[[CUBIN_SM_70]]"], output: "[[BINARY:.+]]"
-// BINDINGS-NEXT: "x86_64-unknown-linux-gnu" - "clang", inputs: ["[[HOST_BC]]", "[[BINARY]]"], output: "[[HOST_OBJ:.+]]"
+// BINDINGS-NEXT: "x86_64-unknown-linux-gnu" - "clang", inputs: ["[[INPUT]]", "[[BINARY]]"], output: "[[HOST_OBJ:.+]]"
// BINDINGS-NEXT: "x86_64-unknown-linux-gnu" - "Offload::Linker", inputs: ["[[HOST_OBJ]]"], output: "a.out"
// RUN: %clang -### -target x86_64-linux-gnu -foffload-via-llvm -ccc-print-bindings \
// RUN: --offload-arch=sm_35 --offload-arch=sm_70 %s 2>&1 \
// RUN: | FileCheck -check-prefix BINDINGS-DEVICE %s
-// BINDINGS-DEVICE: # "nvptx64-nvidia-cuda" - "clang", inputs: ["[[INPUT:.+]]"], output: "[[PTX:.+]]"
-// BINDINGS-DEVICE: # "nvptx64-nvidia-cuda" - "NVPTX::Assembler", inputs: ["[[PTX]]"], output: "[[CUBIN:.+]]"
+// BINDINGS-DEVICE: # "nvptx64-nvidia-cuda-llvm" - "clang", inputs: ["[[INPUT:.+]]"], output: "[[PTX:.+]]"
+// BINDINGS-DEVICE: # "nvptx64-nvidia-cuda-llvm" - "NVPTX::Assembler", inputs: ["[[PTX]]"], output: "[[CUBIN:.+]]"
// RUN: %clang -### -target x86_64-linux-gnu -ccc-print-bindings --offload-link -foffload-via-llvm %s 2>&1 | FileCheck -check-prefix DEVICE-LINK %s
diff --git a/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-image.c b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-image.c
index d05e9d54a108a..083d3340f6f81 100644
--- a/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-image.c
+++ b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-image.c
@@ -94,7 +94,7 @@
// CUDA-NEXT: br i1 %1, label %while.entry, label %while.end
//
// CUDA: while.entry:
-// CUDA-NEXT: %entry1 = phi ptr [ @__start_llvm_offload_entries, %entry ], [ %16, %if.end ]
+// CUDA-NEXT: %entry1 = phi ptr [ @__start_llvm_offload_entries, %entry ], [ %17, %if.end ]
// CUDA-NEXT: %2 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i32 0, i32 4
// CUDA-NEXT: %addr = load ptr, ptr %2, align 8
// CUDA-NEXT: %3 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i32 0, i32 8
@@ -117,15 +117,16 @@
// CUDA-NEXT: %constant = lshr i32 %11, 4
// CUDA-NEXT: %12 = and i32 %flags, 32
// CUDA-NEXT: %normalized = lshr i32 %12, 5
-// CUDA-NEXT: %13 = icmp eq i16 %kind, 2
-// CUDA-NEXT: br i1 %13, label %if.kind, label %if.end
+// CUDA-NEXT: %13 = and i16 %kind, 2
+// CUDA-NEXT: %14 = icmp ne i16 %13, 0
+// CUDA-NEXT: br i1 %14, label %if.kind, label %if.end
//
// CUDA: if.kind:
-// CUDA-NEXT: %14 = icmp eq i64 %size, 0
-// CUDA-NEXT: br i1 %14, label %if.then, label %if.else
+// CUDA-NEXT: %15 = icmp eq i64 %size, 0
+// CUDA-NEXT: br i1 %15, label %if.then, label %if.else
//
// CUDA: if.then:
-// CUDA-NEXT: %15 = call i32 @__cudaRegisterFunction(ptr %0, ptr %addr, ptr %name, ptr %name, i32 -1, ptr null, ptr null, ptr null, ptr null, ptr null)
+// CUDA-NEXT: %16 = call i32 @__cudaRegisterFunction(ptr %0, ptr %addr, ptr %name, ptr %name, i32 -1, ptr null, ptr null, ptr null, ptr null, ptr null)
// CUDA-NEXT: br label %if.end
//
// CUDA: if.else:
@@ -151,9 +152,9 @@
// CUDA-NEXT: br label %if.end
//
// CUDA: if.end:
-// CUDA-NEXT: %16 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i64 1
-// CUDA-NEXT: %17 = icmp eq ptr %16, @__stop_llvm_offload_entries
-// CUDA-NEXT: br i1 %17, label %while.end, label %while.entry
+// CUDA-NEXT: %17 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i64 1
+// CUDA-NEXT: %18 = icmp eq ptr %17, @__stop_llvm_offload_entries
+// CUDA-NEXT: br i1 %18, label %while.end, label %while.entry
//
// CUDA: while.end:
// CUDA-NEXT: ret void
@@ -236,7 +237,7 @@
// HIP-NEXT: br i1 %1, label %while.entry, label %while.end
//
// HIP: while.entry:
-// HIP-NEXT: %entry1 = phi ptr [ @{{.*offload_entries.*}}, %entry ], [ %16, %if.end ]
+// HIP-NEXT: %entry1 = phi ptr [ @{{.*offload_entries.*}}, %entry ], [ %17, %if.end ]
// HIP-NEXT: %2 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i32 0, i32 4
// HIP-NEXT: %addr = load ptr, ptr %2, align 8
// HIP-NEXT: %3 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i32 0, i32 8
@@ -259,15 +260,16 @@
// HIP-NEXT: %constant = lshr i32 %11, 4
// HIP-NEXT: %12 = and i32 %flags, 32
// HIP-NEXT: %normalized = lshr i32 %12, 5
-// HIP-NEXT: %13 = icmp eq i16 %kind, 4
-// HIP-NEXT: br i1 %13, label %if.kind, label %if.end
+// HIP-NEXT: %13 = and i16 %kind, 4
+// HIP-NEXT: %14 = icmp ne i16 %13, 0
+// HIP-NEXT: br i1 %14, label %if.kind, label %if.end
//
// HIP: if.kind:
-// HIP-NEXT: %14 = icmp eq i64 %size, 0
-// HIP-NEXT: br i1 %14, label %if.then, label %if.else
+// HIP-NEXT: %15 = icmp eq i64 %size, 0
+// HIP-NEXT: br i1 %15, label %if.then, label %if.else
//
// HIP: if.then:
-// HIP-NEXT: %15 = call i32 @__hipRegisterFunction(ptr %0, ptr %addr, ptr %name, ptr %name, i32 -1, ptr null, ptr null, ptr null, ptr null, ptr null)
+// HIP-NEXT: %16 = call i32 @__hipRegisterFunction(ptr %0, ptr %addr, ptr %name, ptr %name, i32 -1, ptr null, ptr null, ptr null, ptr null, ptr null)
// HIP-NEXT: br label %if.end
//
// HIP: if.else:
@@ -295,9 +297,9 @@
// HIP-NEXT: br label %if.end
//
// HIP: if.end:
-// HIP-NEXT: %16 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i64 1
-// HIP-NEXT: %17 = icmp eq ptr %16, @{{.*offload_entries.*}}
-// HIP-NEXT: br i1 %17, label %while.end, label %while.entry
+// HIP-NEXT: %17 = getelementptr inbounds %struct.__tgt_offload_entry, ptr %entry1, i64 1
+// HIP-NEXT: %18 = icmp eq ptr %17, @{{.*offload_entries.*}}
+// HIP-NEXT: br i1 %18, label %while.end, label %while.entry
//
// HIP: while.end:
// HIP-NEXT: ret void
diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
index c2de6578773c7..b61c4dae85a86 100644
--- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
+++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
@@ -139,6 +139,13 @@ static bool CanonicalPrefixes = true;
using OffloadingImage = OffloadBinary::OffloadingImage;
+static bool usesLLVMOffloadWrapper(ArrayRef<OffloadingImage> Images) {
+ return llvm::any_of(Images, [](const OffloadingImage &Image) {
+ return Triple(Image.StringData.lookup("triple")).getEnvironment() ==
+ Triple::LLVM;
+ });
+}
+
namespace llvm {
// Provide DenseMapInfo so that OffloadKind can be used in a DenseMap.
template <> struct DenseMapInfo<OffloadKind> {
@@ -971,6 +978,9 @@ Expected<SmallVector<std::unique_ptr<MemoryBuffer>>>
bundleLinkedOutput(ArrayRef<OffloadingImage> Images, const ArgList &Args,
OffloadKind Kind) {
llvm::TimeTraceScope TimeScope("Bundle linked output");
+ if (usesLLVMOffloadWrapper(Images))
+ return bundleOpenMP(Images);
+
switch (Kind) {
case OFK_OpenMP:
return (Verbose && SaveTemps) ? bundleOpenMPVerbose(Images)
@@ -1214,7 +1224,8 @@ linkAndWrapDeviceFiles(ArrayRef<SmallVector<OffloadFile>> LinkerInputFiles,
continue;
}
- auto OutputOrErr = wrapDeviceImages(*BundledImagesOrErr, Args, Kind);
+ OffloadKind WrapperKind = usesLLVMOffloadWrapper(Input) ? OFK_OpenMP : Kind;
+ auto OutputOrErr = wrapDeviceImages(*BundledImagesOrErr, Args, WrapperKind);
if (!OutputOrErr)
return OutputOrErr.takeError();
WrappedOutput.push_back(*OutputOrErr);
diff --git a/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp b/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp
index ded603a1e00e3..037b81a7c42fb 100644
--- a/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp
+++ b/llvm/lib/Frontend/Offloading/OffloadWrapper.cpp
@@ -485,10 +485,12 @@ Function *createRegisterGlobalsFunction(Module &M, bool IsHIP,
llvm::offloading::OffloadGlobalNormalized));
auto *Normalized = Builder.CreateLShr(
NormalizedBit, ConstantInt::get(Type::getInt32Ty(C), 5), "normalized");
- auto *KindCond = Builder.CreateICmpEQ(
+ auto *KindAnd = Builder.CreateAnd(
Kind, ConstantInt::get(Type::getInt16Ty(C),
IsHIP ? object::OffloadKind::OFK_HIP
: object::OffloadKind::OFK_Cuda));
+ auto *KindCond =
+ Builder.CreateICmpNE(KindAnd, ConstantInt::get(Type::getInt16Ty(C), 0));
Builder.CreateCondBr(KindCond, IfKindBB, IfEndBB);
Builder.SetInsertPoint(IfKindBB);
auto *FnCond = Builder.CreateICmpEQ(
>From 47c243d4522c15f3f22c221c62c07e1878bb5dc6 Mon Sep 17 00:00:00 2001
From: Sophia Herrmann <herrmann15 at llnl.gov>
Date: Tue, 28 Jul 2026 17:49:31 -0700
Subject: [PATCH 6/9] add unittests
---
offload/test/lit.cfg | 10 +++-
offload/test/offloading/CUDA/basic_launch.cu | 25 +++++----
.../CUDA/basic_launch_blocks_and_threads.cu | 22 ++++----
.../offloading/CUDA/basic_launch_multi_arg.cu | 36 +++++++------
offload/test/offloading/CUDA/device_api.cu | 45 ++++++++++++++++
.../test/offloading/CUDA/device_properties.cu | 40 +++++++++++++++
offload/test/offloading/CUDA/host_alloc.cu | 40 +++++++++++++++
offload/test/offloading/CUDA/launch_tu.cu | 25 +++++----
offload/test/offloading/CUDA/memcpy_kinds.cu | 51 +++++++++++++++++++
offload/test/offloading/CUDA/stream_api.cu | 46 +++++++++++++++++
offload/test/offloading/CUDA/syncthreads.cu | 40 +++++++++++++++
.../offloading/CUDA/thread_and_block_id.cu | 44 ++++++++++++++++
offload/test/offloading/HIP/basic_launch.hip | 30 +++++++++++
.../HIP/basic_launch_blocks_and_threads.hip | 31 +++++++++++
.../offloading/HIP/basic_launch_multi_arg.hip | 39 ++++++++++++++
offload/test/offloading/HIP/device_api.hip | 45 ++++++++++++++++
.../test/offloading/HIP/device_properties.hip | 40 +++++++++++++++
offload/test/offloading/HIP/host_alloc.hip | 40 +++++++++++++++
offload/test/offloading/HIP/kernel_tu.hip.inc | 1 +
offload/test/offloading/HIP/launch_tu.hip | 30 +++++++++++
offload/test/offloading/HIP/memcpy_kinds.hip | 51 +++++++++++++++++++
offload/test/offloading/HIP/stream_api.hip | 46 +++++++++++++++++
offload/test/offloading/HIP/syncthreads.hip | 40 +++++++++++++++
.../offloading/HIP/thread_and_block_id.hip | 44 ++++++++++++++++
24 files changed, 803 insertions(+), 58 deletions(-)
create mode 100644 offload/test/offloading/CUDA/device_api.cu
create mode 100644 offload/test/offloading/CUDA/device_properties.cu
create mode 100644 offload/test/offloading/CUDA/host_alloc.cu
create mode 100644 offload/test/offloading/CUDA/memcpy_kinds.cu
create mode 100644 offload/test/offloading/CUDA/stream_api.cu
create mode 100644 offload/test/offloading/CUDA/syncthreads.cu
create mode 100644 offload/test/offloading/CUDA/thread_and_block_id.cu
create mode 100644 offload/test/offloading/HIP/basic_launch.hip
create mode 100644 offload/test/offloading/HIP/basic_launch_blocks_and_threads.hip
create mode 100644 offload/test/offloading/HIP/basic_launch_multi_arg.hip
create mode 100644 offload/test/offloading/HIP/device_api.hip
create mode 100644 offload/test/offloading/HIP/device_properties.hip
create mode 100644 offload/test/offloading/HIP/host_alloc.hip
create mode 100644 offload/test/offloading/HIP/kernel_tu.hip.inc
create mode 100644 offload/test/offloading/HIP/launch_tu.hip
create mode 100644 offload/test/offloading/HIP/memcpy_kinds.hip
create mode 100644 offload/test/offloading/HIP/stream_api.hip
create mode 100644 offload/test/offloading/HIP/syncthreads.hip
create mode 100644 offload/test/offloading/HIP/thread_and_block_id.hip
diff --git a/offload/test/lit.cfg b/offload/test/lit.cfg
index ace2b1ea8a749..2981cf5d10fea 100644
--- a/offload/test/lit.cfg
+++ b/offload/test/lit.cfg
@@ -83,7 +83,7 @@ def remove_suffix_if_present(name):
config.name = 'libomptarget :: ' + config.libomptarget_current_target
# suffixes: A list of file extensions to treat as test files.
-config.suffixes = ['.c', '.cpp', '.cc', '.f90', '.cu', '.td']
+config.suffixes = ['.c', '.cpp', '.cc', '.f90', '.cu', '.hip', '.td']
# excludes: A list of directories to exclude from the testuites.
config.excludes = ['Inputs', 'unit']
@@ -91,6 +91,11 @@ config.excludes = ['Inputs', 'unit']
# test_source_root: The root path where tests are located.
config.test_source_root = os.path.dirname(__file__)
+# language includes
+config.test_language_includes = os.path.join(config.test_source_root, "../languages/include")
+config.test_language_cuda_includes = os.path.join(config.test_language_includes, "cuda")
+config.test_language_hip_includes = os.path.join(config.test_language_includes, "hip")
+
# test_exec_root: The root object directory where output is placed
config.test_exec_root = config.libomptarget_obj_root
@@ -100,6 +105,9 @@ config.test_format = lit.formats.ShTest()
# compiler flags
config.test_flags = " -I " + config.test_source_root + \
" -I " + config.omp_header_directory + \
+ " -I " + config.test_language_includes + \
+ " -I " + config.test_language_cuda_includes + \
+ " -I " + config.test_language_hip_includes + \
" -L " + config.library_dir + \
" -L " + config.llvm_library_intdir + \
" -L " + config.llvm_lib_directory
diff --git a/offload/test/offloading/CUDA/basic_launch.cu b/offload/test/offloading/CUDA/basic_launch.cu
index e017241bb9a74..5ecfc3e9d5601 100644
--- a/offload/test/offloading/CUDA/basic_launch.cu
+++ b/offload/test/offloading/CUDA/basic_launch.cu
@@ -7,25 +7,24 @@
// UNSUPPORTED: aarch64-unknown-linux-gnu
// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
// UNSUPPORTED: intelgpu
#include <stdio.h>
-extern "C" {
-void *llvm_omp_target_alloc_shared(size_t Size, int DeviceNum);
-void llvm_omp_target_free_shared(void *DevicePtr, int DeviceNum);
-}
-
__global__ void square(int *A) { *A = 42; }
int main(int argc, char **argv) {
- int DevNo = 0;
- int *Ptr = reinterpret_cast<int *>(llvm_omp_target_alloc_shared(4, DevNo));
- *Ptr = 7;
- printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
- // CHECK: Ptr [[Ptr:0x.*]], *Ptr: 7
+ int *Ptr;
+ cudaMalloc(&Ptr, 4);
+ printf("Ptr %p\n", Ptr);
+ // CHECK: Ptr [[Ptr:0x.*]]
square<<<1, 1>>>(Ptr);
- printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
- // CHECK: Ptr [[Ptr]], *Ptr: 42
- llvm_omp_target_free_shared(Ptr, DevNo);
+ int I = 0;
+ cudaDeviceSynchronize();
+ cudaMemcpy(&I, Ptr, sizeof(int), cudaMemcpyDeviceToHost);
+ printf("I: %i\n", I);
+ // CHECK: I: 42
}
diff --git a/offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu b/offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu
index a428e25d82359..0bfc9e231ffde 100644
--- a/offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu
+++ b/offload/test/offloading/CUDA/basic_launch_blocks_and_threads.cu
@@ -7,27 +7,25 @@
// UNSUPPORTED: aarch64-unknown-linux-gnu
// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
// UNSUPPORTED: intelgpu
#include <stdio.h>
-extern "C" {
-void *llvm_omp_target_alloc_shared(size_t Size, int DeviceNum);
-void llvm_omp_target_free_shared(void *DevicePtr, int DeviceNum);
-}
-
__global__ void square(int *A) {
__scoped_atomic_fetch_add(A, 1, __ATOMIC_SEQ_CST, __MEMORY_SCOPE_DEVICE);
}
int main(int argc, char **argv) {
int DevNo = 0;
- int *Ptr = reinterpret_cast<int *>(llvm_omp_target_alloc_shared(4, DevNo));
- *Ptr = 0;
- printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
- // CHECK: Ptr [[Ptr:0x.*]], *Ptr: 0
+ int *Ptr, I;
+ cudaMalloc(&Ptr, 4);
+ printf("Ptr %p\n", Ptr);
+ // CHECK: Ptr [[Ptr:0x.*]]
square<<<7, 6>>>(Ptr);
- printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
- // CHECK: Ptr [[Ptr]], *Ptr: 42
- llvm_omp_target_free_shared(Ptr, DevNo);
+ cudaMemcpy(&I, Ptr, sizeof(int), cudaMemcpyDeviceToHost);
+ printf("I: %i\n", I);
+ // CHECK: I: 42
}
diff --git a/offload/test/offloading/CUDA/basic_launch_multi_arg.cu b/offload/test/offloading/CUDA/basic_launch_multi_arg.cu
index db2a1e48371b0..7c0c755e71057 100644
--- a/offload/test/offloading/CUDA/basic_launch_multi_arg.cu
+++ b/offload/test/offloading/CUDA/basic_launch_multi_arg.cu
@@ -6,15 +6,13 @@
// clang-format on
// REQUIRES: gpu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
// UNSUPPORTED: intelgpu
#include <stdio.h>
-extern "C" {
-void *llvm_omp_target_alloc_shared(size_t Size, int DeviceNum);
-void llvm_omp_target_free_shared(void *DevicePtr, int DeviceNum);
-}
-
__global__ void square(int *Dst, short Q, int *Src, short P) {
*Dst = (Src[0] + Src[1]) * (Q + P);
Src[0] = Q;
@@ -23,19 +21,19 @@ __global__ void square(int *Dst, short Q, int *Src, short P) {
int main(int argc, char **argv) {
int DevNo = 0;
- int *Ptr = reinterpret_cast<int *>(llvm_omp_target_alloc_shared(4, DevNo));
- int *Src = reinterpret_cast<int *>(llvm_omp_target_alloc_shared(8, DevNo));
- *Ptr = 7;
- Src[0] = -2;
- Src[1] = 8;
- printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
- // CHECK: Ptr [[Ptr:0x.*]], *Ptr: 7
- printf("Src: %i : %i\n", Src[0], Src[1]);
- // CHECK: Src: -2 : 8
+ int *Src, *Ptr;
+ cudaMalloc(&Ptr, 4);
+ cudaMalloc(&Src, 8);
+
+ int I = 7;
+ int HostSrc[2] = {-2,8};
+ cudaMemcpy(Ptr, &I, sizeof(int), cudaMemcpyHostToDevice);
+ cudaMemcpy(Src, &HostSrc[0], 2*sizeof(int), cudaMemcpyHostToDevice);
square<<<1, 1>>>(Ptr, 3, Src, 4);
- printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
- // CHECK: Ptr [[Ptr]], *Ptr: 42
- printf("Src: %i : %i\n", Src[0], Src[1]);
- // CHECK: Src: 3 : 4
- llvm_omp_target_free_shared(Ptr, DevNo);
+ cudaMemcpy(&I, Ptr, sizeof(int), cudaMemcpyDeviceToHost);
+ cudaMemcpy(&HostSrc[0], Src, 2 * sizeof(int), cudaMemcpyDeviceToHost);
+ printf("I: %i\n", I);
+ // CHECK: I: 42
+ printf("Src: %i, %i\n", HostSrc[0], HostSrc[1]);
+ // CHECK: Src: 3, 4
}
diff --git a/offload/test/offloading/CUDA/device_api.cu b/offload/test/offloading/CUDA/device_api.cu
new file mode 100644
index 0000000000000..184167f2d17e4
--- /dev/null
+++ b/offload/test/offloading/CUDA/device_api.cu
@@ -0,0 +1,45 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+int main(int argc, char **argv) {
+ int Count = 0;
+ if (cudaGetDeviceCount(&Count) != cudaSuccess)
+ return 1;
+
+ printf("device count: %d\n", Count);
+ // CHECK: device count: {{[1-9][0-9]*}}
+
+ int Device = -1;
+ if (cudaGetDevice(&Device) != cudaSuccess)
+ return 1;
+
+ printf("device: %d\n", Device);
+ // CHECK: device: {{[0-9]+}}
+
+ if (cudaSetDevice(Device) != cudaSuccess)
+ return 1;
+
+ int After = -1;
+ if (cudaGetDevice(&After) != cudaSuccess)
+ return 1;
+
+ printf("device after set: %d\n", After);
+ // CHECK: device after set: {{[0-9]+}}
+
+ cudaError_t Err = cudaSetDevice(-1);
+ printf("set invalid device: %u\n", Err);
+ // CHECK: set invalid device: 1
+}
diff --git a/offload/test/offloading/CUDA/device_properties.cu b/offload/test/offloading/CUDA/device_properties.cu
new file mode 100644
index 0000000000000..8f625f6ccabe1
--- /dev/null
+++ b/offload/test/offloading/CUDA/device_properties.cu
@@ -0,0 +1,40 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+int main(int argc, char **argv) {
+ cudaDeviceProp Prop = {};
+ cudaError_t Err = cudaGetDeviceProperties(&Prop, 0);
+ if (Err != cudaSuccess) {
+ printf("cudaGetDeviceProperties failed: %u\n", Err);
+ return 1;
+ }
+
+ printf("Device name: %s\n", Prop.name);
+ // CHECK: Device name:
+ printf("Total global memory: %zu\n", Prop.totalGlobalMem);
+ // CHECK: Total global memory:
+ printf("Multiprocessors: %i\n", Prop.multiProcessorCount);
+ // CHECK: Multiprocessors:
+ printf("Warp size: %i\n", Prop.warpSize);
+ // CHECK: Warp size:
+
+ if (!Prop.name[0] || !Prop.totalGlobalMem || !Prop.multiProcessorCount ||
+ !Prop.warpSize)
+ return 1;
+
+ printf("Device properties are populated.\n");
+ // CHECK: Device properties are populated.
+}
diff --git a/offload/test/offloading/CUDA/host_alloc.cu b/offload/test/offloading/CUDA/host_alloc.cu
new file mode 100644
index 0000000000000..f23eac7604582
--- /dev/null
+++ b/offload/test/offloading/CUDA/host_alloc.cu
@@ -0,0 +1,40 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+int main(int argc, char **argv) {
+ int *HostAllocPtr = nullptr;
+ if (cudaHostAlloc(&HostAllocPtr, sizeof(int), cudaHostAllocDefault) !=
+ cudaSuccess)
+ return 1;
+
+ *HostAllocPtr = 17;
+ printf("cudaHostAlloc value: %d\n", *HostAllocPtr);
+ // CHECK: cudaHostAlloc value: 17
+
+ if (cudaFreeHost(HostAllocPtr) != cudaSuccess)
+ return 1;
+
+ int *MallocHostPtr = nullptr;
+ if (cudaMallocHost(&MallocHostPtr, sizeof(int)) != cudaSuccess)
+ return 1;
+
+ *MallocHostPtr = 23;
+ printf("cudaMallocHost value: %d\n", *MallocHostPtr);
+ // CHECK: cudaMallocHost value: 23
+
+ if (cudaFreeHost(MallocHostPtr) != cudaSuccess)
+ return 1;
+}
diff --git a/offload/test/offloading/CUDA/launch_tu.cu b/offload/test/offloading/CUDA/launch_tu.cu
index a46472b514a6c..8b92194ba435e 100644
--- a/offload/test/offloading/CUDA/launch_tu.cu
+++ b/offload/test/offloading/CUDA/launch_tu.cu
@@ -1,31 +1,30 @@
// clang-format off
// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t.launch_tu.o -c
// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x cuda %S/kernel_tu.cu.inc -o %t.kernel_tu.o -c
-// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %t.launch_tu.o %t.kernel_tu.o -o %t
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native --offload-link %t.launch_tu.o %t.kernel_tu.o -o %t
// RUN: %t | %fcheck-generic
// clang-format on
// UNSUPPORTED: aarch64-unknown-linux-gnu
// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
// UNSUPPORTED: intelgpu
#include <stdio.h>
-extern "C" {
-void *llvm_omp_target_alloc_shared(size_t Size, int DeviceNum);
-void llvm_omp_target_free_shared(void *DevicePtr, int DeviceNum);
-}
-
extern __global__ void square(int *A);
int main(int argc, char **argv) {
int DevNo = 0;
- int *Ptr = reinterpret_cast<int *>(llvm_omp_target_alloc_shared(4, DevNo));
- *Ptr = 7;
- printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
- // CHECK: Ptr [[Ptr:0x.*]], *Ptr: 7
+ int *Ptr;
+ cudaMalloc(&Ptr, 4);
+ printf("Ptr %p\n", Ptr);
+ // CHECK: Ptr [[Ptr:0x.*]]
square<<<1, 1>>>(Ptr);
- printf("Ptr %p, *Ptr: %i\n", Ptr, *Ptr);
- // CHECK: Ptr [[Ptr]], *Ptr: 42
- llvm_omp_target_free_shared(Ptr, DevNo);
+ int I;
+ cudaMemcpy(&I, Ptr, sizeof(int), cudaMemcpyDeviceToHost);
+ printf("I: %i\n", I);
+ // CHECK: I: 42
}
diff --git a/offload/test/offloading/CUDA/memcpy_kinds.cu b/offload/test/offloading/CUDA/memcpy_kinds.cu
new file mode 100644
index 0000000000000..a4288ee51ee3b
--- /dev/null
+++ b/offload/test/offloading/CUDA/memcpy_kinds.cu
@@ -0,0 +1,51 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+int main(int argc, char **argv) {
+ int HostSrc = 11;
+ int HostDst = 0;
+ if (cudaMemcpy(&HostDst, &HostSrc, sizeof(int), cudaMemcpyHostToHost) !=
+ cudaSuccess)
+ return 1;
+
+ printf("host to host: %d\n", HostDst);
+ // CHECK: host to host: 11
+
+ int *DevSrc = nullptr;
+ int *DevDst = nullptr;
+ int Result = 0;
+ if (cudaMalloc(&DevSrc, sizeof(int)) != cudaSuccess)
+ return 1;
+ if (cudaMalloc(&DevDst, sizeof(int)) != cudaSuccess)
+ return 1;
+
+ HostSrc = 42;
+ if (cudaMemcpy(DevSrc, &HostSrc, sizeof(int), cudaMemcpyHostToDevice) !=
+ cudaSuccess)
+ return 1;
+ if (cudaMemcpy(DevDst, DevSrc, sizeof(int), cudaMemcpyDeviceToDevice) !=
+ cudaSuccess)
+ return 1;
+ if (cudaMemcpy(&Result, DevDst, sizeof(int), cudaMemcpyDeviceToHost) !=
+ cudaSuccess)
+ return 1;
+
+ printf("device to device: %d\n", Result);
+ // CHECK: device to device: 42
+
+ cudaFree(DevSrc);
+ cudaFree(DevDst);
+}
diff --git a/offload/test/offloading/CUDA/stream_api.cu b/offload/test/offloading/CUDA/stream_api.cu
new file mode 100644
index 0000000000000..7202751f8207e
--- /dev/null
+++ b/offload/test/offloading/CUDA/stream_api.cu
@@ -0,0 +1,46 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+__global__ void setValue(int *Out) { *Out = 42; }
+
+int main(int argc, char **argv) {
+ cudaStream_t Stream = nullptr;
+ if (cudaStreamCreate(&Stream) != cudaSuccess)
+ return 1;
+
+ printf("stream created: %d\n", Stream != nullptr);
+ // CHECK: stream created: 1
+
+ int *DevPtr = nullptr;
+ int Result = 0;
+ if (cudaMalloc(&DevPtr, sizeof(int)) != cudaSuccess)
+ return 1;
+
+ setValue<<<1, 1, 0, Stream>>>(DevPtr);
+
+ if (cudaStreamSynchronize(Stream) != cudaSuccess)
+ return 1;
+ if (cudaMemcpy(&Result, DevPtr, sizeof(int), cudaMemcpyDeviceToHost) !=
+ cudaSuccess)
+ return 1;
+
+ printf("stream result: %d\n", Result);
+ // CHECK: stream result: 42
+
+ if (cudaStreamDestroy(Stream) != cudaSuccess)
+ return 1;
+ cudaFree(DevPtr);
+}
diff --git a/offload/test/offloading/CUDA/syncthreads.cu b/offload/test/offloading/CUDA/syncthreads.cu
new file mode 100644
index 0000000000000..0c6048c32f824
--- /dev/null
+++ b/offload/test/offloading/CUDA/syncthreads.cu
@@ -0,0 +1,40 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+__global__ void reduceBlock(int *Out) {
+ __shared__ int Scratch[64];
+ int Tid = threadIdx.x;
+ Scratch[Tid] = Tid;
+ __syncthreads();
+
+ if (Tid == 0) {
+ int Sum = 0;
+ for (int I = 0; I < 64; ++I)
+ Sum += Scratch[I];
+ Out[0] = Sum;
+ }
+}
+
+int main(int argc, char **argv) {
+ int *DevPtr;
+ int Result = 0;
+ cudaMalloc(&DevPtr, sizeof(int));
+ reduceBlock<<<1, 64>>>(DevPtr);
+ cudaMemcpy(&Result, DevPtr, sizeof(int), cudaMemcpyDeviceToHost);
+
+ printf("sum: %i\n", Result);
+ // CHECK: sum: 2016
+}
diff --git a/offload/test/offloading/CUDA/thread_and_block_id.cu b/offload/test/offloading/CUDA/thread_and_block_id.cu
new file mode 100644
index 0000000000000..074c7e6557d06
--- /dev/null
+++ b/offload/test/offloading/CUDA/thread_and_block_id.cu
@@ -0,0 +1,44 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: aarch64-unknown-linux-gnu-LTO
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu-LTO
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+
+#include <stdio.h>
+#include <stdlib.h>
+
+__global__ void fill(int *A) {
+ int tid = threadIdx.x + blockDim.x * blockIdx.x;
+ A[tid] = 42;
+}
+
+int main(int argc, char **argv) {
+ int NThreads = 128;
+ int NBlocks = 512;
+ int Size = sizeof(int) * NThreads * NBlocks;
+ int *Ptr = (int*)calloc(1, Size);
+ int *DevPtr;
+ cudaMalloc(&DevPtr, Size);
+ cudaMemcpy(DevPtr, Ptr, Size, cudaMemcpyHostToDevice);
+ printf("DevPtr %p\n", DevPtr);
+ // CHECK: DevPtr [[DevPtr:0x.*]]
+ fill<<<NBlocks, NThreads>>>(DevPtr);
+ cudaMemcpy(Ptr, DevPtr, Size, cudaMemcpyDeviceToHost);
+
+ for (int I = 0; I < NBlocks * NThreads; ++I) {
+ if (Ptr[I] == 42)
+ continue;
+ printf("Error at %i: %i vs %i\n", I, Ptr[I], 42);
+ return 1;
+ }
+ return 0;
+}
diff --git a/offload/test/offloading/HIP/basic_launch.hip b/offload/test/offloading/HIP/basic_launch.hip
new file mode 100644
index 0000000000000..bd2f2a6078671
--- /dev/null
+++ b/offload/test/offloading/HIP/basic_launch.hip
@@ -0,0 +1,30 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+__global__ void square(int *A) { *A = 42; }
+
+int main(int argc, char **argv) {
+ int *Ptr;
+ hipMalloc(&Ptr, 4);
+ printf("Ptr %p\n", Ptr);
+ // CHECK: Ptr [[Ptr:0x.*]]
+ square<<<1, 1>>>(Ptr);
+ int I = 0;
+ hipDeviceSynchronize();
+ hipMemcpy(&I, Ptr, sizeof(int), hipMemcpyDeviceToHost);
+ printf("I: %i\n", I);
+ // CHECK: I: 42
+}
diff --git a/offload/test/offloading/HIP/basic_launch_blocks_and_threads.hip b/offload/test/offloading/HIP/basic_launch_blocks_and_threads.hip
new file mode 100644
index 0000000000000..344b98b1636f1
--- /dev/null
+++ b/offload/test/offloading/HIP/basic_launch_blocks_and_threads.hip
@@ -0,0 +1,31 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+__global__ void square(int *A) {
+ __scoped_atomic_fetch_add(A, 1, __ATOMIC_SEQ_CST, __MEMORY_SCOPE_DEVICE);
+}
+
+int main(int argc, char **argv) {
+ int DevNo = 0;
+ int *Ptr, I;
+ hipMalloc(&Ptr, 4);
+ printf("Ptr %p\n", Ptr);
+ // CHECK: Ptr [[Ptr:0x.*]]
+ square<<<7, 6>>>(Ptr);
+ hipMemcpy(&I, Ptr, sizeof(int), hipMemcpyDeviceToHost);
+ printf("I: %i\n", I);
+ // CHECK: I: 42
+}
diff --git a/offload/test/offloading/HIP/basic_launch_multi_arg.hip b/offload/test/offloading/HIP/basic_launch_multi_arg.hip
new file mode 100644
index 0000000000000..6e599d6704598
--- /dev/null
+++ b/offload/test/offloading/HIP/basic_launch_multi_arg.hip
@@ -0,0 +1,39 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// REQUIRES: gpu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+__global__ void square(int *Dst, short Q, int *Src, short P) {
+ *Dst = (Src[0] + Src[1]) * (Q + P);
+ Src[0] = Q;
+ Src[1] = P;
+}
+
+int main(int argc, char **argv) {
+ int DevNo = 0;
+ int *Src, *Ptr;
+ hipMalloc(&Ptr, 4);
+ hipMalloc(&Src, 8);
+
+ int I = 7;
+ int HostSrc[2] = {-2,8};
+ hipMemcpy(Ptr, &I, sizeof(int), hipMemcpyHostToDevice);
+ hipMemcpy(Src, &HostSrc[0], 2*sizeof(int), hipMemcpyHostToDevice);
+ square<<<1, 1>>>(Ptr, 3, Src, 4);
+ hipMemcpy(&I, Ptr, sizeof(int), hipMemcpyDeviceToHost);
+ hipMemcpy(&HostSrc[0], Src, 2 * sizeof(int), hipMemcpyDeviceToHost);
+ printf("I: %i\n", I);
+ // CHECK: I: 42
+ printf("Src: %i, %i\n", HostSrc[0], HostSrc[1]);
+ // CHECK: Src: 3, 4
+}
diff --git a/offload/test/offloading/HIP/device_api.hip b/offload/test/offloading/HIP/device_api.hip
new file mode 100644
index 0000000000000..031e3703e66c1
--- /dev/null
+++ b/offload/test/offloading/HIP/device_api.hip
@@ -0,0 +1,45 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+int main(int argc, char **argv) {
+ int Count = 0;
+ if (hipGetDeviceCount(&Count) != hipSuccess)
+ return 1;
+
+ printf("device count: %d\n", Count);
+ // CHECK: device count: {{[1-9][0-9]*}}
+
+ int Device = -1;
+ if (hipGetDevice(&Device) != hipSuccess)
+ return 1;
+
+ printf("device: %d\n", Device);
+ // CHECK: device: {{[0-9]+}}
+
+ if (hipSetDevice(Device) != hipSuccess)
+ return 1;
+
+ int After = -1;
+ if (hipGetDevice(&After) != hipSuccess)
+ return 1;
+
+ printf("device after set: %d\n", After);
+ // CHECK: device after set: {{[0-9]+}}
+
+ hipError_t Err = hipSetDevice(-1);
+ printf("set invalid device: %u\n", Err);
+ // CHECK: set invalid device: 1
+}
diff --git a/offload/test/offloading/HIP/device_properties.hip b/offload/test/offloading/HIP/device_properties.hip
new file mode 100644
index 0000000000000..1a9b9a70f8ea9
--- /dev/null
+++ b/offload/test/offloading/HIP/device_properties.hip
@@ -0,0 +1,40 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+int main(int argc, char **argv) {
+ hipDeviceProp_t Prop = {};
+ hipError_t Err = hipGetDeviceProperties(&Prop, 0);
+ if (Err != hipSuccess) {
+ printf("hipGetDeviceProperties failed: %u\n", Err);
+ return 1;
+ }
+
+ printf("Device name: %s\n", Prop.name);
+ // CHECK: Device name:
+ printf("Total global memory: %zu\n", Prop.totalGlobalMem);
+ // CHECK: Total global memory:
+ printf("Multiprocessors: %i\n", Prop.multiProcessorCount);
+ // CHECK: Multiprocessors:
+ printf("Warp size: %i\n", Prop.warpSize);
+ // CHECK: Warp size:
+
+ if (!Prop.name[0] || !Prop.totalGlobalMem || !Prop.multiProcessorCount ||
+ !Prop.warpSize)
+ return 1;
+
+ printf("Device properties are populated.\n");
+ // CHECK: Device properties are populated.
+}
diff --git a/offload/test/offloading/HIP/host_alloc.hip b/offload/test/offloading/HIP/host_alloc.hip
new file mode 100644
index 0000000000000..8b067b39f2f81
--- /dev/null
+++ b/offload/test/offloading/HIP/host_alloc.hip
@@ -0,0 +1,40 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+int main(int argc, char **argv) {
+ int *HostAllocPtr = nullptr;
+ if (hipHostAlloc(&HostAllocPtr, sizeof(int), hipHostAllocDefault) !=
+ hipSuccess)
+ return 1;
+
+ *HostAllocPtr = 17;
+ printf("hipHostAlloc value: %d\n", *HostAllocPtr);
+ // CHECK: hipHostAlloc value: 17
+
+ if (hipFreeHost(HostAllocPtr) != hipSuccess)
+ return 1;
+
+ int *MallocHostPtr = nullptr;
+ if (hipMallocHost(&MallocHostPtr, sizeof(int)) != hipSuccess)
+ return 1;
+
+ *MallocHostPtr = 23;
+ printf("hipMallocHost value: %d\n", *MallocHostPtr);
+ // CHECK: hipMallocHost value: 23
+
+ if (hipFreeHost(MallocHostPtr) != hipSuccess)
+ return 1;
+}
diff --git a/offload/test/offloading/HIP/kernel_tu.hip.inc b/offload/test/offloading/HIP/kernel_tu.hip.inc
new file mode 100644
index 0000000000000..d7d28a109dfc5
--- /dev/null
+++ b/offload/test/offloading/HIP/kernel_tu.hip.inc
@@ -0,0 +1 @@
+__global__ void square(int *A) { *A = 42; }
diff --git a/offload/test/offloading/HIP/launch_tu.hip b/offload/test/offloading/HIP/launch_tu.hip
new file mode 100644
index 0000000000000..03073029ca211
--- /dev/null
+++ b/offload/test/offloading/HIP/launch_tu.hip
@@ -0,0 +1,30 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t.launch_tu.o -c
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native -x hip %S/kernel_tu.hip.inc -o %t.kernel_tu.o -c
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native --offload-link %t.launch_tu.o %t.kernel_tu.o -o %t
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+extern __global__ void square(int *A);
+
+int main(int argc, char **argv) {
+ int DevNo = 0;
+ int *Ptr;
+ hipMalloc(&Ptr, 4);
+ printf("Ptr %p\n", Ptr);
+ // CHECK: Ptr [[Ptr:0x.*]]
+ square<<<1, 1>>>(Ptr);
+ int I;
+ hipMemcpy(&I, Ptr, sizeof(int), hipMemcpyDeviceToHost);
+ printf("I: %i\n", I);
+ // CHECK: I: 42
+}
diff --git a/offload/test/offloading/HIP/memcpy_kinds.hip b/offload/test/offloading/HIP/memcpy_kinds.hip
new file mode 100644
index 0000000000000..6755a55aa0794
--- /dev/null
+++ b/offload/test/offloading/HIP/memcpy_kinds.hip
@@ -0,0 +1,51 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+int main(int argc, char **argv) {
+ int HostSrc = 11;
+ int HostDst = 0;
+ if (hipMemcpy(&HostDst, &HostSrc, sizeof(int), hipMemcpyHostToHost) !=
+ hipSuccess)
+ return 1;
+
+ printf("host to host: %d\n", HostDst);
+ // CHECK: host to host: 11
+
+ int *DevSrc = nullptr;
+ int *DevDst = nullptr;
+ int Result = 0;
+ if (hipMalloc(&DevSrc, sizeof(int)) != hipSuccess)
+ return 1;
+ if (hipMalloc(&DevDst, sizeof(int)) != hipSuccess)
+ return 1;
+
+ HostSrc = 42;
+ if (hipMemcpy(DevSrc, &HostSrc, sizeof(int), hipMemcpyHostToDevice) !=
+ hipSuccess)
+ return 1;
+ if (hipMemcpy(DevDst, DevSrc, sizeof(int), hipMemcpyDeviceToDevice) !=
+ hipSuccess)
+ return 1;
+ if (hipMemcpy(&Result, DevDst, sizeof(int), hipMemcpyDeviceToHost) !=
+ hipSuccess)
+ return 1;
+
+ printf("device to device: %d\n", Result);
+ // CHECK: device to device: 42
+
+ hipFree(DevSrc);
+ hipFree(DevDst);
+}
diff --git a/offload/test/offloading/HIP/stream_api.hip b/offload/test/offloading/HIP/stream_api.hip
new file mode 100644
index 0000000000000..c0e2699822814
--- /dev/null
+++ b/offload/test/offloading/HIP/stream_api.hip
@@ -0,0 +1,46 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+__global__ void setValue(int *Out) { *Out = 42; }
+
+int main(int argc, char **argv) {
+ hipStream_t Stream = nullptr;
+ if (hipStreamCreate(&Stream) != hipSuccess)
+ return 1;
+
+ printf("stream created: %d\n", Stream != nullptr);
+ // CHECK: stream created: 1
+
+ int *DevPtr = nullptr;
+ int Result = 0;
+ if (hipMalloc(&DevPtr, sizeof(int)) != hipSuccess)
+ return 1;
+
+ setValue<<<1, 1, 0, Stream>>>(DevPtr);
+
+ if (hipStreamSynchronize(Stream) != hipSuccess)
+ return 1;
+ if (hipMemcpy(&Result, DevPtr, sizeof(int), hipMemcpyDeviceToHost) !=
+ hipSuccess)
+ return 1;
+
+ printf("stream result: %d\n", Result);
+ // CHECK: stream result: 42
+
+ if (hipStreamDestroy(Stream) != hipSuccess)
+ return 1;
+ hipFree(DevPtr);
+}
diff --git a/offload/test/offloading/HIP/syncthreads.hip b/offload/test/offloading/HIP/syncthreads.hip
new file mode 100644
index 0000000000000..5962ab5468b86
--- /dev/null
+++ b/offload/test/offloading/HIP/syncthreads.hip
@@ -0,0 +1,40 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+__global__ void reduceBlock(int *Out) {
+ __shared__ int Scratch[64];
+ int Tid = threadIdx.x;
+ Scratch[Tid] = Tid;
+ __syncthreads();
+
+ if (Tid == 0) {
+ int Sum = 0;
+ for (int I = 0; I < 64; ++I)
+ Sum += Scratch[I];
+ Out[0] = Sum;
+ }
+}
+
+int main(int argc, char **argv) {
+ int *DevPtr;
+ int Result = 0;
+ hipMalloc(&DevPtr, sizeof(int));
+ reduceBlock<<<1, 64>>>(DevPtr);
+ hipMemcpy(&Result, DevPtr, sizeof(int), hipMemcpyDeviceToHost);
+
+ printf("sum: %i\n", Result);
+ // CHECK: sum: 2016
+}
diff --git a/offload/test/offloading/HIP/thread_and_block_id.hip b/offload/test/offloading/HIP/thread_and_block_id.hip
new file mode 100644
index 0000000000000..c9c33c55d72fb
--- /dev/null
+++ b/offload/test/offloading/HIP/thread_and_block_id.hip
@@ -0,0 +1,44 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: aarch64-unknown-linux-gnu-LTO
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu-LTO
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+
+#include <stdio.h>
+#include <stdlib.h>
+
+__global__ void fill(int *A) {
+ int tid = threadIdx.x + blockDim.x * blockIdx.x;
+ A[tid] = 42;
+}
+
+int main(int argc, char **argv) {
+ int NThreads = 128;
+ int NBlocks = 512;
+ int Size = sizeof(int) * NThreads * NBlocks;
+ int *Ptr = (int*)calloc(1, Size);
+ int *DevPtr;
+ hipMalloc(&DevPtr, Size);
+ hipMemcpy(DevPtr, Ptr, Size, hipMemcpyHostToDevice);
+ printf("DevPtr %p\n", DevPtr);
+ // CHECK: DevPtr [[DevPtr:0x.*]]
+ fill<<<NBlocks, NThreads>>>(DevPtr);
+ hipMemcpy(Ptr, DevPtr, Size, hipMemcpyDeviceToHost);
+
+ for (int I = 0; I < NBlocks * NThreads; ++I) {
+ if (Ptr[I] == 42)
+ continue;
+ printf("Error at %i: %i vs %i\n", I, Ptr[I], 42);
+ return 1;
+ }
+ return 0;
+}
>From cb81231d8dc70be2f1e120aea52e9559b9be2faf Mon Sep 17 00:00:00 2001
From: Sophia Herrmann <herrmann15 at llnl.gov>
Date: Wed, 29 Jul 2026 15:12:50 -0700
Subject: [PATCH 7/9] add getErrorName + getErrorString + Test
---
.../include/kernel/DefineLanguageNames.inc | 4 ++
.../include/kernel/LanguageRuntime.h | 6 ++
.../include/kernel/UndefineLanguageNames.inc | 4 ++
.../languages/kernel/include/LanguageUtils.h | 21 ++++++
.../languages/kernel/src/LanguageRuntime.cpp | 33 +++-------
.../languages/kernel/src/LanguageUtils.cpp | 65 +++++++++++++++++++
offload/test/offloading/CUDA/device_api.cu | 2 +-
offload/test/offloading/CUDA/error_kinds.cu | 61 +++++++++++++++++
offload/test/offloading/HIP/device_api.hip | 2 +-
offload/test/offloading/HIP/error_kinds.hip | 61 +++++++++++++++++
10 files changed, 233 insertions(+), 26 deletions(-)
create mode 100644 offload/languages/kernel/include/LanguageUtils.h
create mode 100644 offload/languages/kernel/src/LanguageUtils.cpp
create mode 100644 offload/test/offloading/CUDA/error_kinds.cu
create mode 100644 offload/test/offloading/HIP/error_kinds.hip
diff --git a/offload/languages/include/kernel/DefineLanguageNames.inc b/offload/languages/include/kernel/DefineLanguageNames.inc
index a07ab0ad7e522..04b5bfac65d93 100644
--- a/offload/languages/include/kernel/DefineLanguageNames.inc
+++ b/offload/languages/include/kernel/DefineLanguageNames.inc
@@ -19,6 +19,10 @@
#define DeviceSynchronize COMBINE(LANGUAGE, DeviceSynchronize)
#define Success COMBINE(LANGUAGE, Success)
#define ErrorInvalidValue COMBINE(LANGUAGE, ErrorInvalidValue)
+#define ErrorInvalidDevice COMBINE(LANGUAGE, ErrorInvalidDevice)
+#define ErrorUnknown COMBINE(LANGUAGE, ErrorUnknown)
+#define GetErrorName COMBINE(LANGUAGE, GetErrorName)
+#define GetErrorString COMBINE(LANGUAGE, GetErrorString)
#define MemcpyKind COMBINE(LANGUAGE, MemcpyKind)
#define MemcpyHostToHost COMBINE(LANGUAGE, MemcpyHostToHost)
#define MemcpyHostToDevice COMBINE(LANGUAGE, MemcpyHostToDevice)
diff --git a/offload/languages/include/kernel/LanguageRuntime.h b/offload/languages/include/kernel/LanguageRuntime.h
index 8b114e0d3d054..80c78d7dd6d61 100644
--- a/offload/languages/include/kernel/LanguageRuntime.h
+++ b/offload/languages/include/kernel/LanguageRuntime.h
@@ -17,8 +17,14 @@
enum Error_t : uint32_t {
Success = 0,
ErrorInvalidValue = 1,
+ ErrorInvalidDevice = 2,
+ ErrorUnknown = 3,
};
+const char *GetErrorName(Error_t Error);
+
+const char *GetErrorString(Error_t Error);
+
struct DeviceProp_t {
char name[256];
size_t totalGlobalMem;
diff --git a/offload/languages/include/kernel/UndefineLanguageNames.inc b/offload/languages/include/kernel/UndefineLanguageNames.inc
index 4e6b433bbd8b5..c3c2b7e35b17a 100644
--- a/offload/languages/include/kernel/UndefineLanguageNames.inc
+++ b/offload/languages/include/kernel/UndefineLanguageNames.inc
@@ -16,6 +16,10 @@
#undef DeviceSynchronize
#undef Success
#undef ErrorInvalidValue
+#undef ErrorInvalidDevice
+#undef ErrorUnknown
+#undef GetErrorName
+#undef GetErrorString
#undef MemcpyKind
#undef MemcpyHostToHost
#undef MemcpyHostToDevice
diff --git a/offload/languages/kernel/include/LanguageUtils.h b/offload/languages/kernel/include/LanguageUtils.h
new file mode 100644
index 0000000000000..e064fdb14fea4
--- /dev/null
+++ b/offload/languages/kernel/include/LanguageUtils.h
@@ -0,0 +1,21 @@
+//===-- LanguageUtils.h - Kernel Language utility functions ---------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#ifndef LANGUAGE
+#error This file should be included, or used, with a LANGUAGE macro set.
+#endif
+
+#include "OffloadAPI.h"
+
+/// Convert an ol_result_t to the active language's Error_t.
+static Error_t convertResult(ol_result_t Result);
+
+/// Convert a Stream_t to an ol_queue_handle_t.
+static Error_t getQueueFromStream(Stream_t Stream, ol_queue_handle_t *Queue);
diff --git a/offload/languages/kernel/src/LanguageRuntime.cpp b/offload/languages/kernel/src/LanguageRuntime.cpp
index 0f80c31719176..fa158a2b149d7 100644
--- a/offload/languages/kernel/src/LanguageRuntime.cpp
+++ b/offload/languages/kernel/src/LanguageRuntime.cpp
@@ -20,8 +20,6 @@
#include "OffloadAPI.h"
-#include "DefineLanguageNames.inc"
-
#include <cstdio>
#include <cstdlib>
#include <cstring>
@@ -31,16 +29,8 @@
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;
- }
-}
+#include "DefineLanguageNames.inc"
+#include "LanguageUtils.cpp"
Error_t Malloc(void **DevPtr, size_t Size) {
ol_device_handle_t Device = language_runtime::getDefaultDevice();
@@ -104,7 +94,7 @@ Error_t DeviceSynchronize() {
Error_t GetDevice(int *DeviceNo) {
ol_device_handle_t Device = language_runtime::getDevice(DeviceNo);
if (!Device)
- return ErrorInvalidValue;
+ return ErrorInvalidDevice;
return Success;
}
@@ -116,7 +106,7 @@ Error_t GetDeviceCount(int *Count) {
Error_t SetDevice(int DeviceNo) {
ol_device_handle_t Device = language_runtime::setDefaultDevice(DeviceNo);
if (!Device)
- return ErrorInvalidValue;
+ return ErrorInvalidDevice;
assert(Device == language_runtime::getDefaultDevice() &&
"Set Device is not Default Device");
return Success;
@@ -152,18 +142,13 @@ Error_t GetDeviceProperties(DeviceProp_t *DeviceProp, int DeviceNo) {
return Success;
}
-static Error_t getQueueFromStream(Stream_t Stream, ol_queue_handle_t *Queue) {
- if (!Stream)
- return ErrorInvalidValue;
- *Queue = reinterpret_cast<ol_queue_handle_t>(Stream);
- return Success;
-}
-
Error_t StreamCreate(Stream_t *Stream) {
ol_queue_handle_t Queue;
- olCreateQueue(language_runtime::getDefaultDevice(), &Queue);
- *Stream = reinterpret_cast<Stream_t>(Queue);
- return Success;
+ ol_result_t Result =
+ olCreateQueue(language_runtime::getDefaultDevice(), &Queue);
+ if (Result == OL_SUCCESS)
+ *Stream = reinterpret_cast<Stream_t>(Queue);
+ return convertResult(Result);
}
Error_t StreamDestroy(Stream_t Stream) {
diff --git a/offload/languages/kernel/src/LanguageUtils.cpp b/offload/languages/kernel/src/LanguageUtils.cpp
new file mode 100644
index 0000000000000..64e0f5d7df618
--- /dev/null
+++ b/offload/languages/kernel/src/LanguageUtils.cpp
@@ -0,0 +1,65 @@
+//===-- LanguageUtils.cpp - Kernel language utilities ---------------------===//
+//
+// 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 "LanguageUtils.h"
+#include "LanguageRuntime.h"
+#include "OffloadAPI.h"
+
+static Error_t convertResult(ol_result_t Result) {
+ if (Result == OL_SUCCESS)
+ return Success;
+ switch (Result->Code) {
+ case OL_ERRC_INVALID_VALUE:
+ case OL_ERRC_INVALID_ARGUMENT:
+ case OL_ERRC_INVALID_NULL_POINTER:
+ return ErrorInvalidValue;
+ case OL_ERRC_INVALID_DEVICE:
+ return ErrorInvalidDevice;
+ default:
+ return ErrorUnknown;
+ }
+}
+
+const char *GetErrorName(Error_t Error) {
+ switch (Error) {
+#define LLVM_OFFLOAD_STRINGIFY_IMPL(NAME) #NAME
+#define LLVM_OFFLOAD_STRINGIFY(NAME) LLVM_OFFLOAD_STRINGIFY_IMPL(NAME)
+#define LLVM_OFFLOAD_ERR_STR(NAME) \
+ case NAME: \
+ return LLVM_OFFLOAD_STRINGIFY(NAME);
+ LLVM_OFFLOAD_ERR_STR(Success)
+ LLVM_OFFLOAD_ERR_STR(ErrorInvalidValue)
+ LLVM_OFFLOAD_ERR_STR(ErrorInvalidDevice)
+#undef LLVM_OFFLOAD_ERR_STR
+#undef LLVM_OFFLOAD_STRINGIFY
+#undef LLVM_OFFLOAD_STRINGIFY_IMPL
+ default:
+ return "Unrecognized error";
+ };
+}
+
+const char *GetErrorString(Error_t Error) {
+ switch (Error) {
+ case Success:
+ return "No error";
+ case ErrorInvalidValue:
+ return "Invalid argument value";
+ case ErrorInvalidDevice:
+ return "Invalid device number";
+ case ErrorUnknown:
+ return "Unknown error";
+ }
+ return "Unrecognized error";
+}
+
+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;
+}
diff --git a/offload/test/offloading/CUDA/device_api.cu b/offload/test/offloading/CUDA/device_api.cu
index 184167f2d17e4..af2b046eee397 100644
--- a/offload/test/offloading/CUDA/device_api.cu
+++ b/offload/test/offloading/CUDA/device_api.cu
@@ -41,5 +41,5 @@ int main(int argc, char **argv) {
cudaError_t Err = cudaSetDevice(-1);
printf("set invalid device: %u\n", Err);
- // CHECK: set invalid device: 1
+ // CHECK: set invalid device: 2
}
diff --git a/offload/test/offloading/CUDA/error_kinds.cu b/offload/test/offloading/CUDA/error_kinds.cu
new file mode 100644
index 0000000000000..5fa3035fe52f0
--- /dev/null
+++ b/offload/test/offloading/CUDA/error_kinds.cu
@@ -0,0 +1,61 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+static void print_error(const char *Label, cudaError_t Error) {
+ printf("%s value: %u\n", Label, static_cast<unsigned>(Error));
+ printf("%s name: %s\n", Label, cudaGetErrorName(Error));
+ printf("%s string: %s\n", Label, cudaGetErrorString(Error));
+}
+
+int main() {
+ print_error("success", cudaSuccess);
+ // CHECK: success value: 0
+ // CHECK: success name: cudaSuccess
+ // CHECK: success string: No error
+
+ print_error("invalid value", cudaErrorInvalidValue);
+ // CHECK: invalid value value: 1
+ // CHECK: invalid value name: cudaErrorInvalidValue
+ // CHECK: invalid value string: Invalid argument value
+
+ print_error("invalid device", cudaErrorInvalidDevice);
+ // CHECK: invalid device value: 2
+ // CHECK: invalid device name: cudaErrorInvalidDevice
+ // CHECK: invalid device string: Invalid device number
+
+ print_error("unknown", cudaErrorUnknown);
+ // CHECK: unknown value: 3
+ // CHECK: unknown name: Unrecognized error
+ // CHECK: unknown string: Unknown error
+
+ cudaError_t Unrecognized = static_cast<cudaError_t>(999);
+ print_error("unrecognized", Unrecognized);
+ // CHECK: unrecognized value: 999
+ // CHECK: unrecognized name: Unrecognized error
+ // CHECK: unrecognized string: Unrecognized error
+
+ print_error("set invalid device", cudaSetDevice(-1));
+ // CHECK: set invalid device value: 2
+ // CHECK: set invalid device name: cudaErrorInvalidDevice
+ // CHECK: set invalid device string: Invalid device number
+
+ print_error("null stream destroy", cudaStreamDestroy(nullptr));
+ // CHECK: null stream destroy value: 1
+ // CHECK: null stream destroy name: cudaErrorInvalidValue
+ // CHECK: null stream destroy string: Invalid argument value
+
+ return 0;
+}
diff --git a/offload/test/offloading/HIP/device_api.hip b/offload/test/offloading/HIP/device_api.hip
index 031e3703e66c1..5fb66e6e45eeb 100644
--- a/offload/test/offloading/HIP/device_api.hip
+++ b/offload/test/offloading/HIP/device_api.hip
@@ -41,5 +41,5 @@ int main(int argc, char **argv) {
hipError_t Err = hipSetDevice(-1);
printf("set invalid device: %u\n", Err);
- // CHECK: set invalid device: 1
+ // CHECK: set invalid device: 2
}
diff --git a/offload/test/offloading/HIP/error_kinds.hip b/offload/test/offloading/HIP/error_kinds.hip
new file mode 100644
index 0000000000000..af760a4cd4399
--- /dev/null
+++ b/offload/test/offloading/HIP/error_kinds.hip
@@ -0,0 +1,61 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <stdio.h>
+
+static void print_error(const char *Label, hipError_t Error) {
+ printf("%s value: %u\n", Label, static_cast<unsigned>(Error));
+ printf("%s name: %s\n", Label, hipGetErrorName(Error));
+ printf("%s string: %s\n", Label, hipGetErrorString(Error));
+}
+
+int main() {
+ print_error("success", hipSuccess);
+ // CHECK: success value: 0
+ // CHECK: success name: hipSuccess
+ // CHECK: success string: No error
+
+ print_error("invalid value", hipErrorInvalidValue);
+ // CHECK: invalid value value: 1
+ // CHECK: invalid value name: hipErrorInvalidValue
+ // CHECK: invalid value string: Invalid argument value
+
+ print_error("invalid device", hipErrorInvalidDevice);
+ // CHECK: invalid device value: 2
+ // CHECK: invalid device name: hipErrorInvalidDevice
+ // CHECK: invalid device string: Invalid device number
+
+ print_error("unknown", hipErrorUnknown);
+ // CHECK: unknown value: 3
+ // CHECK: unknown name: Unrecognized error
+ // CHECK: unknown string: Unknown error
+
+ hipError_t Unrecognized = static_cast<hipError_t>(999);
+ print_error("unrecognized", Unrecognized);
+ // CHECK: unrecognized value: 999
+ // CHECK: unrecognized name: Unrecognized error
+ // CHECK: unrecognized string: Unrecognized error
+
+ print_error("set invalid device", hipSetDevice(-1));
+ // CHECK: set invalid device value: 2
+ // CHECK: set invalid device name: hipErrorInvalidDevice
+ // CHECK: set invalid device string: Invalid device number
+
+ print_error("null stream destroy", hipStreamDestroy(nullptr));
+ // CHECK: null stream destroy value: 1
+ // CHECK: null stream destroy name: hipErrorInvalidValue
+ // CHECK: null stream destroy string: Invalid argument value
+
+ return 0;
+}
>From 00befbe26ba5140e5195d279ae8011c38682bfcd Mon Sep 17 00:00:00 2001
From: Sophia Herrmann <herrmann15 at llnl.gov>
Date: Fri, 31 Jul 2026 15:22:08 -0700
Subject: [PATCH 8/9] add new errors and threadlocal last error
---
.../include/kernel/DefineLanguageNames.inc | 4 ++
.../include/kernel/LanguageRuntime.h | 6 +++
.../include/kernel/UndefineLanguageNames.inc | 4 ++
.../languages/kernel/src/LanguageRuntime.cpp | 39 ++++++++++---------
.../languages/kernel/src/LanguageUtils.cpp | 25 ++++++++++++
offload/test/offloading/CUDA/error_kinds.cu | 10 +++++
offload/test/offloading/HIP/error_kinds.hip | 10 +++++
7 files changed, 79 insertions(+), 19 deletions(-)
diff --git a/offload/languages/include/kernel/DefineLanguageNames.inc b/offload/languages/include/kernel/DefineLanguageNames.inc
index 04b5bfac65d93..8f4410bb2cf9a 100644
--- a/offload/languages/include/kernel/DefineLanguageNames.inc
+++ b/offload/languages/include/kernel/DefineLanguageNames.inc
@@ -21,8 +21,12 @@
#define ErrorInvalidValue COMBINE(LANGUAGE, ErrorInvalidValue)
#define ErrorInvalidDevice COMBINE(LANGUAGE, ErrorInvalidDevice)
#define ErrorUnknown COMBINE(LANGUAGE, ErrorUnknown)
+#define ErrorInvalidResourceHandle COMBINE(LANGUAGE, ErrorInvalidResourceHandle)
+#define ErrorInvalidConfiguration COMBINE(LANGUAGE, ErrorInvalidConfiguration)
#define GetErrorName COMBINE(LANGUAGE, GetErrorName)
#define GetErrorString COMBINE(LANGUAGE, GetErrorString)
+#define GetLastError COMBINE(LANGUAGE, GetLastError)
+#define PeekAtLastError COMBINE(LANGUAGE, PeekAtLastError)
#define MemcpyKind COMBINE(LANGUAGE, MemcpyKind)
#define MemcpyHostToHost COMBINE(LANGUAGE, MemcpyHostToHost)
#define MemcpyHostToDevice COMBINE(LANGUAGE, MemcpyHostToDevice)
diff --git a/offload/languages/include/kernel/LanguageRuntime.h b/offload/languages/include/kernel/LanguageRuntime.h
index 80c78d7dd6d61..ea03bcc69faa6 100644
--- a/offload/languages/include/kernel/LanguageRuntime.h
+++ b/offload/languages/include/kernel/LanguageRuntime.h
@@ -19,12 +19,18 @@ enum Error_t : uint32_t {
ErrorInvalidValue = 1,
ErrorInvalidDevice = 2,
ErrorUnknown = 3,
+ ErrorInvalidResourceHandle = 4,
+ ErrorInvalidConfiguration = 5,
};
const char *GetErrorName(Error_t Error);
const char *GetErrorString(Error_t Error);
+Error_t GetLastError();
+
+Error_t PeekAtLastError();
+
struct DeviceProp_t {
char name[256];
size_t totalGlobalMem;
diff --git a/offload/languages/include/kernel/UndefineLanguageNames.inc b/offload/languages/include/kernel/UndefineLanguageNames.inc
index c3c2b7e35b17a..85cb445ab8c28 100644
--- a/offload/languages/include/kernel/UndefineLanguageNames.inc
+++ b/offload/languages/include/kernel/UndefineLanguageNames.inc
@@ -18,8 +18,12 @@
#undef ErrorInvalidValue
#undef ErrorInvalidDevice
#undef ErrorUnknown
+#undef ErrorInvalidResourceHandle
+#undef ErrorInvalidConfiguration
#undef GetErrorName
#undef GetErrorString
+#undef GetLastError
+#undef PeekAtLastError
#undef MemcpyKind
#undef MemcpyHostToHost
#undef MemcpyHostToDevice
diff --git a/offload/languages/kernel/src/LanguageRuntime.cpp b/offload/languages/kernel/src/LanguageRuntime.cpp
index fa158a2b149d7..cbb2040e92a52 100644
--- a/offload/languages/kernel/src/LanguageRuntime.cpp
+++ b/offload/languages/kernel/src/LanguageRuntime.cpp
@@ -35,12 +35,12 @@ namespace language_runtime = llvm::offload::kernel;
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);
+ return LastError = convertResult(Result);
}
Error_t Free(void *DevPtr) {
ol_result_t Result = olMemFree(DevPtr);
- return convertResult(Result);
+ return LastError = convertResult(Result);
}
Error_t Memcpy(void *Dst, const void *Src, size_t Size, MemcpyKind Kind) {
@@ -77,10 +77,11 @@ Error_t Memcpy(void *Dst, const void *Src, size_t Size, MemcpyKind Kind) {
fprintf(stderr, LANGUAGE_STR "MemcpyDefault is not implemented yet");
abort();
};
-
+ if (Result != OL_SUCCESS) {
+ return LastError = convertResult(Result);
+ }
Result = olSyncQueue(Queue);
-
- return convertResult(Result);
+ return LastError = convertResult(Result);
}
Error_t DeviceSynchronize() {
@@ -88,35 +89,35 @@ Error_t DeviceSynchronize() {
// plugins.
ol_queue_handle_t Queue = language_runtime::getDefaultQueue();
ol_result_t Result = olSyncQueue(Queue);
- return convertResult(Result);
+ return LastError = convertResult(Result);
}
Error_t GetDevice(int *DeviceNo) {
ol_device_handle_t Device = language_runtime::getDevice(DeviceNo);
if (!Device)
- return ErrorInvalidDevice;
- return Success;
+ return LastError = ErrorInvalidDevice;
+ return LastError = Success;
}
Error_t GetDeviceCount(int *Count) {
*Count = language_runtime::getDeviceCount();
- return Success;
+ return LastError = Success;
}
Error_t SetDevice(int DeviceNo) {
ol_device_handle_t Device = language_runtime::setDefaultDevice(DeviceNo);
if (!Device)
- return ErrorInvalidDevice;
+ return LastError = ErrorInvalidDevice;
assert(Device == language_runtime::getDefaultDevice() &&
"Set Device is not Default Device");
- return Success;
+ return LastError = Success;
}
Error_t HostAlloc(void **Ptr, size_t Size, unsigned int Flags) {
// TODO:
ol_device_handle_t Device = language_runtime::getDefaultDevice();
ol_result_t Result = olMemAllocHost(Device, Size, Ptr);
- return convertResult(Result);
+ return LastError = convertResult(Result);
}
Error_t MallocHost(void **Ptr, size_t Size) {
@@ -125,7 +126,7 @@ Error_t MallocHost(void **Ptr, size_t Size) {
Error_t FreeHost(void *Ptr) {
ol_result_t Result = olMemFree(Ptr);
- return convertResult(Result);
+ return LastError = convertResult(Result);
}
Error_t GetDeviceProperties(DeviceProp_t *DeviceProp, int DeviceNo) {
@@ -139,7 +140,7 @@ Error_t GetDeviceProperties(DeviceProp_t *DeviceProp, int DeviceNo) {
&DeviceProp->multiProcessorCount);
olGetDeviceInfo(Device, OL_DEVICE_INFO_NUM_LANES, sizeof(uint32_t),
&DeviceProp->warpSize);
- return Success;
+ return LastError = Success;
}
Error_t StreamCreate(Stream_t *Stream) {
@@ -148,23 +149,23 @@ Error_t StreamCreate(Stream_t *Stream) {
olCreateQueue(language_runtime::getDefaultDevice(), &Queue);
if (Result == OL_SUCCESS)
*Stream = reinterpret_cast<Stream_t>(Queue);
- return convertResult(Result);
+ return LastError = convertResult(Result);
}
Error_t StreamDestroy(Stream_t Stream) {
ol_queue_handle_t Queue;
Error_t Err = getQueueFromStream(Stream, &Queue);
if (Err != Success)
- return Err;
+ return LastError = Err;
ol_result_t Result = olDestroyQueue(Queue);
- return convertResult(Result);
+ return LastError = convertResult(Result);
}
Error_t StreamSynchronize(Stream_t Stream) {
ol_queue_handle_t Queue;
Error_t Err = getQueueFromStream(Stream, &Queue);
if (Err != Success)
- return Err;
+ return LastError = Err;
ol_result_t Result = olSyncQueue(Queue);
- return convertResult(Result);
+ return LastError = convertResult(Result);
}
diff --git a/offload/languages/kernel/src/LanguageUtils.cpp b/offload/languages/kernel/src/LanguageUtils.cpp
index 64e0f5d7df618..4433db99438de 100644
--- a/offload/languages/kernel/src/LanguageUtils.cpp
+++ b/offload/languages/kernel/src/LanguageUtils.cpp
@@ -20,11 +20,20 @@ static Error_t convertResult(ol_result_t Result) {
return ErrorInvalidValue;
case OL_ERRC_INVALID_DEVICE:
return ErrorInvalidDevice;
+ case OL_ERRC_INVALID_SIZE:
+ return ErrorInvalidConfiguration;
+ case OL_ERRC_INVALID_NULL_HANDLE:
+ case OL_ERRC_INVALID_QUEUE:
+ case OL_ERRC_INVALID_EVENT:
+ case OL_ERRC_INVALID_CONTEXT:
+ return ErrorInvalidResourceHandle;
default:
return ErrorUnknown;
}
}
+static thread_local Error_t LastError = Success;
+
const char *GetErrorName(Error_t Error) {
switch (Error) {
#define LLVM_OFFLOAD_STRINGIFY_IMPL(NAME) #NAME
@@ -35,6 +44,8 @@ const char *GetErrorName(Error_t Error) {
LLVM_OFFLOAD_ERR_STR(Success)
LLVM_OFFLOAD_ERR_STR(ErrorInvalidValue)
LLVM_OFFLOAD_ERR_STR(ErrorInvalidDevice)
+ LLVM_OFFLOAD_ERR_STR(ErrorInvalidResourceHandle)
+ LLVM_OFFLOAD_ERR_STR(ErrorInvalidConfiguration)
#undef LLVM_OFFLOAD_ERR_STR
#undef LLVM_OFFLOAD_STRINGIFY
#undef LLVM_OFFLOAD_STRINGIFY_IMPL
@@ -53,10 +64,24 @@ const char *GetErrorString(Error_t Error) {
return "Invalid device number";
case ErrorUnknown:
return "Unknown error";
+ case ErrorInvalidResourceHandle:
+ return "Invalid resource handle";
+ case ErrorInvalidConfiguration:
+ return "Invalid configuration argument";
}
return "Unrecognized error";
}
+Error_t GetLastError() {
+ Error_t Error = LastError;
+ LastError = Success;
+ return Error;
+}
+
+Error_t PeekAtLastError() {
+ return LastError;
+}
+
static Error_t getQueueFromStream(Stream_t Stream, ol_queue_handle_t *Queue) {
if (!Stream)
return ErrorInvalidValue;
diff --git a/offload/test/offloading/CUDA/error_kinds.cu b/offload/test/offloading/CUDA/error_kinds.cu
index 5fa3035fe52f0..9789b23d95779 100644
--- a/offload/test/offloading/CUDA/error_kinds.cu
+++ b/offload/test/offloading/CUDA/error_kinds.cu
@@ -41,6 +41,16 @@ int main() {
// CHECK: unknown name: Unrecognized error
// CHECK: unknown string: Unknown error
+ print_error("invalid resource handle", cudaErrorInvalidResourceHandle);
+ // CHECK: invalid resource handle value: 4
+ // CHECK: invalid resource handle name: cudaErrorInvalidResourceHandle
+ // CHECK: invalid resource handle string: Invalid resource handle
+
+ print_error("invalid configuration", cudaErrorInvalidConfiguration);
+ // CHECK: invalid configuration value: 5
+ // CHECK: invalid configuration name: cudaErrorInvalidConfiguration
+ // CHECK: invalid configuration string: Invalid configuration argument
+
cudaError_t Unrecognized = static_cast<cudaError_t>(999);
print_error("unrecognized", Unrecognized);
// CHECK: unrecognized value: 999
diff --git a/offload/test/offloading/HIP/error_kinds.hip b/offload/test/offloading/HIP/error_kinds.hip
index af760a4cd4399..1ba9f827b323a 100644
--- a/offload/test/offloading/HIP/error_kinds.hip
+++ b/offload/test/offloading/HIP/error_kinds.hip
@@ -41,6 +41,16 @@ int main() {
// CHECK: unknown name: Unrecognized error
// CHECK: unknown string: Unknown error
+ print_error("invalid resource handle", hipErrorInvalidResourceHandle);
+ // CHECK: invalid resource handle value: 4
+ // CHECK: invalid resource handle name: hipErrorInvalidResourceHandle
+ // CHECK: invalid resource handle string: Invalid resource handle
+
+ print_error("invalid configuration", hipErrorInvalidConfiguration);
+ // CHECK: invalid configuration value: 5
+ // CHECK: invalid configuration name: hipErrorInvalidConfiguration
+ // CHECK: invalid configuration string: Invalid configuration argument
+
hipError_t Unrecognized = static_cast<hipError_t>(999);
print_error("unrecognized", Unrecognized);
// CHECK: unrecognized value: 999
>From 7e0721430115a9f0a64f785a77825eea097b880e Mon Sep 17 00:00:00 2001
From: Sophia Herrmann <herrmann15 at llnl.gov>
Date: Fri, 31 Jul 2026 15:23:38 -0700
Subject: [PATCH 9/9] add language-based KernelLaunches and preserve error
---
clang/lib/CodeGen/CGCUDANV.cpp | 4 +-
clang/lib/Headers/__clang_gpu_builtin_vars.h | 6 +-
clang/test/CodeGenCUDA/Inputs/cuda.h | 8 +-
clang/test/CodeGenCUDA/offload_via_llvm.cu | 2 +-
offload/languages/cuda/src/cuda_runtime.cpp | 17 ++++
offload/languages/hip/src/hip_runtime.cpp | 16 ++++
offload/languages/include/hip/hip_runtime.h | 4 -
.../include/kernel/DefineLanguageNames.inc | 3 +
.../include/kernel/LanguageRuntime.h | 14 ++++
.../include => include/kernel}/Types.h | 11 ++-
.../include/kernel/UndefineLanguageNames.inc | 3 +
offload/languages/kernel/CMakeLists.txt | 1 +
.../languages/kernel/include/LanguageLaunch.h | 1 -
.../languages/kernel/src/LanguageLaunch.cpp | 43 ++++++++--
offload/test/offloading/CUDA/get_errs.cu | 82 +++++++++++++++++++
offload/test/offloading/HIP/get_errs.hip | 81 ++++++++++++++++++
16 files changed, 275 insertions(+), 21 deletions(-)
rename offload/languages/{kernel/include => include/kernel}/Types.h (73%)
create mode 100644 offload/test/offloading/CUDA/get_errs.cu
create mode 100644 offload/test/offloading/HIP/get_errs.hip
diff --git a/clang/lib/CodeGen/CGCUDANV.cpp b/clang/lib/CodeGen/CGCUDANV.cpp
index 0ea3ed36fae83..f976aac60a062 100644
--- a/clang/lib/CodeGen/CGCUDANV.cpp
+++ b/clang/lib/CodeGen/CGCUDANV.cpp
@@ -445,9 +445,7 @@ void CGNVCUDARuntime::emitDeviceStubBodyNew(CodeGenFunction &CGF,
else if (CGF.getLangOpts().CUDA)
KernelLaunchAPI = KernelLaunchAPI + "_ptsz";
}
- /// Use __llvmLaunchKernel for LLVMOffload.
- auto LaunchKernelName = UsesLLVMOffloading ? "__llvm" + KernelLaunchAPI
- : addPrefixToName(KernelLaunchAPI);
+ auto LaunchKernelName = addPrefixToName(KernelLaunchAPI);
const IdentifierInfo &cudaLaunchKernelII =
CGM.getContext().Idents.get(LaunchKernelName);
FunctionDecl *cudaLaunchKernelFD = nullptr;
diff --git a/clang/lib/Headers/__clang_gpu_builtin_vars.h b/clang/lib/Headers/__clang_gpu_builtin_vars.h
index f3137be0a3181..0cb4a5160ff49 100644
--- a/clang/lib/Headers/__clang_gpu_builtin_vars.h
+++ b/clang/lib/Headers/__clang_gpu_builtin_vars.h
@@ -25,9 +25,9 @@ static inline __attribute__((device)) const struct {
extern "C" {
typedef struct dim3 {
- dim3() {}
- dim3(unsigned x) : x(x) {}
- unsigned x = 0, y = 0, z = 0;
+ constexpr dim3(unsigned x = 1, unsigned y = 1, unsigned z = 1)
+ : x(x), y(y), z(z) {}
+ unsigned x, y, z;
} dim3;
// TODO: For some reason the CUDA device compilation requires this declaration
diff --git a/clang/test/CodeGenCUDA/Inputs/cuda.h b/clang/test/CodeGenCUDA/Inputs/cuda.h
index 83bb7b7bdbb7f..c2d5830dca6ca 100644
--- a/clang/test/CodeGenCUDA/Inputs/cuda.h
+++ b/clang/test/CodeGenCUDA/Inputs/cuda.h
@@ -53,10 +53,14 @@ extern "C" hipError_t hipLaunchKernel_spt(const void *func, dim3 gridDim,
hipStream_t stream);
#endif // __HIP_API_PER_THREAD_DEFAULT_STREAM__
#elif __OFFLOAD_VIA_LLVM__
+typedef struct cudaStream *cudaStream_t;
+typedef enum cudaError {} cudaError_t;
extern "C" unsigned __llvmPushCallConfiguration(dim3 gridDim, dim3 blockDim,
size_t sharedMem = 0, void *stream = 0);
-extern "C" unsigned __llvmLaunchKernel(const void *func, dim3 gridDim, dim3 blockDim,
- void **args, size_t sharedMem = 0, void *stream = 0);
+extern "C" cudaError_t cudaLaunchKernel(const void *func, dim3 gridDim,
+ dim3 blockDim, void **args,
+ size_t sharedMem = 0,
+ cudaStream_t stream = 0);
#else
typedef struct cudaStream *cudaStream_t;
typedef enum cudaError {} cudaError_t;
diff --git a/clang/test/CodeGenCUDA/offload_via_llvm.cu b/clang/test/CodeGenCUDA/offload_via_llvm.cu
index c58e793d11826..b80e453d8a342 100644
--- a/clang/test/CodeGenCUDA/offload_via_llvm.cu
+++ b/clang/test/CodeGenCUDA/offload_via_llvm.cu
@@ -50,7 +50,7 @@
// HST-NEXT: [[TMP15:%.*]] = call i32 @__cudaPopCallConfiguration(ptr [[GRID_DIM]], ptr [[BLOCK_DIM]], ptr [[SHMEM_SIZE]], ptr [[STREAM]])
// HST-NEXT: [[TMP16:%.*]] = load i32, ptr [[SHMEM_SIZE]], align 4
// HST-NEXT: [[TMP17:%.*]] = load ptr, ptr [[STREAM]], align 4
-// HST-NEXT: [[CALL:%.*]] = call noundef i32 @__llvmLaunchKernel(ptr noundef @_Z18__device_stub__fooisPvS_, ptr noundef byval([[STRUCT_DIM3]]) align 4 [[GRID_DIM]], ptr noundef byval([[STRUCT_DIM3]]) align 4 [[BLOCK_DIM]], ptr noundef [[KERNEL_LAUNCH_PARAMS]], i32 noundef [[TMP16]], ptr noundef [[TMP17]]) #[[ATTR3:[0-9]+]]
+// HST-NEXT: [[CALL:%.*]] = call noundef i32 @cudaLaunchKernel(ptr noundef @_Z18__device_stub__fooisPvS_, ptr noundef byval([[STRUCT_DIM3]]) align 4 [[GRID_DIM]], ptr noundef byval([[STRUCT_DIM3]]) align 4 [[BLOCK_DIM]], ptr noundef [[KERNEL_LAUNCH_PARAMS]], i32 noundef [[TMP16]], ptr noundef [[TMP17]]) #[[ATTR3:[0-9]+]]
// HST-NEXT: br label %[[SETUP_END:.*]]
// HST: [[SETUP_END]]:
// HST-NEXT: ret void
diff --git a/offload/languages/cuda/src/cuda_runtime.cpp b/offload/languages/cuda/src/cuda_runtime.cpp
index 536fc8e38bb72..bba072ac4ec0d 100644
--- a/offload/languages/cuda/src/cuda_runtime.cpp
+++ b/offload/languages/cuda/src/cuda_runtime.cpp
@@ -9,8 +9,25 @@
#include "cuda_runtime.h"
+#include "LanguageLaunch.h"
#include "OffloadAPI.h"
#define LANGUAGE cuda
#include "../../kernel/src/LanguageRuntime.cpp"
+
+extern "C" {
+#define CUDA_LAUNCH_KERNEL(SUFFIX) \
+ cudaError_t cudaLaunchKernel##SUFFIX( \
+ const char *KernelID, dim3 GridDim, dim3 BlockDim, void *KernelArgsPtr, \
+ size_t DynamicSharedMem, void *Stream) { \
+ return LastError = convertResult(__llvmLaunchKernelImpl( \
+ KernelID, GridDim, BlockDim, KernelArgsPtr, DynamicSharedMem, \
+ Stream)); \
+ }
+
+CUDA_LAUNCH_KERNEL()
+CUDA_LAUNCH_KERNEL(_ptsz)
+CUDA_LAUNCH_KERNEL(_spt)
+#undef CUDA_LAUNCH_KERNEL
+}
diff --git a/offload/languages/hip/src/hip_runtime.cpp b/offload/languages/hip/src/hip_runtime.cpp
index 881eca894a364..faa9df1e89797 100644
--- a/offload/languages/hip/src/hip_runtime.cpp
+++ b/offload/languages/hip/src/hip_runtime.cpp
@@ -15,3 +15,19 @@
#define LANGUAGE hip
#include "../../kernel/src/LanguageRuntime.cpp"
+
+extern "C" {
+#define HIP_LAUNCH_KERNEL(SUFFIX) \
+ hipError_t hipLaunchKernel##SUFFIX( \
+ const char *KernelID, dim3 GridDim, dim3 BlockDim, void *KernelArgsPtr, \
+ size_t DynamicSharedMem, void *Stream) { \
+ return LastError = convertResult(__llvmLaunchKernelImpl( \
+ KernelID, GridDim, BlockDim, KernelArgsPtr, DynamicSharedMem, \
+ Stream)); \
+ }
+
+HIP_LAUNCH_KERNEL()
+HIP_LAUNCH_KERNEL(_spt)
+HIP_LAUNCH_KERNEL(_ptsz)
+#undef HIP_LAUNCH_KERNEL
+}
diff --git a/offload/languages/include/hip/hip_runtime.h b/offload/languages/include/hip/hip_runtime.h
index 3492fa08e8461..d3cf8cf720bd7 100644
--- a/offload/languages/include/hip/hip_runtime.h
+++ b/offload/languages/include/hip/hip_runtime.h
@@ -46,10 +46,6 @@ template <class T> static inline hipError_t hipHostFree(T *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,
diff --git a/offload/languages/include/kernel/DefineLanguageNames.inc b/offload/languages/include/kernel/DefineLanguageNames.inc
index 8f4410bb2cf9a..a3ef3287fd4f9 100644
--- a/offload/languages/include/kernel/DefineLanguageNames.inc
+++ b/offload/languages/include/kernel/DefineLanguageNames.inc
@@ -48,3 +48,6 @@
#define StreamCreate COMBINE(LANGUAGE, StreamCreate)
#define StreamDestroy COMBINE(LANGUAGE, StreamDestroy)
#define StreamSynchronize COMBINE(LANGUAGE, StreamSynchronize)
+#define LaunchKernel COMBINE(LANGUAGE, LaunchKernel)
+#define LaunchKernel_spt COMBINE(LANGUAGE, LaunchKernel_spt)
+#define LaunchKernel_ptsz COMBINE(LANGUAGE, LaunchKernel_ptsz)
diff --git a/offload/languages/include/kernel/LanguageRuntime.h b/offload/languages/include/kernel/LanguageRuntime.h
index ea03bcc69faa6..6e0e3a68bad61 100644
--- a/offload/languages/include/kernel/LanguageRuntime.h
+++ b/offload/languages/include/kernel/LanguageRuntime.h
@@ -9,6 +9,8 @@
#pragma once
+#include "Types.h"
+
#include <cstddef>
#include <cstdint>
#include <cstdio>
@@ -117,6 +119,18 @@ Error_t StreamDestroy(Stream_t stream);
Error_t StreamSynchronize(Stream_t stream);
+extern "C" Error_t LaunchKernel(const char *KernelID, dim3 GridDim,
+ dim3 BlockDim, void *KernelArgsPtr,
+ size_t DynamicSharedMem, void *Stream);
+
+extern "C" Error_t LaunchKernel_spt(const char *KernelID, dim3 GridDim,
+ dim3 BlockDim, void *KernelArgsPtr,
+ size_t DynamicSharedMem, void *Stream);
+
+extern "C" Error_t LaunchKernel_ptsz(const char *KernelID, dim3 GridDim,
+ dim3 BlockDim, void *KernelArgsPtr,
+ size_t DynamicSharedMem, void *Stream);
+
#if defined(__AMDGPU__) || defined(__NVPTX__)
#include <gpuintrin.h>
diff --git a/offload/languages/kernel/include/Types.h b/offload/languages/include/kernel/Types.h
similarity index 73%
rename from offload/languages/kernel/include/Types.h
rename to offload/languages/include/kernel/Types.h
index 6c5a0ec19a8f7..b2edc1f1c79b9 100644
--- a/offload/languages/kernel/include/Types.h
+++ b/offload/languages/include/kernel/Types.h
@@ -10,15 +10,22 @@
#pragma once
-#include "Types.h"
#include <cstddef>
#include <cstdint>
+#ifdef __CLANG_GPU_BUILTIN_VARS_H__
+using uint3 = dim3;
+#else
struct uint3 {
unsigned x = 0, y = 0, z = 0;
};
-using dim3 = uint3;
+struct dim3 : uint3 {
+ constexpr dim3(unsigned X = 1, unsigned Y = 1, unsigned Z = 1)
+ : uint3{X, Y, Z} {}
+ constexpr dim3(uint3 V) : uint3{V.x, V.y, V.z} {}
+};
+#endif
struct CallConfigurationTy {
dim3 GridSize;
diff --git a/offload/languages/include/kernel/UndefineLanguageNames.inc b/offload/languages/include/kernel/UndefineLanguageNames.inc
index 85cb445ab8c28..6a285946cd158 100644
--- a/offload/languages/include/kernel/UndefineLanguageNames.inc
+++ b/offload/languages/include/kernel/UndefineLanguageNames.inc
@@ -46,3 +46,6 @@
#undef StreamCreate
#undef StreamDestroy
#undef StreamSynchronize
+#undef LaunchKernel
+#undef LaunchKernel_spt
+#undef LaunchKernel_ptsz
diff --git a/offload/languages/kernel/CMakeLists.txt b/offload/languages/kernel/CMakeLists.txt
index 093297184e8e4..b7441cb1a4a80 100644
--- a/offload/languages/kernel/CMakeLists.txt
+++ b/offload/languages/kernel/CMakeLists.txt
@@ -50,5 +50,6 @@ install(TARGETS LLVMOffloadKernel
install(FILES
${CMAKE_CURRENT_SOURCE_DIR}/../include/kernel/DefineLanguageNames.inc
${CMAKE_CURRENT_SOURCE_DIR}/../include/kernel/LanguageRuntime.h
+ ${CMAKE_CURRENT_SOURCE_DIR}/../include/kernel/Types.h
${CMAKE_CURRENT_SOURCE_DIR}/../include/kernel/UndefineLanguageNames.inc
DESTINATION ${CMAKE_INSTALL_PREFIX}/include/offload/kernel/)
diff --git a/offload/languages/kernel/include/LanguageLaunch.h b/offload/languages/kernel/include/LanguageLaunch.h
index 4be6f2499eb04..7df6f1525ba95 100644
--- a/offload/languages/kernel/include/LanguageLaunch.h
+++ b/offload/languages/kernel/include/LanguageLaunch.h
@@ -14,7 +14,6 @@
#include "OffloadAPI.h"
#include "Types.h"
-#include <algorithm> // for std::max
#include <cstddef>
#include <cstdint>
diff --git a/offload/languages/kernel/src/LanguageLaunch.cpp b/offload/languages/kernel/src/LanguageLaunch.cpp
index 9d7fb26d4768a..afe4f0c46cc6e 100644
--- a/offload/languages/kernel/src/LanguageLaunch.cpp
+++ b/offload/languages/kernel/src/LanguageLaunch.cpp
@@ -9,12 +9,25 @@
//===----------------------------------------------------------------------===//
#include "LanguageLaunch.h"
+#include "OffloadAPI.h"
#include "RuntimeAPI.h"
#include <cstdio>
namespace language_launch = llvm::offload::kernel;
+static constexpr ol_error_struct_t InvalidKernelError = {
+ OL_ERRC_INVALID_NULL_HANDLE, "kernel is not registered"};
+
+static constexpr ol_error_struct_t InvalidDeviceError = {OL_ERRC_INVALID_DEVICE,
+ "invalid device"};
+
+static constexpr ol_error_struct_t InvalidArgumentError = {
+ OL_ERRC_INVALID_ARGUMENT, "invalid argument to kernel launch"};
+
+static constexpr ol_error_struct_t InvalidConfigurationError = {
+ OL_ERRC_INVALID_SIZE, "invalid kernel launch configuration"};
+
extern "C" {
/// Push call configuration for kernel launch
@@ -45,19 +58,32 @@ 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();
+ if (!Device)
+ return &InvalidDeviceError;
+ if (!KernelID)
+ return &InvalidArgumentError;
ol_symbol_handle_t Kernel = language_launch::getKernel(KernelID);
+ if (!Kernel)
+ return &InvalidKernelError;
+
+ if (GridDim.x == 0 || GridDim.y == 0 || GridDim.z == 0 || BlockDim.x == 0 ||
+ BlockDim.y == 0 || BlockDim.z == 0)
+ return &InvalidConfigurationError;
+
+ if (!KernelArgsPtr)
+ return &InvalidArgumentError;
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);
+ 1 + (GridDim.y > 1 || BlockDim.y > 1) + (GridDim.z > 1 || BlockDim.z > 1);
GridDimensions.x = GridDim.x;
- GridDimensions.y = std::max(GridDim.y, 1u);
- GridDimensions.z = std::max(GridDim.z, 1u);
+ GridDimensions.y = GridDim.y;
+ GridDimensions.z = GridDim.z;
LaunchSizeArgs.NumGroups = GridDimensions;
BlockDimensions.x = BlockDim.x;
- BlockDimensions.y = std::max(BlockDim.y, 1u);
- BlockDimensions.z = std::max(BlockDim.z, 1u);
+ BlockDimensions.y = BlockDim.y;
+ BlockDimensions.z = BlockDim.z;
LaunchSizeArgs.GroupSize = BlockDimensions;
LaunchSizeArgs.DynSharedMemory = DynamicSharedMem;
@@ -73,6 +99,13 @@ ol_result_t __llvmLaunchKernelImpl(const char *KernelID, dim3 GridDim,
size_t *ArgSizes;
};
OffloadKernelArgs *OKA = reinterpret_cast<OffloadKernelArgs *>(KernelArgsPtr);
+ if ((!OKA->Args) != (!OKA->ArgSizes))
+ return &InvalidArgumentError;
+ if (OKA->NumArgs > 0 && !OKA->Args)
+ return &InvalidArgumentError;
+ for (size_t I = 0; I < OKA->NumArgs; ++I)
+ if (!OKA->Args[I] || OKA->ArgSizes[I] == 0)
+ return &InvalidArgumentError;
ol_result_t Result;
Result = olLaunchKernel(Queue, Device, Kernel, &LaunchSizeArgs, &Properties,
diff --git a/offload/test/offloading/CUDA/get_errs.cu b/offload/test/offloading/CUDA/get_errs.cu
new file mode 100644
index 0000000000000..3e000e7782eb5
--- /dev/null
+++ b/offload/test/offloading/CUDA/get_errs.cu
@@ -0,0 +1,82 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -pthread -std=c++17
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp -pthread -std=c++17
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <cstdio>
+#include <cuda_runtime.h>
+#include <mutex>
+#include <thread>
+
+static std::mutex PrintMutex;
+
+static void printError(int ThreadId, const char *Label, cudaError_t Error) {
+ std::lock_guard<std::mutex> Lock(PrintMutex);
+ printf("thread %d %s: %s\n", ThreadId, Label, cudaGetErrorName(Error));
+ std::fflush(stdout);
+}
+
+__global__ void errorKernel(float *d_out) {
+ int idx = blockIdx.x * blockDim.x + threadIdx.x;
+ d_out[idx] = idx * 0.5f;
+}
+
+void runTask(int thread_id) {
+ const int N = 1 << 20;
+ size_t bytes = N * sizeof(float);
+
+ float *d_data;
+ cudaMalloc(&d_data, bytes);
+ printError(thread_id, "cudaMalloc", cudaGetLastError());
+
+ thread_id == 1 ? errorKernel<<<4096, 256>>>(d_data)
+ : errorKernel<<<4096, 0>>>(d_data);
+ printError(thread_id, "kernel launch", cudaPeekAtLastError());
+
+ printError(thread_id, "kernel launch get", cudaGetLastError());
+
+ printError(thread_id, "kernel launch get again", cudaGetLastError());
+
+ cudaDeviceSynchronize();
+ cudaFree(d_data);
+}
+
+int main() {
+ printError(0, "initial", cudaPeekAtLastError());
+ // CHECK: thread 0 initial: cudaSuccess
+
+ std::thread t1(runTask, 1);
+ std::thread t2(runTask, 2);
+
+ t1.join();
+ t2.join();
+ // CHECK-DAG: thread 1 cudaMalloc: cudaSuccess
+ // CHECK-DAG: thread 2 cudaMalloc: cudaSuccess
+ // CHECK-DAG: thread 1 kernel launch: cudaSuccess
+ // CHECK-DAG: thread 2 kernel launch: cudaErrorInvalidConfiguration
+ // CHECK-DAG: thread 1 kernel launch get: cudaSuccess
+ // CHECK-DAG: thread 2 kernel launch get: cudaErrorInvalidConfiguration
+ // CHECK-DAG: thread 1 kernel launch get again: cudaSuccess
+ // CHECK-DAG: thread 2 kernel launch get again: cudaSuccess
+
+ std::thread t3(runTask, 3);
+ t3.join();
+ // CHECK: thread 3 cudaMalloc: cudaSuccess
+ // CHECK: thread 3 kernel launch: cudaErrorInvalidConfiguration
+ // CHECK: thread 3 kernel launch get: cudaErrorInvalidConfiguration
+ // CHECK: thread 3 kernel launch get again: cudaSuccess
+
+ printError(0, "joined", cudaGetLastError());
+ // CHECK: thread 0 joined: cudaSuccess
+
+ return 0;
+}
diff --git a/offload/test/offloading/HIP/get_errs.hip b/offload/test/offloading/HIP/get_errs.hip
new file mode 100644
index 0000000000000..c400c5e371cda
--- /dev/null
+++ b/offload/test/offloading/HIP/get_errs.hip
@@ -0,0 +1,81 @@
+// clang-format off
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -pthread -std=c++17
+// RUN: %t | %fcheck-generic
+// RUN: %clang++ %flags -foffload-via-llvm --offload-arch=native %s -o %t -fopenmp -pthread -std=c++17
+// RUN: %t | %fcheck-generic
+// clang-format on
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
+// UNSUPPORTED: amdgpu-amd-amdhsa-LTO
+// UNSUPPORTED: intelgpu
+
+#include <cstdio>
+#include <mutex>
+#include <thread>
+
+static std::mutex PrintMutex;
+
+static void printError(int ThreadId, const char *Label, hipError_t Error) {
+ std::lock_guard<std::mutex> Lock(PrintMutex);
+ printf("thread %d %s: %s\n", ThreadId, Label, hipGetErrorName(Error));
+ std::fflush(stdout);
+}
+
+__global__ void errorKernel(float *d_out) {
+ int idx = blockIdx.x * blockDim.x + threadIdx.x;
+ d_out[idx] = idx * 0.5f;
+}
+
+void runTask(int ThreadId) {
+ const int N = 1 << 20;
+ size_t Bytes = N * sizeof(float);
+
+ float *d_data;
+ hipMalloc(&d_data, Bytes);
+ printError(ThreadId, "hipMalloc", hipGetLastError());
+
+ ThreadId == 1 ? errorKernel<<<4096, 256>>>(d_data)
+ : errorKernel<<<4096, 0>>>(d_data);
+ printError(ThreadId, "kernel launch", hipPeekAtLastError());
+
+ printError(ThreadId, "kernel launch get", hipGetLastError());
+
+ printError(ThreadId, "kernel launch get again", hipGetLastError());
+
+ hipDeviceSynchronize();
+ hipFree(d_data);
+}
+
+int main() {
+ printError(0, "initial", hipPeekAtLastError());
+ // CHECK: thread 0 initial: hipSuccess
+
+ std::thread t1(runTask, 1);
+ std::thread t2(runTask, 2);
+
+ t1.join();
+ t2.join();
+ // CHECK-DAG: thread 1 hipMalloc: hipSuccess
+ // CHECK-DAG: thread 2 hipMalloc: hipSuccess
+ // CHECK-DAG: thread 1 kernel launch: hipSuccess
+ // CHECK-DAG: thread 2 kernel launch: hipErrorInvalidConfiguration
+ // CHECK-DAG: thread 1 kernel launch get: hipSuccess
+ // CHECK-DAG: thread 2 kernel launch get: hipErrorInvalidConfiguration
+ // CHECK-DAG: thread 1 kernel launch get again: hipSuccess
+ // CHECK-DAG: thread 2 kernel launch get again: hipSuccess
+
+ std::thread t3(runTask, 3);
+ t3.join();
+ // CHECK: thread 3 hipMalloc: hipSuccess
+ // CHECK: thread 3 kernel launch: hipErrorInvalidConfiguration
+ // CHECK: thread 3 kernel launch get: hipErrorInvalidConfiguration
+ // CHECK: thread 3 kernel launch get again: hipSuccess
+
+ printError(0, "joined", hipGetLastError());
+ // CHECK: thread 0 joined: hipSuccess
+
+ return 0;
+}
More information about the cfe-commits
mailing list