types.hpp

Go to the documentation of this file.

#pragma once
#ifndef CUDA_API_WRAPPERS_COMMON_TYPES_HPP_
#define CUDA_API_WRAPPERS_COMMON_TYPES_HPP_

#ifdef _MSC_VER
// See @url https://stackoverflow.com/q/4913922/1593077
#define NOMINMAX
#endif

#include "detail/preamble.hpp"
#include "detail/optional.hpp"
#include "detail/optional_ref.hpp"
#include "detail/span.hpp"
#include "detail/region.hpp"
#include "detail/type_traits.hpp"

#ifndef __CUDACC__
#include <builtin_types.h>
#endif
#include <cuda.h>

#include <type_traits>
#include <utility>
#include <cassert>
#include <cstddef> // for ::std::size_t
#include <cstdint>
#include <vector>
#ifndef NDEBUG
#include <stdexcept>
#endif

#ifndef __CUDACC__
#ifndef __device__
#define __device__
#define __host__
#endif
#endif


namespace cuda {

// This alias for plain C arrays is required due to an MSVC bug, making it fail to
// accept straight up C array reference parameters to functions under some circumstances;
// see: https://developercommunity.visualstudio.com/t/MSVC-rejects-syntax-of-reference-to-C-ar/10792039
template <typename T, size_t N>
using c_array = T[N];

using status_t = CUresult;

using size_t = ::std::size_t;

using dimensionality_t = size_t;

namespace array {

using dimension_t = size_t;

template<dimensionality_t NumDimensions>
struct dimensions_t;

template<>
struct dimensions_t<3> // this almost-inherits cudaExtent
{
    dimension_t width, height, depth;

    constexpr __host__ __device__ dimensions_t(dimension_t width_, dimension_t height_, dimension_t depth_)
        : width(width_), height(height_), depth(depth_) { }
    constexpr __host__ __device__ dimensions_t(cudaExtent e)
        : dimensions_t(e.width, e.height, e.depth) { }
    constexpr __host__ __device__ dimensions_t(const dimensions_t& other)
        : dimensions_t(other.width, other.height, other.depth) { }
    constexpr __host__ __device__ dimensions_t(dimensions_t&& other)
        : dimensions_t(other.width, other.height, other.depth) { }
    constexpr __host__ __device__ dimensions_t(dimension_t linear_size)
        : dimensions_t(linear_size, 1, 1) { }

    CPP14_CONSTEXPR dimensions_t& operator=(const dimensions_t& other) = default;
    CPP14_CONSTEXPR dimensions_t& operator=(dimensions_t&& other) = default;

    constexpr __host__ __device__ operator cudaExtent() const
    {
        return { width, height, depth };
            // Note: We're not using make_cudaExtent here because:
            // 1. It's not constexpr and
            // 2. It doesn't do anything except construct the plain struct - as of CUDA 10 at least
    }

    constexpr __host__ __device__ size_t volume() const { return width * height * depth; }

    constexpr __host__ __device__ size_t size() const { return volume(); }

    constexpr __host__ __device__ dimensionality_t dimensionality() const
    {
        return (width > 1) + (height> 1) + (depth > 1);
    }

    static constexpr __host__ __device__ dimensions_t cube(dimension_t x)   { return dimensions_t{ x, x, x }; }

    // dimensions
    static constexpr __host__ __device__ dimensions_t zero() { return cube(0); }
};

template<>
struct dimensions_t<2>
{
    dimension_t width, height;

    constexpr __host__ __device__ dimensions_t(dimension_t width_, dimension_t height_)
        : width(width_), height(height_) { }
    constexpr __host__ __device__ dimensions_t(const dimensions_t& other)
        : dimensions_t(other.width, other.height) { }
    constexpr __host__ __device__ dimensions_t(dimensions_t&& other)
        : dimensions_t(other.width, other.height) { }
    constexpr __host__ __device__ dimensions_t(dimension_t linear_size)
        : dimensions_t(linear_size, 1) { }

    CPP14_CONSTEXPR __host__ __device__ dimensions_t& operator=(const dimensions_t& other)
    {
        width = other.width; height = other.height;
        return *this;

    }
    CPP14_CONSTEXPR __host__ __device__ dimensions_t& operator=(dimensions_t&& other)
    {
        width = other.width; height = other.height;
        return *this;
    }

