From a10ca8ed6454361b3889a9d5b41b97d065984648 Mon Sep 17 00:00:00 2001 From: Piotr Balcer Date: Fri, 29 May 2026 19:54:22 +0200 Subject: [PATCH] [offload] use argument pointer array in olLaunchKernel (#194333) 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. ## Mapping the proposed API to backends 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. --- .../llvm/Frontend/Offloading/Utility.h | 5 + llvm/lib/Frontend/Offloading/Utility.cpp | 21 ++ .../tools/llvm-gpu-loader/llvm-gpu-loader.cpp | 46 ++--- 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 | 18 +- offload/plugins-nextgen/amdgpu/src/rtl.cpp | 74 ++++--- .../common/src/PluginInterface.cpp | 6 +- offload/plugins-nextgen/cuda/src/rtl.cpp | 19 +- .../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 | 129 +++++++----- .../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 | 185 ++++++++++++------ .../kernel/olLaunchKernelCooperative.cpp | 16 +- .../unittests/OffloadAPI/memory/olMemcpy.cpp | 13 +- .../OffloadAPI/queue/olLaunchHostFunction.cpp | 9 +- .../OffloadAPI/queue/olWaitEvents.cpp | 39 ++-- 26 files changed, 439 insertions(+), 279 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 a1a8a0291626..4c0bc87786df 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" @@ -154,6 +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, size} in bytes, read from the ".args" array in code + /// object metadata. Explicit user arguments are first, followed by + /// hidden arguments. + SmallVector, 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 909af0237910..c07d276244ee 100644 --- a/llvm/lib/Frontend/Offloading/Utility.cpp +++ b/llvm/lib/Frontend/Offloading/Utility.cpp @@ -315,6 +315,27 @@ private: 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(); + + auto OffsetIt = ArgMap.find(".offset"); + 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"); + + KernelData.ArgMDs.emplace_back(OffsetIt->second.getUInt(), + SizeIt->second.getUInt()); + } } 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 40fed374e067..b17403406c43 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 Args(cl::ConsumeAfter, cl::desc("..."), 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,16 +165,23 @@ ol_device_handle_t getHostDevice() { return Device; } -template +template 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) { + ol_kernel_launch_size_args_t LaunchArgs, + Args &...KernelArgs) { 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 ? 0 : sizeof(Args), - &LaunchArgs, /*Properties=*/nullptr)); + if constexpr (sizeof...(Args) == 0) { + OFFLOAD_ERR(olLaunchKernel(Queue, Device, Kernel, &LaunchArgs, nullptr, 0, + nullptr, nullptr)); + } else { + void *ArgPtrs[] = {static_cast(&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) { @@ -273,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; @@ -282,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/llvm/tools/llvm-gpu-loader/llvm-gpu-loader.h b/llvm/tools/llvm-gpu-loader/llvm-gpu-loader.h index d2fafe8984dc..729ca49b7b32 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 212fb285030f..d1326b433755 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 bc0dcb565f3f..06778105e019 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 66fcbbc264ab..fd2bb703be9a 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,7 @@ Error olLaunchKernel_impl(ol_queue_handle_t Queue, ol_device_handle_t Device, auto *QueueImpl = Queue ? Queue->AsyncInfo : nullptr; KernelArgsTy LaunchArgs{}; + LaunchArgs.NumArgs = static_cast(NumArgs); LaunchArgs.UserNumBlocks[0] = LaunchSizeArgs->NumGroups.x; LaunchArgs.UserNumBlocks[1] = LaunchSizeArgs->NumGroups.y; LaunchArgs.UserNumBlocks[2] = LaunchSizeArgs->NumGroups.z; @@ -1138,12 +1140,10 @@ Error olLaunchKernel_impl(ol_queue_handle_t Queue, ol_device_handle_t Device, } AsyncInfoWrapperTy AsyncInfoWrapper(*DeviceImpl, QueueImpl); - KernelLaunchParamsTy Params; - Params.Data = const_cast(ArgumentsData); - Params.Size = ArgumentsSize; - LaunchArgs.ArgPtrs = reinterpret_cast(&Params); - // Don't do anything with pointer indirection; use arg data as-is - LaunchArgs.Flags.IsCUDA = true; + LaunchArgs.ArgPtrs = ArgPtrs; + LaunchArgs.ArgSizes = + reinterpret_cast(const_cast(ArgSizes)); + LaunchArgs.Flags.IsPtrArgs = true; auto *KernelImpl = std::get(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 9995da9c2aa2..d1c6d0de1128 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 + Expected 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 @@ private: uint32_t ImplicitArgsSize; /// Additional Info for the AMD GPU Kernel - std::optional KernelInfo; + offloading::amdgpu::AMDGPUKernelMetaData KernelInfo; }; /// Class representing an HSA signal. Signals are used to define dependencies @@ -4191,11 +4193,39 @@ 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) { + const auto &ArgMDs = KernelInfo.ArgMDs; + + // 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) exceeds the number of arguments " + "expected by the kernel (%zu)", + KernelArgs.NumArgs, ArgMDs.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); + } + + 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 + // as alternative to copying them twice. + if (LaunchParams.Size) + std::memcpy(AllArgs, LaunchParams.Data, LaunchParams.Size); + ExplicitEnd = LaunchParams.Size; + } AMDGPUDeviceTy &AMDGPUDevice = static_cast(GenericDevice); @@ -4203,8 +4233,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( @@ -4255,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; @@ -4269,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 4379ebd25079..f623a240b137 100644 --- a/offload/plugins-nextgen/common/src/PluginInterface.cpp +++ b/offload/plugins-nextgen/common/src/PluginInterface.cpp @@ -279,11 +279,13 @@ Error GenericKernelTy::launch(GenericDeviceTy &GenericDevice, void **ArgPtrs, KernelLaunchParamsTy LaunchParams; - // Kernel languages don't use indirection. + // 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(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 698bd4544705..51e2bdb0c01d 100644 --- a/offload/plugins-nextgen/cuda/src/rtl.cpp +++ b/offload/plugins-nextgen/cuda/src/rtl.cpp @@ -1472,17 +1472,23 @@ Error CUDAKernelTy::launchImpl(GenericDeviceTy &GenericDevice, AsyncInfoWrapperTy &AsyncInfoWrapper) const { CUDADeviceTy &CUDADevice = static_cast(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) { + 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 only for a contiguous buffer passed through LaunchParams. void *Config[] = {CU_LAUNCH_PARAM_BUFFER_POINTER, LaunchParams.Data, CU_LAUNCH_PARAM_BUFFER_SIZE, reinterpret_cast(&ConfigArgsSize), @@ -1514,7 +1520,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 8643897facac..67cbdca3ab82 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 17e35a37f116..ebbdf565b9ea 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 4d6ba0ee392b..f7134b5c5282 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 dc884e3b753d..c390ae7e4e92 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; + 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(MaxGroupSize, NumThreads[0]); - GroupSizes[1] = std::min(MaxGroupSize, NumThreads[1]); - GroupSizes[2] = std::min(MaxGroupSize, NumThreads[2]); + GroupSizes.groupSizeX = std::min(MaxGroupSize, NumThreads[0]); + GroupSizes.groupSizeY = std::min(MaxGroupSize, NumThreads[1]); + GroupSizes.groupSizeZ = std::min(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(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); + // 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); - 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; + // 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(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(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 diff --git a/offload/unittests/Conformance/include/mathtest/DeviceContext.hpp b/offload/unittests/Conformance/include/mathtest/DeviceContext.hpp index 5c31fc3da53c..17ec7e95fb2b 100644 --- a/offload/unittests/Conformance/include/mathtest/DeviceContext.hpp +++ b/offload/unittests/Conformance/include/mathtest/DeviceContext.hpp @@ -103,16 +103,22 @@ public: "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(Args)...); - static_assert( (std::is_trivially_copyable_v> && ...), "Argument types provided to launchKernel must be trivially copyable"); - launchKernelImpl(Kernel.Handle, NumGroups, GroupSize, &KernelArgs, - sizeof(KernelArgs)); + ProvidedTypes ArgsTuple(std::forward(Args)...); + std::apply( + [&](auto &...A) { + void *ArgPtrs[] = {static_cast(&A)...}; + size_t ArgSizes[] = {sizeof(A)...}; + launchKernelImpl(Kernel.Handle, NumGroups, GroupSize, + sizeof...(Args), ArgPtrs, ArgSizes); + }, + ArgsTuple); } } @@ -126,8 +132,8 @@ private: 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 2d3dceef2a23..6cfb67ce6cbb 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 using KernelSignatureOf_t = typename KernelSignatureOf::type; -//===----------------------------------------------------------------------===// -// Kernel Argument Packing -//===----------------------------------------------------------------------===// - -template struct KernelArgsPack; - -template struct KernelArgsPack { - std::decay_t Arg; - - constexpr KernelArgsPack(ArgType &&Arg) : Arg(std::forward(Arg)) {} -}; - -template -struct KernelArgsPack { - std::decay_t Arg0; - KernelArgsPack Args; - - constexpr KernelArgsPack(ArgType0 &&Arg0, ArgType1 &&Arg1, ArgTypes &&...Args) - : Arg0(std::forward(Arg0)), - Args(std::forward(Arg1), std::forward(Args)...) {} -}; - -template -KernelArgsPack makeKernelArgsPack(ArgTypes &&...Args) { - return KernelArgsPack(std::forward(Args)...); -} - //===----------------------------------------------------------------------===// // Configuration Helpers //===----------------------------------------------------------------------===// diff --git a/offload/unittests/Conformance/lib/DeviceContext.cpp b/offload/unittests/Conformance/lib/DeviceContext.cpp index f209fb3b45bc..638831e40542 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 57f529e1a4ae..5dec8c4cd5a9 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 000000000000..db0b7a6439df --- /dev/null +++ b/offload/unittests/OffloadAPI/device_code/composite.cpp @@ -0,0 +1,10 @@ +#include + +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 265dad124e91..712512b964d2 100644 --- a/offload/unittests/OffloadAPI/device_code/multiargs.cpp +++ b/offload/unittests/OffloadAPI/device_code/multiargs.cpp @@ -1,3 +1,5 @@ #include -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 b3423b02771c..33a62921f4ae 100644 --- a/offload/unittests/OffloadAPI/event/olGetEventElapsedTime.cpp +++ b/offload/unittests/OffloadAPI/event/olGetEventElapsedTime.cpp @@ -41,12 +41,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 DeviceBin; diff --git a/offload/unittests/OffloadAPI/kernel/olLaunchKernel.cpp b/offload/unittests/OffloadAPI/kernel/olLaunchKernel.cpp index 2db80f94dd8c..92fe178b67f5 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)); @@ -86,12 +87,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)); @@ -108,22 +108,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}; + 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(olLaunchKernel(Queue, Device, Kernel, &Args, sizeof(Args), - &LaunchArgs, nullptr)); 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(i)); + + ASSERT_SUCCESS(olMemFree(Mem)); +} + +struct Foo { + uint32_t a; + uint32_t b; +}; + +TEST_P(olLaunchKernelCompositeTest, Success) { + void *Mem; + ASSERT_SUCCESS(olMemAlloc(Device, OL_ALLOC_TYPE_MANAGED, + LaunchArgs.GroupSize.x * sizeof(uint32_t), &Mem)); + + uint8_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) { @@ -131,12 +172,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++) { @@ -146,6 +186,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) { SKIP_KNOWN_FAILURE(LevelZero{"unsupported DynSharedMemory"}); @@ -157,12 +209,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)); @@ -182,12 +234,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)); @@ -207,12 +259,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)); @@ -254,8 +306,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)); @@ -307,21 +363,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; @@ -339,15 +400,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; @@ -362,9 +423,9 @@ 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) { @@ -373,12 +434,12 @@ 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; @@ -393,7 +454,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 3abe3f916693..d7d661fcd3fa 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 8bc81d0f9979..3d210e3d6d01 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 a7ce6df44b4a..387fb308da6d 100644 --- a/offload/unittests/OffloadAPI/queue/olLaunchHostFunction.cpp +++ b/offload/unittests/OffloadAPI/queue/olLaunchHostFunction.cpp @@ -84,11 +84,10 @@ TEST_P(olLaunchHostFunctionKernelTest, SuccessBlocking) { }, const_cast(&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 ded5622d8cde..6396b7cd0a74 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])); }