codekingpro/portable-devtools
114k
1/*
2 * Copyright 1993-2021 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_DEVICE_RUNTIME_API_H__)
51#define __CUDA_DEVICE_RUNTIME_API_H__
52
53#if defined(__CUDACC__) && !defined(__CUDACC_RTC__)
54#include <stdlib.h>
55#endif
56
57/*******************************************************************************
58* *
59* *
60* *
61*******************************************************************************/
62
63#if !defined(CUDA_FORCE_CDP1_IF_SUPPORTED) && !defined(__CUDADEVRT_INTERNAL__) && !defined(_NVHPC_CUDA) && !(defined(_WIN32) && !defined(_WIN64))
64#define __CUDA_INTERNAL_USE_CDP2
65#endif
66
67#if !defined(__CUDACC_RTC__)
68
69#if !defined(__CUDACC_INTERNAL_NO_STUBS__) && !defined(__CUDACC_RDC__) && !defined(__CUDACC_EWP__) && defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 350) && !defined(__CUDADEVRT_INTERNAL__)
70
71#if defined(__cplusplus)
72extern "C" {
73#endif
74
75struct cudaFuncAttributes;
76
77
78#ifndef __CUDA_INTERNAL_USE_CDP2
79inline __device__ cudaError_t CUDARTAPI cudaMalloc(void **p, size_t s)
80{
81 return cudaErrorUnknown;
82}
83
84inline __device__ cudaError_t CUDARTAPI cudaFuncGetAttributes(struct cudaFuncAttributes *p, const void *c)
85{
86 return cudaErrorUnknown;
87}
88
89inline __device__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDeviceAttr attr, int device)
90{
91 return cudaErrorUnknown;
92}
93
94inline __device__ cudaError_t CUDARTAPI cudaGetDevice(int *device)
95{
96 return cudaErrorUnknown;
97}
98
99inline __device__ cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultiprocessor(int *numBlocks, const void *func, int blockSize, size_t dynamicSmemSize)
100{
101 return cudaErrorUnknown;
102}
103
104inline __device__ cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int *numBlocks, const void *func, int blockSize, size_t dynamicSmemSize, unsigned int flags)
105{
106 return cudaErrorUnknown;
107}
108#else // __CUDA_INTERNAL_USE_CDP2
109inline __device__ cudaError_t CUDARTAPI __cudaCDP2Malloc(void **p, size_t s)
110{
111 return cudaErrorUnknown;
112}
113
114inline __device__ cudaError_t CUDARTAPI __cudaCDP2FuncGetAttributes(struct cudaFuncAttributes *p, const void *c)
115{
116 return cudaErrorUnknown;
117}
118
119inline __device__ cudaError_t CUDARTAPI __cudaCDP2DeviceGetAttribute(int *value, enum cudaDeviceAttr attr, int device)
120{
121 return cudaErrorUnknown;
122}
123
124inline __device__ cudaError_t CUDARTAPI __cudaCDP2GetDevice(int *device)
125{
126 return cudaErrorUnknown;
127}
128
129inline __device__ cudaError_t CUDARTAPI __cudaCDP2OccupancyMaxActiveBlocksPerMultiprocessor(int *numBlocks, const void *func, int blockSize, size_t dynamicSmemSize)
130{
131 return cudaErrorUnknown;
132}
133
134inline __device__ cudaError_t CUDARTAPI __cudaCDP2OccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int *numBlocks, const void *func, int blockSize, size_t dynamicSmemSize, unsigned int flags)
135{
136 return cudaErrorUnknown;
137}
138#endif // __CUDA_INTERNAL_USE_CDP2
139
140
141#if defined(__cplusplus)
142}
143#endif
144
145#endif /* !defined(__CUDACC_INTERNAL_NO_STUBS__) && !defined(__CUDACC_RDC__) && !defined(__CUDACC_EWP__) && defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 350) && !defined(__CUDADEVRT_INTERNAL__) */
146
147#endif /* !defined(__CUDACC_RTC__) */
148
149#if defined(__DOXYGEN_ONLY__) || defined(CUDA_ENABLE_DEPRECATED)
150# define __DEPRECATED__(msg)
151#elif defined(_WIN32)
152# define __DEPRECATED__(msg) __declspec(deprecated(msg))
153#elif (defined(__GNUC__) && (__GNUC__ < 4 || (__GNUC__ == 4 && __GNUC_MINOR__ < 5 && !defined(__clang__))))
154# define __DEPRECATED__(msg) __attribute__((deprecated))
155#else
156# define __DEPRECATED__(msg) __attribute__((deprecated(msg)))
157#endif
158
159#if defined(__CUDA_ARCH__) && !defined(__CDPRT_SUPPRESS_SYNC_DEPRECATION_WARNING)
160# define __CDPRT_DEPRECATED(func_name) __DEPRECATED__("Use of "#func_name" from device code is deprecated. Moreover, such use will cause this module to fail to load on sm_90+ devices. If calls to "#func_name" from device code cannot be removed for older devices at this time, you may guard them with __CUDA_ARCH__ macros to remove them only for sm_90+ devices, making sure to generate code for compute_90 for the macros to take effect. Note that this mitigation will no longer work when support for "#func_name" from device code is eventually dropped for all devices. Disable this warning with -D__CDPRT_SUPPRESS_SYNC_DEPRECATION_WARNING.")
161#else
162# define __CDPRT_DEPRECATED(func_name)
163#endif
164
165#if defined(__cplusplus) && defined(__CUDACC__) /* Visible to nvcc front-end only */
166#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 350) // Visible to SM>=3.5 and "__host__ __device__" only
167
168#include "driver_types.h"
169#include "crt/host_defines.h"
170
171#define cudaStreamGraphTailLaunch (cudaStream_t)0x0100000000000000
172#define cudaStreamGraphFireAndForget (cudaStream_t)0x0200000000000000
173#define cudaStreamGraphFireAndForgetAsSibling (cudaStream_t)0x0300000000000000
174
175#ifdef __CUDA_INTERNAL_USE_CDP2
176#define cudaStreamTailLaunch ((cudaStream_t)0x3) /**< Per-grid stream with a tail launch semantics. Only applicable when used with CUDA Dynamic Parallelism. */
177#define cudaStreamFireAndForget ((cudaStream_t)0x4) /**< Per-grid stream with a fire-and-forget synchronization behavior. Only applicable when used with CUDA Dynamic Parallelism. */
178#endif
179
180extern "C"
181{
182
183// Symbols beginning with __cudaCDP* should not be used outside
184// this header file. Instead, compile with -DCUDA_FORCE_CDP1_IF_SUPPORTED if
185// CDP1 support is required.
186
187extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaDeviceSynchronizeDeprecationAvoidance(void);
188
189#ifndef __CUDA_INTERNAL_USE_CDP2
190//// CDP1 endpoints
191extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDeviceAttr attr, int device);
192extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaDeviceGetLimit(size_t *pValue, enum cudaLimit limit);
193extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaDeviceGetCacheConfig(enum cudaFuncCache *pCacheConfig);
194extern __DEPRECATED__("cudaDeviceGetSharedMemConfig deprecated") __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaDeviceGetSharedMemConfig(enum cudaSharedMemConfig *pConfig);
195#if (__CUDA_ARCH__ < 900) && (defined(CUDA_FORCE_CDP1_IF_SUPPORTED) || (defined(_WIN32) && !defined(_WIN64)))
196// cudaDeviceSynchronize is removed on sm_90+
197extern __device__ __cudart_builtin__ __CDPRT_DEPRECATED(cudaDeviceSynchronize) cudaError_t CUDARTAPI cudaDeviceSynchronize(void);
198#endif
199extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaGetLastError(void);
200extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaPeekAtLastError(void);
201extern __device__ __cudart_builtin__ const char* CUDARTAPI cudaGetErrorString(cudaError_t error);
202extern __device__ __cudart_builtin__ const char* CUDARTAPI cudaGetErrorName(cudaError_t error);
203extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaGetDeviceCount(int *count);
204extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaGetDevice(int *device);
205extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaStreamCreateWithFlags(cudaStream_t *pStream, unsigned int flags);
206extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream);
207extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaStreamWaitEvent(cudaStream_t stream, cudaEvent_t event, unsigned int flags);
208extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaStreamWaitEvent_ptsz(cudaStream_t stream, cudaEvent_t event, unsigned int flags);
209extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaEventCreateWithFlags(cudaEvent_t *event, unsigned int flags);
210extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaEventRecord(cudaEvent_t event, cudaStream_t stream);
211extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaEventRecord_ptsz(cudaEvent_t event, cudaStream_t stream);
212extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaEventRecordWithFlags(cudaEvent_t event, cudaStream_t stream, unsigned int flags);
213extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaEventRecordWithFlags_ptsz(cudaEvent_t event, cudaStream_t stream, unsigned int flags);
214extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaEventDestroy(cudaEvent_t event);
215extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaFuncGetAttributes(struct cudaFuncAttributes *attr, const void *func);
216extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaFree(void *devPtr);
217extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size);
218extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMemcpyAsync(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream);
219extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMemcpyAsync_ptsz(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream);
220extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMemcpy2DAsync(void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream);
221extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMemcpy2DAsync_ptsz(void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream);
222extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMemcpy3DAsync(const struct cudaMemcpy3DParms *p, cudaStream_t stream);
223extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMemcpy3DAsync_ptsz(const struct cudaMemcpy3DParms *p, cudaStream_t stream);
224extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMemsetAsync(void *devPtr, int value, size_t count, cudaStream_t stream);
225extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMemsetAsync_ptsz(void *devPtr, int value, size_t count, cudaStream_t stream);
226extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMemset2DAsync(void *devPtr, size_t pitch, int value, size_t width, size_t height, cudaStream_t stream);
227extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMemset2DAsync_ptsz(void *devPtr, size_t pitch, int value, size_t width, size_t height, cudaStream_t stream);
228extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMemset3DAsync(struct cudaPitchedPtr pitchedDevPtr, int value, struct cudaExtent extent, cudaStream_t stream);
229extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMemset3DAsync_ptsz(struct cudaPitchedPtr pitchedDevPtr, int value, struct cudaExtent extent, cudaStream_t stream);
230extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaRuntimeGetVersion(int *runtimeVersion);
231extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultiprocessor(int *numBlocks, const void *func, int blockSize, size_t dynamicSmemSize);
232extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int *numBlocks, const void *func, int blockSize, size_t dynamicSmemSize, unsigned int flags);
233#endif // __CUDA_INTERNAL_USE_CDP2
234
235//// CDP2 endpoints
236extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2DeviceGetAttribute(int *value, enum cudaDeviceAttr attr, int device);
237extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2DeviceGetLimit(size_t *pValue, enum cudaLimit limit);
238extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2DeviceGetCacheConfig(enum cudaFuncCache *pCacheConfig);
239extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2DeviceGetSharedMemConfig(enum cudaSharedMemConfig *pConfig);
240extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2GetLastError(void);
241extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2PeekAtLastError(void);
242extern __device__ __cudart_builtin__ const char* CUDARTAPI __cudaCDP2GetErrorString(cudaError_t error);
243extern __device__ __cudart_builtin__ const char* CUDARTAPI __cudaCDP2GetErrorName(cudaError_t error);
244extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2GetDeviceCount(int *count);
245extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2GetDevice(int *device);
246extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2StreamCreateWithFlags(cudaStream_t *pStream, unsigned int flags);
247extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2StreamDestroy(cudaStream_t stream);
248extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2StreamWaitEvent(cudaStream_t stream, cudaEvent_t event, unsigned int flags);
249extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2StreamWaitEvent_ptsz(cudaStream_t stream, cudaEvent_t event, unsigned int flags);
250extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2EventCreateWithFlags(cudaEvent_t *event, unsigned int flags);
251extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2EventRecord(cudaEvent_t event, cudaStream_t stream);
252extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2EventRecord_ptsz(cudaEvent_t event, cudaStream_t stream);
253extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2EventRecordWithFlags(cudaEvent_t event, cudaStream_t stream, unsigned int flags);
254extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2EventRecordWithFlags_ptsz(cudaEvent_t event, cudaStream_t stream, unsigned int flags);
255extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2EventDestroy(cudaEvent_t event);
256extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2FuncGetAttributes(struct cudaFuncAttributes *attr, const void *func);
257extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2Free(void *devPtr);
258extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2Malloc(void **devPtr, size_t size);
259extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2MemcpyAsync(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream);
260extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2MemcpyAsync_ptsz(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream);
261extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2Memcpy2DAsync(void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream);
262extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2Memcpy2DAsync_ptsz(void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream);
263extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2Memcpy3DAsync(const struct cudaMemcpy3DParms *p, cudaStream_t stream);
264extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2Memcpy3DAsync_ptsz(const struct cudaMemcpy3DParms *p, cudaStream_t stream);
265extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2MemsetAsync(void *devPtr, int value, size_t count, cudaStream_t stream);
266extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2MemsetAsync_ptsz(void *devPtr, int value, size_t count, cudaStream_t stream);
267extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2Memset2DAsync(void *devPtr, size_t pitch, int value, size_t width, size_t height, cudaStream_t stream);
268extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2Memset2DAsync_ptsz(void *devPtr, size_t pitch, int value, size_t width, size_t height, cudaStream_t stream);
269extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2Memset3DAsync(struct cudaPitchedPtr pitchedDevPtr, int value, struct cudaExtent extent, cudaStream_t stream);
270extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2Memset3DAsync_ptsz(struct cudaPitchedPtr pitchedDevPtr, int value, struct cudaExtent extent, cudaStream_t stream);
271extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2RuntimeGetVersion(int *runtimeVersion);
272extern __device__ __cudart_builtin__ void * CUDARTAPI __cudaCDP2GetParameterBuffer(size_t alignment, size_t size);
273extern __device__ __cudart_builtin__ void * CUDARTAPI __cudaCDP2GetParameterBufferV2(void *func, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize);
274extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2LaunchDevice_ptsz(void *func, void *parameterBuffer, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize, cudaStream_t stream);
275extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2LaunchDeviceV2_ptsz(void *parameterBuffer, cudaStream_t stream);
276extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2LaunchDevice(void *func, void *parameterBuffer, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize, cudaStream_t stream);
277extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2LaunchDeviceV2(void *parameterBuffer, cudaStream_t stream);
278extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2OccupancyMaxActiveBlocksPerMultiprocessor(int *numBlocks, const void *func, int blockSize, size_t dynamicSmemSize);
279extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaCDP2OccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int *numBlocks, const void *func, int blockSize, size_t dynamicSmemSize, unsigned int flags);
280
281
282extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaGraphLaunch(cudaGraphExec_t graphExec, cudaStream_t stream);
283#if defined(CUDA_API_PER_THREAD_DEFAULT_STREAM)
284static inline __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaGraphLaunch_ptsz(cudaGraphExec_t graphExec, cudaStream_t stream)
285{
286 if (stream == 0) {
287 stream = cudaStreamPerThread;
288 }
289 return cudaGraphLaunch(graphExec, stream);
290}
291#endif
292
293/**
294 * \ingroup CUDART_GRAPH
295 * \brief Get the currently running device graph id.
296 *
297 * Get the currently running device graph id.
298 * \return Returns the current device graph id, 0 if the call is outside of a device graph.
299 * \sa cudaGraphLaunch
300 */
301static inline __device__ __cudart_builtin__ cudaGraphExec_t CUDARTAPI cudaGetCurrentGraphExec(void)
302{
303 unsigned long long current_graph_exec;
304 asm ("mov.u64 %0, %%current_graph_exec;" : "=l"(current_graph_exec));
305 return (cudaGraphExec_t)current_graph_exec;
306}
307
308/**
309 * \ingroup CUDART_GRAPH
310 * \brief Updates the kernel parameters of the given kernel node
311 *
312 * Updates \p size bytes in the kernel parameters of \p node at \p offset to
313 * the contents of \p value. \p node must be device-updatable, and must reside upon the same
314 * device as the calling kernel.
315 *
316 * If this function is called for the node's immediate dependent and that dependent is configured
317 * for programmatic dependent launch, then a memory fence must be invoked via __threadfence() before
318 * kickoff of the dependent is triggered via ::cudaTriggerProgrammaticLaunchCompletion() to ensure
319 * that the update is visible to that dependent node before it is launched.
320 *
321 * \param node - The node to update
322 * \param offset - The offset into the params at which to make the update
323 * \param value - Buffer containing the params to write
324 * \param size - Size in bytes to update
325 *
326 * \return
327 * cudaSucces,
328 * cudaErrorInvalidValue
329 * \notefnerr
330 *
331 * \sa
332 * ::cudaGraphKernelNodeSetEnabled,
333 * ::cudaGraphKernelNodeSetGridDim,
334 * ::cudaGraphKernelNodeUpdatesApply
335 */
336extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaGraphKernelNodeSetParam(cudaGraphDeviceNode_t node, size_t offset, const void *value , size_t size);
337
338/**
339 * \ingroup CUDART_GRAPH
340 * \brief Enables or disables the given kernel node
341 *
342 * Enables or disables \p node based upon \p enable. If \p enable is true, the node will be enabled;
343 * if it is false, the node will be disabled. Disabled nodes will act as a NOP during execution.
344 * \p node must be device-updatable, and must reside upon the same device as the calling kernel.
345 *
346 * If this function is called for the node's immediate dependent and that dependent is configured
347 * for programmatic dependent launch, then a memory fence must be invoked via __threadfence() before
348 * kickoff of the dependent is triggered via ::cudaTriggerProgrammaticLaunchCompletion() to ensure
349 * that the update is visible to that dependent node before it is launched.
350 *
351 * \param node - The node to update
352 * \param enable - Whether to enable or disable the node
353 *
354 * \return
355 * cudaSucces,
356 * cudaErrorInvalidValue
357 * \notefnerr
358 *
359 * \sa
360 * ::cudaGraphKernelNodeSetParam,
361 * ::cudaGraphKernelNodeSetGridDim,
362 * ::cudaGraphKernelNodeUpdatesApply
363 */
364extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaGraphKernelNodeSetEnabled(cudaGraphDeviceNode_t node, bool enable);
365
366/**
367 * \ingroup CUDART_GRAPH
368 * \brief Updates the grid dimensions of the given kernel node
369 *
370 * Sets the grid dimensions of \p node to \p gridDim. \p node must be device-updatable,
371 * and must reside upon the same device as thecalling kernel.
372 *
373 * If this function is called for the node's immediate dependent and that dependent is configured
374 * for programmatic dependent launch, then a memory fence must be invoked via __threadfence() before
375 * kickoff of the dependent is triggered via ::cudaTriggerProgrammaticLaunchCompletion() to ensure
376 * that the update is visible to that dependent node before it is launched.
377 *
378 * \param node - The node to update
379 * \param gridDim - The grid dimensions to set
380 *
381 * \return
382 * cudaSucces,
383 * cudaErrorInvalidValue
384 * \notefnerr
385 *
386 * \sa
387 * ::cudaGraphKernelNodeSetParam,
388 * ::cudaGraphKernelNodeSetEnabled,
389 * ::cudaGraphKernelNodeUpdatesApply
390 */
391extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaGraphKernelNodeSetGridDim(cudaGraphDeviceNode_t node, dim3 gridDim);
392
393/**
394 * \ingroup CUDART_GRAPH
395 * \brief Batch applies multiple kernel node updates
396 *
397 * Batch applies one or more kernel node updates based on the information provided in \p updates.
398 * \p updateCount specifies the number of updates to apply. Each entry in \p updates must specify
399 * a node to update, the type of update to apply, and the parameters for that type of update. See
400 * the documentation for ::cudaGraphKernelNodeUpdate for more detail.
401 *
402 * If this function is called for the node's immediate dependent and that dependent is configured
403 * for programmatic dependent launch, then a memory fence must be invoked via __threadfence() before
404 * kickoff of the dependent is triggered via ::cudaTriggerProgrammaticLaunchCompletion() to ensure
405 * that the update is visible to that dependent node before it is launched.
406 *
407 * \param updates - The updates to apply
408 * \param updateCount - The number of updates to apply
409 *
410 * \return
411 * cudaSucces,
412 * cudaErrorInvalidValue
413 * \notefnerr
414 *
415 * \sa
416 * ::cudaGraphKernelNodeSetParam,
417 * ::cudaGraphKernelNodeSetEnabled,
418 * ::cudaGraphKernelNodeSetGridDim
419 */
420extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaGraphKernelNodeUpdatesApply(const cudaGraphKernelNodeUpdate *updates, size_t updateCount);
421
422/**
423 * \ingroup CUDART_EXECUTION
424 * \brief Programmatic dependency trigger
425 *
426 * This device function ensures the programmatic launch completion edges /
427 * events are fulfilled. See
428 * ::cudaLaunchAttributeID::cudaLaunchAttributeProgrammaticStreamSerialization
429 * and ::cudaLaunchAttributeID::cudaLaunchAttributeProgrammaticEvent for more
430 * information. The event / edge kick off only happens when every CTAs
431 * in the grid has either exited or called this function at least once,
432 * otherwise the kick off happens automatically after all warps finishes
433 * execution but before the grid completes. The kick off only enables
434 * scheduling of the secondary kernel. It provides no memory visibility
435 * guarantee itself. The user could enforce memory visibility by inserting a
436 * memory fence of the correct scope.
437 */
438static inline __device__ __cudart_builtin__ void CUDARTAPI cudaTriggerProgrammaticLaunchCompletion(void)
439{
440 asm volatile("griddepcontrol.launch_dependents;":::);
441}
442
443/**
444 * \ingroup CUDART_EXECUTION
445 * \brief Programmatic grid dependency synchronization
446 *
447 * This device function will block the thread until all direct grid
448 * dependencies have completed. This API is intended to use in conjuncture with
449 * programmatic / launch event / dependency. See
450 * ::cudaLaunchAttributeID::cudaLaunchAttributeProgrammaticStreamSerialization
451 * and ::cudaLaunchAttributeID::cudaLaunchAttributeProgrammaticEvent for more
452 * information.
453 */
454static inline __device__ __cudart_builtin__ void CUDARTAPI cudaGridDependencySynchronize(void)
455{
456 asm volatile("griddepcontrol.wait;":::"memory");
457}
458
459/**
460 * \ingroup CUDART_GRAPH
461 * \brief Sets the condition value associated with a conditional node.
462 *
463 * Sets the condition value associated with a conditional node.
464 * \sa cudaGraphConditionalHandleCreate
465 */
466extern __device__ __cudart_builtin__ void CUDARTAPI cudaGraphSetConditional(cudaGraphConditionalHandle handle, unsigned int value);
467
468//// CG API
469extern __device__ __cudart_builtin__ unsigned long long CUDARTAPI cudaCGGetIntrinsicHandle(enum cudaCGScope scope);
470extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaCGSynchronize(unsigned long long handle, unsigned int flags);
471extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaCGSynchronizeGrid(unsigned long long handle, unsigned int flags);
472extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaCGGetSize(unsigned int *numThreads, unsigned int *numGrids, unsigned long long handle);
473extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaCGGetRank(unsigned int *threadRank, unsigned int *gridRank, unsigned long long handle);
474
475
476//// CDP API
477
478#ifdef __CUDA_ARCH__
479
480#ifdef __CUDA_INTERNAL_USE_CDP2
481static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDeviceAttr attr, int device)
482{
483 return __cudaCDP2DeviceGetAttribute(value, attr, device);
484}
485
486static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaDeviceGetLimit(size_t *pValue, enum cudaLimit limit)
487{
488 return __cudaCDP2DeviceGetLimit(pValue, limit);
489}
490
491static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaDeviceGetCacheConfig(enum cudaFuncCache *pCacheConfig)
492{
493 return __cudaCDP2DeviceGetCacheConfig(pCacheConfig);
494}
495
496static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaDeviceGetSharedMemConfig(enum cudaSharedMemConfig *pConfig)
497{
498 return __cudaCDP2DeviceGetSharedMemConfig(pConfig);
499}
500
501static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaGetLastError(void)
502{
503 return __cudaCDP2GetLastError();
504}
505
506static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaPeekAtLastError(void)
507{
508 return __cudaCDP2PeekAtLastError();
509}
510
511static __inline__ __device__ __cudart_builtin__ const char* CUDARTAPI cudaGetErrorString(cudaError_t error)
512{
513 return __cudaCDP2GetErrorString(error);
514}
515
516static __inline__ __device__ __cudart_builtin__ const char* CUDARTAPI cudaGetErrorName(cudaError_t error)
517{
518 return __cudaCDP2GetErrorName(error);
519}
520
521static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaGetDeviceCount(int *count)
522{
523 return __cudaCDP2GetDeviceCount(count);
524}
525
526static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaGetDevice(int *device)
527{
528 return __cudaCDP2GetDevice(device);
529}
530
531static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaStreamCreateWithFlags(cudaStream_t *pStream, unsigned int flags)
532{
533 return __cudaCDP2StreamCreateWithFlags(pStream, flags);
534}
535
536static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaStreamDestroy(cudaStream_t stream)
537{
538 return __cudaCDP2StreamDestroy(stream);
539}
540
541static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaStreamWaitEvent(cudaStream_t stream, cudaEvent_t event, unsigned int flags)
542{
543 return __cudaCDP2StreamWaitEvent(stream, event, flags);
544}
545
546static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaStreamWaitEvent_ptsz(cudaStream_t stream, cudaEvent_t event, unsigned int flags)
547{
548 return __cudaCDP2StreamWaitEvent_ptsz(stream, event, flags);
549}
550
551static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaEventCreateWithFlags(cudaEvent_t *event, unsigned int flags)
552{
553 return __cudaCDP2EventCreateWithFlags(event, flags);
554}
555
556static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaEventRecord(cudaEvent_t event, cudaStream_t stream)
557{
558 return __cudaCDP2EventRecord(event, stream);
559}
560
561static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaEventRecord_ptsz(cudaEvent_t event, cudaStream_t stream)
562{
563 return __cudaCDP2EventRecord_ptsz(event, stream);
564}
565
566static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaEventRecordWithFlags(cudaEvent_t event, cudaStream_t stream, unsigned int flags)
567{
568 return __cudaCDP2EventRecordWithFlags(event, stream, flags);
569}
570
571static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaEventRecordWithFlags_ptsz(cudaEvent_t event, cudaStream_t stream, unsigned int flags)
572{
573 return __cudaCDP2EventRecordWithFlags_ptsz(event, stream, flags);
574}
575
576static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaEventDestroy(cudaEvent_t event)
577{
578 return __cudaCDP2EventDestroy(event);
579}
580
581static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaFuncGetAttributes(struct cudaFuncAttributes *attr, const void *func)
582{
583 return __cudaCDP2FuncGetAttributes(attr, func);
584}
585
586static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaFree(void *devPtr)
587{
588 return __cudaCDP2Free(devPtr);
589}
590
591static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMalloc(void **devPtr, size_t size)
592{
593 return __cudaCDP2Malloc(devPtr, size);
594}
595
596static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMemcpyAsync(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream)
597{
598 return __cudaCDP2MemcpyAsync(dst, src, count, kind, stream);
599}
600
601static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMemcpyAsync_ptsz(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind, cudaStream_t stream)
602{
603 return __cudaCDP2MemcpyAsync_ptsz(dst, src, count, kind, stream);
604}
605
606static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMemcpy2DAsync(void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream)
607{
608 return __cudaCDP2Memcpy2DAsync(dst, dpitch, src, spitch, width, height, kind, stream);
609}
610
611static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMemcpy2DAsync_ptsz(void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind, cudaStream_t stream)
612{
613 return __cudaCDP2Memcpy2DAsync_ptsz(dst, dpitch, src, spitch, width, height, kind, stream);
614}
615
616static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMemcpy3DAsync(const struct cudaMemcpy3DParms *p, cudaStream_t stream)
617{
618 return __cudaCDP2Memcpy3DAsync(p, stream);
619}
620
621static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMemcpy3DAsync_ptsz(const struct cudaMemcpy3DParms *p, cudaStream_t stream)
622{
623 return __cudaCDP2Memcpy3DAsync_ptsz(p, stream);
624}
625
626static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMemsetAsync(void *devPtr, int value, size_t count, cudaStream_t stream)
627{
628 return __cudaCDP2MemsetAsync(devPtr, value, count, stream);
629}
630
631static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMemsetAsync_ptsz(void *devPtr, int value, size_t count, cudaStream_t stream)
632{
633 return __cudaCDP2MemsetAsync_ptsz(devPtr, value, count, stream);
634}
635
636static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMemset2DAsync(void *devPtr, size_t pitch, int value, size_t width, size_t height, cudaStream_t stream)
637{
638 return __cudaCDP2Memset2DAsync(devPtr, pitch, value, width, height, stream);
639}
640
641static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMemset2DAsync_ptsz(void *devPtr, size_t pitch, int value, size_t width, size_t height, cudaStream_t stream)
642{
643 return __cudaCDP2Memset2DAsync_ptsz(devPtr, pitch, value, width, height, stream);
644}
645
646static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMemset3DAsync(struct cudaPitchedPtr pitchedDevPtr, int value, struct cudaExtent extent, cudaStream_t stream)
647{
648 return __cudaCDP2Memset3DAsync(pitchedDevPtr, value, extent, stream);
649}
650
651static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMemset3DAsync_ptsz(struct cudaPitchedPtr pitchedDevPtr, int value, struct cudaExtent extent, cudaStream_t stream)
652{
653 return __cudaCDP2Memset3DAsync_ptsz(pitchedDevPtr, value, extent, stream);
654}
655
656static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaRuntimeGetVersion(int *runtimeVersion)
657{
658 return __cudaCDP2RuntimeGetVersion(runtimeVersion);
659}
660
661static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultiprocessor(int *numBlocks, const void *func, int blockSize, size_t dynamicSmemSize)
662{
663 return __cudaCDP2OccupancyMaxActiveBlocksPerMultiprocessor(numBlocks, func, blockSize, dynamicSmemSize);
664}
665
666static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int *numBlocks, const void *func, int blockSize, size_t dynamicSmemSize, unsigned int flags)
667{
668 return __cudaCDP2OccupancyMaxActiveBlocksPerMultiprocessorWithFlags(numBlocks, func, blockSize, dynamicSmemSize, flags);
669}
670#endif // __CUDA_INTERNAL_USE_CDP2
671
672#endif // __CUDA_ARCH__
673
674
675/**
676 * \ingroup CUDART_EXECUTION
677 * \brief Obtains a parameter buffer
678 *
679 * Obtains a parameter buffer which can be filled with parameters for a kernel launch.
680 * Parameters passed to ::cudaLaunchDevice must be allocated via this function.
681 *
682 * This is a low level API and can only be accessed from Parallel Thread Execution (PTX).
683 * CUDA user code should use <<< >>> to launch kernels.
684 *
685 * \param alignment - Specifies alignment requirement of the parameter buffer
686 * \param size - Specifies size requirement in bytes
687 *
688 * \return
689 * Returns pointer to the allocated parameterBuffer
690 * \notefnerr
691 *
692 * \sa cudaLaunchDevice
693 */
694#ifdef __CUDA_INTERNAL_USE_CDP2
695static __inline__ __device__ __cudart_builtin__ void * CUDARTAPI cudaGetParameterBuffer(size_t alignment, size_t size)
696{
697 return __cudaCDP2GetParameterBuffer(alignment, size);
698}
699#else
700extern __device__ __cudart_builtin__ void * CUDARTAPI cudaGetParameterBuffer(size_t alignment, size_t size);
701#endif
702
703
704#ifdef __CUDA_INTERNAL_USE_CDP2
705static __inline__ __device__ __cudart_builtin__ void * CUDARTAPI cudaGetParameterBufferV2(void *func, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize)
706{
707 return __cudaCDP2GetParameterBufferV2(func, gridDimension, blockDimension, sharedMemSize);
708}
709#else
710extern __device__ __cudart_builtin__ void * CUDARTAPI cudaGetParameterBufferV2(void *func, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize);
711#endif
712
713
714#ifdef __CUDA_INTERNAL_USE_CDP2
715static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaLaunchDevice_ptsz(void *func, void *parameterBuffer, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize, cudaStream_t stream)
716{
717 return __cudaCDP2LaunchDevice_ptsz(func, parameterBuffer, gridDimension, blockDimension, sharedMemSize, stream);
718}
719
720static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaLaunchDeviceV2_ptsz(void *parameterBuffer, cudaStream_t stream)
721{
722 return __cudaCDP2LaunchDeviceV2_ptsz(parameterBuffer, stream);
723}
724#else
725extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaLaunchDevice_ptsz(void *func, void *parameterBuffer, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize, cudaStream_t stream);
726extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaLaunchDeviceV2_ptsz(void *parameterBuffer, cudaStream_t stream);
727#endif
728
729
730/**
731 * \ingroup CUDART_EXECUTION
732 * \brief Launches a specified kernel
733 *
734 * Launches a specified kernel with the specified parameter buffer. A parameter buffer can be obtained
735 * by calling ::cudaGetParameterBuffer().
736 *
737 * This is a low level API and can only be accessed from Parallel Thread Execution (PTX).
738 * CUDA user code should use <<< >>> to launch the kernels.
739 *
740 * \param func - Pointer to the kernel to be launched
741 * \param parameterBuffer - Holds the parameters to the launched kernel. parameterBuffer can be NULL. (Optional)
742 * \param gridDimension - Specifies grid dimensions
743 * \param blockDimension - Specifies block dimensions
744 * \param sharedMemSize - Specifies size of shared memory
745 * \param stream - Specifies the stream to be used
746 *
747 * \return
748 * ::cudaSuccess, ::cudaErrorInvalidDevice, ::cudaErrorLaunchMaxDepthExceeded, ::cudaErrorInvalidConfiguration,
749 * ::cudaErrorStartupFailure, ::cudaErrorLaunchPendingCountExceeded, ::cudaErrorLaunchOutOfResources
750 * \notefnerr
751 * \n Please refer to Execution Configuration and Parameter Buffer Layout from the CUDA Programming
752 * Guide for the detailed descriptions of launch configuration and parameter layout respectively.
753 *
754 * \sa cudaGetParameterBuffer
755 */
756#if defined(CUDA_API_PER_THREAD_DEFAULT_STREAM) && defined(__CUDA_ARCH__)
757 // When compiling for the device and per thread default stream is enabled, add
758 // a static inline redirect to the per thread stream entry points.
759
760 static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI
761 cudaLaunchDevice(void *func, void *parameterBuffer, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize, cudaStream_t stream)
762 {
763#ifdef __CUDA_INTERNAL_USE_CDP2
764 return __cudaCDP2LaunchDevice_ptsz(func, parameterBuffer, gridDimension, blockDimension, sharedMemSize, stream);
765#else
766 return cudaLaunchDevice_ptsz(func, parameterBuffer, gridDimension, blockDimension, sharedMemSize, stream);
767#endif
768 }
769
770 static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI
771 cudaLaunchDeviceV2(void *parameterBuffer, cudaStream_t stream)
772 {
773#ifdef __CUDA_INTERNAL_USE_CDP2
774 return __cudaCDP2LaunchDeviceV2_ptsz(parameterBuffer, stream);
775#else
776 return cudaLaunchDeviceV2_ptsz(parameterBuffer, stream);
777#endif
778 }
779#else // defined(CUDA_API_PER_THREAD_DEFAULT_STREAM) && defined(__CUDA_ARCH__)
780#ifdef __CUDA_INTERNAL_USE_CDP2
781 static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaLaunchDevice(void *func, void *parameterBuffer, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize, cudaStream_t stream)
782 {
783 return __cudaCDP2LaunchDevice(func, parameterBuffer, gridDimension, blockDimension, sharedMemSize, stream);
784 }
785
786 static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaLaunchDeviceV2(void *parameterBuffer, cudaStream_t stream)
787 {
788 return __cudaCDP2LaunchDeviceV2(parameterBuffer, stream);
789 }
790#else // __CUDA_INTERNAL_USE_CDP2
791extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaLaunchDevice(void *func, void *parameterBuffer, dim3 gridDimension, dim3 blockDimension, unsigned int sharedMemSize, cudaStream_t stream);
792extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaLaunchDeviceV2(void *parameterBuffer, cudaStream_t stream);
793#endif // __CUDA_INTERNAL_USE_CDP2
794#endif // defined(CUDA_API_PER_THREAD_DEFAULT_STREAM) && defined(__CUDA_ARCH__)
795
796
797// These symbols should not be used outside of this header file.
798#define __cudaCDP2DeviceGetAttribute
799#define __cudaCDP2DeviceGetLimit
800#define __cudaCDP2DeviceGetCacheConfig
801#define __cudaCDP2DeviceGetSharedMemConfig
802#define __cudaCDP2GetLastError
803#define __cudaCDP2PeekAtLastError
804#define __cudaCDP2GetErrorString
805#define __cudaCDP2GetErrorName
806#define __cudaCDP2GetDeviceCount
807#define __cudaCDP2GetDevice
808#define __cudaCDP2StreamCreateWithFlags
809#define __cudaCDP2StreamDestroy
810#define __cudaCDP2StreamWaitEvent
811#define __cudaCDP2StreamWaitEvent_ptsz
812#define __cudaCDP2EventCreateWithFlags
813#define __cudaCDP2EventRecord
814#define __cudaCDP2EventRecord_ptsz
815#define __cudaCDP2EventRecordWithFlags
816#define __cudaCDP2EventRecordWithFlags_ptsz
817#define __cudaCDP2EventDestroy
818#define __cudaCDP2FuncGetAttributes
819#define __cudaCDP2Free
820#define __cudaCDP2Malloc
821#define __cudaCDP2MemcpyAsync
822#define __cudaCDP2MemcpyAsync_ptsz
823#define __cudaCDP2Memcpy2DAsync
824#define __cudaCDP2Memcpy2DAsync_ptsz
825#define __cudaCDP2Memcpy3DAsync
826#define __cudaCDP2Memcpy3DAsync_ptsz
827#define __cudaCDP2MemsetAsync
828#define __cudaCDP2MemsetAsync_ptsz
829#define __cudaCDP2Memset2DAsync
830#define __cudaCDP2Memset2DAsync_ptsz
831#define __cudaCDP2Memset3DAsync
832#define __cudaCDP2Memset3DAsync_ptsz
833#define __cudaCDP2RuntimeGetVersion
834#define __cudaCDP2GetParameterBuffer
835#define __cudaCDP2GetParameterBufferV2
836#define __cudaCDP2LaunchDevice_ptsz
837#define __cudaCDP2LaunchDeviceV2_ptsz
838#define __cudaCDP2LaunchDevice
839#define __cudaCDP2LaunchDeviceV2
840#define __cudaCDP2OccupancyMaxActiveBlocksPerMultiprocessor
841#define __cudaCDP2OccupancyMaxActiveBlocksPerMultiprocessorWithFlags
842
843}
844
845template <typename T> static __inline__ __device__ __cudart_builtin__ cudaError_t cudaMalloc(T **devPtr, size_t size);
846template <typename T> static __inline__ __device__ __cudart_builtin__ cudaError_t cudaFuncGetAttributes(struct cudaFuncAttributes *attr, T *entry);
847template <typename T> static __inline__ __device__ __cudart_builtin__ cudaError_t cudaOccupancyMaxActiveBlocksPerMultiprocessor(int *numBlocks, T func, int blockSize, size_t dynamicSmemSize);
848template <typename T> static __inline__ __device__ __cudart_builtin__ cudaError_t cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int *numBlocks, T func, int blockSize, size_t dynamicSmemSize, unsigned int flags);
849
850/**
851 * \ingroup CUDART_GRAPH
852 * \brief Updates the kernel parameters of the given kernel node
853 *
854 * Updates the kernel parameters of \p node at \p offset to \p value. \p node must be
855 * device-updatable, and must reside upon the same device as the calling kernel.
856 *
857 * If this function is called for the node's immediate dependent and that dependent is configured
858 * for programmatic dependent launch, then a memory fence must be invoked via __threadfence() before
859 * kickoff of the dependent is triggered via ::cudaTriggerProgrammaticLaunchCompletion() to ensure
860 * that the update is visible to that dependent node before it is launched.
861 *
862 * \param node - The node to update
863 * \param offset - The offset into the params at which to make the update
864 * \param value - Parameter value to write
865 *
866 * \return
867 * cudaSucces,
868 * cudaErrorInvalidValue
869 * \notefnerr
870 *
871 * \sa
872 * ::etblGraphKernelNodeSetEnabled,
873 * ::etblGraphKernelNodeSetGridDim,
874 * ::etblGraphKernelNodeUpdatesApply
875 */
876template <typename T>
877static __inline__ __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaGraphKernelNodeSetParam(cudaGraphDeviceNode_t node, size_t offset, const T &value)
878{
879 return cudaGraphKernelNodeSetParam(node, offset, &value, sizeof(T));
880}
881
882#endif // !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 350)
883#endif /* defined(__cplusplus) && defined(__CUDACC__) */
884
885#undef __DEPRECATED__
886#undef __CDPRT_DEPRECATED
887#undef __CUDA_INTERNAL_USE_CDP2
888
889#endif /* !__CUDA_DEVICE_RUNTIME_API_H__ */
890