alpaka
Abstraction Library for Parallel Kernel Acceleration
Loading...
Searching...
No Matches
TaskKernelGpuUniformCudaHipRt.hpp
Go to the documentation of this file.
1/* Copyright 2026 Benjamin Worpitz, Erik Zenker, Matthias Werner, René Widera, Jan Stephan, Andrea Bocci, Bernhard
2 * Manfred Gruber, Antonio Di Pilato, Mehmet Yusufoglu, Maria Michailidi
3 * SPDX-License-Identifier: MPL-2.0
4 */
5
6#pragma once
7
8#include "alpaka/acc/AccGpuUniformCudaHipRt.hpp"
9#include "alpaka/acc/Traits.hpp"
10#include "alpaka/core/Config.hpp"
11#include "alpaka/core/Cuda.hpp"
12#include "alpaka/core/Decay.hpp"
13#include "alpaka/core/DemangleTypeNames.hpp"
14#include "alpaka/core/Hip.hpp"
15#include "alpaka/core/RemoveRestrict.hpp"
16#include "alpaka/dev/DevUniformCudaHipRt.hpp"
17#include "alpaka/dev/Traits.hpp"
18#include "alpaka/dim/Traits.hpp"
19#include "alpaka/idx/Traits.hpp"
20#include "alpaka/kernel/KernelFunctionAttributes.hpp"
21#include "alpaka/kernel/Traits.hpp"
22#include "alpaka/platform/Traits.hpp"
23#include "alpaka/queue/Traits.hpp"
24#include "alpaka/queue/cuda-hip/QueueUniformCudaHipRt.hpp"
25#include "alpaka/workdiv/WorkDivHelpers.hpp"
26#include "alpaka/workdiv/WorkDivMembers.hpp"
27
28#include <stdexcept>
29#include <string>
30#include <tuple>
31#include <type_traits>
32
33#if ALPAKA_DEBUG >= ALPAKA_DEBUG_MINIMAL
34# include <iostream>
35#endif
36
37#if defined(ALPAKA_ACC_GPU_CUDA_ENABLED) || defined(ALPAKA_ACC_GPU_HIP_ENABLED)
38
39# if !defined(ALPAKA_HOST_ONLY)
40
41# include "alpaka/core/Config.hpp"
42
43# if defined(ALPAKA_ACC_GPU_CUDA_ENABLED) && !ALPAKA_LANG_CUDA
44# error If ALPAKA_ACC_GPU_CUDA_ENABLED is set, the compiler has to support CUDA!
45# endif
46
47# if defined(ALPAKA_ACC_GPU_HIP_ENABLED) && !ALPAKA_LANG_HIP
48# error If ALPAKA_ACC_GPU_HIP_ENABLED is set, the compiler has to support HIP!
49# endif
50
51namespace alpaka
52{
53 namespace detail
54 {
55# if ALPAKA_COMP_CLANG
56# pragma clang diagnostic push
57# pragma clang diagnostic ignored "-Wunused-template"
58# endif
59 //! The GPU CUDA/HIP kernel entry point.
60 // \NOTE: 'A __global__ function or function template cannot have a trailing return type.'
61 // We have put the function into a shallow namespace and gave it a short name, so the mangled name in the
62 // profiler (e.g. ncu) is as shorter as possible.
63 template<typename TKernelFnObj, typename TAcc, typename... TArgs>
64 __global__ void gpuKernel(
65 Vec<Dim<TAcc>, Idx<TAcc>> const threadElemExtent,
66 TKernelFnObj const kernelFnObj,
67 TArgs... args)
68 {
69 TAcc const acc(threadElemExtent);
70
71// with clang it is not possible to query std::result_of for a pure device lambda created on the host side
72# if !(ALPAKA_COMP_CLANG_CUDA && ALPAKA_COMP_CLANG)
73 static_assert(
74 std::is_same_v<decltype(kernelFnObj(const_cast<TAcc const&>(acc), args...)), void>,
75 "The TKernelFnObj is required to return void!");
76# endif
77 kernelFnObj(const_cast<TAcc const&>(acc), args...);
78 }
79# if ALPAKA_COMP_CLANG
80# pragma clang diagnostic pop
81# endif
82
83 template<typename TKernelFnObj, typename TAcc, typename... TArgs>
84 inline void (*kernelName)(
85 Vec<Dim<TAcc>, Idx<TAcc>> const,
86 TKernelFnObj const,
87 remove_restrict_t<std::decay_t<TArgs>>...)
88 = gpuKernel<TKernelFnObj, TAcc, TArgs...>;
89
90 } // namespace detail
91
92 namespace uniform_cuda_hip
93 {
94 namespace detail
95 {
96 template<typename TDim, typename TIdx>
97 ALPAKA_FN_HOST auto checkVecOnly3Dim(Vec<TDim, TIdx> const& vec) -> void
98 {
99 if constexpr(TDim::value > 0)
100 {
101 for(auto i = std::min(typename TDim::value_type{3}, TDim::value); i < TDim::value; ++i)
102 {
103 if(vec[TDim::value - 1u - i] != 1)
104 {
105 throw std::runtime_error(
106 "The CUDA/HIP accelerator supports a maximum of 3 dimensions. All "
107 "work division extents of the dimensions higher 3 have to be 1!");
108 }
109 }
110 }
111 }
112
113 template<typename TDim, typename TIdx>
114 ALPAKA_FN_HOST auto convertVecToUniformCudaHipDim(Vec<TDim, TIdx> const& vec) -> dim3
115 {
116 dim3 dim(1, 1, 1);
117 if constexpr(TDim::value >= 1)
118 dim.x = static_cast<unsigned>(vec[TDim::value - 1u]);
119 if constexpr(TDim::value >= 2)
120 dim.y = static_cast<unsigned>(vec[TDim::value - 2u]);
121 if constexpr(TDim::value >= 3)
122 dim.z = static_cast<unsigned>(vec[TDim::value - 3u]);
123 checkVecOnly3Dim(vec);
124 return dim;
125 }
126 } // namespace detail
127 } // namespace uniform_cuda_hip
128
129 //! The GPU CUDA/HIP accelerator execution task.
130 template<
131 typename TApi,
132 typename TAcc,
133 typename TDim,
134 typename TIdx,
135 typename TKernelFnObj,
136 bool TCooperative,
137 typename... TArgs>
138 class TaskKernelGpuUniformCudaHipRt final : public WorkDivMembers<TDim, TIdx>
139 {
140 public:
141 template<typename TWorkDiv>
142 ALPAKA_FN_HOST TaskKernelGpuUniformCudaHipRt(
143 TWorkDiv&& workDiv,
144 TKernelFnObj const& kernelFnObj,
145 TArgs&&... args)
146 : WorkDivMembers<TDim, TIdx>(std::forward<TWorkDiv>(workDiv))
147 , m_kernelFnObj(kernelFnObj)
148 , m_args(std::forward<TArgs>(args)...)
149 {
150 static_assert(
151 Dim<std::decay_t<TWorkDiv>>::value == TDim::value,
152 "The work division and the execution task have to be of the same dimensionality!");
153 }
154
155 TKernelFnObj m_kernelFnObj;
156 std::tuple<remove_restrict_t<std::decay_t<TArgs>>...> m_args;
157 };
158
159 namespace trait
160 {
161 //! The GPU CUDA/HIP execution task accelerator type trait specialization.
162 template<
163 typename TApi,
164 typename TAcc,
165 typename TDim,
166 typename TIdx,
167 typename TKernelFnObj,
168 bool TCooperative,
169 typename... TArgs>
170 struct AccType<TaskKernelGpuUniformCudaHipRt<TApi, TAcc, TDim, TIdx, TKernelFnObj, TCooperative, TArgs...>>
171 {
172 using type = AccGpuUniformCudaHipRt<TApi, TDim, TIdx>;
173 };
174
175 //! The GPU CUDA/HIP execution task device type trait specialization.
176 template<
177 typename TApi,
178 typename TAcc,
179 typename TDim,
180 typename TIdx,
181 typename TKernelFnObj,
182 bool TCooperative,
183 typename... TArgs>
184 struct DevType<TaskKernelGpuUniformCudaHipRt<TApi, TAcc, TDim, TIdx, TKernelFnObj, TCooperative, TArgs...>>
185 {
186 using type = DevUniformCudaHipRt<TApi>;
187 };
188
189 //! The GPU CUDA/HIP execution task dimension getter trait specialization.
190 template<
191 typename TApi,
192 typename TAcc,
193 typename TDim,
194 typename TIdx,
195 typename TKernelFnObj,
196 bool TCooperative,
197 typename... TArgs>
198 struct DimType<TaskKernelGpuUniformCudaHipRt<TApi, TAcc, TDim, TIdx, TKernelFnObj, TCooperative, TArgs...>>
199 {
200 using type = TDim;
201 };
202
203 //! The CPU CUDA/HIP execution task platform type trait specialization.
204 template<
205 typename TApi,
206 typename TAcc,
207 typename TDim,
208 typename TIdx,
209 typename TKernelFnObj,
210 bool TCooperative,
211 typename... TArgs>
212 struct PlatformType<
213 TaskKernelGpuUniformCudaHipRt<TApi, TAcc, TDim, TIdx, TKernelFnObj, TCooperative, TArgs...>>
214 {
215 using type = PlatformUniformCudaHipRt<TApi>;
216 };
217
218 //! The GPU CUDA/HIP execution task idx type trait specialization.
219 template<
220 typename TApi,
221 typename TAcc,
222 typename TDim,
223 typename TIdx,
224 typename TKernelFnObj,
225 bool TCooperative,
226 typename... TArgs>
227 struct IdxType<TaskKernelGpuUniformCudaHipRt<TApi, TAcc, TDim, TIdx, TKernelFnObj, TCooperative, TArgs...>>
228 {
229 using type = TIdx;
230 };
231
232 //! The CUDA/HIP kernel enqueue trait specialization.
233 template<
234 typename TApi,
235 bool TBlocking,
236 typename TAcc,
237 typename TDim,
238 typename TIdx,
239 typename TKernelFnObj,
240 bool TCooperative,
241 typename... TArgs>
242 struct Enqueue<
243 uniform_cuda_hip::detail::QueueUniformCudaHipRt<TApi, TBlocking>,
244 TaskKernelGpuUniformCudaHipRt<TApi, TAcc, TDim, TIdx, TKernelFnObj, TCooperative, TArgs...>>
245 {
246 ALPAKA_FN_HOST static auto enqueue(
247 uniform_cuda_hip::detail::QueueUniformCudaHipRt<TApi, TBlocking>& queue,
248 TaskKernelGpuUniformCudaHipRt<TApi, TAcc, TDim, TIdx, TKernelFnObj, TCooperative, TArgs...> const&
249 task) -> void
250 {
251 ALPAKA_DEBUG_MINIMAL_LOG_SCOPE;
252 // TODO: Check that (sizeof(TKernelFnObj) * m_3uiBlockThreadExtent.prod()) < available memory idx
253
254# if ALPAKA_DEBUG >= ALPAKA_DEBUG_FULL
255 // std::size_t printfFifoSize;
256 // TApi::deviceGetLimit(&printfFifoSize, TApi::limitPrintfFifoSize);
257 // std::cout << __func__ << " INFO: printfFifoSize: " << printfFifoSize << std::endl;
258 // TApi::deviceSetLimit(TApi::limitPrintfFifoSize, printfFifoSize*10);
259 // TApi::deviceGetLimit(&printfFifoSize, TApi::limitPrintfFifoSize);
260 // std::cout << __func__ << " INFO: printfFifoSize: " << printfFifoSize << std::endl;
261# endif
262
263# if ALPAKA_DEBUG >= ALPAKA_DEBUG_MINIMAL
264 // This checks if the device supports cooperative kernel launch
265 if constexpr(TCooperative)
266 {
267 if(!trait::GetAccDevProps<TAcc>::getAccDevProps(getDev(queue)).m_cooperativeLaunch)
268 {
269 throw std::runtime_error("This accelerator doesn't support cooperative groups functionality!");
270 }
271 }
272# endif
273
274 auto const gridBlockExtent = getWorkDiv<Grid, Blocks>(task);
275 auto const blockThreadExtent = getWorkDiv<Block, Threads>(task);
276 auto const threadElemExtent = getWorkDiv<Thread, Elems>(task);
277
278 dim3 const gridDim = uniform_cuda_hip::detail::convertVecToUniformCudaHipDim(gridBlockExtent);
279 dim3 const blockDim = uniform_cuda_hip::detail::convertVecToUniformCudaHipDim(blockThreadExtent);
280 uniform_cuda_hip::detail::checkVecOnly3Dim(threadElemExtent);
281
282# if ALPAKA_DEBUG >= ALPAKA_DEBUG_FULL
283 std::cout << __func__ << " gridDim: (" << gridDim.z << ", " << gridDim.y << ", " << gridDim.x << ")\n";
284 std::cout << __func__ << " blockDim: (" << blockDim.z << ", " << blockDim.y << ", " << blockDim.x
285 << ")\n";
286# endif
287
288# if ALPAKA_DEBUG >= ALPAKA_DEBUG_MINIMAL
289 // This checks for a valid work division that is also compliant with the hardware maxima of the
290 // accelerator.
291 if(!isValidWorkDiv<TAcc>(task, getDev(queue)))
292 {
293 throw std::runtime_error(
294 "The given work division is not valid or not supported by the device of type "
295 + getAccName<AccGpuUniformCudaHipRt<TApi, TDim, TIdx>>() + "!");
296 }
297# endif
298
299 // Get the size of the block shared dynamic memory.
300 auto const blockSharedMemDynSizeBytes = std::apply(
301 [&](remove_restrict_t<std::decay_t<TArgs>> const&... args) {
302 return getBlockSharedMemDynSizeBytes<TAcc>(
303 task.m_kernelFnObj,
304 blockThreadExtent,
305 threadElemExtent,
306 args...);
307 },
308 task.m_args);
309
310# if ALPAKA_DEBUG >= ALPAKA_DEBUG_FULL
311 // Log the block shared memory idx.
312 std::cout << __func__ << " BlockSharedMemDynSizeBytes: " << blockSharedMemDynSizeBytes << " B"
313 << std::endl;
314# endif
315
316 auto kernelName
317 = alpaka::detail::kernelName<TKernelFnObj, TAcc, remove_restrict_t<std::decay_t<TArgs>>...>;
318
319# if ALPAKA_DEBUG >= ALPAKA_DEBUG_MINIMAL
320 if constexpr(TCooperative)
321 {
322 // Get the maximum number of active blocks for the given kernel on the current device.
323 int const maxActiveBlocks = std::apply(
324 [&](remove_restrict_t<std::decay_t<TArgs>> const&... args) {
325 return getMaxActiveBlocks<TAcc>(
326 getDev(queue),
327 task.m_kernelFnObj,
328 blockThreadExtent,
329 threadElemExtent,
330 args...);
331 },
332 task.m_args);
333
334# if ALPAKA_DEBUG >= ALPAKA_DEBUG_FULL
335 std::cout << "cooperative kernel launch\n";
336 std::cout << "maxBlocksPerGrid for the kernel " << core::demangled<TKernelFnObj> << ": "
337 << maxActiveBlocks << std::endl;
338# endif
339
340 if(maxActiveBlocks <= 0)
341 {
342 using namespace std::literals;
343 throw std::runtime_error(
344 "The kernel "s + std::string(core::demangled<TKernelFnObj>)
345 + " cannot be launched as a cooperative kernel with "s
346 + std::to_string(blockThreadExtent.prod()) + " threads per block on the device "s
347 + getAccName<AccGpuUniformCudaHipRt<TApi, TDim, TIdx>>()
348 + ", because not even a single block can be active.\n"s
349 + "Use alpaka::getMaxActiveBlocks(...) to check the block size at runtime."s);
350 }
351
352 int const requestedBlocks = static_cast<int>(gridBlockExtent.prod());
353 if(requestedBlocks < 0 || requestedBlocks > maxActiveBlocks)
354 {
355 using namespace std::literals;
356 throw std::runtime_error(
357 "The requested number of blocks is larger than the device limit for the kernel "s
358 + std::string(core::demangled<TKernelFnObj>) + ":\ndevice: "s
359 + getAccName<AccGpuUniformCudaHipRt<TApi, TDim, TIdx>>() + "\nrequested blocks: "s
360 + std::to_string(gridBlockExtent.prod()) + "\nmaximum allowed: "s
361 + std::to_string(maxActiveBlocks) + "\n"s
362 + "Use alpaka::getMaxActiveBlocks(...) to query the limit at runtime."s);
363 }
364 }
365# endif
366
367# if ALPAKA_DEBUG >= ALPAKA_DEBUG_FULL
368 // Log the function attributes.
369 typename TApi::FuncAttributes_t funcAttrs;
370 ALPAKA_UNIFORM_CUDA_HIP_RT_CHECK(TApi::funcGetAttributes(&funcAttrs, kernelName));
371 std::cout << __func__ << " binaryVersion: " << funcAttrs.binaryVersion
372 << " constSizeBytes: " << funcAttrs.constSizeBytes << " B"
373 << " localSizeBytes: " << funcAttrs.localSizeBytes << " B"
374 << " maxThreadsPerBlock: " << funcAttrs.maxThreadsPerBlock
375 << " numRegs: " << funcAttrs.numRegs << " ptxVersion: " << funcAttrs.ptxVersion
376 << " sharedSizeBytes: " << funcAttrs.sharedSizeBytes << " B" << std::endl;
377# endif
378
379 // Set the current device.
380 ALPAKA_UNIFORM_CUDA_HIP_RT_CHECK(TApi::setDevice(queue.m_spQueueImpl->m_dev.getNativeHandle()));
381
382 // Enqueue the kernel execution.
383 // \NOTE: No const reference (const &) is allowed as the parameter type because the kernel launch
384 // language extension expects the arguments by value. This forces the type of a float argument given
385 // with std::forward to this function to be of type float instead of e.g. "float const & __ptr64"
386 // (MSVC). If not given by value, the kernel launch code does not copy the value but the pointer to the
387 // value location.
388 std::apply(
389 [&](remove_restrict_t<std::decay_t<TArgs>> const&... args)
390 {
391 // checks whether to launch cooperative or non-cooperative kernel
392 if constexpr(TCooperative)
393 {
394 void const* kernelArgs[] = {&threadElemExtent, &task.m_kernelFnObj, &args...};
395 ALPAKA_UNIFORM_CUDA_HIP_RT_CHECK(TApi::launchCooperativeKernel(
396 kernelName,
397 gridDim,
398 blockDim,
399 const_cast<void**>(kernelArgs),
400 static_cast<std::size_t>(blockSharedMemDynSizeBytes),
401 queue.getNativeHandle()));
402 }
403 else
404 {
405 kernelName<<<
406 gridDim,
407 blockDim,
408 static_cast<std::size_t>(blockSharedMemDynSizeBytes),
409 queue.getNativeHandle()>>>(threadElemExtent, task.m_kernelFnObj, args...);
410 }
411 },
412 task.m_args);
413
414 if constexpr(TBlocking || ALPAKA_DEBUG >= ALPAKA_DEBUG_MINIMAL)
415 {
416 // Wait for the kernel execution to finish but do not check error return of this call.
417 // Do not use the alpaka::wait method because it checks the error itself but we want to give a
418 // custom error message.
419 std::ignore = TApi::streamSynchronize(queue.getNativeHandle());
420 }
421 if constexpr(ALPAKA_DEBUG >= ALPAKA_DEBUG_MINIMAL)
422 {
423 using namespace std::literals;
424 std::string const msg
425 = "execution of kernel '"s + std::string(core::demangled<TKernelFnObj>) + "' failed with"s;
426 ::alpaka::uniform_cuda_hip::detail::rtCheckLastError<TApi, true>(msg.c_str(), __FILE__, __LINE__);
427 }
428 }
429 };
430
431 //! \brief Specialisation of the class template FunctionAttributes
432 //! \tparam TApi The type the API of the GPU accelerator backend. Currently Cuda or Hip.
433 //! \tparam TDim The dimensionality of the accelerator device properties.
434 //! \tparam TIdx The idx type of the accelerator device properties.
435 //! \tparam TKernelFn Kernel function object type.
436 //! \tparam TArgs Kernel function object argument types as a parameter pack.
437 template<typename TApi, typename TDev, typename TDim, typename TIdx, typename TKernelFn, typename... TArgs>
438 struct FunctionAttributes<AccGpuUniformCudaHipRt<TApi, TDim, TIdx>, TDev, TKernelFn, TArgs...>
439 {
440 //! \param dev The device instance
441 //! \param kernelFn The kernel function object which should be executed.
442 //! \param args The kernel invocation arguments.
443 //! \return KernelFunctionAttributes instance. The default version always returns an instance with zero
444 //! fields. For CPU, the field of max threads allowed by kernel function for the block is 1.
445 ALPAKA_FN_HOST static auto getFunctionAttributes(
446 [[maybe_unused]] TDev const& dev,
447 [[maybe_unused]] TKernelFn const& kernelFn,
448 [[maybe_unused]] TArgs&&... args) -> alpaka::KernelFunctionAttributes
449 {
450 auto kernelName = alpaka::detail::kernelName<
451 TKernelFn,
452 AccGpuUniformCudaHipRt<TApi, TDim, TIdx>,
453 remove_restrict_t<std::decay_t<TArgs>>...>;
454
455 typename TApi::FuncAttributes_t funcAttrs;
456# if ALPAKA_COMP_GNUC
457 // Disable and enable compile warnings for gcc
458# pragma GCC diagnostic push
459# pragma GCC diagnostic ignored "-Wconditionally-supported"
460# endif
461 ALPAKA_UNIFORM_CUDA_HIP_RT_CHECK(
462 TApi::funcGetAttributes(&funcAttrs, reinterpret_cast<void const*>(kernelName)));
463# if ALPAKA_COMP_GNUC
464# pragma GCC diagnostic pop
465# endif
466
467 alpaka::KernelFunctionAttributes kernelFunctionAttributes;
468 kernelFunctionAttributes.constSizeBytes = funcAttrs.constSizeBytes;
469 kernelFunctionAttributes.localSizeBytes = funcAttrs.localSizeBytes;
470 kernelFunctionAttributes.sharedSizeBytes = funcAttrs.sharedSizeBytes;
471 kernelFunctionAttributes.maxDynamicSharedSizeBytes = funcAttrs.maxDynamicSharedSizeBytes;
472 kernelFunctionAttributes.numRegs = funcAttrs.numRegs;
473 kernelFunctionAttributes.asmVersion = funcAttrs.ptxVersion;
474 kernelFunctionAttributes.maxThreadsPerBlock = static_cast<int>(funcAttrs.maxThreadsPerBlock);
475
476# if ALPAKA_DEBUG >= ALPAKA_DEBUG_FULL
477 printf("Kernel Function Attributes: \n");
478 printf("binaryVersion: %d \n", funcAttrs.binaryVersion);
479 printf(
480 "constSizeBytes: %lu \n localSizeBytes: %lu, sharedSizeBytes %lu maxDynamicSharedSizeBytes: %d "
481 "\n",
482 funcAttrs.constSizeBytes,
483 funcAttrs.localSizeBytes,
484 funcAttrs.sharedSizeBytes,
485 funcAttrs.maxDynamicSharedSizeBytes);
486
487 printf(
488 "numRegs: %d, ptxVersion: %d \n maxThreadsPerBlock: %d .\n ",
489 funcAttrs.numRegs,
490 funcAttrs.ptxVersion,
491 funcAttrs.maxThreadsPerBlock);
492# endif
493 return kernelFunctionAttributes;
494 }
495 };
496
497 //! The CUDA/HIP get max active blocks for cooperative kernel specialization.
498 template<typename TAcc, typename TKernelFnObj, typename TApi, typename TDim, typename TIdx, typename... TArgs>
499 struct MaxActiveBlocks<TAcc, DevUniformCudaHipRt<TApi>, TKernelFnObj, TDim, TIdx, TArgs...>
500 {
501 ALPAKA_FN_HOST static auto getMaxActiveBlocks(
502 TKernelFnObj const& kernelFnObj,
503 DevUniformCudaHipRt<TApi> const& device,
504 alpaka::Vec<TDim, TIdx> const& blockThreadExtent,
505 alpaka::Vec<TDim, TIdx> const& threadElemExtent,
506 TArgs const&... args) -> int
507 {
508 auto const blockSharedMemDynSizeBytes
509 = getBlockSharedMemDynSizeBytes<TAcc>(kernelFnObj, blockThreadExtent, threadElemExtent, args...);
510
511# ifdef __CUDACC_DEBUG__
512 // Empirically, when a CUDA kernel is compiled in device-debug mode, it is not safe to use more than
513 // one block per multiprocessor.
514 int numBlocksPerSm = 1;
515# else
516 // Query the maximum number of blocks per multiprocessor that can be active at the same time.
517 int numBlocksPerSm = 0;
518 ALPAKA_UNIFORM_CUDA_HIP_RT_CHECK(TApi::occupancyMaxActiveBlocksPerMultiprocessor(
519 &numBlocksPerSm,
520 alpaka::detail::gpuKernel<TKernelFnObj, TAcc, remove_restrict_t<std::decay_t<TArgs>>...>,
521 static_cast<int>(blockThreadExtent.prod()),
522 static_cast<std::size_t>(blockSharedMemDynSizeBytes)));
523# endif
524
525 auto multiProcessorCount = trait::GetAccDevProps<TAcc>::getAccDevProps(device).m_multiProcessorCount;
526
527 return numBlocksPerSm * static_cast<int>(multiProcessorCount);
528 }
529 };
530
531 } // namespace trait
532
533} // namespace alpaka
534
535// These macros can be used to give a more readable name to a GPU kernel.
536// KERNEL must be the class or struct whose operator()(acc, ...) implements the kernel.
537// ::NAME (or ::NAMESPACE::NAME) must be a unique name across the whole program.
538
539struct The_ALPAKA_KERNEL_NAME_macro_must_be_called_in_the_global_namespace;
540
541# define ALPAKA_KERNEL_NAME(KERNEL, NAME) \
542 \
543 struct The_ALPAKA_KERNEL_NAME_macro_must_be_called_in_the_global_namespace; \
544 \
545 static_assert( \
546 std::is_same_v< \
547 The_ALPAKA_KERNEL_NAME_macro_must_be_called_in_the_global_namespace, \
548 ::The_ALPAKA_KERNEL_NAME_macro_must_be_called_in_the_global_namespace>, \
549 "The ALPAKA_KERNEL_NAME macro must be called in the global namespace"); \
550 \
551 template<typename TAcc, typename... TArgs> \
552 __global__ void NAME( \
553 alpaka::Vec<alpaka::Dim<TAcc>, alpaka::Idx<TAcc>> const extent, \
554 KERNEL const kernelFnObj, \
555 TArgs... args) \
556 { \
557 TAcc const acc(extent); \
558 kernelFnObj(const_cast<TAcc const&>(acc), args...); \
559 } \
560 \
561 namespace alpaka::detail \
562 { \
563 template<typename TAcc, typename... TArgs> \
564 inline void (*kernelName<KERNEL, TAcc, TArgs...>)( \
565 alpaka::Vec<alpaka::Dim<TAcc>, alpaka::Idx<TAcc>> const, \
566 KERNEL const, \
567 TArgs...) \
568 = ::NAME<TAcc, TArgs...>; \
569 }
570
571struct The_ALPAKA_KERNEL_SCOPED_NAME_macro_must_be_called_in_the_global_namespace;
572
573# define ALPAKA_KERNEL_SCOPED_NAME(KERNEL, NAMESPACE, NAME) \
574 struct The_ALPAKA_KERNEL_SCOPED_NAME_macro_must_be_called_in_the_global_namespace; \
575 \
576 static_assert( \
577 std::is_same_v< \
578 The_ALPAKA_KERNEL_SCOPED_NAME_macro_must_be_called_in_the_global_namespace, \
579 ::The_ALPAKA_KERNEL_SCOPED_NAME_macro_must_be_called_in_the_global_namespace>, \
580 "The ALPAKA_KERNEL_SCOPED_NAME macro must be called in the global namespace"); \
581 \
582 namespace NAMESPACE \
583 { \
584 template<typename TAcc, typename... TArgs> \
585 __global__ void NAME( \
586 alpaka::Vec<alpaka::Dim<TAcc>, alpaka::Idx<TAcc>> const extent, \
587 KERNEL const kernelFnObj, \
588 TArgs... args) \
589 { \
590 TAcc const acc(extent); \
591 kernelFnObj(const_cast<TAcc const&>(acc), args...); \
592 } \
593 } \
594 \
595 namespace alpaka::detail \
596 { \
597 template<typename TAcc, typename... TArgs> \
598 inline void (*kernelName<KERNEL, TAcc, TArgs...>)( \
599 alpaka::Vec<alpaka::Dim<TAcc>, alpaka::Idx<TAcc>> const, \
600 KERNEL const, \
601 TArgs...) \
602 = ::NAMESPACE::NAME<TAcc, TArgs...>; \
603 }
604
605
606# else
607
608// In host-only mode, expand to empty macros
609
610# define ALPAKA_KERNEL_NAME(KERNEL, NAME)
611# define ALPAKA_KERNEL_SCOPED_NAME(KERNEL, NAMESPACE, NAME)
612
613# endif
614
615#else
616
617// If CUDA or HIP are not available, expand to empty macros
618
619# define ALPAKA_KERNEL_NAME(KERNEL, NAME)
620# define ALPAKA_KERNEL_SCOPED_NAME(KERNEL, NAMESPACE, NAME)
621
622#endif
ALPAKA_DEBUG_MINIMAL
#define ALPAKA_DEBUG_MINIMAL
The minimal debug level.
Definition Debug.hpp:16
ALPAKA_DEBUG
#define ALPAKA_DEBUG
Set the minimum log level if it is not defined.
Definition Debug.hpp:22
ALPAKA_DEBUG_MINIMAL_LOG_SCOPE
#define ALPAKA_DEBUG_MINIMAL_LOG_SCOPE
Definition Debug.hpp:55
ALPAKA_UNIFORM_CUDA_HIP_RT_CHECK
#define ALPAKA_UNIFORM_CUDA_HIP_RT_CHECK(cmd)
CUDA/HIP runtime error checking with log and exception.
Definition UniformCudaHip.hpp:105
alpaka::DevUniformCudaHipRt
The CUDA/HIP RT device handle.
Definition QueueUniformCudaHipRt.hpp:35
alpaka::TaskKernelGpuUniformCudaHipRt
The GPU CUDA/HIP accelerator execution task.
Definition TaskKernelGpuUniformCudaHipRt.hpp:139
alpaka::TaskKernelGpuUniformCudaHipRt::TaskKernelGpuUniformCudaHipRt
ALPAKA_FN_HOST TaskKernelGpuUniformCudaHipRt(TWorkDiv &&workDiv, TKernelFnObj const &kernelFnObj, TArgs &&... args)
Definition TaskKernelGpuUniformCudaHipRt.hpp:142
alpaka::TaskKernelGpuUniformCudaHipRt::m_args
std::tuple< remove_restrict_t< std::decay_t< TArgs > >... > m_args
Definition TaskKernelGpuUniformCudaHipRt.hpp:156
alpaka::Vec
A n-dimensional vector.
Definition Vec.hpp:38
alpaka::WorkDivMembers
A basic class holding the work division as grid block extent, block thread and thread element extent.
Definition WorkDivMembers.hpp:20
alpaka::uniform_cuda_hip::detail::QueueUniformCudaHipRt
The CUDA/HIP RT queue.
Definition QueueUniformCudaHipRt.hpp:115
ALPAKA_FN_HOST
#define ALPAKA_FN_HOST
Definition Common.hpp:43
alpaka::core::demangled
constexpr std::string_view demangled
Definition DemangleTypeNames.hpp:111
alpaka::detail::gpuKernel
__global__ void gpuKernel(Vec< Dim< TAcc >, Idx< TAcc > > const threadElemExtent, TKernelFnObj const kernelFnObj, TArgs... args)
The GPU CUDA/HIP kernel entry point.
Definition TaskKernelGpuUniformCudaHipRt.hpp:64
alpaka::detail::kernelName
void(* kernelName)(Vec< Dim< TAcc >, Idx< TAcc > > const, TKernelFnObj const, remove_restrict_t< std::decay_t< TArgs > >...)
Definition TaskKernelGpuUniformCudaHipRt.hpp:84
alpaka::uniform_cuda_hip::detail::convertVecToUniformCudaHipDim
ALPAKA_FN_HOST auto convertVecToUniformCudaHipDim(Vec< TDim, TIdx > const &vec) -> dim3
Definition TaskKernelGpuUniformCudaHipRt.hpp:114
alpaka::uniform_cuda_hip::detail::rtCheckLastError
ALPAKA_FN_HOST void rtCheckLastError(char const *desc, char const *file, int const &line) noexcept(!TThrow)
CUDA/HIP runtime API last error checking with log and exception.
Definition UniformCudaHip.hpp:78
alpaka::uniform_cuda_hip::detail::checkVecOnly3Dim
ALPAKA_FN_HOST auto checkVecOnly3Dim(Vec< TDim, TIdx > const &vec) -> void
Definition TaskKernelGpuUniformCudaHipRt.hpp:97
alpaka
The alpaka accelerator library.
Definition AccCpuOmp2Blocks.hpp:54
alpaka::getWorkDiv
ALPAKA_NO_HOST_ACC_WARNING ALPAKA_FN_HOST_ACC auto getWorkDiv(TWorkDiv const &workDiv) -> Vec< Dim< TWorkDiv >, Idx< TWorkDiv > >
Get the extent requested.
Definition Traits.hpp:33
alpaka::isValidWorkDiv
ALPAKA_FN_HOST auto isValidWorkDiv(TWorkDiv const &workDiv, AccDevProps< TDim, TIdx > const &accDevProps) -> bool
Checks if the work division is supported.
Definition WorkDivHelpers.hpp:436
alpaka::Idx
typename trait::IdxType< T >::type Idx
Definition Traits.hpp:29
alpaka::remove_restrict_t
typename remove_restrict< T >::type remove_restrict_t
Helper to remove restrict from a type.
Definition RemoveRestrict.hpp:34
alpaka::getMaxActiveBlocks
ALPAKA_NO_HOST_ACC_WARNING ALPAKA_FN_HOST_ACC auto getMaxActiveBlocks(TDev const &device, TKernelFnObj const &kernelFnObj, alpaka::Vec< TDim, TIdx > const &blockThreadExtent, alpaka::Vec< TDim, TIdx > const &threadElemExtent, TArgs const &... args) -> int
Definition Traits.hpp:248
alpaka::getDev
ALPAKA_FN_HOST auto getDev(T const &t)
Definition Traits.hpp:75
alpaka::getAccName
ALPAKA_FN_HOST auto getAccName() -> std::string
Definition Traits.hpp:105
alpaka::getBlockSharedMemDynSizeBytes
ALPAKA_NO_HOST_ACC_WARNING ALPAKA_FN_HOST_ACC auto getBlockSharedMemDynSizeBytes(TKernelFnObj const &kernelFnObj, Vec< TDim, Idx< TAcc > > const &blockThreadExtent, Vec< TDim, Idx< TAcc > > const &threadElemExtent, TArgs const &... args) -> std::size_t
Definition Traits.hpp:197
alpaka::Dim
typename trait::DimType< T >::type Dim
The dimension type trait alias template to remove the ::type.
Definition Traits.hpp:19
std
STL namespace.
alpaka::KernelFunctionAttributes
Kernel function attributes struct. Attributes are filled by calling the API of the accelerator using ...
Definition KernelFunctionAttributes.hpp:15
alpaka::trait::AccType
The accelerator type trait.
Definition Traits.hpp:42
alpaka::trait::DevType
The device type trait.
Definition Traits.hpp:23
alpaka::trait::DimType
The dimension getter type trait.
Definition Traits.hpp:14
alpaka::trait::Enqueue< uniform_cuda_hip::detail::QueueUniformCudaHipRt< TApi, TBlocking >, TaskKernelGpuUniformCudaHipRt< TApi, TAcc, TDim, TIdx, TKernelFnObj, TCooperative, TArgs... > >::enqueue
static ALPAKA_FN_HOST auto enqueue(uniform_cuda_hip::detail::QueueUniformCudaHipRt< TApi, TBlocking > &queue, TaskKernelGpuUniformCudaHipRt< TApi, TAcc, TDim, TIdx, TKernelFnObj, TCooperative, TArgs... > const &task) -> void
Definition TaskKernelGpuUniformCudaHipRt.hpp:246
alpaka::trait::Enqueue
The queue enqueue trait.
Definition Traits.hpp:34
alpaka::trait::FunctionAttributes< AccGpuUniformCudaHipRt< TApi, TDim, TIdx >, TDev, TKernelFn, TArgs... >::getFunctionAttributes
static ALPAKA_FN_HOST auto getFunctionAttributes(TDev const &dev, TKernelFn const &kernelFn, TArgs &&... args) -> alpaka::KernelFunctionAttributes
Definition TaskKernelGpuUniformCudaHipRt.hpp:445
alpaka::trait::FunctionAttributes
The structure template to access to the functions attributes of a kernel function object.
Definition Traits.hpp:94
alpaka::trait::GetAccDevProps
The device properties get trait.
Definition Traits.hpp:66
alpaka::trait::IdxType
The idx type trait.
Definition Traits.hpp:25
alpaka::trait::MaxActiveBlocks< TAcc, DevUniformCudaHipRt< TApi >, TKernelFnObj, TDim, TIdx, TArgs... >::getMaxActiveBlocks
static ALPAKA_FN_HOST auto getMaxActiveBlocks(TKernelFnObj const &kernelFnObj, DevUniformCudaHipRt< TApi > const &device, alpaka::Vec< TDim, TIdx > const &blockThreadExtent, alpaka::Vec< TDim, TIdx > const &threadElemExtent, TArgs const &... args) -> int
Definition TaskKernelGpuUniformCudaHipRt.hpp:501
alpaka::trait::MaxActiveBlocks
Get maximum requested blocks for cooperative kernel trait.
Definition Traits.hpp:50
alpaka::trait::PlatformType
The platform type trait.
Definition Traits.hpp:38