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_

#if (__cplusplus < 201103L && (!defined(_MSVC_LANG) || _MSVC_LANG < 201103L))
#error "The CUDA API headers can only be compiled with C++11 or a later version of the C++ language standard"
#endif

#include "detail/optional.hpp"
#include "detail/span.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

#ifndef CPP14_CONSTEXPR
#if __cplusplus >= 201402L
#define CPP14_CONSTEXPR constexpr
#else
#define CPP14_CONSTEXPR
#endif
#endif

#ifdef NDEBUG
#define NOEXCEPT_IF_NDEBUG noexcept(true)
#else
#define NOEXCEPT_IF_NDEBUG noexcept(false)
#endif

#ifdef _MSC_VER
/*
 * Microsoft Visual C++ (upto v2017) does not support the C++
 * keywords `and`, `or` and `not`. Apparently, the following
 * include is a work-around.
 */
#include <ciso646>
#endif

namespace cuda {

namespace detail_ {

template <bool B>
using bool_constant = ::std::integral_constant<bool, B>;

using true_type = bool_constant<true>;
using false_type = bool_constant<false>;

template<bool...> struct bool_pack;

template<bool... bs>
using all_true = ::std::is_same<bool_pack<bs..., true>, bool_pack<true, bs...>>;

// This is available in C++17 as ::std::void_t, but we're only assuming C++11
template<typename...>
using void_t = void;

// This is available in C++14
template<bool B, typename T = void>
using enable_if_t = typename ::std::enable_if<B, T>::type;

template<typename T>
using remove_reference_t = typename ::std::remove_reference<T>::type;

// primary template handles types that have no nested ::type member:
template <typename, typename = void>
struct has_data_method : ::std::false_type { };

// specialization recognizes types that do have a nested ::type member:
template <typename T>
struct has_data_method<T, cuda::detail_::void_t<decltype(::std::declval<T>().data())>> : ::std::true_type { };

template <typename, typename = void>
struct has_value_type_member : ::std::false_type { };

template <typename T>
struct has_value_type_member<T, cuda::detail_::void_t<typename T::value_type>> : ::std::true_type { };

// TODO: Consider either beefing up this type trait or ditching it in favor of something simpler, or
// in the standard library
template <typename T>
struct is_kinda_like_contiguous_container :
    ::std::integral_constant<bool,
        has_data_method<typename ::std::remove_reference<T>::type>::value
            and has_value_type_member<typename ::std::remove_reference<T>::type>::value
    > {};

} // namespace detail_

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));
    }

    // Named constructor idioms

    static constexpr __host__ __device__ dimensions_t cube(dimension_t x)   { return dimensions_t{ x, x, x }; }
    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 }; }
    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
};

namespace detail_ {

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

} // namespace detail_


} // 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)
    : x(x_), y(y_), z(z_) { }

    constexpr __host__ __device__ dimensions_t(const uint3& v) : dimensions_t(v.x, v.y, v.z) { }
    constexpr __host__ __device__ dimensions_t(const dim3& dims) : dimensions_t(dims.x, dims.y, dims.z) { }
    constexpr __host__ __device__ dimensions_t(dim3&& dims) : 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 { return { x, y, z }; }

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

    // Named constructor idioms

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

    static 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 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;

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

    constexpr overall_dimensions_t flatten() const;
    constexpr size_t volume() const;
    constexpr size_t dimensionality() const;

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

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);
}


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 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 },
    };
}

constexpr overall_dimensions_t composite_dimensions_t::flatten() const { return grid * block; }
constexpr size_t composite_dimensions_t::volume() const { return flatten().volume(); }
constexpr size_t composite_dimensions_t::dimensionality() const { return flatten().dimensionality(); }

} // namespace grid

namespace memory {

#if CUDA_VERSION >= 10020
enum : bool {
    read_enabled = true,
    read_disabled = false,
    write_enabled = true,
    write_disabled = false
};

struct access_permissions_t {
    bool read : 1;
    bool write : 1;

    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;
    }

    static access_permissions_t from_access_flags(CUmemAccess_flags access_flags)
    {
        access_permissions_t result;
        result.read = (access_flags & CU_MEM_ACCESS_FLAGS_PROT_READ);
        result.write = (access_flags & CU_MEM_ACCESS_FLAGS_PROT_READWRITE);
        return result;
    }

    static constexpr access_permissions_t read_and_write()
    {
        return access_permissions_t{ read_enabled, write_enabled };
    }
};

namespace physical_allocation {

// TODO: Consider simply aliasing CUmemAllocationHandleType and using constexpr const's or anonymous enums
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 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);
}

} // 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,
};

using range_attribute_t = CUmem_range_attribute;

} // namespace managed

#if CUDA_VERSION >= 11070
enum class barrier_scope_t : typename ::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

using native_word_t = unsigned;

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;

// TODO: Is this really only for kernels, or can any device-side function be
// represented by a CUfunction?
using handle_t = CUfunction;

} // namespace kernel

// The C++ standard library doesn't offer ::std::dynarray (although it almost did),
// and we won't introduce our own here. So...
template <typename T>
using dynarray = ::std::vector<T>;

} // namespace cuda

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

#endif // CUDA_API_WRAPPERS_COMMON_TYPES_HPP_