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 "../device.hpp"
16 #include "../module.hpp"
17 
18 namespace cuda {
19 
20 // Moved over from context.hpp
21 template <typename ContiguousContainer,
22 cuda::detail_::enable_if_t<detail_::is_kinda_like_contiguous_container<ContiguousContainer>::value, bool>>
23 module_t context_t::create_module(ContiguousContainer module_data) const
24 {
25  return module::create<context_t const &>(*this, module_data);
26 }
27 
28 template <typename ContiguousContainer,
29 cuda::detail_::enable_if_t<detail_::is_kinda_like_contiguous_container<ContiguousContainer>::value, bool>>
30 module_t context_t::create_module(ContiguousContainer module_data, const link::options_t& link_options) const
31 {
32  return module::create<context_t const &>(*this, module_data, link_options);
33 }
34 
35 // These API calls are not really the way you want to work.
36 inline cuda::kernel_t module_t::get_kernel(const char* name) const
37 {
38  CAW_SET_SCOPE_CONTEXT(context_handle_);
39  kernel::handle_t kernel_function_handle;
40  auto result = cuModuleGetFunction(&kernel_function_handle, handle_, name);
41  throw_if_error_lazy(result, ::std::string("Failed obtaining function ") + name
42  + " from " + module::detail_::identify(*this));
43  return kernel::wrap(context::detail_::get_device_id(context_handle_), context_handle_, kernel_function_handle);
44 }
45 
46 
47 namespace module {
48 
49 namespace detail_ {
50 
51 template <typename Creator>
52 module_t create(const context_t& context, const void* module_data, Creator creator_function)
53 {
54  CAW_SET_SCOPE_CONTEXT(context.handle());
55  handle_t new_module_handle;
56  auto status = creator_function(new_module_handle, module_data);
57  throw_if_error_lazy(status, ::std::string("Failed loading a module from memory location ")
58  + cuda::detail_::ptr_as_hex(module_data)
59  + " within " + context::detail_::identify(context));
60  bool do_take_ownership { true };
61  bool doesnt_hold_pc_refcount_unit { false };
62  // TODO: Do we want to allow holding a refcount unit here, if context is
63  // the primary context?
64 
65  // TODO: Make sure the default-constructed options correspond to what cuModuleLoadData uses as defaults
66  return detail_::wrap(
67  context.device_id(), context.handle(), new_module_handle,
68  do_take_ownership, doesnt_hold_pc_refcount_unit);
69 }
70 
71 // TODO: Consider adding create_module() methods to context_t
72 inline module_t create(const context_t& context, const void* module_data, const link::options_t& link_options)
73 {
74  auto creator_function =
75  [&link_options](handle_t& new_module_handle, const void* module_data_) {
76  auto marshalled_options = link::detail_::marshal(link_options);
77  return cuModuleLoadDataEx(
78  &new_module_handle,
79  module_data_,
80  marshalled_options.count(),
81  const_cast<link::detail_::option_t *>(marshalled_options.options()),
82  const_cast<void **>(marshalled_options.values())
83  );
84  };
85  return detail_::create(context, module_data, creator_function);
86 }
87 
88 inline module_t create(const context_t& context, const void* module_data)
89 {
90  auto creator_function =
91  [](handle_t& new_module_handle, const void* module_data_) {
92  return cuModuleLoadData(&new_module_handle, module_data_);
93  };
94  return detail_::create(context, module_data, creator_function);
95 }
96 
97 inline device::primary_context_t get_context_for(const device_t& locus) { return locus.primary_context(); }
98 
99 } // namespace detail_
100 
102  const device_t& device,
103  const char* path)
104 {
105  auto pc = device.primary_context();
106  device::primary_context::detail_::increase_refcount(device.id());
107  return load_from_file(pc, path);
108 }
109 
110 inline module_t load_from_file(const char* path)
111 {
112  return load_from_file(device::current::get(), path);
113 }
114 
115 } // namespace module
116 
117 inline context_t module_t::context() const { return context::detail_::from_handle(context_handle_); }
118 inline device_t module_t::device() const { return device::get(context::detail_::get_device_id(context_handle_)); }
119 
120 #if CUDA_VERSION < 12000
121 inline CUsurfref module_t::get_surface(const char* name) const
122 {
123  CAW_SET_SCOPE_CONTEXT(context_handle_);
124  CUsurfref raw_surface_reference;
125  auto status = cuModuleGetSurfRef(&raw_surface_reference, handle_, name);
126  throw_if_error_lazy(status, ::std::string("Failed obtaining a reference to surface \"") + name + "\" from "
127  + module::detail_::identify(*this));
128  return raw_surface_reference;
129 }
130 
131 inline CUtexref module_t::get_texture_reference(const char* name) const
132 {
133  CAW_SET_SCOPE_CONTEXT(context_handle_);
134  CUtexref raw_texture_reference;
135  auto status = cuModuleGetTexRef(&raw_texture_reference, handle_, name);
136  throw_if_error_lazy(status, ::std::string("Failed obtaining a reference to texture \"") + name + "\" from "
137  + module::detail_::identify(*this));
138  return raw_texture_reference;
139 }
140 #endif
141 
142 
143 } // namespace cuda
144 
145 #endif // MULTI_WRAPPER_IMPLS_MODULE_HPP_
146 
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:121
Wrapper class for a CUDA context.
Definition: context.hpp:249
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:122
cuda::kernel_t get_kernel(const char *name) const
Obtains a kernel constituting part of this module.
Definition: module.hpp:36
Wrapper class for a CUDA code module.
Definition: module.hpp:126
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:321
CUtexref get_texture_reference(const char *name) const
A "plug" of a method regarding texture references.
Definition: module.hpp:131
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:832
#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:327
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:30
context_t context() const
Definition: module.hpp:117
device_t device() const
Definition: module.hpp:118
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