GPU::parallelFor() and GPU::transform(), which run a functor over many indices or inputs: in a kernel in a CUDA or HIP build, and in a host loop otherwise. More...
#include <cstddef>#include <cstdint>#include <type_traits>#include <vector>#include "EBGeometry_GPU.hpp"#include "EBGeometry_GPURuntime.hpp"#include "EBGeometry_MemoryResource.hpp"#include "EBGeometry_PODVector.hpp"#include "EBGeometry_GPUTransformImplem.hpp"
Go to the source code of this file.
Namespaces | |
| namespace | EBGeometry |
| Namespace containing all of EBGeometry's functionality. | |
| namespace | EBGeometry::GPU |
| Backend-neutral GPU runtime API surface (CUDA or HIP). | |
Typedefs | |
| template<class Query , class In > | |
| using | EBGeometry::GPU::TransformResult = std::decay_t< std::invoke_result_t< const Query &, const In & > > |
| The result type of a query on an input. | |
Functions | |
| bool | EBGeometry::GPU::deviceAvailable () noexcept |
| Whether a GPU is visible to the runtime. | |
| MemoryResource & | EBGeometry::GPU::memoryResource () noexcept |
| The memory a query's data must live in for parallelFor() and transform() to read it. | |
| template<unsigned int BlockSize = DefaultBlockSize, unsigned int MinResidentThreads = DefaultMinResidentThreads, class F > | |
| void | EBGeometry::GPU::parallelFor (std::size_t a_n, const F &a_f) |
| Run a functor on every index: a_f(i) for i = 0, ..., a_n - 1, in no particular order. | |
| template<unsigned int BlockSize = DefaultBlockSize, unsigned int MinResidentThreads = DefaultMinResidentThreads, class Query , class In , class Out > | |
| void | EBGeometry::GPU::transform (const Query &a_query, PODSpan< In > a_in, PODSpan< Out > a_out) |
| Run a query on every input in place: a_out[i] = a_query(a_in[i]), with no allocation and no copies. | |
| template<unsigned int BlockSize = DefaultBlockSize, unsigned int MinResidentThreads = DefaultMinResidentThreads, class Query , class In , class Out > | |
| void | EBGeometry::GPU::transform (const Query &a_query, const Pool &a_pool, const PODVector< In > &a_in, const PODVector< Out > &a_out) |
| Run a query on every input of a pool-resident array, writing the results into another array in the same pool: out[i] = a_query(in[i]), with no allocation and no copies. | |
| template<unsigned int BlockSize = DefaultBlockSize, unsigned int MinResidentThreads = DefaultMinResidentThreads, class Query , class In , class Out > | |
| void | EBGeometry::GPU::transform (const Query &a_query, const std::vector< In > &a_in, std::vector< Out > &a_out) |
| Run a query on every input, from host memory to host memory: a_out[i] = a_query(a_in[i]). | |
Variables | |
| constexpr bool | EBGeometry::GPU::enabled = false |
| Whether this translation unit is compiled as CUDA or HIP, so that parallelFor() and transform() run a kernel. | |
| constexpr const char * | EBGeometry::GPU::backendName = "host" |
| Where parallelFor() and transform() run in this translation unit: "HIP", "CUDA" or "host". | |
| constexpr unsigned int | EBGeometry::GPU::DefaultBlockSize = 128 |
| The number of threads per block a launch uses unless it is given another. | |
| constexpr unsigned int | EBGeometry::GPU::DefaultMinResidentThreads = 384 |
| The threads a launch asks the CUDA compiler to keep resident on each multiprocessor, at least, unless it is given another: 384, three blocks of 128, which holds a thread to at most 168 registers. | |
Detailed Description
GPU::parallelFor() and GPU::transform(), which run a functor over many indices or inputs: in a kernel in a CUDA or HIP build, and in a host loop otherwise.
GPU::transform() works like std::transform or thrust::transform: it writes a_query(in[i]) to out[i], into an output the caller provides. A query is a trivially copyable functor whose call operator is EBGEOMETRY_HOST_DEVICE and maps one input to one result. It holds what it evaluates by value: an analytic shape, or a pool-resident object rebased onto memory the device can read (a mirror in GPU::memoryResource(), or a frozen managed pool; see the memory model in the user documentation). One source then serves both builds: a CUDA or HIP translation unit launches a kernel, and any other runs the same functor in a host loop, so a plain C++ build needs no GPU toolkit. The GPU:: prefix names where these calls can run, not where they always do. GPU::enabled says which of the two a translation unit gets, and GPU::deviceAvailable() whether a device is visible at run time.
The calls are layered, each built on the one below:
- GPU::parallelFor(n, f) runs f(i) for every i < n, one thread per index.
- GPU::transform(query, in, out) over two PODSpan's: no allocation and no copies, the data is wherever the caller put it (managed memory, a buffer an application already owns on the device, a PODVector bound against a mirrored pool). A span carries no memory location, so nothing checks that the device can reach it. Timing this call times the kernel.
- GPU::transform(query, pool, in, out) over two PODVector's in one pool: the span overload, after checking that the pool is frozen and in memory the device can access. No allocation and no copies either.
- GPU::transform(query, in, out) over two std::vector's: a convenience that also allocates device memory, copies the inputs in and the results back, and frees. Timing it includes those copies.
Every launch takes its block size as a template parameter (GPU::DefaultBlockSize unless given), so that the kernel can carry it as launch bounds. Every call is synchronous: it returns once the kernel has finished. Every transform requires as many results as inputs, and results of exactly the query's result type.
