12 #ifndef MULTI_WRAPPER_IMPLS_MODULE_HPP_ 13 #define MULTI_WRAPPER_IMPLS_MODULE_HPP_ 15 #include "../device.hpp" 16 #include "../module.hpp" 21 template <
typename ContiguousContainer,
22 cuda::detail_::enable_if_t<detail_::is_kinda_like_contiguous_container<ContiguousContainer>::value,
bool>>
25 return module::create<context_t const &>(*
this, module_data);
28 template <
typename ContiguousContainer,
29 cuda::detail_::enable_if_t<detail_::is_kinda_like_contiguous_container<ContiguousContainer>::value,
bool>>
32 return module::create<context_t const &>(*
this, module_data, link_options);
38 CAW_SET_SCOPE_CONTEXT(context_handle_);
39 kernel::handle_t kernel_function_handle;
40 auto result = cuModuleGetFunction(&kernel_function_handle, handle_, name);
42 +
" from " + module::detail_::identify(*
this));
43 return kernel::wrap(context::detail_::get_device_id(context_handle_), context_handle_, kernel_function_handle);
51 template <
typename Creator>
52 module_t create(
const context_t& context,
const void* module_data, Creator creator_function)
54 CAW_SET_SCOPE_CONTEXT(context.handle());
55 handle_t new_module_handle;
56 auto status = creator_function(new_module_handle, module_data);
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 };
67 context.device_id(), context.handle(), new_module_handle,
68 do_take_ownership, doesnt_hold_pc_refcount_unit);
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(
80 marshalled_options.count(),
81 const_cast<link::detail_::option_t *
>(marshalled_options.options()),
82 const_cast<void **>(marshalled_options.values())
85 return detail_::create(context, module_data, creator_function);
90 auto creator_function =
91 [](handle_t& new_module_handle,
const void* module_data_) {
92 return cuModuleLoadData(&new_module_handle, module_data_);
94 return detail_::create(context, module_data, creator_function);
106 device::primary_context::detail_::increase_refcount(device.
id());
120 #if CUDA_VERSION < 12000 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;
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;
145 #endif // MULTI_WRAPPER_IMPLS_MODULE_HPP_ CUsurfref get_surface(const char *name) const
A "plug" of a method regarding surfaces, which modules support but our wrappers don't really cater to...
Definition: module.hpp:121
A convenience class for holding, setting and inspecting options for a CUDA binary code linking proces...
Definition: link_options.hpp:130
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'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'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'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