    constexpr __host__ __device__ size_t area() const { return width * height; }

    constexpr __host__ __device__ size_t size() const { return area(); }

    constexpr __host__ __device__ dimensionality_t dimensionality() const
    {
        return (width > 1) + (height> 1);
    }

    // Named constructor idioms

    static constexpr __host__ __device__ dimensions_t square(dimension_t x)   { return dimensions_t{ x, x }; }

    // dimensions
    static constexpr __host__ __device__ dimensions_t zero() { return square(0); }
};

} // namespace array

namespace event {

using handle_t = CUevent;

namespace ipc {

using handle_t = CUipcEventHandle;

} // namespace ipc

} // namespace event

namespace stream {

using handle_t = CUstream;

using priority_t       = int;
enum : priority_t {
    default_priority   = 0
};


#if CUDA_VERSION >= 10000
using callback_t = CUhostFn;
#else
using callback_t = CUstreamCallback;
#endif

#if CUDA_VERSION >= 10000

namespace capture {

enum class mode_t : ::std::underlying_type<CUstreamCaptureMode>::type {
    global        = CU_STREAM_CAPTURE_MODE_GLOBAL,
    thread        = CU_STREAM_CAPTURE_MODE_THREAD_LOCAL,
    thread_local_ = thread,
    relaxed       = CU_STREAM_CAPTURE_MODE_RELAXED
};

enum class state_t : ::std::underlying_type<CUstreamCaptureStatus>::type {
    active        = CU_STREAM_CAPTURE_STATUS_ACTIVE,
    capturing     = active,
    invalidated   = CU_STREAM_CAPTURE_STATUS_INVALIDATED,
    none          = CU_STREAM_CAPTURE_STATUS_NONE,
    not_capturing = none
};

} // namespace capture

inline bool is_capturing(capture::state_t status) noexcept
{
    return status == capture::state_t::active;
}

#endif // CUDA_VERSION >= 10000

} // namespace stream

namespace grid {

using dimension_t        = decltype(dim3::x);

using block_dimension_t  = dimension_t;


struct dimensions_t // this almost-inherits dim3
{
    dimension_t x, y, z;
    constexpr __host__ __device__ dimensions_t(dimension_t x_ = 1, dimension_t y_ = 1, dimension_t z_ = 1) noexcept
        : x(x_), y(y_), z(z_) { }

    constexpr __host__ __device__ dimensions_t(const uint3& v) noexcept : dimensions_t(v.x, v.y, v.z) { }
    constexpr __host__ __device__ dimensions_t(const dim3& dims) noexcept : dimensions_t(dims.x, dims.y, dims.z) { }
    constexpr __host__ __device__ dimensions_t(dim3&& dims) noexcept : dimensions_t(dims.x, dims.y, dims.z) { }

    constexpr __host__ __device__ operator uint3(void) const { return { x, y, z }; }

    // This _should_ have been constexpr, but nVIDIA have not marked the dim3 constructors
    // as constexpr, so it isn't
    __host__ __device__ operator dim3(void) const noexcept { return { x, y, z }; }

    constexpr __host__ __device__ size_t volume() const noexcept { return static_cast<size_t>(x) * y * z; }

    constexpr __host__ __device__ dimensionality_t dimensionality() const noexcept
    {
        return (z > 1) + (y > 1) + (x > 1);
    }

    // Named constructor idioms

    static constexpr __host__ __device__ dimensions_t cube(dimension_t x) noexcept   { return dimensions_t{ x, x, x }; }

    static constexpr __host__ __device__ dimensions_t square(dimension_t x) noexcept { return dimensions_t{ x, x, 1 }; }

    static constexpr __host__ __device__ dimensions_t line(dimension_t x) noexcept   { return dimensions_t{ x, 1, 1 }; }

    static constexpr __host__ __device__ dimensions_t point() noexcept               { return dimensions_t{ 1, 1, 1 }; }

    static constexpr bool divides(dimensions_t lhs, dimensions_t rhs)
    {
        return
            (rhs.x % lhs.x == 0) and
            (rhs.y % lhs.y == 0) and
            (rhs.z % lhs.z == 0);
    }

