[clang] [llvm] [LLVMOffload] Add more error types + GetLastError/PeekAtLastError (PR #213389)
Sophia Herrmann via cfe-commits
cfe-commits at lists.llvm.org
Fri Jul 31 17:50:51 PDT 2026
https://github.com/jellytabby updated https://github.com/llvm/llvm-project/pull/213389
>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/8] 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/8] 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/8] 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/8] 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/8] 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/8] 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/8] 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 f4016fad1af9e650e6c6872dd20e85235a5577b7 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/8] 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 | 23 +++++++++++
offload/test/offloading/CUDA/error_kinds.cu | 10 +++++
offload/test/offloading/HIP/error_kinds.hip | 10 +++++
7 files changed, 77 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..f4fdea178b2e2 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,22 @@ 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
More information about the cfe-commits
mailing list