cuda-api-wrappers
Thin C++-flavored wrappers for the CUDA Runtime API
module.hpp
Go to the documentation of this file.
1 
11 #pragma once
12 #ifndef MULTI_WRAPPER_IMPLS_MODULE_HPP_
13 #define MULTI_WRAPPER_IMPLS_MODULE_HPP_
14 
15 #include "context.hpp"
16 #include "../device.hpp"
17 #include "../module.hpp"
18 
19 namespace cuda {
20 
21 // Moved over from context.hpp
22 template <typename ContiguousContainer,
23 cuda::detail_::enable_if_t<detail_::is_kinda_like_contiguous_container<ContiguousContainer>::value, bool>>
24 module_t context_t::create_module(ContiguousContainer module_data) const
25 {
26  return module::create<context_t const &>(*this, module_data);
27 }
28 
29 template <typename ContiguousContainer,
30 cuda::detail_::enable_if_t<detail_::is_kinda_like_contiguous_container<ContiguousContainer>::value, bool>>
31 module_t context_t::create_module(ContiguousContainer module_data, const link::options_t& link_options) const
32 {
33  return module::create<context_t const &>(*this, module_data, link_options);
34 }
35 
36 // These API calls are not really the way you want to work.
37 inline cuda::kernel_t module_t::get_kernel(const char* name) const
38 {
39  CAW_SET_SCOPE_CONTEXT(context_handle_);
40  kernel::handle_t kernel_function_handle;
41  auto result = cuModuleGetFunction(&kernel_function_handle, handle_, name);
42  throw_if_error_lazy(result, ::std::string("Failed obtaining function ") + name
43  + " from " + module::detail_::identify(*this));
44  return kernel::wrap(context::detail_::get_device_id(context_handle_), context_handle_, kernel_function_handle);
45 }
46 
47 
48 namespace module {
49 
50 namespace detail_ {
51 
52 template <typename Creator>
53 module_t create(const context_t& context, const void* module_data, Creator creator_function)
54 {
55  CAW_SET_SCOPE_CONTEXT(context.handle());
56  handle_t new_module_handle;
57  auto status = creator_function(new_module_handle, module_data);
58  throw_if_error_lazy(status, ::std::string("Failed loading a module from memory location ")
59  + cuda::detail_::ptr_as_hex(module_data)
60  + " within " + context::detail_::identify(context));
61  bool do_take_ownership { true };
62  bool doesnt_hold_pc_refcount_unit { false };
63  // TODO: Do we want to allow holding a refcount unit here, if context is
64  // the primary context?
65 
66  // TODO: Make sure the default-constructed options correspond to what cuModuleLoadData uses as defaults
67  return detail_::wrap(
68  context.device_id(), context.handle(), new_module_handle,
69  do_take_ownership, doesnt_hold_pc_refcount_unit);
70 }
71 
72 // TODO: Consider adding create_module() methods to context_t
73 inline module_t create(const context_t& context, const void* module_data, const link::options_t& link_options)
74 {
75  auto creator_function =
76  [&link_options](handle_t& new_module_handle, const void* module_data_) {
77  auto marshalled_options = link::detail_::marshal(link_options);
78  return cuModuleLoadDataEx(
79  &new_module_handle,
80  module_data_,
81  marshalled_options.count(),
82  const_cast<link::detail_::option_t *>(marshalled_options.options()),
83  const_cast<void **>(marshalled_options.values())
84  );
85  };
86  return detail_::create(context, module_data, creator_function);
87 }
88 
89 inline module_t create(const context_t& context, const void* module_data)
90 {
91  auto creator_function =
92  [](handle_t& new_module_handle, const void* module_data_) {
93  return cuModuleLoadData(&new_module_handle, module_data_);
94  };
95  return detail_::create(context, module_data, creator_function);
96 }
97 
98 
99 
100 inline device::primary_context_t get_context_for(device_t& locus) { return locus.primary_context(); }
101 
102 } // namespace detail_
103 
105  const device_t& device,
106  const char* path)
107 {
108  auto pc = device.primary_context();
109  device::primary_context::detail_::increase_refcount(device.id());
110  return load_from_file(pc, path);
111 }
112 
113 inline module_t load_from_file(const char* path)
114 {
115  return load_from_file(device::current::get(), path);
116 }
117 
118 } // namespace module
119 
120 inline context_t module_t::context() const { return context::detail_::from_handle(context_handle_); }
121 inline device_t module_t::device() const { return device::get(context::detail_::get_device_id(context_handle_)); }
122 
123 #if CUDA_VERSION < 12000
124 inline CUsurfref module_t::get_surface(const char* name) const
125 {
126  CAW_SET_SCOPE_CONTEXT(context_handle_);
127  CUsurfref raw_surface_reference;
128  auto status = cuModuleGetSurfRef(&raw_surface_reference, handle_, name);
129  throw_if_error_lazy(status, ::std::string("Failed obtaining a reference to surface \"") + name + "\" from "
130  + module::detail_::identify(*this));
131  return raw_surface_reference;
132 }
133 
134 inline CUtexref module_t::get_texture_reference(const char* name) const
135 {
136  CAW_SET_SCOPE_CONTEXT(context_handle_);
137  CUtexref raw_texture_reference;
138  auto status = cuModuleGetTexRef(&raw_texture_reference, handle_, name);
139  throw_if_error_lazy(status, ::std::string("Failed obtaining a reference to texture \"") + name + "\" from "
140  + module::detail_::identify(*this));
141  return raw_texture_reference;
142 }
143 #endif
144 
145 
146 } // namespace cuda
147 
148 #endif // MULTI_WRAPPER_IMPLS_MODULE_HPP_
149 
CUsurfref get_surface(const char *name) const
A "plug" of a method regarding surfaces, which modules support but our wrappers don&#39;t really cater to...
Definition: module.hpp:124
Wrapper class for a CUDA context.
Definition: context.hpp:244
Definitions and functionality wrapping CUDA APIs.
Definition: array.hpp:22
A class for holding the primary context of a CUDA device.
Definition: primary_context.hpp:112
cuda::kernel_t get_kernel(const char *name) const
Obtains a kernel constituting part of this module.
Definition: module.hpp:37
Wrapper class for a CUDA code module.
Definition: module.hpp:123
device::id_t id() const noexcept
Return the proxied device&#39;s ID.
Definition: device.hpp:594
module_t load_from_file(const context_t &context, const char *path)
Load a module from an appropriate compiled or semi-compiled file, allocating all relevant resources f...
Definition: module.hpp:317
CUtexref get_texture_reference(const char *name) const
A "plug" of a method regarding texture references.
Definition: module.hpp:134
Implementations requiring the definitions of multiple CUDA entity proxy classes, and which regard con...
device::primary_context_t primary_context(bool hold_pc_refcount_unit=false) const
Produce a proxy for the device&#39;s primary context - the one used by runtime API calls.
Definition: device.hpp:152
device_t get(id_t id)
Returns a proxy for the CUDA device with a given id.
Definition: device.hpp:837
#define throw_if_error_lazy(status__,...)
A macro for only throwing an error if we&#39;ve failed - which also ensures no string is constructed unle...
Definition: error.hpp:316
module_t create_module(ContiguousContainer module_data, const link::options_t &link_options) const
Create a new module of kernels and global memory regions within this context; see also cuda::module::...
Definition: module.hpp:31
context_t context() const
Definition: module.hpp:120
device_t device() const
Definition: module.hpp:121
A non-owning wrapper for CUDA kernels - whether they be __global__ functions compiled apriori...
Definition: kernel.hpp:159
Wrapper class for a CUDA device.
Definition: device.hpp:135