    constexpr dimension_t operator[](int i) const noexcept {
        return (i == 0) ? x :
               (i == 1) ? y :
               z;
    }
    CPP14_CONSTEXPR dimension_t& operator[](int i) noexcept {
        return (i == 0) ? x :
               (i == 1) ? y :
               z;
    }
};

constexpr inline bool operator==(const dim3& lhs, const dim3& rhs) noexcept
{
    return lhs.x == rhs.x and lhs.y == rhs.y and lhs.z == rhs.z;
}
constexpr inline bool operator!=(const dim3& lhs, const dim3& rhs) noexcept
{
    return not (lhs == rhs);
}
constexpr inline bool operator==(const dimensions_t& lhs, const dimensions_t& rhs) noexcept
{
    return lhs.x == rhs.x and lhs.y == rhs.y and lhs.z == rhs.z;
}
constexpr inline bool operator!=(const dimensions_t& lhs, const dimensions_t& rhs) noexcept
{
    return not (lhs == rhs);
}


using block_dimensions_t = dimensions_t;

using overall_dimension_t = size_t;

struct overall_dimensions_t
{
    using dimension_type = overall_dimension_t;
    dimension_type x, y, z;

    constexpr __host__ __device__ overall_dimensions_t(
        dimension_type width_, dimension_type height_, dimension_type depth_) noexcept
        : x(width_), y(height_), z(depth_) { }

    constexpr __host__ __device__ overall_dimensions_t(const dim3& dims) noexcept
        : x(dims.x), y(dims.y), z(dims.z) { }

    constexpr __host__ __device__ overall_dimensions_t(dim3&& dims) noexcept
        : x(dims.x), y(dims.y), z(dims.z) { }

    constexpr __host__ __device__ overall_dimensions_t(const overall_dimensions_t& other) noexcept
        : overall_dimensions_t(other.x, other.y, other.z) { }

    constexpr __host__ __device__ overall_dimensions_t(overall_dimensions_t&& other) noexcept
        : overall_dimensions_t(other.x, other.y, other.z) { }

    explicit constexpr __host__ __device__ overall_dimensions_t(dimensions_t dims) noexcept
        : overall_dimensions_t(dims.x, dims.y, dims.z) { }

    CPP14_CONSTEXPR overall_dimensions_t& operator=(const overall_dimensions_t& other) noexcept = default;
    CPP14_CONSTEXPR overall_dimensions_t& operator=(overall_dimensions_t&& other) noexcept = default;

    constexpr __host__ __device__ size_t volume() const noexcept { return x * y * z; }
    constexpr __host__ __device__ size_t size() const noexcept { return volume(); }
    constexpr __host__ __device__ dimensionality_t dimensionality() const noexcept
    {
        return ((x > 1) + (y > 1) + (z > 1));
    }

    constexpr dimension_type operator[](int i) const noexcept {
        return (i == 0) ? x :
               (i == 1) ? y :
               z;
    }
    CPP14_CONSTEXPR dimension_type& operator[](int i) noexcept {
        return (i == 0) ? x :
               (i == 1) ? y :
               z;
    }
};

constexpr bool operator==(overall_dimensions_t lhs, overall_dimensions_t rhs) noexcept
{
    return (lhs.x == rhs.x) and (lhs.y == rhs.y) and (lhs.z == rhs.z);
}

constexpr bool operator!=(overall_dimensions_t lhs, overall_dimensions_t rhs) noexcept
{
    return not (lhs == rhs);
}

constexpr overall_dimensions_t operator*(dimensions_t grid_dims, block_dimensions_t block_dims) noexcept
{
    return overall_dimensions_t {
        grid_dims.x * overall_dimension_t { block_dims.x },
        grid_dims.y * overall_dimension_t { block_dims.y },
        grid_dims.z * overall_dimension_t { block_dims.z },
    };
}

struct composite_dimensions_t {
    grid::dimensions_t       grid;
    grid::block_dimensions_t block;

    constexpr overall_dimensions_t flatten() const noexcept { return grid * block; }

    constexpr size_t volume() const noexcept { return flatten().volume(); }

    constexpr size_t dimensionality() const noexcept { return flatten().dimensionality(); }

