codekingpro/portable-devtools
114k
1/*
2 * Copyright 1993-2023 NVIDIA Corporation. All rights reserved.
3 *
4 * NOTICE TO LICENSEE:
5 *
6 * This source code and/or documentation ("Licensed Deliverables") are
7 * subject to NVIDIA intellectual property rights under U.S. and
8 * international Copyright laws.
9 *
10 * These Licensed Deliverables contained herein is PROPRIETARY and
11 * CONFIDENTIAL to NVIDIA and is being provided under the terms and
12 * conditions of a form of NVIDIA software license agreement by and
13 * between NVIDIA and Licensee ("License Agreement") or electronically
14 * accepted by Licensee. Notwithstanding any terms or conditions to
15 * the contrary in the License Agreement, reproduction or disclosure
16 * of the Licensed Deliverables to any third party without the express
17 * written consent of NVIDIA is prohibited.
18 *
19 * NOTWITHSTANDING ANY TERMS OR CONDITIONS TO THE CONTRARY IN THE
20 * LICENSE AGREEMENT, NVIDIA MAKES NO REPRESENTATION ABOUT THE
21 * SUITABILITY OF THESE LICENSED DELIVERABLES FOR ANY PURPOSE. IT IS
22 * PROVIDED "AS IS" WITHOUT EXPRESS OR IMPLIED WARRANTY OF ANY KIND.
23 * NVIDIA DISCLAIMS ALL WARRANTIES WITH REGARD TO THESE LICENSED
24 * DELIVERABLES, INCLUDING ALL IMPLIED WARRANTIES OF MERCHANTABILITY,
25 * NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE.
26 * NOTWITHSTANDING ANY TERMS OR CONDITIONS TO THE CONTRARY IN THE
27 * LICENSE AGREEMENT, IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY
28 * SPECIAL, INDIRECT, INCIDENTAL, OR CONSEQUENTIAL DAMAGES, OR ANY
29 * DAMAGES WHATSOEVER RESULTING FROM LOSS OF USE, DATA OR PROFITS,
30 * WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE OR OTHER TORTIOUS
31 * ACTION, ARISING OUT OF OR IN CONNECTION WITH THE USE OR PERFORMANCE
32 * OF THESE LICENSED DELIVERABLES.
33 *
34 * U.S. Government End Users. These Licensed Deliverables are a
35 * "commercial item" as that term is defined at 48 C.F.R. 2.101 (OCT
36 * 1995), consisting of "commercial computer software" and "commercial
37 * computer software documentation" as such terms are used in 48
38 * C.F.R. 12.212 (SEPT 1995) and is provided to the U.S. Government
39 * only as a commercial end item. Consistent with 48 C.F.R.12.212 and
40 * 48 C.F.R. 227.7202-1 through 227.7202-4 (JUNE 1995), all
41 * U.S. Government End Users acquire the Licensed Deliverables with
42 * only those rights set forth herein.
43 *
44 * Any use of the Licensed Deliverables in individual and commercial
45 * software must include, in the user documentation and internal
46 * comments to the code, the above Disclaimer and U.S. Government End
47 * Users Notice.
48 */
49
50#if !defined(__CUDA_RUNTIME_H__)
51#define __CUDA_RUNTIME_H__
52
53#if !defined(__CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__)
54#define __CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__
55#define __UNDEF_CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS_CUDA_RUNTIME_H__
56#endif
57
58#define EXCLUDE_FROM_RTC
59#if defined(__GNUC__)
60#if defined(__clang__) || (!defined(__PGIC__) && (__GNUC__ > 4 || (__GNUC__ == 4 && __GNUC_MINOR__ >= 6)))
61#pragma GCC diagnostic push
62#endif
63#if defined(__clang__) || (!defined(__PGIC__) && (__GNUC__ > 4 || (__GNUC__ == 4 && __GNUC_MINOR__ >= 2)))
64#pragma GCC diagnostic ignored "-Wunused-function"
65#endif
66#elif defined(_MSC_VER)
67#pragma warning(push)
68#pragma warning(disable: 4820)
69#endif
70#ifdef __QNX__
71#if (__GNUC__ == 4 && __GNUC_MINOR__ >= 7)
72typedef unsigned size_t;
73#endif
74#endif
75#undef EXCLUDE_FROM_RTC
76/*******************************************************************************
77* *
78* *
79* *
80*******************************************************************************/
81
82#include "crt/host_config.h"
83
84/*******************************************************************************
85* *
86* *
87* *
88*******************************************************************************/
89
90#include "builtin_types.h"
91#include "library_types.h"
92#if !defined(__CUDACC_RTC__)
93#define EXCLUDE_FROM_RTC
94#include "channel_descriptor.h"
95#include "cuda_runtime_api.h"
96#include "driver_functions.h"
97#undef EXCLUDE_FROM_RTC
98#endif /* !__CUDACC_RTC__ */
99#include "crt/host_defines.h"
100#ifdef __CUDACC_RTC__
101#include "target"
102#endif /* defined(__CUDACC_RTC__) */
103
104
105#include "vector_functions.h"
106
107#if defined(__CUDACC__)
108
109#if defined(__CUDACC_RTC__)
110#include "nvrtc_device_runtime.h"
111#include "crt/device_functions.h"
112#include "crt/common_functions.h"
113#include "device_launch_parameters.h"
114
115#else /* !__CUDACC_RTC__ */
116#define EXCLUDE_FROM_RTC
117#include "crt/common_functions.h"
118#include "crt/device_functions.h"
119#include "device_launch_parameters.h"
120
121#if defined(__CUDACC_EXTENDED_LAMBDA__)
122#include <functional>
123#include <utility>
124struct __device_builtin__ __nv_lambda_preheader_injection { };
125#endif /* defined(__CUDACC_EXTENDED_LAMBDA__) */
126
127#undef EXCLUDE_FROM_RTC
128#endif /* __CUDACC_RTC__ */
129
130#endif /* __CUDACC__ */
131
132/** \cond impl_private */
133#if defined(__DOXYGEN_ONLY__) || defined(CUDA_ENABLE_DEPRECATED)
134#define __CUDA_DEPRECATED
135#elif defined(_MSC_VER)
136#define __CUDA_DEPRECATED __declspec(deprecated)
137#elif defined(__GNUC__)
138#define __CUDA_DEPRECATED __attribute__((deprecated))
139#else
140#define __CUDA_DEPRECATED
141#endif
142/** \endcond impl_private */
143
144#define EXCLUDE_FROM_RTC
145#if defined(__cplusplus) && !defined(__CUDACC_RTC__)
146
147#if __cplusplus >= 201103L || (defined(_MSC_VER) && (_MSC_VER >= 1900))
148#include <utility>
149#endif
150
151/*******************************************************************************
152* *
153* *
154* *
155*******************************************************************************/
156
157/**
158 * \addtogroup CUDART_HIGHLEVEL
159 * @{
160 */
161
162/**
163 *\brief Launches a device function
164 *
165 * The function invokes kernel \p func on \p gridDim (\p gridDim.x × \p gridDim.y
166 * × \p gridDim.z) grid of blocks. Each block contains \p blockDim (\p blockDim.x ×
167 * \p blockDim.y × \p blockDim.z) threads.
168 *
169 * If the kernel has N parameters the \p args should point to array of N pointers.
170 * Each pointer, from <tt>args[0]</tt> to <tt>args[N - 1]</tt>, point to the region
171 * of memory from which the actual parameter will be copied.
172 *
173 * \p sharedMem sets the amount of dynamic shared memory that will be available to
174 * each thread block.
175 *
176 * \p stream specifies a stream the invocation is associated to.
177 *
178 * \param func - Device function symbol
179 * \param gridDim - Grid dimentions
180 * \param blockDim - Block dimentions
181 * \param args - Arguments
182 * \param sharedMem - Shared memory (defaults to 0)
183 * \param stream - Stream identifier (defaults to NULL)
184 *
185 * \return
186 * ::cudaSuccess,
187 * ::cudaErrorInvalidDeviceFunction,
188 * ::cudaErrorInvalidConfiguration,
189 * ::cudaErrorLaunchFailure,
190 * ::cudaErrorLaunchTimeout,
191 * ::cudaErrorLaunchOutOfResources,
192 * ::cudaErrorSharedObjectInitFailed,
193 * ::cudaErrorInvalidPtx,
194 * ::cudaErrorUnsupportedPtxVersion,
195 * ::cudaErrorNoKernelImageForDevice,
196 * ::cudaErrorJitCompilerNotFound,
197 * ::cudaErrorJitCompilationDisabled
198 * \notefnerr
199 * \note_async
200 * \note_null_stream
201 * \note_init_rt
202 * \note_callback
203 *
204 * \ref ::cudaLaunchKernel(const void *func, dim3 gridDim, dim3 blockDim, void **args, size_t sharedMem, cudaStream_t stream) "cudaLaunchKernel (C API)"
205 */
206template<class T>
207static __inline__ __host__ cudaError_t cudaLaunchKernel(
208 const T *func,
209 dim3 gridDim,
210 dim3 blockDim,
211 void **args,
212 size_t sharedMem = 0,
213 cudaStream_t stream = 0
214)
215{
216 return ::cudaLaunchKernel((const void *)func, gridDim, blockDim, args, sharedMem, stream);
217}
218
219
220#if __cplusplus >= 201103L || (defined(_MSC_VER) && (_MSC_VER >= 1900)) || defined(__DOXYGEN_ONLY__)
221/**
222 * \brief Launches a CUDA function with launch-time configuration
223 *
224 * Invokes the kernel \p func on \p config->gridDim (\p config->gridDim.x
225 * × \p config->gridDim.y × \p config->gridDim.z) grid of blocks.
226 * Each block contains \p config->blockDim (\p config->blockDim.x ×
227 * \p config->blockDim.y × \p config->blockDim.z) threads.
228 *
229 * \p config->dynamicSmemBytes sets the amount of dynamic shared memory that
230 * will be available to each thread block.
231 *
232 * \p config->stream specifies a stream the invocation is associated to.
233 *
234 * Configuration beyond grid and block dimensions, dynamic shared memory size,
235 * and stream can be provided with the following two fields of \p config:
236 *
237 * \p config->attrs is an array of \p config->numAttrs contiguous
238 * ::cudaLaunchAttribute elements. The value of this pointer is not considered
239 * if \p config->numAttrs is zero. However, in that case, it is recommended to
240 * set the pointer to NULL.
241 * \p config->numAttrs is the number of attributes populating the first
242 * \p config->numAttrs positions of the \p config->attrs array.
243 *
244 * The kernel arguments should be passed as arguments to this function via the
245 * \p args parameter pack.
246 *
247 * The C API version of this function, \p cudaLaunchKernelExC, is also available
248 * for pre-C++11 compilers and for use cases where the ability to pass kernel
249 * parameters via void* array is preferable.
250 *
251 * \param config - Launch configuration
252 * \param func - Kernel to launch
253 * \param args - Parameter pack of kernel parameters
254 *
255 * \return
256 * ::cudaSuccess,
257 * ::cudaErrorInvalidDeviceFunction,
258 * ::cudaErrorInvalidConfiguration,
259 * ::cudaErrorLaunchFailure,
260 * ::cudaErrorLaunchTimeout,
261 * ::cudaErrorLaunchOutOfResources,
262 * ::cudaErrorSharedObjectInitFailed,
263 * ::cudaErrorInvalidPtx,
264 * ::cudaErrorUnsupportedPtxVersion,
265 * ::cudaErrorNoKernelImageForDevice,
266 * ::cudaErrorJitCompilerNotFound,
267 * ::cudaErrorJitCompilationDisabled
268 * \note_null_stream
269 * \notefnerr
270 * \note_init_rt
271 * \note_callback
272 *
273 * \sa
274 * \ref ::cudaLaunchKernelExC(const cudaLaunchConfig_t *config, const void *func, void **args) "cudaLaunchKernelEx (C API)",
275 * ::cuLaunchKernelEx
276 */
277template<typename... ExpTypes, typename... ActTypes>
278static __inline__ __host__ cudaError_t cudaLaunchKernelEx(
279 const cudaLaunchConfig_t *config,
280 void (*kernel)(ExpTypes...),
281 ActTypes &&... args
282)
283{
284 return [&](ExpTypes... coercedArgs){
285 void *pArgs[] = { &coercedArgs... };
286 return ::cudaLaunchKernelExC(config, (const void *)kernel, pArgs);
287 }(std::forward<ActTypes>(args)...);
288}
289#endif
290
291/**
292 *\brief Launches a device function
293 *
294 * The function invokes kernel \p func on \p gridDim (\p gridDim.x × \p gridDim.y
295 * × \p gridDim.z) grid of blocks. Each block contains \p blockDim (\p blockDim.x ×
296 * \p blockDim.y × \p blockDim.z) threads.
297 *
298 * The device on which this kernel is invoked must have a non-zero value for
299 * the device attribute ::cudaDevAttrCooperativeLaunch.
300 *
301 * The total number of blocks launched cannot exceed the maximum number of blocks per
302 * multiprocessor as returned by ::cudaOccupancyMaxActiveBlocksPerMultiprocessor (or
303 * ::cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags) times the number of multiprocessors
304 * as specified by the device attribute ::cudaDevAttrMultiProcessorCount.
305 *
306 * The kernel cannot make use of CUDA dynamic parallelism.
307 *
308 * If the kernel has N parameters the \p args should point to array of N pointers.
309 * Each pointer, from <tt>args[0]</tt> to <tt>args[N - 1]</tt>, point to the region
310 * of memory from which the actual parameter will be copied.
311 *
312 * \p sharedMem sets the amount of dynamic shared memory that will be available to
313 * each thread block.
314 *
315 * \p stream specifies a stream the invocation is associated to.
316 *
317 * \param func - Device function symbol
318 * \param gridDim - Grid dimentions
319 * \param blockDim - Block dimentions
320 * \param args - Arguments
321 * \param sharedMem - Shared memory (defaults to 0)
322 * \param stream - Stream identifier (defaults to NULL)
323 *
324 * \return
325 * ::cudaSuccess,
326 * ::cudaErrorInvalidDeviceFunction,
327 * ::cudaErrorInvalidConfiguration,
328 * ::cudaErrorLaunchFailure,
329 * ::cudaErrorLaunchTimeout,
330 * ::cudaErrorLaunchOutOfResources,
331 * ::cudaErrorSharedObjectInitFailed
332 * \notefnerr
333 * \note_async
334 * \note_null_stream
335 * \note_init_rt
336 * \note_callback
337 *
338 * \ref ::cudaLaunchCooperativeKernel(const void *func, dim3 gridDim, dim3 blockDim, void **args, size_t sharedMem, cudaStream_t stream) "cudaLaunchCooperativeKernel (C API)"
339 */
340template<class T>
341static __inline__ __host__ cudaError_t cudaLaunchCooperativeKernel(
342 const T *func,
343 dim3 gridDim,
344 dim3 blockDim,
345 void **args,
346 size_t sharedMem = 0,
347 cudaStream_t stream = 0
348)
349{
350 return ::cudaLaunchCooperativeKernel((const void *)func, gridDim, blockDim, args, sharedMem, stream);
351}
352
353/**
354 * \brief \hl Creates an event object with the specified flags
355 *
356 * Creates an event object with the specified flags. Valid flags include:
357 * - ::cudaEventDefault: Default event creation flag.
358 * - ::cudaEventBlockingSync: Specifies that event should use blocking
359 * synchronization. A host thread that uses ::cudaEventSynchronize() to wait
360 * on an event created with this flag will block until the event actually
361 * completes.
362 * - ::cudaEventDisableTiming: Specifies that the created event does not need
363 * to record timing data. Events created with this flag specified and
364 * the ::cudaEventBlockingSync flag not specified will provide the best
365 * performance when used with ::cudaStreamWaitEvent() and ::cudaEventQuery().
366 *
367 * \param event - Newly created event
368 * \param flags - Flags for new event
369 *
370 * \return
371 * ::cudaSuccess,
372 * ::cudaErrorInvalidValue,
373 * ::cudaErrorLaunchFailure,
374 * ::cudaErrorMemoryAllocation
375 * \notefnerr
376 * \note_init_rt
377 * \note_callback
378 *
379 * \sa \ref ::cudaEventCreate(cudaEvent_t*) "cudaEventCreate (C API)",
380 * ::cudaEventCreateWithFlags, ::cudaEventRecord, ::cudaEventQuery,
381 * ::cudaEventSynchronize, ::cudaEventDestroy, ::cudaEventElapsedTime,
382 * ::cudaStreamWaitEvent
383 */
384static __inline__ __host__ cudaError_t cudaEventCreate(
385 cudaEvent_t *event,
386 unsigned int flags
387)
388{
389 return ::cudaEventCreateWithFlags(event, flags);
390}
391
392/**
393 * \brief Creates an executable graph from a graph
394 *
395 * Instantiates \p graph as an executable graph. The graph is validated for any
396 * structural constraints or intra-node constraints which were not previously
397 * validated. If instantiation is successful, a handle to the instantiated graph
398 * is returned in \p pGraphExec.
399 *
400 * If there are any errors, diagnostic information may be returned in \p pErrorNode and
401 * \p pLogBuffer. This is the primary way to inspect instantiation errors. The output
402 * will be null terminated unless the diagnostics overflow
403 * the buffer. In this case, they will be truncated, and the last byte can be
404 * inspected to determine if truncation occurred.
405 *
406 * \param pGraphExec - Returns instantiated graph
407 * \param graph - Graph to instantiate
408 * \param pErrorNode - In case of an instantiation error, this may be modified to
409 * indicate a node contributing to the error
410 * \param pLogBuffer - A character buffer to store diagnostic messages
411 * \param bufferSize - Size of the log buffer in bytes
412 *
413 * \return
414 * ::cudaSuccess,
415 * ::cudaErrorInvalidValue
416 * \note_graph_thread_safety
417 * \notefnerr
418 * \note_init_rt
419 * \note_callback
420 *
421 * \sa
422 * ::cudaGraphInstantiateWithFlags,
423 * ::cudaGraphCreate,
424 * ::cudaGraphUpload,
425 * ::cudaGraphLaunch,
426 * ::cudaGraphExecDestroy
427 */
428static __inline__ __host__ cudaError_t cudaGraphInstantiate(
429 cudaGraphExec_t *pGraphExec,
430 cudaGraph_t graph,
431 cudaGraphNode_t *pErrorNode,
432 char *pLogBuffer,
433 size_t bufferSize
434)
435{
436 (void)pErrorNode;
437 (void)pLogBuffer;
438 (void)bufferSize;
439 return ::cudaGraphInstantiate(pGraphExec, graph, 0);
440}
441
442/**
443 * \brief \hl Allocates page-locked memory on the host
444 *
445 * Allocates \p size bytes of host memory that is page-locked and accessible
446 * to the device. The driver tracks the virtual memory ranges allocated with
447 * this function and automatically accelerates calls to functions such as
448 * ::cudaMemcpy(). Since the memory can be accessed directly by the device, it
449 * can be read or written with much higher bandwidth than pageable memory
450 * obtained with functions such as ::malloc(). Allocating excessive amounts of
451 * pinned memory may degrade system performance, since it reduces the amount
452 * of memory available to the system for paging. As a result, this function is
453 * best used sparingly to allocate staging areas for data exchange between host
454 * and device.
455 *
456 * The \p flags parameter enables different options to be specified that affect
457 * the allocation, as follows.
458 * - ::cudaHostAllocDefault: This flag's value is defined to be 0.
459 * - ::cudaHostAllocPortable: The memory returned by this call will be
460 * considered as pinned memory by all CUDA contexts, not just the one that
461 * performed the allocation.
462 * - ::cudaHostAllocMapped: Maps the allocation into the CUDA address space.
463 * The device pointer to the memory may be obtained by calling
464 * ::cudaHostGetDevicePointer().
465 * - ::cudaHostAllocWriteCombined: Allocates the memory as write-combined (WC).
466 * WC memory can be transferred across the PCI Express bus more quickly on some
467 * system configurations, but cannot be read efficiently by most CPUs. WC
468 * memory is a good option for buffers that will be written by the CPU and read
469 * by the device via mapped pinned memory or host->device transfers.
470 *
471 * All of these flags are orthogonal to one another: a developer may allocate
472 * memory that is portable, mapped and/or write-combined with no restrictions.
473 *
474 * ::cudaSetDeviceFlags() must have been called with the ::cudaDeviceMapHost
475 * flag in order for the ::cudaHostAllocMapped flag to have any effect.
476 *
477 * The ::cudaHostAllocMapped flag may be specified on CUDA contexts for devices
478 * that do not support mapped pinned memory. The failure is deferred to
479 * ::cudaHostGetDevicePointer() because the memory may be mapped into other
480 * CUDA contexts via the ::cudaHostAllocPortable flag.
481 *
482 * Memory allocated by this function must be freed with ::cudaFreeHost().
483 *
484 * \param ptr - Device pointer to allocated memory
485 * \param size - Requested allocation size in bytes
486 * \param flags - Requested properties of allocated memory
487 *
488 * \return
489 * ::cudaSuccess,
490 * ::cudaErrorMemoryAllocation
491 * \notefnerr
492 * \note_init_rt
493 * \note_callback
494 *
495 * \sa ::cudaSetDeviceFlags,
496 * \ref ::cudaMallocHost(void**, size_t) "cudaMallocHost (C API)",
497 * ::cudaFreeHost, ::cudaHostAlloc
498 */
499static __inline__ __host__ cudaError_t cudaMallocHost(
500 void **ptr,
501 size_t size,
502 unsigned int flags
503)
504{
505 return ::cudaHostAlloc(ptr, size, flags);
506}
507
508template<class T>
509static __inline__ __host__ cudaError_t cudaHostAlloc(
510 T **ptr,
511 size_t size,
512 unsigned int flags
513)
514{
515 return ::cudaHostAlloc((void**)(void*)ptr, size, flags);
516}
517
518template<class T>
519static __inline__ __host__ cudaError_t cudaHostGetDevicePointer(
520 T **pDevice,
521 void *pHost,
522 unsigned int flags
523)
524{
525 return ::cudaHostGetDevicePointer((void**)(void*)pDevice, pHost, flags);
526}
527
528/**
529 * \brief Allocates memory that will be automatically managed by the Unified Memory system
530 *
531 * Allocates \p size bytes of managed memory on the device and returns in
532 * \p *devPtr a pointer to the allocated memory. If the device doesn't support
533 * allocating managed memory, ::cudaErrorNotSupported is returned. Support
534 * for managed memory can be queried using the device attribute
535 * ::cudaDevAttrManagedMemory. The allocated memory is suitably
536 * aligned for any kind of variable. The memory is not cleared. If \p size
537 * is 0, ::cudaMallocManaged returns ::cudaErrorInvalidValue. The pointer
538 * is valid on the CPU and on all GPUs in the system that support managed memory.
539 * All accesses to this pointer must obey the Unified Memory programming model.
540 *
541 * \p flags specifies the default stream association for this allocation.
542 * \p flags must be one of ::cudaMemAttachGlobal or ::cudaMemAttachHost. The
543 * default value for \p flags is ::cudaMemAttachGlobal.
544 * If ::cudaMemAttachGlobal is specified, then this memory is accessible from
545 * any stream on any device. If ::cudaMemAttachHost is specified, then the
546 * allocation should not be accessed from devices that have a zero value for the
547 * device attribute ::cudaDevAttrConcurrentManagedAccess; an explicit call to
548 * ::cudaStreamAttachMemAsync will be required to enable access on such devices.
549 *
550 * If the association is later changed via ::cudaStreamAttachMemAsync to
551 * a single stream, the default association, as specifed during ::cudaMallocManaged,
552 * is restored when that stream is destroyed. For __managed__ variables, the
553 * default association is always ::cudaMemAttachGlobal. Note that destroying a
554 * stream is an asynchronous operation, and as a result, the change to default
555 * association won't happen until all work in the stream has completed.
556 *
557 * Memory allocated with ::cudaMallocManaged should be released with ::cudaFree.
558 *
559 * Device memory oversubscription is possible for GPUs that have a non-zero value for the
560 * device attribute ::cudaDevAttrConcurrentManagedAccess. Managed memory on
561 * such GPUs may be evicted from device memory to host memory at any time by the Unified
562 * Memory driver in order to make room for other allocations.
563 *
564 * In a multi-GPU system where all GPUs have a non-zero value for the device attribute
565 * ::cudaDevAttrConcurrentManagedAccess, managed memory may not be populated when this
566 * API returns and instead may be populated on access. In such systems, managed memory can
567 * migrate to any processor's memory at any time. The Unified Memory driver will employ heuristics to
568 * maintain data locality and prevent excessive page faults to the extent possible. The application
569 * can also guide the driver about memory usage patterns via ::cudaMemAdvise. The application
570 * can also explicitly migrate memory to a desired processor's memory via
571 * ::cudaMemPrefetchAsync.
572 *
573 * In a multi-GPU system where all of the GPUs have a zero value for the device attribute
574 * ::cudaDevAttrConcurrentManagedAccess and all the GPUs have peer-to-peer support
575 * with each other, the physical storage for managed memory is created on the GPU which is active
576 * at the time ::cudaMallocManaged is called. All other GPUs will reference the data at reduced
577 * bandwidth via peer mappings over the PCIe bus. The Unified Memory driver does not migrate
578 * memory among such GPUs.
579 *
580 * In a multi-GPU system where not all GPUs have peer-to-peer support with each other and
581 * where the value of the device attribute ::cudaDevAttrConcurrentManagedAccess
582 * is zero for at least one of those GPUs, the location chosen for physical storage of managed
583 * memory is system-dependent.
584 * - On Linux, the location chosen will be device memory as long as the current set of active
585 * contexts are on devices that either have peer-to-peer support with each other or have a
586 * non-zero value for the device attribute ::cudaDevAttrConcurrentManagedAccess.
587 * If there is an active context on a GPU that does not have a non-zero value for that device
588 * attribute and it does not have peer-to-peer support with the other devices that have active
589 * contexts on them, then the location for physical storage will be 'zero-copy' or host memory.
590 * Note that this means that managed memory that is located in device memory is migrated to
591 * host memory if a new context is created on a GPU that doesn't have a non-zero value for
592 * the device attribute and does not support peer-to-peer with at least one of the other devices
593 * that has an active context. This in turn implies that context creation may fail if there is
594 * insufficient host memory to migrate all managed allocations.
595 * - On Windows, the physical storage is always created in 'zero-copy' or host memory.
596 * All GPUs will reference the data at reduced bandwidth over the PCIe bus. In these
597 * circumstances, use of the environment variable CUDA_VISIBLE_DEVICES is recommended to
598 * restrict CUDA to only use those GPUs that have peer-to-peer support.
599 * Alternatively, users can also set CUDA_MANAGED_FORCE_DEVICE_ALLOC to a non-zero
600 * value to force the driver to always use device memory for physical storage.
601 * When this environment variable is set to a non-zero value, all devices used in
602 * that process that support managed memory have to be peer-to-peer compatible
603 * with each other. The error ::cudaErrorInvalidDevice will be returned if a device
604 * that supports managed memory is used and it is not peer-to-peer compatible with
605 * any of the other managed memory supporting devices that were previously used in
606 * that process, even if ::cudaDeviceReset has been called on those devices. These
607 * environment variables are described in the CUDA programming guide under the
608 * "CUDA environment variables" section.
609 * - On ARM, managed memory is not available on discrete gpu with Drive PX-2.
610 *
611 * \param devPtr - Pointer to allocated device memory
612 * \param size - Requested allocation size in bytes
613 * \param flags - Must be either ::cudaMemAttachGlobal or ::cudaMemAttachHost (defaults to ::cudaMemAttachGlobal)
614 *
615 * \return
616 * ::cudaSuccess,
617 * ::cudaErrorMemoryAllocation,
618 * ::cudaErrorNotSupported,
619 * ::cudaErrorInvalidValue
620 * \note_init_rt
621 * \note_callback
622 *
623 * \sa ::cudaMallocPitch, ::cudaFree, ::cudaMallocArray, ::cudaFreeArray,
624 * ::cudaMalloc3D, ::cudaMalloc3DArray,
625 * \ref ::cudaMallocHost(void**, size_t) "cudaMallocHost (C API)",
626 * ::cudaFreeHost, ::cudaHostAlloc, ::cudaDeviceGetAttribute, ::cudaStreamAttachMemAsync
627 */
628template<class T>
629static __inline__ __host__ cudaError_t cudaMallocManaged(
630 T **devPtr,
631 size_t size,
632 unsigned int flags = cudaMemAttachGlobal
633)
634{
635 return ::cudaMallocManaged((void**)(void*)devPtr, size, flags);
636}
637
638/**
639 * \brief Advise about the usage of a given memory range.
640 *
641 * This is an alternate spelling for cudaMemAdvise made available through operator overloading.
642 *
643 * \sa ::cudaMemAdvise,
644 * \ref ::cudaMemAdvise(const void* devPtr, size_t count, enum cudaMemoryAdvise advice, struct cudaMemLocation location) "cudaMemAdvise (C API)"
645 */
646template<class T>
647cudaError_t cudaMemAdvise(
648 T *devPtr,
649 size_t count,
650 enum cudaMemoryAdvise advice,
651 struct cudaMemLocation location
652)
653{
654 return ::cudaMemAdvise_v2((const void *)devPtr, count, advice, location);
655}
656
657template<class T>
658static __inline__ __host__ cudaError_t cudaMemPrefetchAsync(
659 T *devPtr,
660 size_t count,
661 struct cudaMemLocation location,
662 unsigned int flags,
663 cudaStream_t stream = 0
664)
665{
666 return ::cudaMemPrefetchAsync_v2((const void *)devPtr, count, location, flags, stream);
667}
668
669/**
670 * \brief Attach memory to a stream asynchronously
671 *
672 * Enqueues an operation in \p stream to specify stream association of
673 * \p length bytes of memory starting from \p devPtr. This function is a
674 * stream-ordered operation, meaning that it is dependent on, and will
675 * only take effect when, previous work in stream has completed. Any
676 * previous association is automatically replaced.
677 *
678 * \p devPtr must point to an one of the following types of memories:
679 * - managed memory declared using the __managed__ keyword or allocated with
680 * ::cudaMallocManaged.
681 * - a valid host-accessible region of system-allocated pageable memory. This
682 * type of memory may only be specified if the device associated with the
683 * stream reports a non-zero value for the device attribute
684 * ::cudaDevAttrPageableMemoryAccess.
685 *
686 * For managed allocations, \p length must be either zero or the entire
687 * allocation's size. Both indicate that the entire allocation's stream
688 * association is being changed. Currently, it is not possible to change stream
689 * association for a portion of a managed allocation.
690 *
691 * For pageable allocations, \p length must be non-zero.
692 *
693 * The stream association is specified using \p flags which must be
694 * one of ::cudaMemAttachGlobal, ::cudaMemAttachHost or ::cudaMemAttachSingle.
695 * The default value for \p flags is ::cudaMemAttachSingle
696 * If the ::cudaMemAttachGlobal flag is specified, the memory can be accessed
697 * by any stream on any device.
698 * If the ::cudaMemAttachHost flag is specified, the program makes a guarantee
699 * that it won't access the memory on the device from any stream on a device that
700 * has a zero value for the device attribute ::cudaDevAttrConcurrentManagedAccess.
701 * If the ::cudaMemAttachSingle flag is specified and \p stream is associated with
702 * a device that has a zero value for the device attribute ::cudaDevAttrConcurrentManagedAccess,
703 * the program makes a guarantee that it will only access the memory on the device
704 * from \p stream. It is illegal to attach singly to the NULL stream, because the
705 * NULL stream is a virtual global stream and not a specific stream. An error will
706 * be returned in this case.
707 *
708 * When memory is associated with a single stream, the Unified Memory system will
709 * allow CPU access to this memory region so long as all operations in \p stream
710 * have completed, regardless of whether other streams are active. In effect,
711 * this constrains exclusive ownership of the managed memory region by
712 * an active GPU to per-stream activity instead of whole-GPU activity.
713 *
714 * Accessing memory on the device from streams that are not associated with
715 * it will produce undefined results. No error checking is performed by the
716 * Unified Memory system to ensure that kernels launched into other streams
717 * do not access this region.
718 *
719 * It is a program's responsibility to order calls to ::cudaStreamAttachMemAsync
720 * via events, synchronization or other means to ensure legal access to memory
721 * at all times. Data visibility and coherency will be changed appropriately
722 * for all kernels which follow a stream-association change.
723 *
724 * If \p stream is destroyed while data is associated with it, the association is
725 * removed and the association reverts to the default visibility of the allocation
726 * as specified at ::cudaMallocManaged. For __managed__ variables, the default
727 * association is always ::cudaMemAttachGlobal. Note that destroying a stream is an
728 * asynchronous operation, and as a result, the change to default association won't
729 * happen until all work in the stream has completed.
730 *
731 * \param stream - Stream in which to enqueue the attach operation
732 * \param devPtr - Pointer to memory (must be a pointer to managed memory or
733 * to a valid host-accessible region of system-allocated
734 * memory)
735 * \param length - Length of memory (defaults to zero)
736 * \param flags - Must be one of ::cudaMemAttachGlobal, ::cudaMemAttachHost or ::cudaMemAttachSingle (defaults to ::cudaMemAttachSingle)
737 *
738 * \return
739 * ::cudaSuccess,
740 * ::cudaErrorNotReady,
741 * ::cudaErrorInvalidValue,
742 * ::cudaErrorInvalidResourceHandle
743 * \notefnerr
744 * \note_init_rt
745 * \note_callback
746 *
747 * \sa ::cudaStreamCreate, ::cudaStreamCreateWithFlags, ::cudaStreamWaitEvent, ::cudaStreamSynchronize, ::cudaStreamAddCallback, ::cudaStreamDestroy, ::cudaMallocManaged
748 */
749template<class T>
750static __inline__ __host__ cudaError_t cudaStreamAttachMemAsync(
751 cudaStream_t stream,
752 T *devPtr,
753 size_t length = 0,
754 unsigned int flags = cudaMemAttachSingle
755)
756{
757 return ::cudaStreamAttachMemAsync(stream, (void*)devPtr, length, flags);
758}
759
760template<class T>
761static __inline__ __host__ cudaError_t cudaMalloc(
762 T **devPtr,
763 size_t size
764)
765{
766 return ::cudaMalloc((void**)(void*)devPtr, size);
767}
768
769template<class T>
770static __inline__ __host__ cudaError_t cudaMallocHost(
771 T **ptr,
772 size_t size,
773 unsigned int flags = 0
774)
775{
776 return cudaMallocHost((void**)(void*)ptr, size, flags);
777}
778
779template<class T>
780static __inline__ __host__ cudaError_t cudaMallocPitch(
781 T **devPtr,
782 size_t *pitch,
783 size_t width,
784 size_t height
785)
786{
787 return ::cudaMallocPitch((void**)(void*)devPtr, pitch, width, height);
788}
789
790/**
791 * \brief Allocate from a pool
792 *
793 * This is an alternate spelling for cudaMallocFromPoolAsync
794 * made available through operator overloading.
795 *
796 * \sa ::cudaMallocFromPoolAsync,
797 * \ref ::cudaMallocAsync(void** ptr, size_t size, cudaStream_t hStream) "cudaMallocAsync (C API)"
798 */
799static __inline__ __host__ cudaError_t cudaMallocAsync(
800 void **ptr,
801 size_t size,
802 cudaMemPool_t memPool,
803 cudaStream_t stream
804)
805{
806 return ::cudaMallocFromPoolAsync(ptr, size, memPool, stream);
807}
808
809template<class T>
810static __inline__ __host__ cudaError_t cudaMallocAsync(
811 T **ptr,
812 size_t size,
813 cudaMemPool_t memPool,
814 cudaStream_t stream
815)
816{
817 return ::cudaMallocFromPoolAsync((void**)(void*)ptr, size, memPool, stream);
818}
819
820template<class T>
821static __inline__ __host__ cudaError_t cudaMallocAsync(
822 T **ptr,
823 size_t size,
824 cudaStream_t stream
825)
826{
827 return ::cudaMallocAsync((void**)(void*)ptr, size, stream);
828}
829
830template<class T>
831static __inline__ __host__ cudaError_t cudaMallocFromPoolAsync(
832 T **ptr,
833 size_t size,
834 cudaMemPool_t memPool,
835 cudaStream_t stream
836)
837{
838 return ::cudaMallocFromPoolAsync((void**)(void*)ptr, size, memPool, stream);
839}
840
841#if defined(__CUDACC__)
842
843/**
844 * \brief \hl Copies data to the given symbol on the device
845 *
846 * Copies \p count bytes from the memory area pointed to by \p src
847 * to the memory area \p offset bytes from the start of symbol
848 * \p symbol. The memory areas may not overlap. \p symbol is a variable that
849 * resides in global or constant memory space. \p kind can be either
850 * ::cudaMemcpyHostToDevice or ::cudaMemcpyDeviceToDevice.
851 *
852 * \param symbol - Device symbol reference
853 * \param src - Source memory address
854 * \param count - Size in bytes to copy
855 * \param offset - Offset from start of symbol in bytes
856 * \param kind - Type of transfer
857 *
858 * \return
859 * ::cudaSuccess,
860 * ::cudaErrorInvalidValue,
861 * ::cudaErrorInvalidSymbol,
862 * ::cudaErrorInvalidMemcpyDirection,
863 * ::cudaErrorNoKernelImageForDevice
864 * \notefnerr
865 * \note_sync
866 * \note_string_api_deprecation
867 * \note_init_rt
868 * \note_callback
869 *
870 * \sa ::cudaMemcpy, ::cudaMemcpy2D,
871 * ::cudaMemcpy2DToArray, ::cudaMemcpy2DFromArray,
872 * ::cudaMemcpy2DArrayToArray,
873 * ::cudaMemcpyFromSymbol, ::cudaMemcpyAsync, ::cudaMemcpy2DAsync,
874 * ::cudaMemcpy2DToArrayAsync,
875 * ::cudaMemcpy2DFromArrayAsync,
876 * ::cudaMemcpyToSymbolAsync, ::cudaMemcpyFromSymbolAsync
877 */
878template<class T>
879static __inline__ __host__ cudaError_t cudaMemcpyToSymbol(
880 const T &symbol,
881 const void *src,
882 size_t count,
883 size_t offset = 0,
884 enum cudaMemcpyKind kind = cudaMemcpyHostToDevice
885)
886{
887 return ::cudaMemcpyToSymbol((const void*)&symbol, src, count, offset, kind);
888}
889
890/**
891 * \brief \hl Copies data to the given symbol on the device
892 *
893 * Copies \p count bytes from the memory area pointed to by \p src
894 * to the memory area \p offset bytes from the start of symbol
895 * \p symbol. The memory areas may not overlap. \p symbol is a variable that
896 * resides in global or constant memory space. \p kind can be either
897 * ::cudaMemcpyHostToDevice or ::cudaMemcpyDeviceToDevice.
898 *
899 * ::cudaMemcpyToSymbolAsync() is asynchronous with respect to the host, so
900 * the call may return before the copy is complete. The copy can optionally
901 * be associated to a stream by passing a non-zero \p stream argument. If
902 * \p kind is ::cudaMemcpyHostToDevice and \p stream is non-zero, the copy
903 * may overlap with operations in other streams.
904 *
905 * \param symbol - Device symbol reference
906 * \param src - Source memory address
907 * \param count - Size in bytes to copy
908 * \param offset - Offset from start of symbol in bytes
909 * \param kind - Type of transfer
910 * \param stream - Stream identifier
911 *
912 * \return
913 * ::cudaSuccess,
914 * ::cudaErrorInvalidValue,
915 * ::cudaErrorInvalidSymbol,
916 * ::cudaErrorInvalidMemcpyDirection,
917 * ::cudaErrorNoKernelImageForDevice
918 * \notefnerr
919 * \note_async
920 * \note_string_api_deprecation
921 * \note_init_rt
922 * \note_callback
923 *
924 * \sa ::cudaMemcpy, ::cudaMemcpy2D,
925 * ::cudaMemcpy2DToArray, ::cudaMemcpy2DFromArray,
926 * ::cudaMemcpy2DArrayToArray, ::cudaMemcpyToSymbol,
927 * ::cudaMemcpyFromSymbol, ::cudaMemcpyAsync, ::cudaMemcpy2DAsync,
928 * ::cudaMemcpy2DToArrayAsync,
929 * ::cudaMemcpy2DFromArrayAsync,
930 * ::cudaMemcpyFromSymbolAsync
931 */
932template<class T>
933static __inline__ __host__ cudaError_t cudaMemcpyToSymbolAsync(
934 const T &symbol,
935 const void *src,
936 size_t count,
937 size_t offset = 0,
938 enum cudaMemcpyKind kind = cudaMemcpyHostToDevice,
939 cudaStream_t stream = 0
940)
941{
942 return ::cudaMemcpyToSymbolAsync((const void*)&symbol, src, count, offset, kind, stream);
943}
944
945/**
946 * \brief \hl Copies data from the given symbol on the device
947 *
948 * Copies \p count bytes from the memory area \p offset bytes
949 * from the start of symbol \p symbol to the memory area pointed to by \p dst.
950 * The memory areas may not overlap. \p symbol is a variable that
951 * resides in global or constant memory space. \p kind can be either
952 * ::cudaMemcpyDeviceToHost or ::cudaMemcpyDeviceToDevice.
953 *
954 * \param dst - Destination memory address
955 * \param symbol - Device symbol reference
956 * \param count - Size in bytes to copy
957 * \param offset - Offset from start of symbol in bytes
958 * \param kind - Type of transfer
959 *
960 * \return
961 * ::cudaSuccess,
962 * ::cudaErrorInvalidValue,
963 * ::cudaErrorInvalidSymbol,
964 * ::cudaErrorInvalidMemcpyDirection,
965 * ::cudaErrorNoKernelImageForDevice
966 * \notefnerr
967 * \note_sync
968 * \note_string_api_deprecation
969 * \note_init_rt
970 * \note_callback
971 *
972 * \sa ::cudaMemcpy, ::cudaMemcpy2D,
973 * ::cudaMemcpy2DToArray, ::cudaMemcpy2DFromArray,
974 * ::cudaMemcpy2DArrayToArray, ::cudaMemcpyToSymbol,
975 * ::cudaMemcpyAsync, ::cudaMemcpy2DAsync,
976 * ::cudaMemcpy2DToArrayAsync,
977 * ::cudaMemcpy2DFromArrayAsync,
978 * ::cudaMemcpyToSymbolAsync, ::cudaMemcpyFromSymbolAsync
979 */
980template<class T>
981static __inline__ __host__ cudaError_t cudaMemcpyFromSymbol(
982 void *dst,
983 const T &symbol,
984 size_t count,
985 size_t offset = 0,
986 enum cudaMemcpyKind kind = cudaMemcpyDeviceToHost
987)
988{
989 return ::cudaMemcpyFromSymbol(dst, (const void*)&symbol, count, offset, kind);
990}
991
992/**
993 * \brief \hl Copies data from the given symbol on the device
994 *
995 * Copies \p count bytes from the memory area \p offset bytes
996 * from the start of symbol \p symbol to the memory area pointed to by \p dst.
997 * The memory areas may not overlap. \p symbol is a variable that resides in
998 * global or constant memory space. \p kind can be either
999 * ::cudaMemcpyDeviceToHost or ::cudaMemcpyDeviceToDevice.
1000 *
1001 * ::cudaMemcpyFromSymbolAsync() is asynchronous with respect to the host, so
1002 * the call may return before the copy is complete. The copy can optionally be
1003 * associated to a stream by passing a non-zero \p stream argument. If \p kind
1004 * is ::cudaMemcpyDeviceToHost and \p stream is non-zero, the copy may overlap
1005 * with operations in other streams.
1006 *
1007 * \param dst - Destination memory address
1008 * \param symbol - Device symbol reference
1009 * \param count - Size in bytes to copy
1010 * \param offset - Offset from start of symbol in bytes
1011 * \param kind - Type of transfer
1012 * \param stream - Stream identifier
1013 *
1014 * \return
1015 * ::cudaSuccess,
1016 * ::cudaErrorInvalidValue,
1017 * ::cudaErrorInvalidSymbol,
1018 * ::cudaErrorInvalidMemcpyDirection,
1019 * ::cudaErrorNoKernelImageForDevice
1020 * \notefnerr
1021 * \note_async
1022 * \note_string_api_deprecation
1023 * \note_init_rt
1024 * \note_callback
1025 *
1026 * \sa ::cudaMemcpy, ::cudaMemcpy2D,
1027 * ::cudaMemcpy2DToArray, ::cudaMemcpy2DFromArray,
1028 * ::cudaMemcpy2DArrayToArray, ::cudaMemcpyToSymbol,
1029 * ::cudaMemcpyFromSymbol, ::cudaMemcpyAsync, ::cudaMemcpy2DAsync,
1030 * ::cudaMemcpy2DToArrayAsync,
1031 * ::cudaMemcpy2DFromArrayAsync,
1032 * ::cudaMemcpyToSymbolAsync
1033 */
1034template<class T>
1035static __inline__ __host__ cudaError_t cudaMemcpyFromSymbolAsync(
1036 void *dst,
1037 const T &symbol,
1038 size_t count,
1039 size_t offset = 0,
1040 enum cudaMemcpyKind kind = cudaMemcpyDeviceToHost,
1041 cudaStream_t stream = 0
1042)
1043{
1044 return ::cudaMemcpyFromSymbolAsync(dst, (const void*)&symbol, count, offset, kind, stream);
1045}
1046
1047/**
1048 * \brief Creates a memcpy node to copy to a symbol on the device and adds it to a graph
1049 *
1050 * Creates a new memcpy node to copy to \p symbol and adds it to \p graph with
1051 * \p numDependencies dependencies specified via \p pDependencies.
1052 * It is possible for \p numDependencies to be 0, in which case the node will be placed
1053 * at the root of the graph. \p pDependencies may not have any duplicate entries.
1054 * A handle to the new node will be returned in \p pGraphNode.
1055 *
1056 * When the graph is launched, the node will copy \p count bytes from the memory area
1057 * pointed to by \p src to the memory area pointed to by \p offset bytes from the start
1058 * of symbol \p symbol. The memory areas may not overlap. \p symbol is a variable that
1059 * resides in global or constant memory space. \p kind can be either
1060 * ::cudaMemcpyHostToDevice, ::cudaMemcpyDeviceToDevice, or ::cudaMemcpyDefault.
1061 * Passing ::cudaMemcpyDefault is recommended, in which case the type of
1062 * transfer is inferred from the pointer values. However, ::cudaMemcpyDefault
1063 * is only allowed on systems that support unified virtual addressing.
1064 *
1065 * Memcpy nodes have some additional restrictions with regards to managed memory, if the
1066 * system contains at least one device which has a zero value for the device attribute
1067 * ::cudaDevAttrConcurrentManagedAccess.
1068 *
1069 * \param pGraphNode - Returns newly created node
1070 * \param graph - Graph to which to add the node
1071 * \param pDependencies - Dependencies of the node
1072 * \param numDependencies - Number of dependencies
1073 * \param symbol - Device symbol address
1074 * \param src - Source memory address
1075 * \param count - Size in bytes to copy
1076 * \param offset - Offset from start of symbol in bytes
1077 * \param kind - Type of transfer
1078 *
1079 * \return
1080 * ::cudaSuccess,
1081 * ::cudaErrorInvalidValue
1082 * \note_graph_thread_safety
1083 * \notefnerr
1084 * \note_init_rt
1085 * \note_callback
1086 *
1087 * \sa
1088 * ::cudaMemcpyToSymbol,
1089 * ::cudaGraphAddMemcpyNode,
1090 * ::cudaGraphAddMemcpyNodeFromSymbol,
1091 * ::cudaGraphMemcpyNodeGetParams,
1092 * ::cudaGraphMemcpyNodeSetParams,
1093 * ::cudaGraphMemcpyNodeSetParamsToSymbol,
1094 * ::cudaGraphMemcpyNodeSetParamsFromSymbol,
1095 * ::cudaGraphCreate,
1096 * ::cudaGraphDestroyNode,
1097 * ::cudaGraphAddChildGraphNode,
1098 * ::cudaGraphAddEmptyNode,
1099 * ::cudaGraphAddKernelNode,
1100 * ::cudaGraphAddHostNode,
1101 * ::cudaGraphAddMemsetNode
1102 */
1103template<class T>
1104static __inline__ __host__ cudaError_t cudaGraphAddMemcpyNodeToSymbol(
1105 cudaGraphNode_t *pGraphNode,
1106 cudaGraph_t graph,
1107 const cudaGraphNode_t *pDependencies,
1108 size_t numDependencies,
1109 const T &symbol,
1110 const void* src,
1111 size_t count,
1112 size_t offset,
1113 enum cudaMemcpyKind kind)
1114{
1115 return ::cudaGraphAddMemcpyNodeToSymbol(pGraphNode, graph, pDependencies, numDependencies, (const void*)&symbol, src, count, offset, kind);
1116}
1117
1118/**
1119 * \brief Creates a memcpy node to copy from a symbol on the device and adds it to a graph
1120 *
1121 * Creates a new memcpy node to copy from \p symbol and adds it to \p graph with
1122 * \p numDependencies dependencies specified via \p pDependencies.
1123 * It is possible for \p numDependencies to be 0, in which case the node will be placed
1124 * at the root of the graph. \p pDependencies may not have any duplicate entries.
1125 * A handle to the new node will be returned in \p pGraphNode.
1126 *
1127 * When the graph is launched, the node will copy \p count bytes from the memory area
1128 * pointed to by \p offset bytes from the start of symbol \p symbol to the memory area
1129 * pointed to by \p dst. The memory areas may not overlap. \p symbol is a variable
1130 * that resides in global or constant memory space. \p kind can be either
1131 * ::cudaMemcpyDeviceToHost, ::cudaMemcpyDeviceToDevice, or ::cudaMemcpyDefault.
1132 * Passing ::cudaMemcpyDefault is recommended, in which case the type of transfer
1133 * is inferred from the pointer values. However, ::cudaMemcpyDefault is only
1134 * allowed on systems that support unified virtual addressing.
1135 *
1136 * Memcpy nodes have some additional restrictions with regards to managed memory, if the
1137 * system contains at least one device which has a zero value for the device attribute
1138 * ::cudaDevAttrConcurrentManagedAccess.
1139 *
1140 * \param pGraphNode - Returns newly created node
1141 * \param graph - Graph to which to add the node
1142 * \param pDependencies - Dependencies of the node
1143 * \param numDependencies - Number of dependencies
1144 * \param dst - Destination memory address
1145 * \param symbol - Device symbol address
1146 * \param count - Size in bytes to copy
1147 * \param offset - Offset from start of symbol in bytes
1148 * \param kind - Type of transfer
1149 *
1150 * \return
1151 * ::cudaSuccess,
1152 * ::cudaErrorInvalidValue
1153 * \note_graph_thread_safety
1154 * \notefnerr
1155 * \note_init_rt
1156 * \note_callback
1157 *
1158 * \sa
1159 * ::cudaMemcpyFromSymbol,
1160 * ::cudaGraphAddMemcpyNode,
1161 * ::cudaGraphAddMemcpyNodeToSymbol,
1162 * ::cudaGraphMemcpyNodeGetParams,
1163 * ::cudaGraphMemcpyNodeSetParams,
1164 * ::cudaGraphMemcpyNodeSetParamsFromSymbol,
1165 * ::cudaGraphMemcpyNodeSetParamsToSymbol,
1166 * ::cudaGraphCreate,
1167 * ::cudaGraphDestroyNode,
1168 * ::cudaGraphAddChildGraphNode,
1169 * ::cudaGraphAddEmptyNode,
1170 * ::cudaGraphAddKernelNode,
1171 * ::cudaGraphAddHostNode,
1172 * ::cudaGraphAddMemsetNode
1173 */
1174template<class T>
1175static __inline__ __host__ cudaError_t cudaGraphAddMemcpyNodeFromSymbol(
1176 cudaGraphNode_t* pGraphNode,
1177 cudaGraph_t graph,
1178 const cudaGraphNode_t* pDependencies,
1179 size_t numDependencies,
1180 void* dst,
1181 const T &symbol,
1182 size_t count,
1183 size_t offset,
1184 enum cudaMemcpyKind kind)
1185{
1186 return ::cudaGraphAddMemcpyNodeFromSymbol(pGraphNode, graph, pDependencies, numDependencies, dst, (const void*)&symbol, count, offset, kind);
1187}
1188
1189/**
1190 * \brief Sets a memcpy node's parameters to copy to a symbol on the device
1191 *
1192 * Sets the parameters of memcpy node \p node to the copy described by the provided parameters.
1193 *
1194 * When the graph is launched, the node will copy \p count bytes from the memory area
1195 * pointed to by \p src to the memory area pointed to by \p offset bytes from the start
1196 * of symbol \p symbol. The memory areas may not overlap. \p symbol is a variable that
1197 * resides in global or constant memory space. \p kind can be either
1198 * ::cudaMemcpyHostToDevice, ::cudaMemcpyDeviceToDevice, or ::cudaMemcpyDefault.
1199 * Passing ::cudaMemcpyDefault is recommended, in which case the type of
1200 * transfer is inferred from the pointer values. However, ::cudaMemcpyDefault