[llvm] [offload] use argument pointer array in olLaunchKernel (PR #194333)
Piotr Balcer via llvm-commits
llvm-commits at lists.llvm.org
Thu May 28 07:39:06 PDT 2026
https://github.com/pbalcer updated https://github.com/llvm/llvm-project/pull/194333
>From fbac7f4645269e2e7077bc8c7dd6d08d63673941 Mon Sep 17 00:00:00 2001
From: Piotr Balcer <piotr.balcer at intel.com>
Date: Mon, 20 Apr 2026 09:41:23 +0000
Subject: [PATCH 1/2] [offload] use argument pointer array in olLaunchKernel
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit
The existing olLaunchKernel function accepts arguments through a single
contiguous packaged buffer, with an unspecified, implementation-defined,
layout. In practice, the existing AMDGPU and CUDA backend plugins require
the layout to follow C/C++ ABI rules.
This makes for a simple API when the arguments are already packaged in
an existing structure:
```
struct {
char A;
int *B;
short C;
} Args{0, nullptr, 0};
olLaunchKernel(Queue, Device, Kernel, &Args, sizeof(Args), &LaunchArgs);
```
In practice, this has a number of hard to workaround downsides.
The example above shows the simplest scenario, where the olLaunchKernel
function is invoked directly by a program. In practice, a language
runtime that uses liboffload may need to construct the argument buffer
at runtime. This may involve allocating a variably-sized buffer and
populating it correctly, accounting for alignment and padding,
with user-provided arguments.
Constructing such a buffer may be non-trivial, or even impossible,
in scenarios where the language runtime doesn't have full argument
metadata containing size or alignment information about arguments.
In addition to the above issue, constructing such a buffer, at runtime
during kernel launch, adds an otherwise avoidable host copy before the
plugin's, or driver's, own copy. Each existing (AMDGPU/HSA, CUDA, Level-Zero)
and anticipated (OpenCL) backend already needs to create its own
GPU-accessible copy of the kernel argument buffer, extending its lifetime
past the kernel launch function.
Another issue with an argument buffer is that its layout is currently
implementation-defined. While they may be compatible in practice,
relying on C/C++ ABI rules, the AMDGPU kernarg and NVPTX .param are
two distinct ABIs.
Allowing the layout to be implementation-defined may require the caller
to construct the buffer differently based on the underlying backend.
In most scenarios, liboffload users should not need to be aware of the
choice of the backend, instead relying on the provided abstractions.
Similarily, the compiler can't optimize code that produces the buffer
if the layout can't be determined at compile-time.
A target independent API accepting a buffer would have to define its own
layout compatible with, or at least convertable to, all existing
and future backends. Documenting the layout format would be a non-trivial
endeavour, likely involving collaboration with various other entities,
such as the Khronos Group, if we wanted for that layout to be natively
supported by all the backends.
The only two backends that can natively implement the existing method of
passing arguments are CUDA and AMDGPU. For OpenCL and Level-Zero, the
plugins would need to internally translate the contiguous buffer into
either calls to clSetKernelArg / zeKernelSetArgumentValue or, in the case
of Level Zero, an array of pointers to arguments, for use in
zeCommandListAppendLaunchKernelWithArguments
(https://oneapi-src.github.io/level-zero-spec/level-zero/latest/core/api.html#zecommandlistappendlaunchkernelwitharguments).
Doing this translation is non-trivial, since it requires the knowledge
of size and alignment of arguments within the buffer. Information which
is not provided through the olLaunchKernel API. There's currently no API in
Level-Zero or OpenCL that can be reliably used to retrieve argument
metadata from a program. Previously, we've attempted to solve this
by adding explicit argument metadata to olLaunchKernel:
```
struct {
char A;
int *B;
short C;
} Args{0, nullptr, 0};
uint64_t ArgSizes[] = {sizeof(char), sizeof(int *), sizeof(short)};
ol_kernel_launch_prop_t Props[] = {
{OL_KERNEL_LAUNCH_PROP_TYPE_SIZE, ArgSizes},
OL_KERNEL_LAUNCH_PROP_END
};
olLaunchKernel(Queue, Device, Kernel, &Args, sizeof(Args), …, Props);
```
This makes the API harder to use, and, most importantly, is
insufficient if we want the API to be target independent. The API above,
with just an array of sizes for each argument, does not allow the plugin
to account for padding inside of the buffer. OpenCL and Level-Zero plugins
would need to see a tightly-packed buffer, while CUDA and AMDGPU plugins need
a C/C++ ABI-padded one.
One solution to retrieving kernel argument metadata could be to read
them from the program image. This is similar to how the AMDGPU plugin
already reads general kernel metadata (readAMDGPUMetaDataFromImage).
However, this would require that the program image format accepted by
liboffload is formalized, and that the OpenCL and Level-Zero plugins
parse and read that format to extract the required metadata. While for
many initial use cases, it may be enough to define the program image
format to be SPIR-V for Level-Zero and OpenCL, it would not allow us to
support creating programs from an existing implementation-specific binary.
Such functionality will be required to support some SYCL features,
such as kernel module interop (where the kernel is created from
an existing backend binary), runtime compilation, and kernel caching.
The equivalent OpenCL API is clCreateProgramWithBinary, where the
documentation states the following:
> The program binary can consist of either or both:
> - Device-specific code and/or,
> - Implementation-specific intermediate representation (IR) which
> will be converted to the device-specific code.
https://registry.khronos.org/OpenCL/sdk/3.0/docs/man/html/clCreateProgramWithBinary.html
Extracting kernel argument metadata from such binaries would require
the generic OpenCL backend to parse vendor-specific binary formats,
one per implementation it intends to support.
This PR changes olLaunchKernel to accept an array of pointers to arguments:
```
void *ArgPtrs[] = {&A, &B, &C};
size_t ArgSizes[] = {sizeof(A), sizeof(B), sizeof(C)};
olLaunchKernel(Queue, Device, Kernel, &LaunchArgs, std::size(ArgPtrs), ArgPtrs, ArgSizes);
```
The newly proposed interface is implementable by existing and anticipated
backends, is familiar to CUDA programmers, eliminates the extraneous
construction of a contiguous arguments buffer, replacing it with constructing
an array of pointers, sidesteps the alignment requirements, does not
require reading program image metadata where it's impractical, and enables
a compliant SYCL implementation to be built on top of it.
The ArgSizes array is required to support OpenCL, which does not have
native support for launching a kernel with an argument pointer array, or
a reliable way of retrieving argument sizes for a kernel.
CUDA and Level-Zero both support accepting an array of pointers to
kernel arguments, through `cuLaunchKernel`
(https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__EXEC.html#group__CUDA__EXEC_1gb8f3dc3031b40da29d5f9a7139e52e15)
and `zeCommandListAppendLaunchKernelWithArguments`
(https://oneapi-src.github.io/level-zero-spec/level-zero/latest/core/api.html#zecommandlistappendlaunchkernelwitharguments)
respectively.
For OpenCL, which requires the kernel arguments to be set separately
from the kernel launch, a potential implementation can extract the
number of arguments from the OpenCL API, and then iterate over the
argument pointer and size arrays:
```
cl_uint num_args = 0;
cl_int err = clGetKernelInfo(kernel, CL_KERNEL_NUM_ARGS,
sizeof(num_args), &num_args, NULL);
for (cl_uint i = 0; i < num_args; ++i) {
clSetKernelArg(kernel, i, ArgSizes[i], ArgPtrs[i]);
}
```
The AMDGPU plugin needs to construct a contiguous buffer using the array
of argument pointers. To do this conversion, we need to have offsets
at which to place the arguments. Here, luckily, as mentioned before,
the AMD plugin already reads kernel metadata from the program image.
The implementation simply retrieves size and offset for each argument
from the kernel image, and then uses it to populate a buffer.
SYCL 2020, and Intel's intel/llvm SYCL implementation, uses 6
fundamental kinds of arguments:
1. USM pointers
2. Buffers
3. Images
4. Samplers
5. Values
6. Local Arguments
Which map to 5 kinds of arguments in Unified Runtime:
```
typedef enum ur_exp_kernel_arg_type_t {
UR_EXP_KERNEL_ARG_TYPE_VALUE = 0,
UR_EXP_KERNEL_ARG_TYPE_POINTER = 1,
UR_EXP_KERNEL_ARG_TYPE_MEM_OBJ = 2,
UR_EXP_KERNEL_ARG_TYPE_LOCAL = 3,
UR_EXP_KERNEL_ARG_TYPE_SAMPLER = 4,
...
} ur_exp_kernel_arg_type_t;
```
The following SYCL 2020 program demonstrates their use:
```
int main() {
sycl::queue q;
int *usm_ptr = sycl::malloc_host<int>(1, q);
*usm_ptr = -1;
int host_data[1] = {5};
sycl::buffer<int, 1> buf(host_data, sycl::range<1>(1));
sycl::image<2> img(sycl::image_channel_order::rgba,
sycl::image_channel_type::fp32,
sycl::range<2>(4, 4));
sycl::sampler samp(
sycl::coordinate_normalization_mode::unnormalized,
sycl::addressing_mode::clamp,
sycl::filtering_mode::nearest);
int val = 42;
q.submit([&](sycl::handler &h) {
auto buf_acc =
buf.get_access<sycl::access::mode::read_write>(h);
auto img_acc =
img.get_access<sycl::float4,
sycl::access::mode::read>(h);
sycl::local_accessor<int, 1> local_mem(
sycl::range<1>(64), h);
sycl::local_accessor<int, 1> local_mem2(
sycl::range<1>(64), h);
h.parallel_for(sycl::nd_range<1>(1, 1),
[=](sycl::nd_item<1> item) {
local_mem[0] = val;
local_mem2[0] = val + 1;
buf_acc[0] = buf_acc[0] + local_mem[0] + local_mem2[0];
usm_ptr[0] = buf_acc[0];
// just to keep the args
(void)img_acc;
(void)samp;
});
});
q.wait();
int expected = 5 + 42 + 43;
printf("usm_ptr[0] = %d (expected %d)\n", *usm_ptr, expected);
sycl::free(usm_ptr, q);
return 0;
}
```
The kernel launch uses the following UR call (simplified for
legibility):
```
urEnqueueKernelLaunchWithArgsExp(
.hQueue = 0xca2f090,
.hKernel = 0xc7640f0,
.workDim = 1,
.pGlobalWorkOffset = nullptr,
.pGlobalWorkSize = 0xc6e1c38 (1),
.pLocalWorkSize = 0xc6e1c50 (1),
.numArgs = 8,
.pArgs = 0xc754710 {
{.type = UR_EXP_KERNEL_ARG_TYPE_LOCAL, .index = 0, .size = 256,
.value = (union ur_exp_kernel_arg_value_t){<unknown>}},
{.type = UR_EXP_KERNEL_ARG_TYPE_VALUE, .index = 1, .size = 4,
.value = (union ur_exp_kernel_arg_value_t){.value = 0xc6e23e8}},
{.type = UR_EXP_KERNEL_ARG_TYPE_LOCAL, .index = 2, .size = 256,
.value = (union ur_exp_kernel_arg_value_t){<unknown>}},
{.type = UR_EXP_KERNEL_ARG_TYPE_MEM_OBJ, .index = 3, .size = 8,
.value = (union ur_exp_kernel_arg_value_t){
.memObjTuple = (struct ur_exp_kernel_arg_mem_obj_tuple_t){
.hMem = 0xc6e60f0, .flags = UR_MEM_FLAG_READ_WRITE}}},
{.type = UR_EXP_KERNEL_ARG_TYPE_VALUE, .index = 4, .size = 8,
.value = (union ur_exp_kernel_arg_value_t){.value = 0xc6966f0}},
{.type = UR_EXP_KERNEL_ARG_TYPE_POINTER, .index = 5, .size = 8,
.value = (union ur_exp_kernel_arg_value_t){.pointer = ...}},
{.type = UR_EXP_KERNEL_ARG_TYPE_MEM_OBJ, .index = 6, .size = 8,
.value = (union ur_exp_kernel_arg_value_t){
.memObjTuple = (struct ur_exp_kernel_arg_mem_obj_tuple_t){
.hMem = 0xc6ebac0, .flags = UR_MEM_FLAG_READ_ONLY}}},
{.type = UR_EXP_KERNEL_ARG_TYPE_SAMPLER, .index = 7, .size = 8,
.value = (union ur_exp_kernel_arg_value_t){.sampler = 0xca50750}}
},
...
) -> UR_RESULT_SUCCESS;
```
In order to support a compliant SYCL 2020 specification, liboffload
launch kernel API needs to be able to support all these kinds of
arguments. The proposed API makes this straightforward. To launch the
kernel above, the language runtime could use the liboffload API in the
following way:
```
ol_kernel_launch_size_args_t LaunchArgs;
LaunchArgs.Dimensions = 1;
LaunchArgs.GroupSize = {1, 1, 1};
LaunchArgs.NumGroups = {1, 1, 1};
// CUDA, HSA, Level Zero, and OpenCL have a differing model of
// shared / local memory. The lowest common denominator is setting
// only a single shared memory block to contain all memory required
// by SYCL's local accessors, with the language runtime additionally
// setting offsets into that block as value arguments.
// Level Zero doesn't currently implement this feature.
// The SPIR-V kernel is expected to declare a single SLM argument
// at the last position when this field is set.
LaunchArgs.DynSharedMemory = 512;
SmallVector<void *, 8> ArgPtrs;
SmallVector<size_t, 8> ArgSizes;
// [0] local_mem: first local accessor, offset 0 into the shared
// memory block.
uint32_t LocalOffset0 = 0;
ArgPtrs.push_back(&LocalOffset0);
ArgSizes.push_back(sizeof(uint32_t));
// [1] val: plain value argument (int val = 42).
int Val = 42;
ArgPtrs.push_back(&Val);
ArgSizes.push_back(sizeof(int));
// [2] local_mem2: second local accessor, offset 256 into the shared
// memory block.
uint32_t LocalOffset1 = 256;
ArgPtrs.push_back(&LocalOffset1);
ArgSizes.push_back(sizeof(uint32_t));
// [3] buf_acc: buffer accessor. Buffers are not supported in
// liboffload. The language runtimes can use USM memory instead.
void *BufPtr = /* usm pointer obtained from olMemAlloc */
ArgPtrs.push_back(&BufPtr);
ArgSizes.push_back(sizeof(void *));
// [4] buffer accessor metadata, passed by value.
size_t buf_accessor_md = 0;
ArgPtrs.push_back(&buf_accessor_md);
ArgSizes.push_back(sizeof(size_t));
// [5] usm_ptr: USM pointer argument.
void *UsmPtr = /* usm pointer obtained from olMemAlloc */
ArgPtrs.push_back(&UsmPtr);
ArgSizes.push_back(sizeof(void *));
// All argument kinds are ultimately just pointers to values,
// including images and samplers.
// [6] img_acc: Images can be opaque handles in
// liboffload. This API doesn't exist yet.
ol_image_handle_t Image = olCreateImage(...);
ArgPtrs.push_back(olImageGetArgPtr(Image));
ArgSizes.push_back(olImageGetArgSize(Image));
// [7] sampler: Samplers can be opaque handles in liboffload.
// This API doesn't exist yet.
ol_sampler_handle_t Sampler = olCreateSampler(...);
ArgPtrs.push_back(olSamplerGetArgPtr(Sampler));
// Sampler size can be different depending on the platform
ArgSizes.push_back(olSamplerGetArgSize(Sampler));
olLaunchKernel(Queue, Device, Kernel, &LaunchArgs,
ArgPtrs.size(), ArgPtrs.data(), ArgSizes.data());
```
This API provides flexibility and target independence, while allowing
efficient kernel dispatch from all backends we expect to support.
---
.../llvm/Frontend/Offloading/Utility.h | 14 ++
llvm/lib/Frontend/Offloading/Utility.cpp | 32 +++
.../tools/llvm-gpu-loader/llvm-gpu-loader.cpp | 44 ++++-
llvm/tools/llvm-gpu-loader/llvm-gpu-loader.h | 5 +-
offload/include/Shared/APITypes.h | 5 +-
offload/liboffload/API/Kernel.td | 20 +-
offload/liboffload/src/OffloadImpl.cpp | 20 +-
offload/plugins-nextgen/amdgpu/src/rtl.cpp | 35 +++-
.../common/src/PluginInterface.cpp | 7 +-
offload/plugins-nextgen/cuda/src/rtl.cpp | 38 +++-
.../level_zero/dynamic_l0/L0DynWrapper.cpp | 1 +
.../level_zero/dynamic_l0/level_zero/ze_api.h | 19 ++
.../level_zero/include/L0Kernel.h | 5 +-
.../level_zero/src/L0Kernel.cpp | 131 +++++++-----
.../include/mathtest/DeviceContext.hpp | 20 +-
.../Conformance/include/mathtest/Support.hpp | 27 ---
.../Conformance/lib/DeviceContext.cpp | 11 +-
.../OffloadAPI/device_code/CMakeLists.txt | 2 +
.../OffloadAPI/device_code/composite.cpp | 10 +
.../OffloadAPI/device_code/multiargs.cpp | 4 +-
.../event/olGetEventElapsedTime.cpp | 9 +-
.../OffloadAPI/kernel/olLaunchKernel.cpp | 187 ++++++++++++------
.../kernel/olLaunchKernelCooperative.cpp | 16 +-
.../unittests/OffloadAPI/memory/olMemcpy.cpp | 13 +-
.../OffloadAPI/queue/olLaunchHostFunction.cpp | 9 +-
.../OffloadAPI/queue/olWaitEvents.cpp | 39 ++--
26 files changed, 479 insertions(+), 244 deletions(-)
create mode 100644 offload/unittests/OffloadAPI/device_code/composite.cpp
diff --git a/llvm/include/llvm/Frontend/Offloading/Utility.h b/llvm/include/llvm/Frontend/Offloading/Utility.h
index a1a8a02916265..812bbb3234246 100644
--- a/llvm/include/llvm/Frontend/Offloading/Utility.h
+++ b/llvm/include/llvm/Frontend/Offloading/Utility.h
@@ -15,6 +15,7 @@
#include "llvm/ADT/ArrayRef.h"
#include "llvm/ADT/SmallString.h"
+#include "llvm/ADT/SmallVector.h"
#include "llvm/ADT/StringMap.h"
#include "llvm/ADT/StringRef.h"
#include "llvm/IR/Module.h"
@@ -121,6 +122,16 @@ namespace amdgpu {
LLVM_ABI bool isImageCompatibleWithEnv(StringRef ImageArch, uint32_t ImageFlags,
StringRef EnvTargetID);
+/// Struct for holding AMDGPU Kernel Argument Metadata, see:
+/// https://llvm.org/docs/AMDGPUUsage.html#amdgpu-amdhsa-code-object-kernel-argument-metadata-map-table-v3
+struct AMDGPUKernelArgMetaData {
+ /// Kernel argument offset in bytes. The offset must be a multiple
+ /// of the alignment required by the argument.
+ uint32_t Offset = 0;
+ /// Kernel argument size in bytes.
+ uint32_t Size = 0;
+};
+
/// Struct for holding metadata related to AMDGPU kernels, for more information
/// about the metadata and its meaning see:
/// https://llvm.org/docs/AMDGPUUsage.html#code-object-v3
@@ -154,6 +165,9 @@ struct AMDGPUKernelMetaData {
uint32_t WavefrontSize = KInvalidValue;
/// Maximum flat work-group size supported by the kernel in work-items.
uint32_t MaxFlatWorkgroupSize = KInvalidValue;
+ /// Per-argument offset and size, read from the ".args" array in code object
+ /// metadata. Includes only explicit user arguments.
+ SmallVector<AMDGPUKernelArgMetaData, 8> ExplicitArgMDs;
};
/// Reads AMDGPU specific metadata from the ELF file and propagates the
diff --git a/llvm/lib/Frontend/Offloading/Utility.cpp b/llvm/lib/Frontend/Offloading/Utility.cpp
index 909af0237910f..6648e8c3d1d5a 100644
--- a/llvm/lib/Frontend/Offloading/Utility.cpp
+++ b/llvm/lib/Frontend/Offloading/Utility.cpp
@@ -315,6 +315,38 @@ class KernelInfoReader {
KernelData.WavefrontSize = V.second.getUInt();
} else if (IsKey(V.first, ".max_flat_workgroup_size")) {
KernelData.MaxFlatWorkgroupSize = V.second.getUInt();
+ } else if (IsKey(V.first, ".args")) {
+ auto ArgsArray = V.second.getArray();
+ for (auto ArgIt = ArgsArray.begin(), ArgEnd = ArgsArray.end();
+ ArgIt != ArgEnd; ++ArgIt) {
+ auto ArgMap = ArgIt->getMap();
+
+ // Skip hidden arguments
+ auto VKIt = ArgMap.find(".value_kind");
+ if (VKIt != ArgMap.end() &&
+ VKIt->second.getString().starts_with("hidden_"))
+ continue;
+
+ auto OffsetIt = ArgMap.find(".offset");
+ // TODO: should these be asserts?
+ // if .offset or .size isn't found, that means the kernel
+ // is malformed.
+ if (OffsetIt == ArgMap.end())
+ return createStringError(
+ inconvertibleErrorCode(),
+ "Missing required .offset key in kernel argument metadata map");
+ auto SizeIt = ArgMap.find(".size");
+ if (SizeIt == ArgMap.end())
+ return createStringError(
+ inconvertibleErrorCode(),
+ "Missing required .size key in kernel argument metadata map");
+
+ amdgpu::AMDGPUKernelArgMetaData ArgMD;
+ ArgMD.Offset = OffsetIt->second.getUInt();
+ ArgMD.Size = SizeIt->second.getUInt();
+
+ KernelData.ExplicitArgMDs.push_back(ArgMD);
+ }
}
return Error::success();
diff --git a/llvm/tools/llvm-gpu-loader/llvm-gpu-loader.cpp b/llvm/tools/llvm-gpu-loader/llvm-gpu-loader.cpp
index 40fed374e0676..341ae82c60c60 100644
--- a/llvm/tools/llvm-gpu-loader/llvm-gpu-loader.cpp
+++ b/llvm/tools/llvm-gpu-loader/llvm-gpu-loader.cpp
@@ -183,16 +183,46 @@ ol_device_handle_t getHostDevice() {
return Device;
}
-template <typename Args>
-void launchKernel(ol_queue_handle_t Queue, ol_device_handle_t Device,
- ol_program_handle_t Program, const char *Name,
- ol_kernel_launch_size_args_t LaunchArgs, Args &KernelArgs) {
+static void launchKernelImpl(ol_queue_handle_t Queue, ol_device_handle_t Device,
+ ol_program_handle_t Program, const char *Name,
+ ol_kernel_launch_size_args_t LaunchArgs,
+ size_t NumArgs, void **ArgPtrs,
+ const size_t *ArgSizes) {
ol_symbol_handle_t Kernel;
OFFLOAD_ERR(olGetSymbol(Program, Name, OL_SYMBOL_KIND_KERNEL, &Kernel));
- OFFLOAD_ERR(olLaunchKernel(Queue, Device, Kernel, &KernelArgs,
- std::is_empty_v<Args> ? 0 : sizeof(Args),
- &LaunchArgs, /*Properties=*/nullptr));
+ OFFLOAD_ERR(olLaunchKernel(Queue, Device, Kernel, &LaunchArgs, nullptr,
+ NumArgs, ArgPtrs, ArgSizes));
+}
+
+void launchKernel(ol_queue_handle_t Queue, ol_device_handle_t Device,
+ ol_program_handle_t Program, const char *Name,
+ ol_kernel_launch_size_args_t LaunchArgs,
+ BeginArgs &KernelArgs) {
+ void *ArgPtrs[] = {&KernelArgs.Argc, &KernelArgs.Argv, &KernelArgs.Envp};
+ size_t ArgSizes[] = {sizeof(KernelArgs.Argc), sizeof(KernelArgs.Argv),
+ sizeof(KernelArgs.Envp)};
+ launchKernelImpl(Queue, Device, Program, Name, LaunchArgs, 3, ArgPtrs,
+ ArgSizes);
+}
+
+void launchKernel(ol_queue_handle_t Queue, ol_device_handle_t Device,
+ ol_program_handle_t Program, const char *Name,
+ ol_kernel_launch_size_args_t LaunchArgs,
+ StartArgs &KernelArgs) {
+ void *ArgPtrs[] = {&KernelArgs.Argc, &KernelArgs.Argv, &KernelArgs.Envp,
+ &KernelArgs.Ret};
+ size_t ArgSizes[] = {sizeof(KernelArgs.Argc), sizeof(KernelArgs.Argv),
+ sizeof(KernelArgs.Envp), sizeof(KernelArgs.Ret)};
+ launchKernelImpl(Queue, Device, Program, Name, LaunchArgs, 4, ArgPtrs,
+ ArgSizes);
+}
+
+void launchKernel(ol_queue_handle_t Queue, ol_device_handle_t Device,
+ ol_program_handle_t Program, const char *Name,
+ ol_kernel_launch_size_args_t LaunchArgs, EndArgs &) {
+ launchKernelImpl(Queue, Device, Program, Name, LaunchArgs, 0, nullptr,
+ nullptr);
}
int main(int argc, const char **argv, const char **envp) {
diff --git a/llvm/tools/llvm-gpu-loader/llvm-gpu-loader.h b/llvm/tools/llvm-gpu-loader/llvm-gpu-loader.h
index d2fafe8984dce..729ca49b7b32c 100644
--- a/llvm/tools/llvm-gpu-loader/llvm-gpu-loader.h
+++ b/llvm/tools/llvm-gpu-loader/llvm-gpu-loader.h
@@ -133,9 +133,10 @@ ol_result_t (*olGetSymbol)(ol_program_handle_t Program, const char *Name,
ol_result_t (*olLaunchKernel)(
ol_queue_handle_t Queue, ol_device_handle_t Device,
- ol_symbol_handle_t Kernel, const void *ArgumentsData, size_t ArgumentsSize,
+ ol_symbol_handle_t Kernel,
const ol_kernel_launch_size_args_t *LaunchSizeArgs,
- const ol_kernel_launch_prop_t *Properties);
+ const ol_kernel_launch_prop_t *Properties, size_t NumArgs, void **ArgPtrs,
+ const size_t *ArgSizes);
ol_result_t (*olCreateQueue)(ol_device_handle_t Device,
ol_queue_handle_t *Queue);
diff --git a/offload/include/Shared/APITypes.h b/offload/include/Shared/APITypes.h
index 212fb285030fb..d1326b4337551 100644
--- a/offload/include/Shared/APITypes.h
+++ b/offload/include/Shared/APITypes.h
@@ -105,8 +105,9 @@ struct KernelArgsTy {
uint64_t IsCUDA : 1; // Was this kernel spawned via CUDA.
uint64_t DynCGroupMemFallback : 2; // The fallback for dynamic cgroup mem.
uint64_t Cooperative : 1; // Was this kernel spawned as cooperative.
- uint64_t Unused : 59;
- } Flags = {0, 0, 0, 0, 0};
+ uint64_t IsPtrArgs : 1; // Arguments are laid out as an array of pointers.
+ uint64_t Unused : 58;
+ } Flags = {0, 0, 0, 0, 0, 0};
// User-requested number of blocks (for x,y,z dimension).
uint32_t UserNumBlocks[3] = {0, 0, 0};
// User-requested number of threads (for x,y,z dimension).
diff --git a/offload/liboffload/API/Kernel.td b/offload/liboffload/API/Kernel.td
index bc0dcb565f3f3..06778105e0194 100644
--- a/offload/liboffload/API/Kernel.td
+++ b/offload/liboffload/API/Kernel.td
@@ -42,22 +42,30 @@ def ol_kernel_launch_prop_t : Struct {
}
def olLaunchKernel : Function {
- let desc = "Enqueue a kernel launch with the specified size and parameters.";
+ let desc = "Enqueue a kernel launch with arguments specified as an array of pointers and sizes.";
let details = [
"If a queue is not specified, kernel execution happens synchronously",
- "ArgumentsData may be set to NULL (to indicate no parameters)"
+ "Each element of ArgPtrs points to the value of the corresponding kernel argument",
+ "Each element of ArgSizes specifies the size in bytes of the corresponding argument",
+ "NumArgs must match the number of arguments expected by the kernel",
+ "ArgPtrs and ArgSizes must both be NULL (when NumArgs == 0) or both be non-NULL"
];
let params = [
Param<"ol_queue_handle_t", "Queue", "handle of the queue", PARAM_IN_OPTIONAL>,
Param<"ol_device_handle_t", "Device", "handle of the device to execute on", PARAM_IN>,
Param<"ol_symbol_handle_t", "Kernel", "handle of the kernel", PARAM_IN>,
- Param<"const void*", "ArgumentsData", "pointer to the kernel argument struct", PARAM_IN_OPTIONAL>,
- Param<"size_t", "ArgumentsSize", "size of the kernel argument struct", PARAM_IN>,
Param<"const ol_kernel_launch_size_args_t*", "LaunchSizeArgs", "pointer to the struct containing launch size parameters", PARAM_IN>,
- Param<"const ol_kernel_launch_prop_t *", "Properties", "Array of optional properties, last element must be OL_KERNEL_LAUNCH_PROP_END", PARAM_IN_OPTIONAL>
+ Param<"const ol_kernel_launch_prop_t *", "Properties", "Array of optional properties, last element must be OL_KERNEL_LAUNCH_PROP_END", PARAM_IN_OPTIONAL>,
+ Param<"size_t", "NumArgs", "number of kernel arguments", PARAM_IN>,
+ Param<"void**", "ArgPtrs", "array of pointers, each pointing to a kernel argument value", PARAM_IN_OPTIONAL>,
+ Param<"const size_t*", "ArgSizes", "array of kernel argument sizes", PARAM_IN_OPTIONAL>,
];
let returns = [
- Return<"OL_ERRC_INVALID_ARGUMENT", ["`ArgumentsSize > 0 && ArgumentsData == NULL`"]>,
+ Return<"OL_ERRC_INVALID_ARGUMENT", [
+ "`(ArgPtrs == NULL) != (ArgSizes == NULL)`",
+ "`NumArgs > 0 && (ArgPtrs == NULL || ArgSizes == NULL)`",
+ "NumArgs does not match the number of arguments expected by the kernel",
+ ]>,
Return<"OL_ERRC_INVALID_DEVICE", ["If Queue is non-null but does not belong to Device"]>,
Return<"OL_ERRC_SYMBOL_KIND", ["The provided symbol is not a kernel"]>,
];
diff --git a/offload/liboffload/src/OffloadImpl.cpp b/offload/liboffload/src/OffloadImpl.cpp
index 66fcbbc264ab4..6e174fc77c2af 100644
--- a/offload/liboffload/src/OffloadImpl.cpp
+++ b/offload/liboffload/src/OffloadImpl.cpp
@@ -1098,10 +1098,11 @@ Error olGetKernelMaxCooperativeGroupCount_impl(
}
Error olLaunchKernel_impl(ol_queue_handle_t Queue, ol_device_handle_t Device,
- ol_symbol_handle_t Kernel, const void *ArgumentsData,
- size_t ArgumentsSize,
+ ol_symbol_handle_t Kernel,
const ol_kernel_launch_size_args_t *LaunchSizeArgs,
- const ol_kernel_launch_prop_t *Properties) {
+ const ol_kernel_launch_prop_t *Properties,
+ size_t NumArgs, void **ArgPtrs,
+ const size_t *ArgSizes) {
auto *DeviceImpl = Device->Device;
if (Queue && Device != Queue->Device) {
return createOffloadError(
@@ -1115,6 +1116,8 @@ Error olLaunchKernel_impl(ol_queue_handle_t Queue, ol_device_handle_t Device,
auto *QueueImpl = Queue ? Queue->AsyncInfo : nullptr;
KernelArgsTy LaunchArgs{};
+ // TODO: change LaunchArgs.NumArgs to size_t to match the API type
+ LaunchArgs.NumArgs = static_cast<uint32_t>(NumArgs);
LaunchArgs.UserNumBlocks[0] = LaunchSizeArgs->NumGroups.x;
LaunchArgs.UserNumBlocks[1] = LaunchSizeArgs->NumGroups.y;
LaunchArgs.UserNumBlocks[2] = LaunchSizeArgs->NumGroups.z;
@@ -1138,12 +1141,11 @@ Error olLaunchKernel_impl(ol_queue_handle_t Queue, ol_device_handle_t Device,
}
AsyncInfoWrapperTy AsyncInfoWrapper(*DeviceImpl, QueueImpl);
- KernelLaunchParamsTy Params;
- Params.Data = const_cast<void *>(ArgumentsData);
- Params.Size = ArgumentsSize;
- LaunchArgs.ArgPtrs = reinterpret_cast<void **>(&Params);
- // Don't do anything with pointer indirection; use arg data as-is
- LaunchArgs.Flags.IsCUDA = true;
+ LaunchArgs.ArgPtrs = ArgPtrs;
+ // TODO: Either change ArgSizes to const size_t * or add a new variable.
+ LaunchArgs.ArgSizes =
+ reinterpret_cast<int64_t *>(const_cast<size_t *>(ArgSizes));
+ LaunchArgs.Flags.IsPtrArgs = true;
auto *KernelImpl = std::get<GenericKernelTy *>(Kernel->PluginImpl);
auto Err = KernelImpl->launch(*DeviceImpl, LaunchArgs.ArgPtrs, nullptr,
diff --git a/offload/plugins-nextgen/amdgpu/src/rtl.cpp b/offload/plugins-nextgen/amdgpu/src/rtl.cpp
index 9995da9c2aa2b..ebf3573efd87f 100644
--- a/offload/plugins-nextgen/amdgpu/src/rtl.cpp
+++ b/offload/plugins-nextgen/amdgpu/src/rtl.cpp
@@ -4191,11 +4191,32 @@ Error AMDGPUKernelTy::launchImpl(GenericDeviceTy &GenericDevice,
if (auto Err = GenericDevice.getDeviceStackSize(StackSize))
return Err;
- // Copy the explicit arguments.
- // TODO: We should expose the args memory manager alloc to the common part as
- // alternative to copying them twice.
- if (LaunchParams.Size)
- std::memcpy(AllArgs, LaunchParams.Data, LaunchParams.Size);
+ // Copy explicit arguments.
+ size_t ExplicitEnd = 0;
+ if (KernelArgs.Flags.IsPtrArgs) {
+ if (KernelArgs.ArgPtrs) { // may be null
+ const auto &ArgMDs = KernelInfo->ExplicitArgMDs;
+
+ if (KernelArgs.NumArgs != ArgMDs.size())
+ return Plugin::error(
+ ErrorCode::INVALID_ARGUMENT,
+ "Number of arguments (%u) does not match the number of arguments "
+ "expected by the kernel (%zu)",
+ KernelArgs.NumArgs, ArgMDs.size());
+
+ for (size_t I = 0; I < ArgMDs.size(); I++)
+ std::memcpy(utils::advancePtr(AllArgs, ArgMDs[I].Offset),
+ KernelArgs.ArgPtrs[I], ArgMDs[I].Size);
+
+ ExplicitEnd = ArgMDs.back().Offset + ArgMDs.back().Size;
+ }
+ } else {
+ // TODO: We should expose the args memory manager alloc to the common part
+ // as alternative to copying them twice.
+ if (LaunchParams.Size)
+ std::memcpy(AllArgs, LaunchParams.Data, LaunchParams.Size);
+ ExplicitEnd = LaunchParams.Size;
+ }
AMDGPUDeviceTy &AMDGPUDevice = static_cast<AMDGPUDeviceTy &>(GenericDevice);
@@ -4203,8 +4224,8 @@ Error AMDGPUKernelTy::launchImpl(GenericDeviceTy &GenericDevice,
if (auto Err = AMDGPUDevice.getStream(AsyncInfoWrapper, Stream))
return Err;
- uint64_t ImplArgsOffset = llvm::alignTo(
- LaunchParams.Size, alignof(hsa_utils::AMDGPUImplicitArgsTy));
+ uint64_t ImplArgsOffset =
+ llvm::alignTo(ExplicitEnd, alignof(hsa_utils::AMDGPUImplicitArgsTy));
if (ArgsSize > ImplArgsOffset) {
hsa_utils::AMDGPUImplicitArgsTy *ImplArgs =
reinterpret_cast<hsa_utils::AMDGPUImplicitArgsTy *>(
diff --git a/offload/plugins-nextgen/common/src/PluginInterface.cpp b/offload/plugins-nextgen/common/src/PluginInterface.cpp
index 4379ebd250794..d43f9268d3cc3 100644
--- a/offload/plugins-nextgen/common/src/PluginInterface.cpp
+++ b/offload/plugins-nextgen/common/src/PluginInterface.cpp
@@ -277,13 +277,14 @@ Error GenericKernelTy::launch(GenericDeviceTy &GenericDevice, void **ArgPtrs,
if (!KernelLaunchEnvOrErr)
return KernelLaunchEnvOrErr.takeError();
- KernelLaunchParamsTy LaunchParams;
-
// Kernel languages don't use indirection.
+ // IsPtrArgs bypasses LaunchParms entirely,
+ // plugins read KernelArgs.ArgPtrs/ArgSizes directly.
+ KernelLaunchParamsTy LaunchParams;
if (KernelArgs.Flags.IsCUDA) {
LaunchParams =
*reinterpret_cast<KernelLaunchParamsTy *>(KernelArgs.ArgPtrs);
- } else {
+ } else if (!KernelArgs.Flags.IsPtrArgs) {
LaunchParams =
prepareArgs(GenericDevice, ArgPtrs, ArgOffsets, KernelArgs.NumArgs,
Args, Ptrs, *KernelLaunchEnvOrErr, KernelArgs.Version);
diff --git a/offload/plugins-nextgen/cuda/src/rtl.cpp b/offload/plugins-nextgen/cuda/src/rtl.cpp
index 698bd4544705f..564a4c0e3fd56 100644
--- a/offload/plugins-nextgen/cuda/src/rtl.cpp
+++ b/offload/plugins-nextgen/cuda/src/rtl.cpp
@@ -171,11 +171,14 @@ struct CUDAKernelTy : public GenericKernelTy {
size_t Arg = 0;
ArgsSize = 0;
+ NumArgs = 0;
// Find the last argument to know the total size of the arguments.
while ((Res = cuFuncGetParamInfo(Func, Arg++, &ArgOffset, &ArgSize)) ==
- CUDA_SUCCESS)
+ CUDA_SUCCESS) {
ArgsSize = ArgOffset + ArgSize;
+ ++NumArgs;
+ }
if (Res != CUDA_ERROR_INVALID_VALUE)
return Plugin::check(Res, "error in cuFuncGetParamInfo: %s");
@@ -190,6 +193,13 @@ struct CUDAKernelTy : public GenericKernelTy {
/// The size of the kernel arguments.
size_t ArgsSize;
+
+ /// The number of kernel arguments.
+ size_t NumArgs = 0;
+
+public:
+ /// Return the number of kernel arguments reported by the driver.
+ size_t getNumArgs() const { return NumArgs; }
};
/// Class wrapping a CUDA stream reference. These are the objects handled by the
@@ -1472,17 +1482,30 @@ Error CUDAKernelTy::launchImpl(GenericDeviceTy &GenericDevice,
AsyncInfoWrapperTy &AsyncInfoWrapper) const {
CUDADeviceTy &CUDADevice = static_cast<CUDADeviceTy &>(GenericDevice);
- // The args size passed in LaunchParams may have tail padding, which is not
- // accepted by the CUDA driver.
- if (ArgsSize > LaunchParams.Size)
- return Plugin::error(ErrorCode::INVALID_ARGUMENT,
- "mismatch in kernel arguments");
+ void **KernelParams = nullptr;
+ if (KernelArgs.Flags.IsPtrArgs) {
+ if (KernelArgs.NumArgs != NumArgs)
+ return Plugin::error(
+ ErrorCode::INVALID_ARGUMENT,
+ "Number of arguments (%u) does not match the number of arguments "
+ "expected by the kernel (%zu)",
+ KernelArgs.NumArgs, NumArgs);
+
+ KernelParams = KernelArgs.ArgPtrs;
+ } else {
+ // The args size passed in LaunchParams may have tail padding,
+ // which is not accepted by the CUDA driver.
+ if (ArgsSize > LaunchParams.Size)
+ return Plugin::error(ErrorCode::INVALID_ARGUMENT,
+ "mismatch in kernel arguments");
+ }
CUstream Stream;
if (auto Err = CUDADevice.getStream(AsyncInfoWrapper, Stream))
return Err;
size_t ConfigArgsSize = ArgsSize;
+ // valid for buffer passed through LaunchParams
void *Config[] = {CU_LAUNCH_PARAM_BUFFER_POINTER, LaunchParams.Data,
CU_LAUNCH_PARAM_BUFFER_SIZE,
reinterpret_cast<void *>(&ConfigArgsSize),
@@ -1514,7 +1537,8 @@ Error CUDAKernelTy::launchImpl(GenericDeviceTy &GenericDevice,
DynBlockMemSize, Stream,
&CoopAttr, 1};
- CUresult Res = cuLaunchKernelEx(&LaunchConfig, Func, nullptr, Config);
+ CUresult Res = cuLaunchKernelEx(&LaunchConfig, Func, KernelParams,
+ KernelParams ? nullptr : Config);
// Register a callback to indicate when the kernel is complete.
if (GenericDevice.getRPCServer())
diff --git a/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp b/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
index 8643897facacd..67cbdca3ab82b 100644
--- a/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
+++ b/offload/plugins-nextgen/level_zero/dynamic_l0/L0DynWrapper.cpp
@@ -29,6 +29,7 @@ DLWRAP(zeModuleDestroy, 1)
DLWRAP(zeCommandListAppendBarrier, 4)
DLWRAP(zeCommandListAppendLaunchKernel, 6)
DLWRAP(zeCommandListAppendLaunchCooperativeKernel, 6)
+DLWRAP(zeCommandListAppendLaunchKernelWithArguments, 9)
DLWRAP(zeCommandListAppendMemoryCopy, 7)
DLWRAP(zeCommandListAppendMemoryCopyRegion, 12)
DLWRAP(zeCommandListAppendMemoryFill, 8)
diff --git a/offload/plugins-nextgen/level_zero/dynamic_l0/level_zero/ze_api.h b/offload/plugins-nextgen/level_zero/dynamic_l0/level_zero/ze_api.h
index 17e35a37f116e..ebbdf565b9ea9 100644
--- a/offload/plugins-nextgen/level_zero/dynamic_l0/level_zero/ze_api.h
+++ b/offload/plugins-nextgen/level_zero/dynamic_l0/level_zero/ze_api.h
@@ -167,6 +167,7 @@ typedef enum _ze_structure_type_t {
ZE_STRUCTURE_TYPE_KERNEL_PROPERTIES = 0x1e,
ZE_STRUCTURE_TYPE_KERNEL_PREFERRED_GROUP_SIZE_PROPERTIES = 0x21,
ZE_STRUCTURE_TYPE_DEVICE_IP_VERSION_EXT = 0x1000f,
+ ZE_STRUCTURE_TYPE_COMMAND_LIST_APPEND_PARAM_COOPERATIVE_DESC = 0x00020036,
ZE_STRUCTURE_TYPE_RELAXED_ALLOCATION_LIMITS_EXP_DESC = 0x00020001,
ZE_STRUCTURE_TYPE_FORCE_UINT32 = 0x7fffffff
} ze_structure_type_t;
@@ -489,6 +490,12 @@ typedef struct _ze_group_count_t {
uint32_t groupCountZ;
} ze_group_count_t;
+typedef struct _ze_group_size_t {
+ uint32_t groupSizeX;
+ uint32_t groupSizeY;
+ uint32_t groupSizeZ;
+} ze_group_size_t;
+
/* Memory allocation properties */
typedef struct _ze_memory_allocation_properties_t {
ze_structure_type_t stype;
@@ -731,6 +738,18 @@ ZE_APIEXPORT ze_result_t ZE_APICALL zeCommandListAppendLaunchCooperativeKernel(
ze_command_list_handle_t hCommandList, ze_kernel_handle_t hKernel,
const ze_group_count_t *pLaunchFuncArgs, ze_event_handle_t hSignalEvent,
uint32_t numWaitEvents, ze_event_handle_t *phWaitEvents);
+typedef struct _ze_command_list_append_launch_kernel_param_cooperative_desc_t {
+ ze_structure_type_t stype;
+ const void *pNext;
+ ze_bool_t isCooperative;
+} ze_command_list_append_launch_kernel_param_cooperative_desc_t;
+
+ZE_APIEXPORT ze_result_t ZE_APICALL
+zeCommandListAppendLaunchKernelWithArguments(
+ ze_command_list_handle_t hCommandList, ze_kernel_handle_t hKernel,
+ const ze_group_count_t groupCounts, const ze_group_size_t groupSizes,
+ void **pArguments, const void *pNext, ze_event_handle_t hSignalEvent,
+ uint32_t numWaitEvents, ze_event_handle_t *phWaitEvents);
ZE_APIEXPORT ze_result_t ZE_APICALL zeCommandListAppendMemoryCopy(
ze_command_list_handle_t hCommandList, void *dstptr, const void *srcptr,
size_t size, ze_event_handle_t hSignalEvent, uint32_t numWaitEvents,
diff --git a/offload/plugins-nextgen/level_zero/include/L0Kernel.h b/offload/plugins-nextgen/level_zero/include/L0Kernel.h
index 4d6ba0ee392b4..f7134b5c5282a 100644
--- a/offload/plugins-nextgen/level_zero/include/L0Kernel.h
+++ b/offload/plugins-nextgen/level_zero/include/L0Kernel.h
@@ -63,8 +63,9 @@ class L0KernelTy : public GenericKernelTy {
Error buildKernel(L0ProgramTy &Program);
Error readKernelProperties(L0ProgramTy &Program);
- Error setKernelGroups(L0DeviceTy &l0Device, L0LaunchEnvTy &KEnv,
- uint32_t NumThreads[3], uint32_t NumBlocks[3]) const;
+ ze_group_size_t createKernelGroups(L0DeviceTy &l0Device, L0LaunchEnvTy &KEnv,
+ uint32_t NumThreads[3],
+ uint32_t NumBlocks[3]) const;
Error setIndirectFlags(L0DeviceTy &l0Device, L0LaunchEnvTy &KEnv) const;
public:
diff --git a/offload/plugins-nextgen/level_zero/src/L0Kernel.cpp b/offload/plugins-nextgen/level_zero/src/L0Kernel.cpp
index dc884e3b753dc..c390ae7e4e92e 100644
--- a/offload/plugins-nextgen/level_zero/src/L0Kernel.cpp
+++ b/offload/plugins-nextgen/level_zero/src/L0Kernel.cpp
@@ -15,8 +15,6 @@
#include "L0Plugin.h"
#include "L0Program.h"
-#include "llvm/ADT/ScopeExit.h"
-
namespace llvm::omp::target::plugin {
Error L0KernelTy::readKernelProperties(L0ProgramTy &Program) {
@@ -85,17 +83,21 @@ Error L0KernelTy::initImpl(GenericDeviceTy &GenericDevice,
return Plugin::success();
}
+using AppendLaunchFnTy = llvm::function_ref<Error(
+ ze_command_list_handle_t CmdList, ze_event_handle_t Event,
+ uint32_t NumWaitEvents, ze_event_handle_t *WaitEvents)>;
+
static Error launchKernelWithImmCmdList(L0DeviceTy &l0Device,
ze_kernel_handle_t zeKernel,
L0LaunchEnvTy &KEnv,
- CommandModeTy CommandMode) {
+ CommandModeTy CommandMode,
+ AppendLaunchFnTy AppendLaunch) {
const auto DeviceId = l0Device.getDeviceId();
auto *IdStr = l0Device.getZeIdCStr();
auto CmdListOrErr = l0Device.getImmCmdList();
if (!CmdListOrErr)
return CmdListOrErr.takeError();
const ze_command_list_handle_t CmdList = *CmdListOrErr;
- // Command queue is not used with immediate command list.
INFO(OMP_INFOTYPE_PLUGIN_KERNEL, DeviceId,
"Using immediate command list for kernel submission.\n");
@@ -120,17 +122,8 @@ static Error launchKernelWithImmCmdList(L0DeviceTy &l0Device,
Error AllErrors = Error::success();
- if (KEnv.IsCooperative) {
- INFO(OMP_INFOTYPE_PLUGIN_KERNEL, DeviceId,
- "Launching cooperative kernel " DPxMOD "\n", DPxPTR(zeKernel));
- CALL_ZE_ACCUM_ERROR(AllErrors, zeCommandListAppendLaunchCooperativeKernel,
- CmdList, zeKernel, &KEnv.GroupCounts, Event,
- NumWaitEvents, WaitEvents);
- } else {
- CALL_ZE_ACCUM_ERROR(AllErrors, zeCommandListAppendLaunchKernel, CmdList,
- zeKernel, &KEnv.GroupCounts, Event, NumWaitEvents,
- WaitEvents);
- }
+ if (auto Err = AppendLaunch(CmdList, Event, NumWaitEvents, WaitEvents))
+ AllErrors = joinErrors(std::move(AllErrors), std::move(Err));
KEnv.Lock.unlock();
if (AllErrors) {
@@ -159,35 +152,33 @@ static Error launchKernelWithImmCmdList(L0DeviceTy &l0Device,
return Plugin::success();
}
-Error L0KernelTy::setKernelGroups(L0DeviceTy &l0Device, L0LaunchEnvTy &KEnv,
- uint32_t NumThreads[3],
- uint32_t NumBlocks[3]) const {
+ze_group_size_t L0KernelTy::createKernelGroups(L0DeviceTy &l0Device,
+ L0LaunchEnvTy &KEnv,
+ uint32_t NumThreads[3],
+ uint32_t NumBlocks[3]) const {
+ ze_group_size_t GroupSizes;
assert(NumThreads[0] > 0 && NumThreads[1] > 0 && NumThreads[2] > 0 &&
"Pre-computed ThreadLimit values must be non-zero");
assert(NumBlocks[0] > 0 && NumBlocks[1] > 0 && NumBlocks[2] > 0 &&
"Pre-computed NumTeams values must be non-zero");
- uint32_t GroupSizes[3];
KEnv.GroupCounts = {NumBlocks[0], NumBlocks[1], NumBlocks[2]};
// Respect max group size attribute in the kernel.
uint32_t MaxGroupSize = KEnv.KernelPR.MaxThreadGroupSize;
- GroupSizes[0] = std::min<uint32_t>(MaxGroupSize, NumThreads[0]);
- GroupSizes[1] = std::min<uint32_t>(MaxGroupSize, NumThreads[1]);
- GroupSizes[2] = std::min<uint32_t>(MaxGroupSize, NumThreads[2]);
+ GroupSizes.groupSizeX = std::min<uint32_t>(MaxGroupSize, NumThreads[0]);
+ GroupSizes.groupSizeY = std::min<uint32_t>(MaxGroupSize, NumThreads[1]);
+ GroupSizes.groupSizeZ = std::min<uint32_t>(MaxGroupSize, NumThreads[2]);
auto DeviceId = l0Device.getDeviceId();
INFO(OMP_INFOTYPE_PLUGIN_KERNEL, DeviceId,
- "Team sizes = {%" PRIu32 ", %" PRIu32 ", %" PRIu32 "}\n", GroupSizes[0],
- GroupSizes[1], GroupSizes[2]);
+ "Team sizes = {%" PRIu32 ", %" PRIu32 ", %" PRIu32 "}\n",
+ GroupSizes.groupSizeX, GroupSizes.groupSizeY, GroupSizes.groupSizeZ);
INFO(OMP_INFOTYPE_PLUGIN_KERNEL, DeviceId,
"Number of teams = {%" PRIu32 ", %" PRIu32 ", %" PRIu32 "}\n",
KEnv.GroupCounts.groupCountX, KEnv.GroupCounts.groupCountY,
KEnv.GroupCounts.groupCountZ);
- CALL_ZE_RET_ERROR(zeKernelSetGroupSize, getZeKernel(), GroupSizes[0],
- GroupSizes[1], GroupSizes[2]);
-
- return Plugin::success();
+ return GroupSizes;
}
Error L0KernelTy::setIndirectFlags(L0DeviceTy &l0Device,
@@ -251,8 +242,8 @@ Error L0KernelTy::launchImpl(GenericDeviceTy &GenericDevice,
// Protect from kernel preparation to submission as kernels are shared.
KEnv.Lock.lock();
- if (auto Err = setKernelGroups(l0Device, KEnv, NumThreads, NumBlocks))
- return Err;
+ ze_group_size_t GroupSizes =
+ createKernelGroups(l0Device, KEnv, NumThreads, NumBlocks);
// Validate cooperative kernel launch constraints
if (IsCooperative) {
@@ -278,34 +269,74 @@ Error L0KernelTy::launchImpl(GenericDeviceTy &GenericDevice,
TotalGroupCount, MaxCooperativeGroupCount);
}
- // Set kernel arguments.
- uint32_t NumKernelArgs = KernelPR.NumKernelArgs;
- if (NumKernelArgs > 0) {
- if (!KernelPR.ArgSizes)
- return Plugin::error(ErrorCode::INVALID_ARGUMENT,
- "level zero plugin requires kernel argument sizes.");
- // Use sizes from kernel properties.
- // TODO: This is temporary workaround it will not work if there is
- // padding/alignment between arguments.
- char *Arg = static_cast<char *>(LaunchParams.Data);
- for (uint32_t I = 0; I < NumKernelArgs; I++) {
- uint32_t ArgSize = KernelPR.ArgSizes[I];
- CALL_ZE_RET_ERROR(zeKernelSetArgumentValue, zeKernel, I, ArgSize, Arg);
-
- INFO(OMP_INFOTYPE_PLUGIN_KERNEL, DeviceId,
- "Kernel Pointer argument %" PRIu32 " (value: " DPxMOD
- ") was set successfully for device %s.\n",
- I, DPxPTR(Arg), IdStr);
- Arg += ArgSize;
+ // With pointer-array arguments, zeCommandListAppendLaunchKernelWithArguments
+ // folds group-size, per-argument set, and launch into a single call.
+ const bool IsPtrArgs = KernelArgs.Flags.IsPtrArgs;
+ if (IsPtrArgs) {
+ if (KernelArgs.NumArgs != KernelPR.NumKernelArgs)
+ return Plugin::error(
+ ErrorCode::INVALID_ARGUMENT,
+ "Number of arguments (%u) does not match the number of arguments "
+ "expected by the kernel (%u)",
+ KernelArgs.NumArgs, KernelPR.NumKernelArgs);
+ } else {
+ CALL_ZE_RET_ERROR(zeKernelSetGroupSize, zeKernel, GroupSizes.groupSizeX,
+ GroupSizes.groupSizeY, GroupSizes.groupSizeZ);
+
+ // Set kernel arguments.
+ uint32_t NumKernelArgs = KernelPR.NumKernelArgs;
+ if (NumKernelArgs > 0) {
+ if (!KernelPR.ArgSizes)
+ return Plugin::error(
+ ErrorCode::INVALID_ARGUMENT,
+ "level zero plugin requires kernel argument sizes.");
+ // Use sizes from kernel properties.
+ // TODO: This is temporary workaround it will not work if there is
+ // padding/alignment between arguments.
+ char *Arg = static_cast<char *>(LaunchParams.Data);
+ for (uint32_t I = 0; I < NumKernelArgs; I++) {
+ uint32_t ArgSize = KernelPR.ArgSizes[I];
+ CALL_ZE_RET_ERROR(zeKernelSetArgumentValue, zeKernel, I, ArgSize, Arg);
+
+ INFO(OMP_INFOTYPE_PLUGIN_KERNEL, DeviceId,
+ "Kernel Pointer argument %" PRIu32 " (value: " DPxMOD
+ ") was set successfully for device %s.\n",
+ I, DPxPTR(Arg), IdStr);
+ Arg += ArgSize;
+ }
}
}
if (auto Err = setIndirectFlags(l0Device, KEnv))
return Err;
+ auto AppendLaunch = [&](ze_command_list_handle_t CmdList,
+ ze_event_handle_t Event, uint32_t NumWaitEvents,
+ ze_event_handle_t *WaitEvents) -> Error {
+ if (IsPtrArgs) {
+ ze_command_list_append_launch_kernel_param_cooperative_desc_t CoopDesc = {
+ ZE_STRUCTURE_TYPE_COMMAND_LIST_APPEND_PARAM_COOPERATIVE_DESC, nullptr,
+ static_cast<ze_bool_t>(IsCooperative)};
+ CALL_ZE_RET_ERROR(zeCommandListAppendLaunchKernelWithArguments, CmdList,
+ zeKernel, KEnv.GroupCounts, GroupSizes,
+ KernelArgs.ArgPtrs, IsCooperative ? &CoopDesc : nullptr,
+ Event, NumWaitEvents, WaitEvents);
+ } else if (IsCooperative) {
+ INFO(OMP_INFOTYPE_PLUGIN_KERNEL, DeviceId,
+ "Launching cooperative kernel " DPxMOD "\n", DPxPTR(zeKernel));
+ CALL_ZE_RET_ERROR(zeCommandListAppendLaunchCooperativeKernel, CmdList,
+ zeKernel, &KEnv.GroupCounts, Event, NumWaitEvents,
+ WaitEvents);
+ } else {
+ CALL_ZE_RET_ERROR(zeCommandListAppendLaunchKernel, CmdList, zeKernel,
+ &KEnv.GroupCounts, Event, NumWaitEvents, WaitEvents);
+ }
+ return Plugin::success();
+ };
+
// The next call should unlock the KernelLock internally.
return launchKernelWithImmCmdList(l0Device, zeKernel, KEnv,
- Options.CommandMode);
+ Options.CommandMode, AppendLaunch);
}
Expected<uint32_t>
diff --git a/offload/unittests/Conformance/include/mathtest/DeviceContext.hpp b/offload/unittests/Conformance/include/mathtest/DeviceContext.hpp
index 5c31fc3da53cd..17ec7e95fb2be 100644
--- a/offload/unittests/Conformance/include/mathtest/DeviceContext.hpp
+++ b/offload/unittests/Conformance/include/mathtest/DeviceContext.hpp
@@ -103,16 +103,22 @@ class DeviceContext {
"device");
if constexpr (sizeof...(Args) == 0) {
- launchKernelImpl(Kernel.Handle, NumGroups, GroupSize, nullptr, 0);
+ launchKernelImpl(Kernel.Handle, NumGroups, GroupSize, 0, nullptr,
+ nullptr);
} else {
- auto KernelArgs = makeKernelArgsPack(std::forward<ArgTypes>(Args)...);
-
static_assert(
(std::is_trivially_copyable_v<std::decay_t<ArgTypes>> && ...),
"Argument types provided to launchKernel must be trivially copyable");
- launchKernelImpl(Kernel.Handle, NumGroups, GroupSize, &KernelArgs,
- sizeof(KernelArgs));
+ ProvidedTypes ArgsTuple(std::forward<ArgTypes>(Args)...);
+ std::apply(
+ [&](auto &...A) {
+ void *ArgPtrs[] = {static_cast<void *>(&A)...};
+ size_t ArgSizes[] = {sizeof(A)...};
+ launchKernelImpl(Kernel.Handle, NumGroups, GroupSize,
+ sizeof...(Args), ArgPtrs, ArgSizes);
+ },
+ ArgsTuple);
}
}
@@ -126,8 +132,8 @@ class DeviceContext {
llvm::StringRef KernelName) const noexcept;
void launchKernelImpl(ol_symbol_handle_t KernelHandle, uint32_t NumGroups,
- uint32_t GroupSize, const void *KernelArgs,
- std::size_t KernelArgsSize) const noexcept;
+ uint32_t GroupSize, size_t NumArgs, void **ArgPtrs,
+ const size_t *ArgSizes) const noexcept;
std::size_t GlobalDeviceId;
ol_device_handle_t DeviceHandle;
diff --git a/offload/unittests/Conformance/include/mathtest/Support.hpp b/offload/unittests/Conformance/include/mathtest/Support.hpp
index 2d3dceef2a230..6cfb67ce6cbb7 100644
--- a/offload/unittests/Conformance/include/mathtest/Support.hpp
+++ b/offload/unittests/Conformance/include/mathtest/Support.hpp
@@ -87,33 +87,6 @@ using KernelSignatureOf = detail::KernelSignatureOfImpl<
template <auto Func>
using KernelSignatureOf_t = typename KernelSignatureOf<Func>::type;
-//===----------------------------------------------------------------------===//
-// Kernel Argument Packing
-//===----------------------------------------------------------------------===//
-
-template <typename... ArgTypes> struct KernelArgsPack;
-
-template <typename ArgType> struct KernelArgsPack<ArgType> {
- std::decay_t<ArgType> Arg;
-
- constexpr KernelArgsPack(ArgType &&Arg) : Arg(std::forward<ArgType>(Arg)) {}
-};
-
-template <typename ArgType0, typename ArgType1, typename... ArgTypes>
-struct KernelArgsPack<ArgType0, ArgType1, ArgTypes...> {
- std::decay_t<ArgType0> Arg0;
- KernelArgsPack<ArgType1, ArgTypes...> Args;
-
- constexpr KernelArgsPack(ArgType0 &&Arg0, ArgType1 &&Arg1, ArgTypes &&...Args)
- : Arg0(std::forward<ArgType0>(Arg0)),
- Args(std::forward<ArgType1>(Arg1), std::forward<ArgTypes>(Args)...) {}
-};
-
-template <typename... ArgTypes>
-KernelArgsPack<ArgTypes...> makeKernelArgsPack(ArgTypes &&...Args) {
- return KernelArgsPack<ArgTypes...>(std::forward<ArgTypes>(Args)...);
-}
-
//===----------------------------------------------------------------------===//
// Configuration Helpers
//===----------------------------------------------------------------------===//
diff --git a/offload/unittests/Conformance/lib/DeviceContext.cpp b/offload/unittests/Conformance/lib/DeviceContext.cpp
index f209fb3b45bc4..638831e405424 100644
--- a/offload/unittests/Conformance/lib/DeviceContext.cpp
+++ b/offload/unittests/Conformance/lib/DeviceContext.cpp
@@ -286,17 +286,18 @@ DeviceContext::getKernelHandle(ol_program_handle_t ProgramHandle,
return Handle;
}
-void DeviceContext::launchKernelImpl(
- ol_symbol_handle_t KernelHandle, uint32_t NumGroups, uint32_t GroupSize,
- const void *KernelArgs, std::size_t KernelArgsSize) const noexcept {
+void DeviceContext::launchKernelImpl(ol_symbol_handle_t KernelHandle,
+ uint32_t NumGroups, uint32_t GroupSize,
+ size_t NumArgs, void **ArgPtrs,
+ const size_t *ArgSizes) const noexcept {
ol_kernel_launch_size_args_t LaunchSizeArgs;
LaunchSizeArgs.Dimensions = 1;
LaunchSizeArgs.NumGroups = {NumGroups, 1, 1};
LaunchSizeArgs.GroupSize = {GroupSize, 1, 1};
LaunchSizeArgs.DynSharedMemory = 0;
- OL_CHECK(olLaunchKernel(nullptr, DeviceHandle, KernelHandle, KernelArgs,
- KernelArgsSize, &LaunchSizeArgs, nullptr));
+ OL_CHECK(olLaunchKernel(nullptr, DeviceHandle, KernelHandle, &LaunchSizeArgs,
+ nullptr, NumArgs, ArgPtrs, ArgSizes));
}
[[nodiscard]] llvm::StringRef DeviceContext::getName() const noexcept {
diff --git a/offload/unittests/OffloadAPI/device_code/CMakeLists.txt b/offload/unittests/OffloadAPI/device_code/CMakeLists.txt
index 57f529e1a4ae6..5dec8c4cd5a9d 100644
--- a/offload/unittests/OffloadAPI/device_code/CMakeLists.txt
+++ b/offload/unittests/OffloadAPI/device_code/CMakeLists.txt
@@ -3,6 +3,7 @@ add_offload_test_device_code(bar.cpp bar)
# Compile with optimizations to eliminate AMDGPU implicit arguments.
add_offload_test_device_code(noargs.cpp noargs -O3)
add_offload_test_device_code(multiargs.cpp multiargs -O3)
+add_offload_test_device_code(composite.cpp composite -O3)
add_offload_test_device_code(byte.cpp byte)
add_offload_test_device_code(localmem.cpp localmem)
add_offload_test_device_code(localmem_reduction.cpp localmem_reduction)
@@ -18,6 +19,7 @@ add_custom_target(offload_device_binaries DEPENDS
bar.bin
noargs.bin
multiargs.bin
+ composite.bin
byte.bin
localmem.bin
localmem_reduction.bin
diff --git a/offload/unittests/OffloadAPI/device_code/composite.cpp b/offload/unittests/OffloadAPI/device_code/composite.cpp
new file mode 100644
index 0000000000000..db0b7a6439df2
--- /dev/null
+++ b/offload/unittests/OffloadAPI/device_code/composite.cpp
@@ -0,0 +1,10 @@
+#include <gpuintrin.h>
+
+struct Foo {
+ uint32_t a;
+ uint32_t b;
+};
+
+extern "C" __gpu_kernel void composite(uint8_t N, Foo F, uint32_t *Out) {
+ Out[__gpu_thread_id(0)] = N + F.a + F.b + __gpu_thread_id(0);
+}
diff --git a/offload/unittests/OffloadAPI/device_code/multiargs.cpp b/offload/unittests/OffloadAPI/device_code/multiargs.cpp
index 265dad124e91e..712512b964d27 100644
--- a/offload/unittests/OffloadAPI/device_code/multiargs.cpp
+++ b/offload/unittests/OffloadAPI/device_code/multiargs.cpp
@@ -1,3 +1,5 @@
#include <gpuintrin.h>
-extern "C" __gpu_kernel void multiargs(char, int *, short) { (void)0; }
+extern "C" __gpu_kernel void multiargs(char A, int *B, short C) {
+ B[__gpu_thread_id(0)] = A + C + __gpu_thread_id(0);
+}
diff --git a/offload/unittests/OffloadAPI/event/olGetEventElapsedTime.cpp b/offload/unittests/OffloadAPI/event/olGetEventElapsedTime.cpp
index b8f3f4d8f0ea1..5e53707a3d3bb 100644
--- a/offload/unittests/OffloadAPI/event/olGetEventElapsedTime.cpp
+++ b/offload/unittests/OffloadAPI/event/olGetEventElapsedTime.cpp
@@ -40,12 +40,11 @@ struct olGetEventElapsedTimeTest : OffloadQueueTest {
}
void launchFoo() {
- struct {
- void *Mem;
- } Args{Mem};
+ void *ArgPtrs[] = {&Mem};
+ size_t ArgSizes[] = {sizeof(Mem)};
- ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernel, &Args, sizeof(Args),
- &LaunchArgs, nullptr));
+ ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernel, &LaunchArgs, nullptr,
+ 1, ArgPtrs, ArgSizes));
}
std::unique_ptr<llvm::MemoryBuffer> DeviceBin;
diff --git a/offload/unittests/OffloadAPI/kernel/olLaunchKernel.cpp b/offload/unittests/OffloadAPI/kernel/olLaunchKernel.cpp
index e2dfa1a9c6b64..00f32e0acf0a8 100644
--- a/offload/unittests/OffloadAPI/kernel/olLaunchKernel.cpp
+++ b/offload/unittests/OffloadAPI/kernel/olLaunchKernel.cpp
@@ -19,6 +19,7 @@
KERNEL_TEST(Foo, foo)
KERNEL_TEST(NoArgs, noargs)
KERNEL_TEST(MultiArgs, multiargs)
+KERNEL_TEST(Composite, composite)
KERNEL_TEST(Byte, byte)
KERNEL_TEST(LocalMem, localmem)
KERNEL_TEST(LocalMemReduction, localmem_reduction)
@@ -60,12 +61,12 @@ TEST_P(olLaunchKernelFooTest, Success) {
void *Mem;
ASSERT_SUCCESS(olMemAlloc(Device, OL_ALLOC_TYPE_MANAGED,
LaunchArgs.GroupSize.x * sizeof(uint32_t), &Mem));
- struct {
- void *Mem;
- } Args{Mem};
- ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernel, &Args, sizeof(Args),
- &LaunchArgs, nullptr));
+ void *ArgPtrs[] = {&Mem};
+ size_t ArgSizes[] = {sizeof(Mem)};
+
+ ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernel, &LaunchArgs, nullptr,
+ std::size(ArgPtrs), ArgPtrs, ArgSizes));
ASSERT_SUCCESS(olSyncQueue(Queue));
@@ -84,12 +85,11 @@ TEST_P(olLaunchKernelFooTest, SuccessThreaded) {
ASSERT_SUCCESS(olMemAlloc(Device, OL_ALLOC_TYPE_DEVICE, Size, &DevAlloc));
ASSERT_SUCCESS(olMemAlloc(Device, OL_ALLOC_TYPE_HOST, Size, &HstAlloc));
- struct {
- void *Mem;
- } Args{DevAlloc};
+ void *ArgPtrs[] = {&DevAlloc};
+ size_t ArgSizes[] = {sizeof(DevAlloc)};
- ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernel, &Args, sizeof(Args),
- &LaunchArgs, nullptr));
+ ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernel, &LaunchArgs, nullptr,
+ std::size(ArgPtrs), ArgPtrs, ArgSizes));
ASSERT_SUCCESS(olMemcpy(Queue, HstAlloc, Host, DevAlloc, Device, Size));
@@ -106,22 +106,63 @@ TEST_P(olLaunchKernelFooTest, SuccessThreaded) {
}
TEST_P(olLaunchKernelNoArgsTest, Success) {
- ASSERT_SUCCESS(
- olLaunchKernel(Queue, Device, Kernel, nullptr, 0, &LaunchArgs, nullptr));
+ ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernel, &LaunchArgs, nullptr, 0,
+ nullptr, nullptr));
ASSERT_SUCCESS(olSyncQueue(Queue));
}
TEST_P(olLaunchKernelMultiArgsTest, Success) {
- struct {
- char A;
- int *B;
- short C;
- } Args{0, nullptr, 0};
-
- ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernel, &Args, sizeof(Args),
- &LaunchArgs, nullptr));
+ void *Mem;
+ ASSERT_SUCCESS(olMemAlloc(Device, OL_ALLOC_TYPE_MANAGED,
+ LaunchArgs.GroupSize.x * sizeof(int), &Mem));
+
+ char A = 3;
+ int *B = (int *)Mem;
+ short C = 5;
+
+ void *ArgPtrs[] = {&A, &B, &C};
+ size_t ArgSizes[] = {sizeof(A), sizeof(B), sizeof(C)};
+
+ ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernel, &LaunchArgs, nullptr,
+ std::size(ArgPtrs), ArgPtrs, ArgSizes));
+
+ ASSERT_SUCCESS(olSyncQueue(Queue));
+
+ int *Data = (int *)Mem;
+ for (uint32_t i = 0; i < LaunchArgs.GroupSize.x; i++)
+ ASSERT_EQ(Data[i], A + C + static_cast<int>(i));
+
+ ASSERT_SUCCESS(olMemFree(Mem));
+}
+
+struct Foo {
+ int a;
+ int b;
+};
+
+TEST_P(olLaunchKernelCompositeTest, Success) {
+ void *Mem;
+ ASSERT_SUCCESS(olMemAlloc(Device, OL_ALLOC_TYPE_MANAGED,
+ LaunchArgs.GroupSize.x * sizeof(uint32_t), &Mem));
+
+ uint32_t N = 1;
+ Foo F{2, 3};
+ uint32_t *Out = (uint32_t *)Mem;
+
+ void *ArgPtrs[] = {&N, &F, &Out};
+ size_t ArgSizes[] = {sizeof(N), sizeof(F), sizeof(Out)};
+
+ ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernel, &LaunchArgs, nullptr,
+ std::size(ArgPtrs), ArgPtrs, ArgSizes));
+
ASSERT_SUCCESS(olSyncQueue(Queue));
+
+ uint32_t *Data = (uint32_t *)Mem;
+ for (uint32_t i = 0; i < LaunchArgs.GroupSize.x; i++)
+ ASSERT_EQ(Data[i], N + F.a + F.b + i);
+
+ ASSERT_SUCCESS(olMemFree(Mem));
}
TEST_P(olLaunchKernelFooTest, SuccessSynchronous) {
@@ -129,12 +170,11 @@ TEST_P(olLaunchKernelFooTest, SuccessSynchronous) {
ASSERT_SUCCESS(olMemAlloc(Device, OL_ALLOC_TYPE_MANAGED,
LaunchArgs.GroupSize.x * sizeof(uint32_t), &Mem));
- struct {
- void *Mem;
- } Args{Mem};
+ void *ArgPtrs[] = {&Mem};
+ size_t ArgSizes[] = {sizeof(Mem)};
- ASSERT_SUCCESS(olLaunchKernel(nullptr, Device, Kernel, &Args, sizeof(Args),
- &LaunchArgs, nullptr));
+ ASSERT_SUCCESS(olLaunchKernel(nullptr, Device, Kernel, &LaunchArgs, nullptr,
+ std::size(ArgPtrs), ArgPtrs, ArgSizes));
uint32_t *Data = (uint32_t *)Mem;
for (uint32_t i = 0; i < 64; i++) {
@@ -144,6 +184,18 @@ TEST_P(olLaunchKernelFooTest, SuccessSynchronous) {
ASSERT_SUCCESS(olMemFree(Mem));
}
+TEST_P(olLaunchKernelByteTest, Success) {
+ unsigned char C = 42;
+
+ void *ArgPtrs[] = {&C};
+ size_t ArgSizes[] = {sizeof(C)};
+
+ ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernel, &LaunchArgs, nullptr,
+ std::size(ArgPtrs), ArgPtrs, ArgSizes));
+
+ ASSERT_SUCCESS(olSyncQueue(Queue));
+}
+
TEST_P(olLaunchKernelLocalMemTest, Success) {
LaunchArgs.NumGroups.x = 4;
LaunchArgs.DynSharedMemory = 64 * sizeof(uint32_t);
@@ -153,12 +205,12 @@ TEST_P(olLaunchKernelLocalMemTest, Success) {
LaunchArgs.GroupSize.x * LaunchArgs.NumGroups.x *
sizeof(uint32_t),
&Mem));
- struct {
- void *Mem;
- } Args{Mem};
- ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernel, &Args, sizeof(Args),
- &LaunchArgs, nullptr));
+ void *ArgPtrs[] = {&Mem};
+ size_t ArgSizes[] = {sizeof(Mem)};
+
+ ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernel, &LaunchArgs, nullptr,
+ std::size(ArgPtrs), ArgPtrs, ArgSizes));
ASSERT_SUCCESS(olSyncQueue(Queue));
@@ -176,12 +228,12 @@ TEST_P(olLaunchKernelLocalMemReductionTest, Success) {
void *Mem;
ASSERT_SUCCESS(olMemAlloc(Device, OL_ALLOC_TYPE_MANAGED,
LaunchArgs.NumGroups.x * sizeof(uint32_t), &Mem));
- struct {
- void *Mem;
- } Args{Mem};
- ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernel, &Args, sizeof(Args),
- &LaunchArgs, nullptr));
+ void *ArgPtrs[] = {&Mem};
+ size_t ArgSizes[] = {sizeof(Mem)};
+
+ ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernel, &LaunchArgs, nullptr,
+ std::size(ArgPtrs), ArgPtrs, ArgSizes));
ASSERT_SUCCESS(olSyncQueue(Queue));
@@ -199,12 +251,12 @@ TEST_P(olLaunchKernelLocalMemStaticTest, Success) {
void *Mem;
ASSERT_SUCCESS(olMemAlloc(Device, OL_ALLOC_TYPE_MANAGED,
LaunchArgs.NumGroups.x * sizeof(uint32_t), &Mem));
- struct {
- void *Mem;
- } Args{Mem};
- ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernel, &Args, sizeof(Args),
- &LaunchArgs, nullptr));
+ void *ArgPtrs[] = {&Mem};
+ size_t ArgSizes[] = {sizeof(Mem)};
+
+ ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernel, &LaunchArgs, nullptr,
+ std::size(ArgPtrs), ArgPtrs, ArgSizes));
ASSERT_SUCCESS(olSyncQueue(Queue));
@@ -246,8 +298,12 @@ TEST_P(olLaunchKernelSingleCounterSyncEventTest, SuccessSyncEvent) {
ArgsSingleCounter Args{LoopRange, NumberToAdd, (uint32_t *)InitValuePassed,
(uint32_t *)ResNum};
- ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernel, &Args, sizeof(Args),
- &LaunchArgs, nullptr));
+ void *ArgPtrs[] = {&Args.InitLoop, &Args.Addend, &Args.InitVal, &Args.Out};
+ size_t ArgSizes[] = {sizeof(Args.InitLoop), sizeof(Args.Addend),
+ sizeof(Args.InitVal), sizeof(Args.Out)};
+
+ ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernel, &LaunchArgs, nullptr,
+ std::size(ArgPtrs), ArgPtrs, ArgSizes));
uint32_t FinalResVal = 0;
ASSERT_SUCCESS(olMemcpy(Queue, &FinalResVal, Host, ResNum, Device, Size));
@@ -299,21 +355,26 @@ TEST_P(olLaunchKernelSingleCounterSyncEventTest, SuccessTwoQueues) {
ArgsSingleCounter Args{LoopRange, NumberToAdd, (uint32_t *)InitValuePassed,
(uint32_t *)ResNum1};
- ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernel, &Args, sizeof(Args),
- &LaunchArgs, nullptr));
+ void *ArgPtrs[] = {&Args.InitLoop, &Args.Addend, &Args.InitVal, &Args.Out};
+ size_t ArgSizes[] = {sizeof(Args.InitLoop), sizeof(Args.Addend),
+ sizeof(Args.InitVal), sizeof(Args.Out)};
+ ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernel, &LaunchArgs, nullptr,
+ std::size(ArgPtrs), ArgPtrs, ArgSizes));
ol_event_handle_t Event = nullptr;
ASSERT_SUCCESS(olCreateEvent(Queue, OL_EVENT_FLAGS_NONE, &Event));
ASSERT_SUCCESS(olWaitEvents(Queue2, &Event, 1));
ArgsSingleCounter Args2{LoopRange, NumberToAdd, (uint32_t *)ResNum1,
(uint32_t *)ResNum2};
+ void *ArgPtrs2[] = {&Args2.InitLoop, &Args2.Addend, &Args2.InitVal,
+ &Args2.Out};
// At the beginning of the kernel, ResNum from the first queue is saved
// locally as the initial value. If operations enqueued before Event have not
// been completed by the time the kernel is executed using the second queue,
// the FinalResVal would be incorrect.
- ASSERT_SUCCESS(olLaunchKernel(Queue2, Device, Kernel, &Args2, sizeof(Args2),
- &LaunchArgs, nullptr));
+ ASSERT_SUCCESS(olLaunchKernel(Queue2, Device, Kernel, &LaunchArgs, nullptr,
+ std::size(ArgPtrs2), ArgPtrs2, ArgSizes));
ASSERT_SUCCESS(olSyncQueue(Queue2));
uint32_t FinalResVal = 0;
@@ -331,15 +392,15 @@ TEST_P(olLaunchKernelGlobalTest, Success) {
void *Mem;
ASSERT_SUCCESS(olMemAlloc(Device, OL_ALLOC_TYPE_MANAGED,
LaunchArgs.GroupSize.x * sizeof(uint32_t), &Mem));
- struct {
- void *Mem;
- } Args{Mem};
- ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernels[0], nullptr, 0,
- &LaunchArgs, nullptr));
+ void *ArgPtrs[] = {&Mem};
+ size_t ArgSizes[] = {sizeof(Mem)};
+
+ ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernels[0], &LaunchArgs, nullptr,
+ 0, nullptr, nullptr));
ASSERT_SUCCESS(olSyncQueue(Queue));
- ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernels[1], &Args, sizeof(Args),
- &LaunchArgs, nullptr));
+ ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernels[1], &LaunchArgs, nullptr,
+ std::size(ArgPtrs), ArgPtrs, ArgSizes));
ASSERT_SUCCESS(olSyncQueue(Queue));
uint32_t *Data = (uint32_t *)Mem;
@@ -354,21 +415,21 @@ TEST_P(olLaunchKernelGlobalTest, InvalidNotAKernel) {
ol_symbol_handle_t Global = nullptr;
ASSERT_SUCCESS(
olGetSymbol(Program, "global", OL_SYMBOL_KIND_GLOBAL_VARIABLE, &Global));
- ASSERT_ERROR(
- OL_ERRC_SYMBOL_KIND,
- olLaunchKernel(Queue, Device, Global, nullptr, 0, &LaunchArgs, nullptr));
+ ASSERT_ERROR(OL_ERRC_SYMBOL_KIND,
+ olLaunchKernel(Queue, Device, Global, &LaunchArgs, nullptr, 0,
+ nullptr, nullptr));
}
TEST_P(olLaunchKernelGlobalCtorTest, Success) {
void *Mem;
ASSERT_SUCCESS(olMemAlloc(Device, OL_ALLOC_TYPE_MANAGED,
LaunchArgs.GroupSize.x * sizeof(uint32_t), &Mem));
- struct {
- void *Mem;
- } Args{Mem};
- ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernel, &Args, sizeof(Args),
- &LaunchArgs, nullptr));
+ void *ArgPtrs[] = {&Mem};
+ size_t ArgSizes[] = {sizeof(Mem)};
+
+ ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernel, &LaunchArgs, nullptr,
+ std::size(ArgPtrs), ArgPtrs, ArgSizes));
ASSERT_SUCCESS(olSyncQueue(Queue));
uint32_t *Data = (uint32_t *)Mem;
@@ -383,7 +444,7 @@ TEST_P(olLaunchKernelGlobalDtorTest, Success) {
// TODO: We can't inspect the result of a destructor yet, once we
// find/implement a way, update this test. For now we just check that nothing
// crashes
- ASSERT_SUCCESS(
- olLaunchKernel(Queue, Device, Kernel, nullptr, 0, &LaunchArgs, nullptr));
+ ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernel, &LaunchArgs, nullptr, 0,
+ nullptr, nullptr));
ASSERT_SUCCESS(olSyncQueue(Queue));
}
diff --git a/offload/unittests/OffloadAPI/kernel/olLaunchKernelCooperative.cpp b/offload/unittests/OffloadAPI/kernel/olLaunchKernelCooperative.cpp
index 3abe3f916693f..d7d661fcd3fab 100644
--- a/offload/unittests/OffloadAPI/kernel/olLaunchKernelCooperative.cpp
+++ b/offload/unittests/OffloadAPI/kernel/olLaunchKernelCooperative.cpp
@@ -56,14 +56,14 @@ TEST_P(olLaunchKernelCooperativeTest, SuccessCooperative) {
{OL_KERNEL_LAUNCH_PROP_TYPE_IS_COOPERATIVE, &IsCooperative},
OL_KERNEL_LAUNCH_PROP_END};
- ASSERT_SUCCESS(
- olLaunchKernel(Queue, Device, Kernel, nullptr, 0, &LaunchArgs, Props));
+ ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernel, &LaunchArgs, Props, 0,
+ nullptr, nullptr));
ASSERT_SUCCESS(olSyncQueue(Queue));
}
TEST_P(olLaunchKernelCooperativeTest, SuccessNonCooperative) {
- ASSERT_SUCCESS(
- olLaunchKernel(Queue, Device, Kernel, nullptr, 0, &LaunchArgs, nullptr));
+ ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernel, &LaunchArgs, nullptr, 0,
+ nullptr, nullptr));
ASSERT_SUCCESS(olSyncQueue(Queue));
}
@@ -79,8 +79,8 @@ TEST_P(olLaunchKernelCooperativeTest, TooManyGroups) {
{OL_KERNEL_LAUNCH_PROP_TYPE_IS_COOPERATIVE, &IsCooperative},
OL_KERNEL_LAUNCH_PROP_END};
- ASSERT_ANY_ERROR(
- olLaunchKernel(Queue, Device, Kernel, nullptr, 0, &LaunchArgs, Props));
+ ASSERT_ANY_ERROR(olLaunchKernel(Queue, Device, Kernel, &LaunchArgs, Props, 0,
+ nullptr, nullptr));
}
TEST_P(olLaunchKernelCooperativeTest, SynchronousLaunch) {
@@ -95,8 +95,8 @@ TEST_P(olLaunchKernelCooperativeTest, SynchronousLaunch) {
{OL_KERNEL_LAUNCH_PROP_TYPE_IS_COOPERATIVE, &IsCooperative},
OL_KERNEL_LAUNCH_PROP_END};
- ASSERT_SUCCESS(
- olLaunchKernel(nullptr, Device, Kernel, nullptr, 0, &LaunchArgs, Props));
+ ASSERT_SUCCESS(olLaunchKernel(nullptr, Device, Kernel, &LaunchArgs, Props, 0,
+ nullptr, nullptr));
}
TEST_P(olLaunchKernelCooperativeTest, InvalidNullHandleKernel) {
diff --git a/offload/unittests/OffloadAPI/memory/olMemcpy.cpp b/offload/unittests/OffloadAPI/memory/olMemcpy.cpp
index 8bc81d0f99796..3d210e3d6d015 100644
--- a/offload/unittests/OffloadAPI/memory/olMemcpy.cpp
+++ b/offload/unittests/OffloadAPI/memory/olMemcpy.cpp
@@ -163,15 +163,14 @@ TEST_P(olMemcpyGlobalTest, SuccessWrite) {
ASSERT_SUCCESS(olMemAlloc(Device, OL_ALLOC_TYPE_MANAGED,
LaunchArgs.GroupSize.x * sizeof(uint32_t),
&DestMem));
- struct {
- void *Mem;
- } Args{DestMem};
+ void *ArgPtrs[] = {&DestMem};
+ size_t ArgSizes[] = {sizeof(DestMem)};
ASSERT_SUCCESS(
olMemcpy(Queue, Addr, Device, SourceMem, Host, 64 * sizeof(uint32_t)));
ASSERT_SUCCESS(olSyncQueue(Queue));
- ASSERT_SUCCESS(olLaunchKernel(Queue, Device, ReadKernel, &Args, sizeof(Args),
- &LaunchArgs, nullptr));
+ ASSERT_SUCCESS(olLaunchKernel(Queue, Device, ReadKernel, &LaunchArgs, nullptr,
+ 1, ArgPtrs, ArgSizes));
ASSERT_SUCCESS(olSyncQueue(Queue));
uint32_t *DestData = (uint32_t *)DestMem;
@@ -188,8 +187,8 @@ TEST_P(olMemcpyGlobalTest, SuccessRead) {
LaunchArgs.GroupSize.x * sizeof(uint32_t),
&DestMem));
- ASSERT_SUCCESS(olLaunchKernel(Queue, Device, WriteKernel, nullptr, 0,
- &LaunchArgs, nullptr));
+ ASSERT_SUCCESS(olLaunchKernel(Queue, Device, WriteKernel, &LaunchArgs,
+ nullptr, 0, nullptr, nullptr));
ASSERT_SUCCESS(olSyncQueue(Queue));
ASSERT_SUCCESS(
olMemcpy(Queue, DestMem, Host, Addr, Device, 64 * sizeof(uint32_t)));
diff --git a/offload/unittests/OffloadAPI/queue/olLaunchHostFunction.cpp b/offload/unittests/OffloadAPI/queue/olLaunchHostFunction.cpp
index 0ad5b58774feb..6219580648abd 100644
--- a/offload/unittests/OffloadAPI/queue/olLaunchHostFunction.cpp
+++ b/offload/unittests/OffloadAPI/queue/olLaunchHostFunction.cpp
@@ -74,11 +74,10 @@ TEST_P(olLaunchHostFunctionKernelTest, SuccessBlocking) {
},
const_cast<bool *>(&Block)));
- struct {
- void *Mem;
- } Args{Mem};
- ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernel, &Args, sizeof(Args),
- &LaunchArgs, nullptr));
+ void *ArgPtrs[] = {&Mem};
+ size_t ArgSizes[] = {sizeof(Mem)};
+ ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernel, &LaunchArgs, nullptr, 1,
+ ArgPtrs, ArgSizes));
std::this_thread::sleep_for(std::chrono::milliseconds(500));
for (uint32_t i = 0; i < 64; i++) {
diff --git a/offload/unittests/OffloadAPI/queue/olWaitEvents.cpp b/offload/unittests/OffloadAPI/queue/olWaitEvents.cpp
index ded5622d8cde1..6396b7cd0a748 100644
--- a/offload/unittests/OffloadAPI/queue/olWaitEvents.cpp
+++ b/offload/unittests/OffloadAPI/queue/olWaitEvents.cpp
@@ -38,21 +38,20 @@ TEST_P(olWaitEventsTest, Success) {
void *Mem;
ASSERT_SUCCESS(olMemAlloc(Device, OL_ALLOC_TYPE_MANAGED,
NUM_KERNELS * sizeof(uint32_t), &Mem));
- struct {
- uint32_t Idx;
- void *Mem;
- } Args{0, Mem};
+ uint32_t Idx = 0;
+ void *ArgPtrs[] = {&Idx, &Mem};
+ size_t ArgSizes[] = {sizeof(Idx), sizeof(Mem)};
for (size_t I = 0; I < NUM_KERNELS; I++) {
- Args.Idx = I;
+ Idx = I;
ASSERT_SUCCESS(olCreateQueue(Device, &Queues[I]));
if (I > 0)
ASSERT_SUCCESS(olWaitEvents(Queues[I], &Events[I - 1], 1));
- ASSERT_SUCCESS(olLaunchKernel(Queues[I], Device, Kernel, &Args,
- sizeof(Args), &LaunchArgs, nullptr));
+ ASSERT_SUCCESS(olLaunchKernel(Queues[I], Device, Kernel, &LaunchArgs,
+ nullptr, 2, ArgPtrs, ArgSizes));
ASSERT_SUCCESS(olCreateEvent(Queues[I], OL_EVENT_FLAGS_NONE, &Events[I]));
}
@@ -74,19 +73,18 @@ TEST_P(olWaitEventsTest, SuccessSingleQueue) {
void *Mem;
ASSERT_SUCCESS(olMemAlloc(Device, OL_ALLOC_TYPE_MANAGED,
NUM_KERNELS * sizeof(uint32_t), &Mem));
- struct {
- uint32_t Idx;
- void *Mem;
- } Args{0, Mem};
+ uint32_t Idx = 0;
+ void *ArgPtrs[] = {&Idx, &Mem};
+ size_t ArgSizes[] = {sizeof(Idx), sizeof(Mem)};
for (size_t I = 0; I < NUM_KERNELS; I++) {
- Args.Idx = I;
+ Idx = I;
if (I > 0)
ASSERT_SUCCESS(olWaitEvents(Queue, &Events[I - 1], 1));
- ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernel, &Args, sizeof(Args),
- &LaunchArgs, nullptr));
+ ASSERT_SUCCESS(olLaunchKernel(Queue, Device, Kernel, &LaunchArgs, nullptr,
+ 2, ArgPtrs, ArgSizes));
ASSERT_SUCCESS(olCreateEvent(Queue, OL_EVENT_FLAGS_NONE, &Events[I]));
}
@@ -106,21 +104,20 @@ TEST_P(olWaitEventsTest, SuccessMultipleEvents) {
void *Mem;
ASSERT_SUCCESS(olMemAlloc(Device, OL_ALLOC_TYPE_MANAGED,
NUM_KERNELS * sizeof(uint32_t), &Mem));
- struct {
- uint32_t Idx;
- void *Mem;
- } Args{0, Mem};
+ uint32_t Idx = 0;
+ void *ArgPtrs[] = {&Idx, &Mem};
+ size_t ArgSizes[] = {sizeof(Idx), sizeof(Mem)};
for (size_t I = 0; I < NUM_KERNELS; I++) {
- Args.Idx = I;
+ Idx = I;
ASSERT_SUCCESS(olCreateQueue(Device, &Queues[I]));
if (I > 0)
ASSERT_SUCCESS(olWaitEvents(Queues[I], Events, I));
- ASSERT_SUCCESS(olLaunchKernel(Queues[I], Device, Kernel, &Args,
- sizeof(Args), &LaunchArgs, nullptr));
+ ASSERT_SUCCESS(olLaunchKernel(Queues[I], Device, Kernel, &LaunchArgs,
+ nullptr, 2, ArgPtrs, ArgSizes));
ASSERT_SUCCESS(olCreateEvent(Queues[I], OL_EVENT_FLAGS_NONE, &Events[I]));
}
>From 43325240896a52492b6932c5eac104a8ee5b152b Mon Sep 17 00:00:00 2001
From: Piotr Balcer <piotr.balcer at intel.com>
Date: Thu, 28 May 2026 13:38:34 +0000
Subject: [PATCH 2/2] fixup! [offload] use argument pointer array in
olLaunchKernel
---
.../llvm/Frontend/Offloading/Utility.h | 17 +---
llvm/lib/Frontend/Offloading/Utility.cpp | 17 +---
.../tools/llvm-gpu-loader/llvm-gpu-loader.cpp | 78 +++++--------------
offload/liboffload/src/OffloadImpl.cpp | 2 -
offload/plugins-nextgen/amdgpu/src/rtl.cpp | 55 +++++++------
.../common/src/PluginInterface.cpp | 7 +-
offload/plugins-nextgen/cuda/src/rtl.cpp | 21 +----
.../OffloadAPI/kernel/olLaunchKernel.cpp | 6 +-
8 files changed, 64 insertions(+), 139 deletions(-)
diff --git a/llvm/include/llvm/Frontend/Offloading/Utility.h b/llvm/include/llvm/Frontend/Offloading/Utility.h
index 812bbb3234246..4c0bc87786dfb 100644
--- a/llvm/include/llvm/Frontend/Offloading/Utility.h
+++ b/llvm/include/llvm/Frontend/Offloading/Utility.h
@@ -122,16 +122,6 @@ namespace amdgpu {
LLVM_ABI bool isImageCompatibleWithEnv(StringRef ImageArch, uint32_t ImageFlags,
StringRef EnvTargetID);
-/// Struct for holding AMDGPU Kernel Argument Metadata, see:
-/// https://llvm.org/docs/AMDGPUUsage.html#amdgpu-amdhsa-code-object-kernel-argument-metadata-map-table-v3
-struct AMDGPUKernelArgMetaData {
- /// Kernel argument offset in bytes. The offset must be a multiple
- /// of the alignment required by the argument.
- uint32_t Offset = 0;
- /// Kernel argument size in bytes.
- uint32_t Size = 0;
-};
-
/// Struct for holding metadata related to AMDGPU kernels, for more information
/// about the metadata and its meaning see:
/// https://llvm.org/docs/AMDGPUUsage.html#code-object-v3
@@ -165,9 +155,10 @@ struct AMDGPUKernelMetaData {
uint32_t WavefrontSize = KInvalidValue;
/// Maximum flat work-group size supported by the kernel in work-items.
uint32_t MaxFlatWorkgroupSize = KInvalidValue;
- /// Per-argument offset and size, read from the ".args" array in code object
- /// metadata. Includes only explicit user arguments.
- SmallVector<AMDGPUKernelArgMetaData, 8> ExplicitArgMDs;
+ /// Per-argument {offset, size} in bytes, read from the ".args" array in code
+ /// object metadata. Explicit user arguments are first, followed by
+ /// hidden arguments.
+ SmallVector<std::pair<uint32_t, uint32_t>, 8> ArgMDs;
};
/// Reads AMDGPU specific metadata from the ELF file and propagates the
diff --git a/llvm/lib/Frontend/Offloading/Utility.cpp b/llvm/lib/Frontend/Offloading/Utility.cpp
index 6648e8c3d1d5a..c07d276244ee1 100644
--- a/llvm/lib/Frontend/Offloading/Utility.cpp
+++ b/llvm/lib/Frontend/Offloading/Utility.cpp
@@ -321,31 +321,20 @@ class KernelInfoReader {
ArgIt != ArgEnd; ++ArgIt) {
auto ArgMap = ArgIt->getMap();
- // Skip hidden arguments
- auto VKIt = ArgMap.find(".value_kind");
- if (VKIt != ArgMap.end() &&
- VKIt->second.getString().starts_with("hidden_"))
- continue;
-
auto OffsetIt = ArgMap.find(".offset");
- // TODO: should these be asserts?
- // if .offset or .size isn't found, that means the kernel
- // is malformed.
if (OffsetIt == ArgMap.end())
return createStringError(
inconvertibleErrorCode(),
"Missing required .offset key in kernel argument metadata map");
+
auto SizeIt = ArgMap.find(".size");
if (SizeIt == ArgMap.end())
return createStringError(
inconvertibleErrorCode(),
"Missing required .size key in kernel argument metadata map");
- amdgpu::AMDGPUKernelArgMetaData ArgMD;
- ArgMD.Offset = OffsetIt->second.getUInt();
- ArgMD.Size = SizeIt->second.getUInt();
-
- KernelData.ExplicitArgMDs.push_back(ArgMD);
+ KernelData.ArgMDs.emplace_back(OffsetIt->second.getUInt(),
+ SizeIt->second.getUInt());
}
}
diff --git a/llvm/tools/llvm-gpu-loader/llvm-gpu-loader.cpp b/llvm/tools/llvm-gpu-loader/llvm-gpu-loader.cpp
index 341ae82c60c60..b17403406c43f 100644
--- a/llvm/tools/llvm-gpu-loader/llvm-gpu-loader.cpp
+++ b/llvm/tools/llvm-gpu-loader/llvm-gpu-loader.cpp
@@ -72,24 +72,6 @@ static cl::list<std::string> Args(cl::ConsumeAfter,
cl::desc("<program arguments>..."),
cl::cat(LoaderCategory));
-// The arguments to the '_begin' kernel.
-struct BeginArgs {
- int Argc;
- void *Argv;
- void *Envp;
-};
-
-// The arguments to the '_start' kernel.
-struct StartArgs {
- int Argc;
- void *Argv;
- void *Envp;
- void *Ret;
-};
-
-// The arguments to the '_end' kernel.
-struct EndArgs {};
-
[[noreturn]] static void handleError(Error E) {
outs().flush();
logAllUnhandledErrors(std::move(E), WithColor::error(errs(), "loader"));
@@ -183,46 +165,23 @@ ol_device_handle_t getHostDevice() {
return Device;
}
-static void launchKernelImpl(ol_queue_handle_t Queue, ol_device_handle_t Device,
- ol_program_handle_t Program, const char *Name,
- ol_kernel_launch_size_args_t LaunchArgs,
- size_t NumArgs, void **ArgPtrs,
- const size_t *ArgSizes) {
- ol_symbol_handle_t Kernel;
- OFFLOAD_ERR(olGetSymbol(Program, Name, OL_SYMBOL_KIND_KERNEL, &Kernel));
-
- OFFLOAD_ERR(olLaunchKernel(Queue, Device, Kernel, &LaunchArgs, nullptr,
- NumArgs, ArgPtrs, ArgSizes));
-}
-
+template <typename... Args>
void launchKernel(ol_queue_handle_t Queue, ol_device_handle_t Device,
ol_program_handle_t Program, const char *Name,
ol_kernel_launch_size_args_t LaunchArgs,
- BeginArgs &KernelArgs) {
- void *ArgPtrs[] = {&KernelArgs.Argc, &KernelArgs.Argv, &KernelArgs.Envp};
- size_t ArgSizes[] = {sizeof(KernelArgs.Argc), sizeof(KernelArgs.Argv),
- sizeof(KernelArgs.Envp)};
- launchKernelImpl(Queue, Device, Program, Name, LaunchArgs, 3, ArgPtrs,
- ArgSizes);
-}
-
-void launchKernel(ol_queue_handle_t Queue, ol_device_handle_t Device,
- ol_program_handle_t Program, const char *Name,
- ol_kernel_launch_size_args_t LaunchArgs,
- StartArgs &KernelArgs) {
- void *ArgPtrs[] = {&KernelArgs.Argc, &KernelArgs.Argv, &KernelArgs.Envp,
- &KernelArgs.Ret};
- size_t ArgSizes[] = {sizeof(KernelArgs.Argc), sizeof(KernelArgs.Argv),
- sizeof(KernelArgs.Envp), sizeof(KernelArgs.Ret)};
- launchKernelImpl(Queue, Device, Program, Name, LaunchArgs, 4, ArgPtrs,
- ArgSizes);
-}
+ Args &...KernelArgs) {
+ ol_symbol_handle_t Kernel;
+ OFFLOAD_ERR(olGetSymbol(Program, Name, OL_SYMBOL_KIND_KERNEL, &Kernel));
-void launchKernel(ol_queue_handle_t Queue, ol_device_handle_t Device,
- ol_program_handle_t Program, const char *Name,
- ol_kernel_launch_size_args_t LaunchArgs, EndArgs &) {
- launchKernelImpl(Queue, Device, Program, Name, LaunchArgs, 0, nullptr,
- nullptr);
+ if constexpr (sizeof...(Args) == 0) {
+ OFFLOAD_ERR(olLaunchKernel(Queue, Device, Kernel, &LaunchArgs, nullptr, 0,
+ nullptr, nullptr));
+ } else {
+ void *ArgPtrs[] = {static_cast<void *>(&KernelArgs)...};
+ size_t ArgSizes[] = {sizeof(KernelArgs)...};
+ OFFLOAD_ERR(olLaunchKernel(Queue, Device, Kernel, &LaunchArgs, nullptr,
+ sizeof...(Args), ArgPtrs, ArgSizes));
+ }
}
int main(int argc, const char **argv, const char **envp) {
@@ -303,8 +262,8 @@ int main(int argc, const char **argv, const char **envp) {
OFFLOAD_ERR(olMemcpy(Queue, DevRet, Device, &Zero, Host, sizeof(int)));
ol_kernel_launch_size_args_t BeginLaunch{1, {1, 1, 1}, {1, 1, 1}, 0};
- BeginArgs BeginArgs = {DevArgc, DevArgv, DevEnvp};
- launchKernel(Queue, Device, Program, "_begin", BeginLaunch, BeginArgs);
+ launchKernel(Queue, Device, Program, "_begin", BeginLaunch, DevArgc, DevArgv,
+ DevEnvp);
OFFLOAD_ERR(olSyncQueue(Queue));
uint32_t Dims = (BlocksZ > 1) ? 3 : (BlocksY > 1) ? 2 : 1;
@@ -312,12 +271,11 @@ int main(int argc, const char **argv, const char **envp) {
{BlocksX, BlocksY, BlocksZ},
{ThreadsX, ThreadsY, ThreadsZ},
/*SharedMemBytes=*/0};
- StartArgs StartArgs = {DevArgc, DevArgv, DevEnvp, DevRet};
- launchKernel(Queue, Device, Program, "_start", StartLaunch, StartArgs);
+ launchKernel(Queue, Device, Program, "_start", StartLaunch, DevArgc, DevArgv,
+ DevEnvp, DevRet);
ol_kernel_launch_size_args_t EndLaunch{1, {1, 1, 1}, {1, 1, 1}, 0};
- EndArgs EndArgs = {};
- launchKernel(Queue, Device, Program, "_end", EndLaunch, EndArgs);
+ launchKernel(Queue, Device, Program, "_end", EndLaunch);
int Ret;
OFFLOAD_ERR(olMemcpy(Queue, &Ret, Host, DevRet, Device, sizeof(int)));
diff --git a/offload/liboffload/src/OffloadImpl.cpp b/offload/liboffload/src/OffloadImpl.cpp
index 6e174fc77c2af..fd2bb703be9a4 100644
--- a/offload/liboffload/src/OffloadImpl.cpp
+++ b/offload/liboffload/src/OffloadImpl.cpp
@@ -1116,7 +1116,6 @@ Error olLaunchKernel_impl(ol_queue_handle_t Queue, ol_device_handle_t Device,
auto *QueueImpl = Queue ? Queue->AsyncInfo : nullptr;
KernelArgsTy LaunchArgs{};
- // TODO: change LaunchArgs.NumArgs to size_t to match the API type
LaunchArgs.NumArgs = static_cast<uint32_t>(NumArgs);
LaunchArgs.UserNumBlocks[0] = LaunchSizeArgs->NumGroups.x;
LaunchArgs.UserNumBlocks[1] = LaunchSizeArgs->NumGroups.y;
@@ -1142,7 +1141,6 @@ Error olLaunchKernel_impl(ol_queue_handle_t Queue, ol_device_handle_t Device,
AsyncInfoWrapperTy AsyncInfoWrapper(*DeviceImpl, QueueImpl);
LaunchArgs.ArgPtrs = ArgPtrs;
- // TODO: Either change ArgSizes to const size_t * or add a new variable.
LaunchArgs.ArgSizes =
reinterpret_cast<int64_t *>(const_cast<size_t *>(ArgSizes));
LaunchArgs.Flags.IsPtrArgs = true;
diff --git a/offload/plugins-nextgen/amdgpu/src/rtl.cpp b/offload/plugins-nextgen/amdgpu/src/rtl.cpp
index ebf3573efd87f..5be5c87c9824a 100644
--- a/offload/plugins-nextgen/amdgpu/src/rtl.cpp
+++ b/offload/plugins-nextgen/amdgpu/src/rtl.cpp
@@ -511,12 +511,14 @@ struct AMDGPUDeviceImageTy : public DeviceImageTy {
findDeviceSymbol(GenericDeviceTy &Device, StringRef SymbolName) const;
/// Get additional info for kernel, e.g., register spill counts
- std::optional<offloading::amdgpu::AMDGPUKernelMetaData>
+ Expected<offloading::amdgpu::AMDGPUKernelMetaData>
getKernelInfo(StringRef Identifier) const {
auto It = KernelInfoMap.find(Identifier);
if (It == KernelInfoMap.end())
- return {};
+ return Plugin::error(ErrorCode::INVALID_BINARY,
+ "Could not find metadata for kernel %s",
+ Identifier.str().c_str());
return It->second;
}
@@ -593,10 +595,10 @@ struct AMDGPUKernelTy : public GenericKernelTy {
ODBG(OLDT_Module) << "ELFABIVersion: " << AMDImage.getELFABIVersion();
// Get additional kernel info read from image
- KernelInfo = AMDImage.getKernelInfo(getName());
- if (!KernelInfo.has_value())
- INFO(OMP_INFOTYPE_PLUGIN_KERNEL, Device.getDeviceId(),
- "Could not read extra information for kernel %s.", getName());
+ auto KernelInfoOrErr = AMDImage.getKernelInfo(getName());
+ if (!KernelInfoOrErr)
+ return KernelInfoOrErr.takeError();
+ KernelInfo = std::move(*KernelInfoOrErr);
return Plugin::success();
}
@@ -649,7 +651,7 @@ struct AMDGPUKernelTy : public GenericKernelTy {
uint32_t ImplicitArgsSize;
/// Additional Info for the AMD GPU Kernel
- std::optional<offloading::amdgpu::AMDGPUKernelMetaData> KernelInfo;
+ offloading::amdgpu::AMDGPUKernelMetaData KernelInfo;
};
/// Class representing an HSA signal. Signals are used to define dependencies
@@ -4194,21 +4196,28 @@ Error AMDGPUKernelTy::launchImpl(GenericDeviceTy &GenericDevice,
// Copy explicit arguments.
size_t ExplicitEnd = 0;
if (KernelArgs.Flags.IsPtrArgs) {
- if (KernelArgs.ArgPtrs) { // may be null
- const auto &ArgMDs = KernelInfo->ExplicitArgMDs;
+ if (KernelArgs.ArgPtrs) {
+ const auto &ArgMDs = KernelInfo.ArgMDs;
- if (KernelArgs.NumArgs != ArgMDs.size())
+ // ArgMDs might also contain hidden implicit arguments, so we can't check
+ // if user-provided NumArgs matches exactly.
+ if (KernelArgs.NumArgs > ArgMDs.size())
return Plugin::error(
ErrorCode::INVALID_ARGUMENT,
- "Number of arguments (%u) does not match the number of arguments "
+ "Number of arguments (%u) exceeds the number of arguments "
"expected by the kernel (%zu)",
KernelArgs.NumArgs, ArgMDs.size());
- for (size_t I = 0; I < ArgMDs.size(); I++)
- std::memcpy(utils::advancePtr(AllArgs, ArgMDs[I].Offset),
- KernelArgs.ArgPtrs[I], ArgMDs[I].Size);
+ for (size_t I = 0; I < KernelArgs.NumArgs; I++) {
+ auto [Offset, Size] = ArgMDs[I];
+ std::memcpy(utils::advancePtr(AllArgs, Offset), KernelArgs.ArgPtrs[I],
+ Size);
+ }
- ExplicitEnd = ArgMDs.back().Offset + ArgMDs.back().Size;
+ if (KernelArgs.NumArgs) {
+ auto [Offset, Size] = ArgMDs[KernelArgs.NumArgs - 1];
+ ExplicitEnd = Offset + Size;
+ }
}
} else {
// TODO: We should expose the args memory manager alloc to the common part
@@ -4276,10 +4285,6 @@ Error AMDGPUKernelTy::printLaunchInfoDetails(GenericDeviceTy &GenericDevice,
if (!(getInfoLevel() & OMP_INFOTYPE_PLUGIN_KERNEL))
return Plugin::success();
- // We don't have data to print additional info, but no hard error
- if (!KernelInfo.has_value())
- return Plugin::success();
-
// General Info
auto *NumGroups = NumBlocks;
auto *ThreadsPerGroup = NumThreads;
@@ -4290,12 +4295,12 @@ Error AMDGPUKernelTy::printLaunchInfoDetails(GenericDeviceTy &GenericDevice,
// Details for AMDGPU kernels (read from image)
// https://www.llvm.org/docs/AMDGPUUsage.html#code-object-v4-metadata
- auto GroupSegmentSize = (*KernelInfo).GroupSegmentList;
- auto SGPRCount = (*KernelInfo).SGPRCount;
- auto VGPRCount = (*KernelInfo).VGPRCount;
- auto SGPRSpillCount = (*KernelInfo).SGPRSpillCount;
- auto VGPRSpillCount = (*KernelInfo).VGPRSpillCount;
- auto MaxFlatWorkgroupSize = (*KernelInfo).MaxFlatWorkgroupSize;
+ auto GroupSegmentSize = KernelInfo.GroupSegmentList;
+ auto SGPRCount = KernelInfo.SGPRCount;
+ auto VGPRCount = KernelInfo.VGPRCount;
+ auto SGPRSpillCount = KernelInfo.SGPRSpillCount;
+ auto VGPRSpillCount = KernelInfo.VGPRSpillCount;
+ auto MaxFlatWorkgroupSize = KernelInfo.MaxFlatWorkgroupSize;
// Prints additional launch info that contains the following.
// Num Args: The number of kernel arguments
diff --git a/offload/plugins-nextgen/common/src/PluginInterface.cpp b/offload/plugins-nextgen/common/src/PluginInterface.cpp
index d43f9268d3cc3..f623a240b137e 100644
--- a/offload/plugins-nextgen/common/src/PluginInterface.cpp
+++ b/offload/plugins-nextgen/common/src/PluginInterface.cpp
@@ -277,10 +277,11 @@ Error GenericKernelTy::launch(GenericDeviceTy &GenericDevice, void **ArgPtrs,
if (!KernelLaunchEnvOrErr)
return KernelLaunchEnvOrErr.takeError();
- // Kernel languages don't use indirection.
- // IsPtrArgs bypasses LaunchParms entirely,
- // plugins read KernelArgs.ArgPtrs/ArgSizes directly.
KernelLaunchParamsTy LaunchParams;
+
+ // Kernel languages (.IsCUDA) don't use indirection, whereas dispatching with
+ // an array of kernel argument pointers (.IsPtrArgs) uses KernelArgs.ArgPtrs
+ // and KernelArgs.ArgSizes directly.
if (KernelArgs.Flags.IsCUDA) {
LaunchParams =
*reinterpret_cast<KernelLaunchParamsTy *>(KernelArgs.ArgPtrs);
diff --git a/offload/plugins-nextgen/cuda/src/rtl.cpp b/offload/plugins-nextgen/cuda/src/rtl.cpp
index 564a4c0e3fd56..51e2bdb0c01dc 100644
--- a/offload/plugins-nextgen/cuda/src/rtl.cpp
+++ b/offload/plugins-nextgen/cuda/src/rtl.cpp
@@ -171,14 +171,11 @@ struct CUDAKernelTy : public GenericKernelTy {
size_t Arg = 0;
ArgsSize = 0;
- NumArgs = 0;
// Find the last argument to know the total size of the arguments.
while ((Res = cuFuncGetParamInfo(Func, Arg++, &ArgOffset, &ArgSize)) ==
- CUDA_SUCCESS) {
+ CUDA_SUCCESS)
ArgsSize = ArgOffset + ArgSize;
- ++NumArgs;
- }
if (Res != CUDA_ERROR_INVALID_VALUE)
return Plugin::check(Res, "error in cuFuncGetParamInfo: %s");
@@ -193,13 +190,6 @@ struct CUDAKernelTy : public GenericKernelTy {
/// The size of the kernel arguments.
size_t ArgsSize;
-
- /// The number of kernel arguments.
- size_t NumArgs = 0;
-
-public:
- /// Return the number of kernel arguments reported by the driver.
- size_t getNumArgs() const { return NumArgs; }
};
/// Class wrapping a CUDA stream reference. These are the objects handled by the
@@ -1484,13 +1474,6 @@ Error CUDAKernelTy::launchImpl(GenericDeviceTy &GenericDevice,
void **KernelParams = nullptr;
if (KernelArgs.Flags.IsPtrArgs) {
- if (KernelArgs.NumArgs != NumArgs)
- return Plugin::error(
- ErrorCode::INVALID_ARGUMENT,
- "Number of arguments (%u) does not match the number of arguments "
- "expected by the kernel (%zu)",
- KernelArgs.NumArgs, NumArgs);
-
KernelParams = KernelArgs.ArgPtrs;
} else {
// The args size passed in LaunchParams may have tail padding,
@@ -1505,7 +1488,7 @@ Error CUDAKernelTy::launchImpl(GenericDeviceTy &GenericDevice,
return Err;
size_t ConfigArgsSize = ArgsSize;
- // valid for buffer passed through LaunchParams
+ // Valid only for a contiguous buffer passed through LaunchParams.
void *Config[] = {CU_LAUNCH_PARAM_BUFFER_POINTER, LaunchParams.Data,
CU_LAUNCH_PARAM_BUFFER_SIZE,
reinterpret_cast<void *>(&ConfigArgsSize),
diff --git a/offload/unittests/OffloadAPI/kernel/olLaunchKernel.cpp b/offload/unittests/OffloadAPI/kernel/olLaunchKernel.cpp
index 00f32e0acf0a8..7ea2991c3c711 100644
--- a/offload/unittests/OffloadAPI/kernel/olLaunchKernel.cpp
+++ b/offload/unittests/OffloadAPI/kernel/olLaunchKernel.cpp
@@ -137,8 +137,8 @@ TEST_P(olLaunchKernelMultiArgsTest, Success) {
}
struct Foo {
- int a;
- int b;
+ uint32_t a;
+ uint32_t b;
};
TEST_P(olLaunchKernelCompositeTest, Success) {
@@ -146,7 +146,7 @@ TEST_P(olLaunchKernelCompositeTest, Success) {
ASSERT_SUCCESS(olMemAlloc(Device, OL_ALLOC_TYPE_MANAGED,
LaunchArgs.GroupSize.x * sizeof(uint32_t), &Mem));
- uint32_t N = 1;
+ uint8_t N = 1;
Foo F{2, 3};
uint32_t *Out = (uint32_t *)Mem;
More information about the llvm-commits
mailing list