    static constexpr composite_dimensions_t point() noexcept
    {
        return { dimensions_t::point(), block_dimensions_t::point() };
    }

#if __cplusplus >= 202002L
    constexpr bool operator==(const composite_dimensions_t&) const noexcept = default;
    constexpr bool operator!=(const composite_dimensions_t&) const noexcept = default;
#endif
};

#if __cplusplus < 202002L
constexpr bool operator==(composite_dimensions_t lhs, composite_dimensions_t rhs) noexcept
{
    return (lhs.grid == rhs.grid) and (lhs.block == rhs.block);
}

constexpr bool operator!=(composite_dimensions_t lhs, composite_dimensions_t rhs) noexcept
{
    return not (lhs == rhs);
}
#endif // __cplusplus < 202002L

} // namespace grid

namespace memory {

using location_t = CUmemLocation;

#if CUDA_VERSION >= 10020

struct permissions_t {
    bool read;
    bool write;

    operator CUmemAccess_flags() const noexcept
    {
        return read ?
               (write ? CU_MEM_ACCESS_FLAGS_PROT_READWRITE : CU_MEM_ACCESS_FLAGS_PROT_READ) :
               CU_MEM_ACCESS_FLAGS_PROT_NONE;
    }

};

namespace permissions {

constexpr inline permissions_t none()           { return permissions_t{ false, false }; }
constexpr inline permissions_t read_only()      { return permissions_t{ true,  false }; }
constexpr inline permissions_t write_only()     { return permissions_t{ false, true  }; }
constexpr inline permissions_t read_and_write() { return permissions_t{ true,  true  }; }

namespace detail_ {

inline permissions_t from_flags(CUmemAccess_flags access_flags)
{
    bool read = (access_flags & CU_MEM_ACCESS_FLAGS_PROT_READ);
    bool write = (access_flags & CU_MEM_ACCESS_FLAGS_PROT_READWRITE);
    return permissions_t{read, write};
}

} // namespace detail_

} // namespace permissions


namespace physical_allocation {

enum class shared_handle_kind_t : ::std::underlying_type<CUmemAllocationHandleType>::type {
#if CUDA_VERSION >= 11020
    no_export             = CU_MEM_HANDLE_TYPE_NONE,
#endif
    posix_file_descriptor = CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR,
    win32_handle          = CU_MEM_HANDLE_TYPE_WIN32,
    win32_kmt             = CU_MEM_HANDLE_TYPE_WIN32_KMT,
};

namespace detail_ {

template<shared_handle_kind_t SharedHandleKind> struct shared_handle_type_helper;

template <> struct shared_handle_type_helper<shared_handle_kind_t::posix_file_descriptor> { using type = int; };
#if defined(WIN32) || defined(_WIN32) || defined(WIN64) || defined(_WIN64)
template <> struct shared_handle_type_helper<shared_handle_kind_t::win32_handle> { using type = void *; };
#endif
// TODO: What about WIN32_KMT?

} // namespace detail_

template<shared_handle_kind_t SharedHandleKind>
using shared_handle_t = typename detail_::shared_handle_type_helper<SharedHandleKind>::type;

} // namespace physical_allocation
#endif // CUDA_VERSION >= 10020
#if CUDA_VERSION >= 11020

namespace pool {

using handle_t = CUmemoryPool;

using shared_handle_kind_t = physical_allocation::shared_handle_kind_t;

using physical_allocation::shared_handle_t;

namespace ipc {

using ptr_handle_t = CUmemPoolPtrExportData;

} // namespace ipc

} // namespace pool
#endif // CUDA_VERSION >= 11020

namespace pointer {

using attribute_t = CUpointer_attribute;

} // namespace pointer

namespace device {


using address_t = CUdeviceptr;

static_assert(sizeof(void *) == sizeof(device::address_t), "Unexpected address size");

inline address_t address(const void* device_ptr) noexcept
{
    static_assert(sizeof(void*) == sizeof(address_t), "Incompatible sizes for a void pointer and memory::device::address_t");
    return reinterpret_cast<address_t>(device_ptr);
}

inline address_t address(memory::const_region_t region) noexcept { return address(region.start()); }

} // namespace device

inline void* as_pointer(device::address_t address) noexcept
{
    static_assert(sizeof(void*) == sizeof(device::address_t), "Incompatible sizes for a void pointer and memory::device::address_t");
    return reinterpret_cast<void*>(address);
}

namespace shared {

using size_t = unsigned;

} // namespace shared

namespace managed {

enum class initial_visibility_t {
    to_all_devices,
    to_supporters_of_concurrent_managed_access,
};

} // namespace managed

#if CUDA_VERSION >= 11070
enum class barrier_scope_t : ::std::underlying_type<CUstreamMemoryBarrier_flags>::type {
    device = CU_STREAM_MEMORY_BARRIER_TYPE_GPU,
    system = CU_STREAM_MEMORY_BARRIER_TYPE_SYS
};
#endif // CUDA_VERSION >= 11700

#if CUDA_VERSION >= 10000
namespace external {

using handle_t = CUexternalMemory;

struct subregion_spec_t {
    size_t offset;
    size_t size;
};

} // namespace external

#endif // CUDA_VERSION >= 10000

} // namespace memory

struct launch_configuration_t;

enum class multiprocessor_cache_preference_t : ::std::underlying_type<CUfunc_cache_enum>::type {
    no_preference                 = CU_FUNC_CACHE_PREFER_NONE,
    equal_l1_and_shared_memory    = CU_FUNC_CACHE_PREFER_EQUAL,
    prefer_shared_memory_over_l1  = CU_FUNC_CACHE_PREFER_SHARED,
    prefer_l1_over_shared_memory  = CU_FUNC_CACHE_PREFER_L1,
    // aliases
    none                          = no_preference,
    equal                         = equal_l1_and_shared_memory,
    prefer_shared                 = prefer_shared_memory_over_l1,
    prefer_l1                     = prefer_l1_over_shared_memory,
};

enum multiprocessor_shared_memory_bank_size_option_t
    : ::std::underlying_type<CUsharedconfig>::type
{
    device_default       = CU_SHARED_MEM_CONFIG_DEFAULT_BANK_SIZE,
    four_bytes_per_bank  = CU_SHARED_MEM_CONFIG_FOUR_BYTE_BANK_SIZE,
    eight_bytes_per_bank = CU_SHARED_MEM_CONFIG_EIGHT_BYTE_BANK_SIZE
};

namespace device {

using id_t               = CUdevice;

using attribute_t        = CUdevice_attribute;
using attribute_value_t  = int;

namespace peer_to_peer {

using attribute_t = CUdevice_P2PAttribute;

} // namespace peer_to_peer

} // namespace device

namespace context {

using handle_t = CUcontext;

using flags_t = unsigned;

enum host_thread_sync_scheduling_policy_t : unsigned int {

