12 #ifndef MULTI_WRAPPER_IMPLS_MODULE_HPP_ 13 #define MULTI_WRAPPER_IMPLS_MODULE_HPP_ 16 #include "../device.hpp" 17 #include "../module.hpp" 22 template <
typename ContiguousContainer,
23 cuda::detail_::enable_if_t<detail_::is_kinda_like_contiguous_container<ContiguousContainer>::value,
bool>>
26 return module::create<context_t const &>(*
this, module_data);
29 template <
typename ContiguousContainer,
30 cuda::detail_::enable_if_t<detail_::is_kinda_like_contiguous_container<ContiguousContainer>::value,
bool>>
33 return module::create<context_t const &>(*
this, module_data, link_options);
39 CAW_SET_SCOPE_CONTEXT(context_handle_);
40 kernel::handle_t kernel_function_handle;
41 auto result = cuModuleGetFunction(&kernel_function_handle, handle_, name);
43 +
" from " + module::detail_::identify(*
this));
44 return kernel::wrap(context::detail_::get_device_id(context_handle_), context_handle_, kernel_function_handle);
52 template <
typename Creator>
53 module_t create(
const context_t& context,
const void* module_data, Creator creator_function)
55 CAW_SET_SCOPE_CONTEXT(context.handle());
56 handle_t new_module_handle;
57 auto status = creator_function(new_module_handle, module_data);
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 };
68 context.device_id(), context.handle(), new_module_handle,
69 do_take_ownership, doesnt_hold_pc_refcount_unit);
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(
81 marshalled_options.count(),
82 const_cast<link::detail_::option_t *
>(marshalled_options.options()),
83 const_cast<void **>(marshalled_options.values())
86 return detail_::create(context, module_data, creator_function);
91 auto creator_function =
92 [](handle_t& new_module_handle,
const void* module_data_) {
93 return cuModuleLoadData(&new_module_handle, module_data_);
95 return detail_::create(context, module_data, creator_function);
109 device::primary_context::detail_::increase_refcount(device.
id());
123 #if CUDA_VERSION < 12000 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;
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;
148 #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:124
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: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'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'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'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