    heuristic = CU_CTX_SCHED_AUTO,

    default_ = heuristic,

    spin      = CU_CTX_SCHED_SPIN,

    block     = CU_CTX_SCHED_BLOCKING_SYNC,

    yield     = CU_CTX_SCHED_YIELD,

    automatic = heuristic,
};

} // namespace context

namespace device {

using flags_t = context::flags_t;

namespace primary_context {

using handle_t = cuda::context::handle_t;

} // namespace primary_context

using host_thread_sync_scheduling_policy_t = context::host_thread_sync_scheduling_policy_t;

} // namespace device

namespace detail_ {

template <typename T, typename U>
inline T identity_cast(U&& x)
{
    static_assert(::std::is_same<
            typename ::std::remove_reference<T>::type,
            typename ::std::remove_reference<U>::type
        >::value,
        "Casting to a different type - don't use identity_cast");
    return static_cast<T>(::std::forward<U>(x));
}

} // namespace detail_

using uuid_t = CUuuid;

namespace module {

using handle_t = CUmodule;

} // namespace module

namespace kernel {

using attribute_t = CUfunction_attribute;

using attribute_value_t = int;

// A raw CUDA driver handle for a kernel; prefer using the @ref cuda::kernel_t type.
using handle_t = CUfunction;

} // namespace kernel

#if CUDA_VERSION >= 10000

namespace graph {

#if CUDA_VERSION >= 13010
using id_t = unsigned;
#endif

namespace node {

using handle_t = CUgraphNode;
using const_handle_t = CUgraphNode_st const *;

constexpr const const_handle_t no_handle = nullptr;

} // namespace node

namespace template_ {

using handle_t = CUgraph;
constexpr const handle_t null_handle = nullptr;

} // namespace template_

namespace instance {

using handle_t = CUgraphExec;

} // namespace instance

} // namespace graph

#endif // CUDA_VERSION >= 10000

} // namespace cuda

#endif // CUDA_API_WRAPPERS_COMMON_TYPES_HPP_