diff --git a/CHANGELOG b/CHANGELOG new file mode 100644 index 0000000..f046421 --- /dev/null +++ b/CHANGELOG @@ -0,0 +1,8 @@ +Changelog (v0.6.5.2) + - AI Green Screen filter + - Quality mode is updated for higher quality + - Added CUDA graph optimization which may improve overall latency under a GPU-intensive workload. + - Super Resolution filter + - 4K input support for 4/3x(~1.33x), 1.5x and 2x super resolution + - Migrated to TensorRT 8.0.1.6 + - Migrated to CUDA 11.3u1 diff --git a/CMakeLists.txt b/CMakeLists.txt index 4309db4..efaa07f 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -21,45 +21,9 @@ set(CMAKE_LIBRARY_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}) set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}) set(CMAKE_MODULE_PATH "${PROJECT_SOURCE_DIR}/cmake" ${CMAKE_MODULE_PATH}) -if(MSVC) - - set(SDK_INCLUDES_PATH ${CMAKE_CURRENT_SOURCE_DIR}/nvvfx/include) - # Add target for NVVideoEffects - add_library(NVVideoEffects INTERFACE) - target_include_directories(NVVideoEffects INTERFACE ${SDK_INCLUDES_PATH}) - -else() - # Add target for NVVideoEffects - add_library(NVVideoEffects INTERFACE) - - # found in different locations depending on type of package - find_path(VideoFX_INCLUDES - NAMES nvVideoEffects.h - PATHS - /usr/local/VideoFX/include - /usr/include/x86_64-linux-gnu - /usr/include - REQUIRED - ) - - target_include_directories(NVVideoEffects INTERFACE ${VideoFX_INCLUDES}) - set(SDK_INCLUDES_PATH ${VideoFX_INCLUDES}) - - find_library(VideoFX_LIB - NAMES libVideoFX.so - PATHS - /usr/local/VideoFX/lib - /usr/lib/x86_64-linux-gnu - /usr/lib64 - /usr/lib - REQUIRED - NO_DEFAULT_PATH) - - target_link_libraries(NVVideoEffects INTERFACE "${VideoFX_LIB}") - - message(STATUS "VideoFX_LIB: ${VideoFX_LIB}") - message(STATUS "SDK_INCLUDES_PATH: ${SDK_INCLUDES_PATH}") - -endif() +set(SDK_INCLUDES_PATH ${CMAKE_CURRENT_SOURCE_DIR}/nvvfx/include) +# Add target for NVVideoEffects +add_library(NVVideoEffects INTERFACE) +target_include_directories(NVVideoEffects INTERFACE ${SDK_INCLUDES_PATH}) add_subdirectory(samples) diff --git a/README.MD b/README.MD index 253e009..9140cfc 100644 --- a/README.MD +++ b/README.MD @@ -27,9 +27,11 @@ The SDK provides several sample applications that demonstrate the features liste - **DenoiseEffect App**, which is a sample app that demonstrates the webcam denoising feature. The input and output resolutions supported by the features of the SDK are listed below. -- The Super Resolution and Encoder Artifact Reduction features support between 90p to 1080p as input resolutions. +- The Encoder Artifact Reduction feature supports between 90p to 1080p as input resolutions. +- The Super Resolution feature supports between 90p to 2160p as input resolutions. - Super Resolution supports the following scaling factors: 4/3x (~1.33x), 1.5x, 2x, 3x and 4x. - - The maximum output resolution for the Super Resolution feature is 2160p. + - 2160p input is only supported for the following scaling factors: 4/3x (~1.33x), 1.5x and 2x + - The maximum output resolution for the Super Resolution feature is 4320p. - The Upscale feature supports any input resolution, and the following scaling factors: 4/3x (~1.33x), 1.5x, 2x, 3x and 4x. - The Webcam Denoising feature supports between 80p to 1080p as input resolutions. - The AI Green Screen and Background Blur features require that an input image/video be at least 288 pixels high. @@ -37,17 +39,17 @@ The input and output resolutions supported by the features of the SDK are listed NVIDIA MAXINE VideoEffects SDK is distributed in the following parts: - This open source repository that includes the [SDK API and proxy linking source code](https://github.com/NVIDIA/MAXINE-VFX-SDK/tree/master/nvvfx), and [sample applications and their dependency libraries](https://github.com/NVIDIA/MAXINE-VFX-SDK/tree/master/samples). -- An installer hosted on [NVIDIA Maxine developer page](https://www.nvidia.com/broadcast-sdk-resources) that installs the SDK DLLs, the models, and the SDK dependency libraries. +- An installer hosted on [NVIDIA Maxine End-user Redistributables page](https://www.nvidia.com/broadcast-sdk-resources) that installs the SDK DLLs, the models, and the SDK dependency libraries. -Please refer to [SDK programming guide](https://github.com/NVIDIA/MAXINE-VFX-SDK/blob/master/docs/NVIDIA%20Video%20Effects%20SDK%20Programming%20Guide.pdf) for configuring and integrating the SDK, compiling and running the sample applications. Please visit the [NVIDIA MAXINE Video Effects SDK](https://developer.nvidia.com/maxine-getting-started) webpage for more information about the SDK. +Please refer to the [SDK System guide](https://docs.nvidia.com/deeplearning/maxine/vfx-sdk-system-guide/) for configuring and integrating the SDK, compiling and running the sample applications. Please visit the [NVIDIA MAXINE Video Effects SDK](https://developer.nvidia.com/maxine-getting-started) webpage for more information about the SDK. ## System requirements The SDK is supported on NVIDIA GPUs that are based on the NVIDIA® Turing™ or Ampere™ architecture and have Tensor Cores. -* Windows OS supported: 64-bit Windows 10 +* Windows OS supported: 64-bit Windows 10 or later * Microsoft Visual Studio: 2017 (MSVC15.0) or later * CMake: v3.12 or later -* NVIDIA Graphics Driver for Windows: 455.57 or later +* NVIDIA Graphics Driver for Windows: 465.89 or later ## NVIDIA MAXINE Branding Guidelines If you integrate an NVIDIA MAXINE SDK within your product, please follow the required branding guidelines that are available [here](https://www.nvidia.com/maxine-sdk-guidelines/) @@ -73,3 +75,14 @@ The open source repository includes the source code to build the sample applicat 3. Use Visual Studio to generate the application binary .exe file from the NvVideoEffects_SDK.sln file. * In CMake, to open Visual Studio, click Open Project. * In Visual Studio, select Build > Build Solution. + +## Documentation +Please refer to the online documentation guides - +* [NVIDIA Video Effects SDK Programming Guide](https://docs.nvidia.com/deeplearning/maxine/vfx-sdk-programming-guide/index.html) +* [NVIDIA Video Effects SDK System Guide](https://docs.nvidia.com/deeplearning/maxine/vfx-sdk-system-guide/index.html) +* [NvCVImage API Guide](https://docs.nvidia.com/deeplearning/maxine/nvcvimage-api-guide/index.html) + +PDF versions of these guides are also available at the following locations - +* [NVIDIA Video Effects SDK Programming Guide](https://docs.nvidia.com/deeplearning/maxine/pdf/vfx-sdk-programming-guide.pdf) +* [NVIDIA Video Effects SDK System Guide](https://docs.nvidia.com/deeplearning/maxine/pdf/vfx-sdk-system-guide.pdf) +* [NvCVImage API Guide](https://docs.nvidia.com/deeplearning/maxine/pdf/nvcvimage-api-guide.pdf) diff --git a/docs/NVIDIA Video Effects SDK Programming Guide.pdf b/docs/NVIDIA Video Effects SDK Programming Guide.pdf deleted file mode 100644 index c779dc4..0000000 Binary files a/docs/NVIDIA Video Effects SDK Programming Guide.pdf and /dev/null differ diff --git a/nvvfx/include/nvCVImage.h b/nvvfx/include/nvCVImage.h index c5ed9ed..014cd94 100644 --- a/nvvfx/include/nvCVImage.h +++ b/nvvfx/include/nvCVImage.h @@ -204,21 +204,24 @@ NvCVImage { //! \param[in] dstY The top coordinate of the dst rectangle. //! \param[in] width The width of the rectangle to be copied, in pixels. //! \param[in] height The height of the rectangle to be copied, in pixels. + //! \param[in] stream the CUDA stream. //! \note NvCVImage_Transfer() can handle more cases. //! \return NVCV_SUCCESS if successful //! \return NVCV_ERR_MISMATCH if the formats are different //! \return NVCV_ERR_CUDA if a CUDA error occurred //! \return NVCV_ERR_PIXELFORMAT if the pixel format is not yet accommodated. - inline NvCV_Status copyFrom(const NvCVImage *src, int srcX, int srcY, int dstX, int dstY, unsigned width, unsigned height); + inline NvCV_Status copyFrom(const NvCVImage *src, int srcX, int srcY, int dstX, int dstY, + unsigned width, unsigned height, struct CUstream_st* stream = 0); //! Copy from one image to another. This works for CPU->CPU, CPU->GPU, GPU->GPU, and GPU->CPU. //! \param[in] src The source image from which to copy. + //! \param[in] stream the CUDA stream. //! \note NvCVImage_Transfer() can handle more cases. //! \return NVCV_SUCCESS if successful //! \return NVCV_ERR_MISMATCH if the formats are different //! \return NVCV_ERR_CUDA if a CUDA error occurred //! \return NVCV_ERR_PIXELFORMAT if the pixel format is not yet accommodated. - inline NvCV_Status copyFrom(const NvCVImage *src); + inline NvCV_Status copyFrom(const NvCVImage *src, struct CUstream_st* stream = 0); #endif // ___cplusplus } NvCVImage; @@ -466,6 +469,8 @@ NvCV_Status NvCV_API NvCVImage_TransferRect( //! \param[in] tmp a staging image. //! \return NVCV_SUCCESS if the operation was completed successfully. //! \note The actual transfer region may be smaller, because the rects are clipped against the images. +//! \note This is supplied for use with YUV buffers that do not have the standard structure +//! that are expected for NvCVImage_Transfer() and NvCVImage_TransferRect. NvCV_Status NvCV_API NvCVImage_TransferFromYUV( const void *y, int yPixBytes, int yPitch, const void *u, const void *v, int uvPixBytes, int uvPitch, @@ -492,6 +497,8 @@ NvCV_Status NvCV_API NvCVImage_TransferFromYUV( //! \param[in] tmp a staging image. //! \return NVCV_SUCCESS if the operation was completed successfully. //! \note The actual transfer region may be smaller, because the rects are clipped against the images. +//! \note This is supplied for use with YUV buffers that do not have the standard structure +//! that are expected for NvCVImage_Transfer() and NvCVImage_TransferRect. NvCV_Status NvCV_API NvCVImage_TransferToYUV( const NvCVImage *src, const NvCVRect2i *srcRect, const void *y, int yPixBytes, int yPitch, @@ -507,7 +514,9 @@ NvCV_Status NvCV_API NvCVImage_TransferToYUV( //! \param[in,out] im the image to be mapped. //! \param[in] stream the stream on which the mapping is to be performed. //! \return NVCV_SUCCESS is the operation was completed successfully. -NvCV_Status NvCV_API NvCVImage_MapResource(NvCVImage *im, struct CUstream_st *stream); +//! \note This is an experimental API. If you find it useful, please respond to XXX@YYY.com, +//! otherwise we may drop support. +/* EXPERIMENTAL */ NvCV_Status NvCV_API NvCVImage_MapResource(NvCVImage *im, struct CUstream_st *stream); //! After transfer by CUDA, the texture resource must be unmapped in order to be used by the graphics system again. @@ -516,7 +525,9 @@ NvCV_Status NvCV_API NvCVImage_MapResource(NvCVImage *im, struct CUstream_st *st //! \param[in,out] im the image to be mapped. //! \param[in] stream the CUDA stream on which the mapping is to be performed. //! \return NVCV_SUCCESS is the operation was completed successfully. -NvCV_Status NvCV_API NvCVImage_UnmapResource(NvCVImage *im, struct CUstream_st *stream); +//! \note This is an experimental API. If you find it useful, please respond to XXX@YYY.com, +//! otherwise we may drop support. +/* EXPERIMENTAL */ NvCV_Status NvCV_API NvCVImage_UnmapResource(NvCVImage *im, struct CUstream_st *stream); //! Composite one source image over another using the given matte. @@ -530,6 +541,7 @@ NvCV_Status NvCV_API NvCVImage_UnmapResource(NvCVImage *im, struct CUstream_st * //! \return NVCV_ERR_PIXELFORMAT if the pixel format is not accommodated. //! \return NVCV_ERR_MISMATCH if either the fg & bg & dst formats do not match, or if fg & bg & dst & mat are not //! in the same address space (CPU or GPU). +//! \bug Though RGBA destinations are accommodated, the A channel is not updated at all. #if RTX_CAMERA_IMAGE == 0 NvCV_Status NvCV_API NvCVImage_Composite(const NvCVImage *fg, const NvCVImage *bg, const NvCVImage *mat, NvCVImage *dst, struct CUstream_st *stream); @@ -537,8 +549,9 @@ NvCV_Status NvCV_API NvCVImage_Composite(const NvCVImage *fg, const NvCVImage *b NvCV_Status NvCV_API NvCVImage_Composite(const NvCVImage *fg, const NvCVImage *bg, const NvCVImage *mat, NvCVImage *dst); #endif // RTX_CAMERA_IMAGE == 1 + //! Composite one source image over another using the given matte. -//! Not all pixel format combinations are accommodated. +//! This accommodates all RGB and RGBA formats, with u8 and f32 components. //! \param[in] fg the foreground source image. //! \param[in] fgOrg the upper-left corner of the fg image to be composited (NULL implies (0,0)). //! \param[in] bg the background source image. @@ -566,17 +579,27 @@ NvCV_Status NvCV_API NvCVImage_CompositeRect( NvCVImage *dst, const NvCVPoint2i *dstOrg, struct CUstream_st *stream); -//! Composite a BGRu8 source image over a constant color field using the given matte. -//! \param[in] src the source BGRu8 (or RGBu8) image. -//! \param[in] mat the matte Yu8 (or Au8) image, indicating where the src should come through. -//! \param[in] bgColor the desired flat background color, with the same component ordering as the src and dst. -//! \param[in,out] dst the destination BGRu8 (or RGBu8) image. May be the same as src. + +//! Composite a source image over a constant color field using the given matte. +//! \param[in] src the source image. +//! \param[in] mat the matte image, indicating where the src should come through. +//! \param[in] bgColor pointer to a location holding the desired flat background color, with the same format +//! and component ordering as the dst. This acts as a 1x1 background pixel buffer, +//! so should reside in the same memory space (CUDA or CPU) as the other buffers. +//! \param[in,out] dst the destination image. May be the same as src. //! \return NVCV_SUCCESS if the operation was successful. //! \return NVCV_ERR_PIXELFORMAT if the pixel format is not accommodated. -//! \bug This is only implemented for 3-component u8 src and dst, and 1-component mat, -//! where all images are resident on the CPU. +//! \return NVCV_ERR_MISMATCH if fg & mat & dst & bgColor are not in the same address space (CPU or GPU). +//! \note The bgColor must remain valid until complete; this is an important consideration especially if +//! the buffers are on the GPU and NvCVImage_CompositeOverConstant() runs asynchronously. +//! \bug Though RGBA destinations are accommodated, the A channel is not updated at all. NvCV_Status NvCV_API NvCVImage_CompositeOverConstant( - const NvCVImage *src, const NvCVImage *mat, const unsigned char bgColor[3], NvCVImage *dst); +#if RTX_CAMERA_IMAGE == 0 + const NvCVImage *src, const NvCVImage *mat, const void *bgColor, NvCVImage *dst, struct CUstream_st *stream +#else // RTX_CAMERA_IMAGE == 1 + const NvCVImage *src, const NvCVImage *mat, const unsigned char bgColor[3], NvCVImage *dst +#endif // RTX_CAMERA_IMAGE +); //! Flip the image vertically. @@ -649,16 +672,16 @@ NvCVImage::~NvCVImage() { NvCVImage_Dealloc(this); } ********************************************************************************/ NvCV_Status NvCVImage::copyFrom(const NvCVImage *src, int srcX, int srcY, int dstX, int dstY, unsigned wd, - unsigned ht) { + unsigned ht, struct CUstream_st* stream) { #if RTX_CAMERA_IMAGE // This only works for chunky images NvCVImage srcView, dstView; NvCVImage_InitView(&srcView, const_cast(src), srcX, srcY, wd, ht); NvCVImage_InitView(&dstView, this, dstX, dstY, wd, ht); - return NvCVImage_Transfer(&srcView, &dstView, 1.f, 0, nullptr); + return NvCVImage_Transfer(&srcView, &dstView, 1.f, stream, nullptr); #else // !RTX_CAMERA_IMAGE bug fix for non-chunky images NvCVRect2i srcRect = { (int)srcX, (int)srcY, (int)wd, (int)ht }; NvCVPoint2i dstPt = { (int)dstX, (int)dstY }; - return NvCVImage_TransferRect(src, &srcRect, this, &dstPt, 1.f, 0, nullptr); + return NvCVImage_TransferRect(src, &srcRect, this, &dstPt, 1.f, stream, nullptr); #endif // RTX_CAMERA_IMAGE } @@ -666,7 +689,9 @@ NvCV_Status NvCVImage::copyFrom(const NvCVImage *src, int srcX, int srcY, int ds * copy image ********************************************************************************/ -NvCV_Status NvCVImage::copyFrom(const NvCVImage *src) { return NvCVImage_Transfer(src, this, 1.f, 0, nullptr); } +NvCV_Status NvCVImage::copyFrom(const NvCVImage *src, struct CUstream_st* stream) { + return NvCVImage_Transfer(src, this, 1.f, stream, nullptr); +} #endif // ___cplusplus diff --git a/nvvfx/include/nvTransferD3D.h b/nvvfx/include/nvTransferD3D.h index e914eb5..b560322 100644 --- a/nvvfx/include/nvTransferD3D.h +++ b/nvvfx/include/nvTransferD3D.h @@ -32,7 +32,8 @@ extern "C" { //! \param[in] layout the layout. //! \param[out] d3dFormat a place to store the corresponding D3D format. //! \return NVCV_SUCCESS if successful. -NvCV_Status NvCV_API NvCVImage_ToD3DFormat(NvCVImage_PixelFormat format, NvCVImage_ComponentType type, unsigned layout, DXGI_FORMAT *d3dFormat); +//! \note This is an experimental API. If you find it useful, please respond to XXX@YYY.com, otherwise we may drop support. +/* EXPERIMENTAL */ NvCV_Status NvCV_API NvCVImage_ToD3DFormat(NvCVImage_PixelFormat format, NvCVImage_ComponentType type, unsigned layout, DXGI_FORMAT *d3dFormat); //! Utility to determine the NvCVImage format, component type and layout from a D3D format. @@ -41,7 +42,8 @@ NvCV_Status NvCV_API NvCVImage_ToD3DFormat(NvCVImage_PixelFormat format, NvCVIma //! \param[out] type a place to store the NvCVImage component type. //! \param[out] layout a place to store the NvCVImage layout. //! \return NVCV_SUCCESS if successful. -NvCV_Status NvCV_API NvCVImage_FromD3DFormat(DXGI_FORMAT d3dFormat, NvCVImage_PixelFormat *format, NvCVImage_ComponentType *type, unsigned char *layout); +//! \note This is an experimental API. If you find it useful, please respond to XXX@YYY.com, otherwise we may drop support. +/* EXPERIMENTAL */ NvCV_Status NvCV_API NvCVImage_FromD3DFormat(DXGI_FORMAT d3dFormat, NvCVImage_PixelFormat *format, NvCVImage_ComponentType *type, unsigned char *layout); #ifdef __dxgicommon_h__ @@ -51,7 +53,8 @@ NvCV_Status NvCV_API NvCVImage_FromD3DFormat(DXGI_FORMAT d3dFormat, NvCVImage_Pi //! \param[out] pD3dColorSpace a place to store the resultant D3D color space. //! \return NVCV_SUCCESS if successful. //! \return NVCV_ERR_PIXELFORMAT if there is no equivalent color space. -NvCV_Status NvCV_API NvCVImage_ToD3DColorSpace(unsigned char nvcvColorSpace, DXGI_COLOR_SPACE_TYPE *pD3dColorSpace); +//! \note This is an experimental API. If you find it useful, please respond to XXX@YYY.com, otherwise we may drop support. +/* EXPERIMENTAL */ NvCV_Status NvCV_API NvCVImage_ToD3DColorSpace(unsigned char nvcvColorSpace, DXGI_COLOR_SPACE_TYPE *pD3dColorSpace); //! Utility to determine the NvCVImage color space from the D3D color space. @@ -59,7 +62,8 @@ NvCV_Status NvCV_API NvCVImage_ToD3DColorSpace(unsigned char nvcvColorSpace, DXG //! \param[out] pNvcvColorSpace a place to store the resultant NvCVImage color space. //! \return NVCV_SUCCESS if successful. //! \return NVCV_ERR_PIXELFORMAT if there is no equivalent color space. -NvCV_Status NvCV_API NvCVImage_FromD3DColorSpace(DXGI_COLOR_SPACE_TYPE d3dColorSpace, unsigned char *pNvcvColorSpace); +//! \note This is an experimental API. If you find it useful, please respond to XXX@YYY.com, otherwise we may drop support. +/* EXPERIMENTAL */ NvCV_Status NvCV_API NvCVImage_FromD3DColorSpace(DXGI_COLOR_SPACE_TYPE d3dColorSpace, unsigned char *pNvcvColorSpace); #endif // __dxgicommon_h__ diff --git a/nvvfx/include/nvTransferD3D11.h b/nvvfx/include/nvTransferD3D11.h index fabf067..b827ed1 100644 --- a/nvvfx/include/nvTransferD3D11.h +++ b/nvvfx/include/nvTransferD3D11.h @@ -26,13 +26,14 @@ extern "C" { //! Initialize an NvCVImage from a D3D11 texture. //! The pixelFormat and component types with be transferred over, and a cudaGraphicsResource will be registered; //! the NvCVImage destructor will unregister the resource. -//! This is designed to work with NvCVImage_TransferFromArray() (and eventually NvCVImage_Transfer()); -//! however it is necessary to call NvCVImage_MapResource beforehand, and NvCVImage_UnmapResource -//! before allowing D3D to render into it. +//! It is necessary to call NvCVImage_MapResource() after rendering D3D and before calling NvCVImage_Transfer(), +//! and to call NvCVImage_UnmapResource() before rendering in D3D again. //! \param[in,out] im the image to be initialized. //! \param[in] tx the texture to be used for initialization. //! \return NVCV_SUCCESS if successful. -NvCV_Status NvCV_API NvCVImage_InitFromD3D11Texture(NvCVImage *im, struct ID3D11Texture2D *tx); +//! \note This is an experimental API. If you find it useful, please respond to XXX@YYY.com, +//! otherwise we may drop support. +/* EXPERIMENTAL */ NvCV_Status NvCV_API NvCVImage_InitFromD3D11Texture(NvCVImage *im, struct ID3D11Texture2D *tx); diff --git a/nvvfx/include/nvVideoEffects.h b/nvvfx/include/nvVideoEffects.h index 9d461bc..7434175 100644 --- a/nvvfx/include/nvVideoEffects.h +++ b/nvvfx/include/nvVideoEffects.h @@ -207,6 +207,7 @@ NvCV_Status NvVFX_API NvVFX_CudaStreamDestroy(CUstream stream); #define NVVFX_OUTPUT_IMAGE NVVFX_OUTPUT_IMAGE_0 //!< but there is usually only one output image #define NVVFX_MODEL_DIRECTORY "ModelDir" //!< The directory where the model may be found #define NVVFX_CUDA_STREAM "CudaStream" //!< The CUDA stream to use +#define NVVFX_CUDA_GRAPH "CudaGraph" //!< Enable CUDA graph to use #define NVVFX_INFO "Info" //!< Get info about the effects #define NVVFX_SCALE "Scale" //!< Scale factor #define NVVFX_STRENGTH "Strength" //!< Strength for different filters diff --git a/nvvfx/src/NVVideoEffectsProxy.cpp b/nvvfx/src/NVVideoEffectsProxy.cpp index b7fca4b..631e81e 100644 --- a/nvvfx/src/NVVideoEffectsProxy.cpp +++ b/nvvfx/src/NVVideoEffectsProxy.cpp @@ -1,3 +1,6 @@ +#if defined(linux) || defined(unix) || defined(__linux) +#warning nvVideoEffectsProxy.cpp not ported +#else // _WIN32_ /*############################################################################### # # Copyright (c) 2020 NVIDIA Corporation @@ -37,9 +40,9 @@ // Parameter string does not include the file extension #ifdef _WIN32 -#define nvLoadLibrary(library) LoadLibrary(TEXT(library ".dll")) + #define nvLoadLibrary(library) LoadLibrary(TEXT(library ".dll")) #else // !_WIN32 -#define nvLoadLibrary(library) dlopen("lib" library ".so", RTLD_LAZY) + #define nvLoadLibrary(library) dlopen("lib" library ".so", RTLD_LAZY) #endif // _WIN32 @@ -55,27 +58,44 @@ inline void* nvGetProcAddress(HINSTANCE handle, const char* proc) { inline int nvFreeLibrary(HINSTANCE handle) { #ifdef _WIN32 return FreeLibrary(handle); -#else +#else // !_WIN32 return dlclose(handle); -#endif +#endif // _WIN32 } HINSTANCE getNvVfxLib() { TCHAR path[MAX_PATH], fullPath[MAX_PATH]; + bool bSDKPathSet = false; - // There can be multiple apps on the system, - // some might include the SDK in the app package and - // others might expect the SDK to be installed in Program Files - GetEnvironmentVariable(TEXT("NV_VIDEO_EFFECTS_PATH"), path, MAX_PATH); - if (_tcscmp(path, TEXT("USE_APP_PATH"))) { - // App has not set environment variable to "USE_APP_PATH" - // So pick up the SDK dll and dependencies from Program Files - GetEnvironmentVariable(TEXT("ProgramFiles"), path, MAX_PATH); - size_t max_len = sizeof(fullPath)/sizeof(TCHAR); - _stprintf_s(fullPath, max_len, TEXT("%s\\NVIDIA Corporation\\NVIDIA Video Effects\\"), path); + extern char* g_nvVFXSDKPath; + if (g_nvVFXSDKPath && g_nvVFXSDKPath[0]) { +#ifndef UNICODE + strncpy_s(fullPath, MAX_PATH, g_nvVFXSDKPath, MAX_PATH); + +#else // !UNICODE + size_t res = 0; + mbstowcs_s(&res, fullPath, MAX_PATH, g_nvVFXSDKPath, MAX_PATH); +#endif // UNICODE SetDllDirectory(fullPath); + bSDKPathSet = true; } + + if (!bSDKPathSet) { + // There can be multiple apps on the system, + // some might include the SDK in the app package and + // others might expect the SDK to be installed in Program Files + GetEnvironmentVariable(TEXT("NV_VIDEO_EFFECTS_PATH"), path, MAX_PATH); + if (_tcscmp(path, TEXT("USE_APP_PATH"))) { + // App has not set environment variable to "USE_APP_PATH" + // So pick up the SDK dll and dependencies from Program Files + GetEnvironmentVariable(TEXT("ProgramFiles"), path, MAX_PATH); + size_t max_len = sizeof(fullPath) / sizeof(TCHAR); + _stprintf_s(fullPath, max_len, TEXT("%s\\NVIDIA Corporation\\NVIDIA Video Effects\\"), path); + SetDllDirectory(fullPath); + } + } + static const HINSTANCE NvVfxLib = nvLoadLibrary("NVVideoEffects"); return NvVfxLib; } @@ -256,3 +276,4 @@ NvCV_Status NvVFX_API NvVFX_CudaStreamDestroy(CUstream stream) { return funcPtr(stream); } +#endif // enabling for this file diff --git a/nvvfx/src/nvCVImageProxy.cpp b/nvvfx/src/nvCVImageProxy.cpp index f724d7a..4f5199a 100644 --- a/nvvfx/src/nvCVImageProxy.cpp +++ b/nvvfx/src/nvCVImageProxy.cpp @@ -67,7 +67,11 @@ inline int nvFreeLibrary(HINSTANCE handle) { HINSTANCE getNvCVImageLib() { TCHAR path[MAX_PATH], tmpPath[MAX_PATH], fullPath[MAX_PATH]; static HINSTANCE nvCVImageLib = NULL; - static bool bSDKPathSet = false; + static bool bSDKPathSet = false; + if (!bSDKPathSet) { + nvCVImageLib = nvLoadLibrary("NVCVImage"); + if (nvCVImageLib) bSDKPathSet = true; + } if (!bSDKPathSet) { // There can be multiple apps on the system, // some might include the SDK in the app package and @@ -234,16 +238,28 @@ NvCV_Status NvCV_API NvCVImage_CompositeRect( return funcPtr(fg, fgOrg, bg, bgOrg, mat, mode, dst, dstOrg, stream); } -NvCV_Status NvCV_API NvCVImage_CompositeOverConstant(const NvCVImage* src, const NvCVImage* mat, - const unsigned char bgColor[3], NvCVImage* dst) { +#if RTX_CAMERA_IMAGE == 0 +NvCV_Status NvCV_API NvCVImage_CompositeOverConstant(const NvCVImage *src, const NvCVImage *mat, + const void *bgColor, NvCVImage *dst, struct CUstream_st *stream) { + static const auto funcPtr = + (decltype(NvCVImage_CompositeOverConstant)*)nvGetProcAddress(getNvCVImageLib(), "NvCVImage_CompositeOverConstant"); + + if (nullptr == funcPtr) return NVCV_ERR_LIBRARY; + return funcPtr(src, mat, bgColor, dst, stream); +} +#else // RTX_CAMERA_IMAGE == 1 +NvCV_Status NvCV_API NvCVImage_CompositeOverConstant(const NvCVImage *src, const NvCVImage *mat, + const unsigned char bgColor[3], NvCVImage *dst) { static const auto funcPtr = (decltype(NvCVImage_CompositeOverConstant)*)nvGetProcAddress(getNvCVImageLib(), "NvCVImage_CompositeOverConstant"); if (nullptr == funcPtr) return NVCV_ERR_LIBRARY; return funcPtr(src, mat, bgColor, dst); } +#endif // RTX_CAMERA_IMAGE -NvCV_Status NvCV_API NvCVImage_FlipY(const NvCVImage* src, NvCVImage* dst) { + +NvCV_Status NvCV_API NvCVImage_FlipY(const NvCVImage *src, NvCVImage *dst) { static const auto funcPtr = (decltype(NvCVImage_FlipY)*)nvGetProcAddress(getNvCVImageLib(), "NvCVImage_FlipY"); if (nullptr == funcPtr) return NVCV_ERR_LIBRARY; diff --git a/resources/Denoise.gif b/resources/Denoise.gif index e09fb1a..9d9c011 100644 Binary files a/resources/Denoise.gif and b/resources/Denoise.gif differ diff --git a/resources/SR.gif b/resources/SR.gif index b22a5bb..190b70a 100644 Binary files a/resources/SR.gif and b/resources/SR.gif differ diff --git a/samples/AigsEffectApp/AigsEffectApp.cpp b/samples/AigsEffectApp/AigsEffectApp.cpp index 4019738..981a44d 100644 --- a/samples/AigsEffectApp/AigsEffectApp.cpp +++ b/samples/AigsEffectApp/AigsEffectApp.cpp @@ -63,9 +63,9 @@ bool FLAG_progress = false; bool FLAG_show = false; -bool FLAG_useOTAU = false; bool FLAG_verbose = false; bool FLAG_webcam = false; +bool FLAG_cudaGraph = false; int FLAG_compMode = 3 /*compWhite*/; int FLAG_mode = 0; float FLAG_blurStrength = 0.5; @@ -75,6 +75,7 @@ std::string FLAG_inFile; std::string FLAG_modelDir; std::string FLAG_outDir; std::string FLAG_outFile; +std::string FLAG_bgFile; static bool GetFlagArgVal(const char *flag, const char *arg, const char **val) { if (*arg != '-') return false; @@ -135,6 +136,7 @@ static void Usage() { " where args is:\n" " --in_file= input file to be processed\n" " --out_file= output file to be written\n" + " --bg_file= background file for composition\n" " --webcam use a webcam as input\n" " --cam_res=[WWWx]HHH specify resolution as height or width x height\n" " --model_dir= the path to the directory that contains the models\n" @@ -144,8 +146,15 @@ static void Usage() { " --mode=(0|1) pick one of the green screen modes\n" " 0 - Best quality\n" " 1 - Best performance\n" - " --comp_mode choose the composition mode - { compMatte = 0, compLight = 1, compGreen = 2, compWhite = 3, compNone = 4, compBG = 5, compBlur = 6}\n" - " --blur_strength change the blur strength, range is [0, 1]" + " --comp_mode choose the composition mode - {\n" + " 0 (show matte - compMatte),\n" + " 1 (overlay mask on foreground - compLight),\n" + " 2 (composite over green - compGreen),\n" + " 3 (composite over white - compWhite),\n" + " 4 (show input - compNone),\n" + " 5 (composite over a specified background image - compBG),\n" + " 6 (blur the background of the image - compBlur) }\n" + " --cuda_graph Enable cuda graph.\n" ); } @@ -160,10 +169,12 @@ static int ParseMyArgs(int argc, char **argv) { (GetFlagArgVal("verbose", arg, &FLAG_verbose) || GetFlagArgVal("in", arg, &FLAG_inFile) || GetFlagArgVal("in_file", arg, &FLAG_inFile) || GetFlagArgVal("out", arg, &FLAG_outFile) || GetFlagArgVal("out_file", arg, &FLAG_outFile) || GetFlagArgVal("model_dir", arg, &FLAG_modelDir) || + GetFlagArgVal("bg_file", arg, &FLAG_bgFile) || GetFlagArgVal("codec", arg, &FLAG_codec) || GetFlagArgVal("webcam", arg, &FLAG_webcam) || GetFlagArgVal("cam_res", arg, &FLAG_camRes) || GetFlagArgVal("mode", arg, &FLAG_mode) || GetFlagArgVal("progress", arg, &FLAG_progress) || GetFlagArgVal("show", arg, &FLAG_show) || - GetFlagArgVal("comp_mode", arg, &FLAG_compMode) || GetFlagArgVal("blur_strength", arg, &FLAG_blurStrength))) { + GetFlagArgVal("comp_mode", arg, &FLAG_compMode) || GetFlagArgVal("blur_strength", arg, &FLAG_blurStrength) || + GetFlagArgVal("cuda_graph", arg, &FLAG_cudaGraph) )) { continue; } else if (GetFlagArgVal("help", arg, &help)) { return NVCV_ERR_HELP; @@ -208,6 +219,10 @@ static bool IsImageFile(const char *str) { return HasOneOfTheseSuffixes(str, ".bmp", ".jpg", ".jpeg", ".png", nullptr); } +static bool IsLossyImageFile(const char *str) { + return HasOneOfTheseSuffixes(str, ".jpg", ".jpeg", nullptr); +} + static const char *DurationString(double sc) { static char buf[16]; int hr, mn; @@ -246,9 +261,9 @@ static void GetVideoInfo(cv::VideoCapture &reader, const char *fileName, VideoIn info->height = (int)reader.get(cv::CAP_PROP_FRAME_HEIGHT); info->frameRate = (double)reader.get(cv::CAP_PROP_FPS); if(!strcmp(fileName,"webcam")) - info->frameCount = 0; + info->frameCount = 0; else - info->frameCount = (long long)reader.get(cv::CAP_PROP_FRAME_COUNT); + info->frameCount = (long long)reader.get(cv::CAP_PROP_FRAME_COUNT); if (FLAG_verbose) PrintVideoInfo(info, fileName); } @@ -282,8 +297,8 @@ struct FXApp { errLibrary = NVCV_ERR_LIBRARY, errInitialization = NVCV_ERR_INITIALIZATION, errFileNotFound = NVCV_ERR_FILE, - errFeatureNotFound = NVCV_ERR_FEATURENOTFOUND, - errMissingInput = NVCV_ERR_MISSINGINPUT, + errFeatureNotFound = NVCV_ERR_FEATURENOTFOUND, + errMissingInput = NVCV_ERR_MISSINGINPUT, errResolution = NVCV_ERR_RESOLUTION, errUnsupportedGPU = NVCV_ERR_UNSUPPORTEDGPU, errWrongGPU = NVCV_ERR_WRONGGPU, @@ -341,6 +356,8 @@ struct FXApp { NvVFX_Handle _eff, _bgblurEff; cv::Mat _srcImg; cv::Mat _dstImg; + cv::Mat _bgImg; + cv::Mat _resizedCroppedBgImg; NvCVImage _srcVFX; NvCVImage _dstVFX; bool _show; @@ -401,24 +418,26 @@ void FXApp::drawFrameRate(cv::Mat &img) { void FXApp::nextCompMode() { switch (_compMode) { default: - case compBG: + case compMatte: + _compMode = compLight; + break; case compLight: _compMode = compGreen; break; - case compMatte: - _compMode = compNone; - break; case compGreen: _compMode = compWhite; break; case compWhite: - _compMode = compMatte; + _compMode = compNone; break; case compNone: + _compMode = compBG; + break; + case compBG: _compMode = compBlur; break; case compBlur: - _compMode = compLight; + _compMode = compMatte; break; } } @@ -473,7 +492,7 @@ NvCV_Status FXApp::createAigsEffect() { if (!FLAG_modelDir.empty()) { vfxErr = NvVFX_SetString(_eff, NVVFX_MODEL_DIRECTORY, FLAG_modelDir.c_str()); - } + } if (vfxErr != NVCV_SUCCESS) { std::cerr << "Error setting the model path to \"" << FLAG_modelDir << "\"\n"; return vfxErr; @@ -493,6 +512,12 @@ NvCV_Status FXApp::createAigsEffect() { return vfxErr; } + vfxErr = NvVFX_SetU32(_eff, NVVFX_CUDA_GRAPH, FLAG_cudaGraph?1u:0u); + if (vfxErr != NVCV_SUCCESS) { + std::cerr << "Error enabling cuda graph \n"; + return vfxErr; + } + vfxErr = NvVFX_CudaStreamCreate(&_stream); if (vfxErr != NVCV_SUCCESS) { std::cerr << "Error creating CUDA stream " << std::endl; @@ -570,6 +595,8 @@ FXApp::Err FXApp::processImage(const char *inFile, const char *outFile) { overlay(_srcImg, _dstImg, 0.5, result); if (!std::string(outFile).empty()) { + if(IsLossyImageFile(outFile)) + fprintf(stderr, "WARNING: JPEG output file format will reduce image quality\n"); ok = cv::imwrite(outFile, result); if (!ok) { printf("Error writing: \"%s\"\n", outFile); @@ -647,6 +674,30 @@ FXApp::Err FXApp::processMovie(const char *inFile, const char *outFile) { unsigned int width = (unsigned int)reader.get(cv::CAP_PROP_FRAME_WIDTH); unsigned int height = (unsigned int)reader.get(cv::CAP_PROP_FRAME_HEIGHT); + if (!FLAG_bgFile.empty()) + { + _bgImg = cv::imread(FLAG_bgFile); + if (!_bgImg.data) + { + return errRead; + } + else + { + // Find the scale to resize background such that image can fit into background + float scale = float(height) / float(_bgImg.rows); + if ((scale * _bgImg.cols) < float(width)) + { + scale = float(width) / float(_bgImg.cols); + } + cv::Mat resizedBg; + cv::resize(_bgImg, resizedBg, cv::Size(), scale, scale, cv::INTER_AREA); + + // Always crop from top left of background. + cv::Rect rect(0, 0, width, height); + _resizedCroppedBgImg = resizedBg(rect); + } + } + // allocate src for GPU if (!_srcNvVFXImage.pixels) BAIL_IF_ERR(vfxErr = @@ -695,6 +746,25 @@ FXApp::Err FXApp::processMovie(const char *inFile, const char *outFile) { case compNone: _srcImg.copyTo(result); break; + case compBG: { + if (FLAG_bgFile.empty()) + { + _resizedCroppedBgImg = cv::Mat(_srcImg.rows, _srcImg.cols, CV_8UC3, cv::Scalar(118, 185, 0)); + size_t startX = _resizedCroppedBgImg.cols/20; + size_t offsetY = _resizedCroppedBgImg.rows/20; + std::string text = "No Background Image!"; + for (size_t startY = offsetY; startY < _resizedCroppedBgImg.rows; startY += offsetY) + { + cv::putText(_resizedCroppedBgImg, text, cv::Point(startX, startY), + cv::FONT_HERSHEY_DUPLEX, 1.0, CV_RGB(0, 0, 0), 1); + } + } + NvCVImage bgVFX; + (void)NVWrapperForCVMat(&_resizedCroppedBgImg, &bgVFX); + NvCVImage matVFX; + (void)NVWrapperForCVMat(&result, &matVFX); + NvCVImage_Composite(&_srcVFX, &bgVFX, &_dstVFX, &matVFX, _stream); + } break; case compLight: if (inFile) { overlay(_srcImg, _dstImg, 0.5, result); @@ -708,13 +778,13 @@ FXApp::Err FXApp::processMovie(const char *inFile, const char *outFile) { const unsigned char bgColor[3] = {0, 255, 0}; NvCVImage matVFX; (void)NVWrapperForCVMat(&result, &matVFX); - NvCVImage_CompositeOverConstant(&_srcVFX, &_dstVFX, bgColor, &matVFX); + NvCVImage_CompositeOverConstant(&_srcVFX, &_dstVFX, bgColor, &matVFX, _stream); } break; case compWhite: { const unsigned char bgColor[3] = {255, 255, 255}; NvCVImage matVFX; (void)NVWrapperForCVMat(&result, &matVFX); - NvCVImage_CompositeOverConstant(&_srcVFX, &_dstVFX, bgColor, &matVFX); + NvCVImage_CompositeOverConstant(&_srcVFX, &_dstVFX, bgColor, &matVFX, _stream); } break; case compMatte: cv::cvtColor(_dstImg, result, cv::COLOR_GRAY2BGR); @@ -726,7 +796,7 @@ FXApp::Err FXApp::processMovie(const char *inFile, const char *outFile) { BAIL_IF_ERR(vfxErr = NvVFX_SetImage(_bgblurEff, NVVFX_OUTPUT_IMAGE, &_blurNvVFXImage)); BAIL_IF_ERR(vfxErr = NvVFX_Load(_bgblurEff)); BAIL_IF_ERR(vfxErr = NvVFX_Run(_bgblurEff, 0)); - + NvCVImage matVFX; (void)NVWrapperForCVMat(&result, &matVFX); BAIL_IF_ERR(vfxErr = NvCVImage_Transfer(&_blurNvVFXImage, &matVFX, 1.0f, _stream, NULL)); @@ -750,12 +820,11 @@ FXApp::Err FXApp::processMovie(const char *inFile, const char *outFile) { if (errQuit == appErr) break; } } - if (_progress) { - if(info.frameCount == 0) // no progress for a webcam - fprintf(stderr, "\b\b\b\b???%%"); - else - fprintf(stderr, "\b\b\b\b%3.0f%%", 100.f * frameNum / info.frameCount); - } + if (_progress) + if(info.frameCount == 0) // no progress for a webcam + fprintf(stderr, "\b\b\b\b???%%"); + else + fprintf(stderr, "\b\b\b\b%3.0f%%", 100.f * frameNum / info.frameCount); } if (_progress) fprintf(stderr, "\n"); @@ -769,17 +838,44 @@ bail: return appErrFromVfxStatus(vfxErr); } +// This path is used by nvVideoEffectsProxy.cpp to load the SDK dll +char *g_nvVFXSDKPath = NULL; + +int chooseGPU() { + // If the system has multiple supported GPUs then the application + // should use CUDA driver APIs or CUDA runtime APIs to enumerate + // the GPUs and select one based on the application's requirements + + // Cuda device 0 + return 0; +} + +bool isCompModeEnumValid(const FXApp::CompMode& mode) +{ + if (mode != FXApp::CompMode::compMatte && + mode != FXApp::CompMode::compLight && + mode != FXApp::CompMode::compGreen && + mode != FXApp::CompMode::compWhite && + mode != FXApp::CompMode::compNone && + mode != FXApp::CompMode::compBG && + mode != FXApp::CompMode::compBlur) + { + return false; + } + return true; +} + int main(int argc, char **argv) { int nErrs = 0; - FXApp::Err fxErr = FXApp::errNone; - FXApp app; - nErrs = ParseMyArgs(argc, argv); if (nErrs) { Usage(); return nErrs; } + FXApp::Err fxErr = FXApp::errNone; + FXApp app; + if (FLAG_inFile.empty() && !FLAG_webcam) { std::cerr << "Please specify --in_file=XXX or --webcam\n"; ++nErrs; @@ -793,6 +889,12 @@ int main(int argc, char **argv) { app.setShow(FLAG_show); app._compMode = static_cast(FLAG_compMode); + if (!isCompModeEnumValid(app._compMode)) + { + std::cerr << "Please specify a valid --comp_mode=XXX, valid range is [0,6] check help section\n"; + ++nErrs; + } + app._blurStrength = FLAG_blurStrength; if (app._blurStrength < 0) { app._blurStrength = 0; @@ -823,4 +925,4 @@ int main(int argc, char **argv) { if (fxErr) std::cerr << "Error: " << app.errorStringFromCode(fxErr) << std::endl; return (int)fxErr; -} +} \ No newline at end of file diff --git a/samples/AigsEffectApp/AigsEffectApp.exe b/samples/AigsEffectApp/AigsEffectApp.exe index a1feb4c..02274be 100644 Binary files a/samples/AigsEffectApp/AigsEffectApp.exe and b/samples/AigsEffectApp/AigsEffectApp.exe differ diff --git a/samples/AigsEffectApp/CMakeLists.txt b/samples/AigsEffectApp/CMakeLists.txt index 3bcce30..fcb804d 100644 --- a/samples/AigsEffectApp/CMakeLists.txt +++ b/samples/AigsEffectApp/CMakeLists.txt @@ -22,11 +22,10 @@ target_link_libraries(AigsEffectApp PUBLIC ) set(OPENCV_PATH_STR ${CMAKE_CURRENT_SOURCE_DIR}/../external/opencv/bin) -set(VFXSDK_PATH_STR ${CMAKE_CURRENT_SOURCE_DIR}/../../bin) # Also the location for CUDA/NVTRT/libcrypto -set(PATH_STR "PATH=%PATH%" ${VFXSDK_PATH_STR} ${OPENCV_PATH_STR}) -set(CMD_ARG_STR "--show --in_file=\"${CMAKE_CURRENT_SOURCE_DIR}/../input/input_003054.jpg\"") +set(PATH_STR "PATH=%PATH%" ${OPENCV_PATH_STR}) +set(CMD_ARG_STR "--show --in_file=\"${CMAKE_CURRENT_SOURCE_DIR}/../input/input_003054.jpg\" ") set_target_properties(AigsEffectApp PROPERTIES FOLDER SampleApps VS_DEBUGGER_ENVIRONMENT "${PATH_STR}" VS_DEBUGGER_COMMAND_ARGUMENTS "${CMD_ARG_STR}" -) + ) \ No newline at end of file diff --git a/samples/BatchEffectApp/BatchDenoiseEffectApp.exe b/samples/BatchEffectApp/BatchDenoiseEffectApp.exe index cb9c803..7a056fc 100644 Binary files a/samples/BatchEffectApp/BatchDenoiseEffectApp.exe and b/samples/BatchEffectApp/BatchDenoiseEffectApp.exe differ diff --git a/samples/BatchEffectApp/BatchEffectApp.cpp b/samples/BatchEffectApp/BatchEffectApp.cpp index adae384..2e67dcc 100644 --- a/samples/BatchEffectApp/BatchEffectApp.cpp +++ b/samples/BatchEffectApp/BatchEffectApp.cpp @@ -21,6 +21,7 @@ # ###############################################################################*/ +#include #include #include @@ -127,8 +128,7 @@ static void Usage() { " where flags is:\n" " --out_file= output image files to be written, default \"BatchOut_%%02u.png\"\n" " --effect= the effect to apply\n" - " --strength= strength of an effect, 0 or 1 for super res and artifact reduction,\n" - " and [0.0, 1.0] for upscaling\n" + " --strength= strength of the upscaling effect, [0.0, 1.0]\n" " --scale= scale factor to be applied: 1.5, 2, 3, maybe 1.3333333\n" " --resolution= the desired height (either --scale or --resolution may be used)\n" " --mode= mode 0 or 1\n" @@ -183,6 +183,32 @@ static int ParseMyArgs(int argc, char **argv) { return errs; } +static bool HasSuffix(const char *str, const char *suf) { + size_t strSize = strlen(str), + sufSize = strlen(suf); + if (strSize < sufSize) + return false; + return (0 == strcasecmp(suf, str + strSize - sufSize)); +} + +static bool HasOneOfTheseSuffixes(const char *str, ...) { + bool matches = false; + const char *suf; + va_list ap; + va_start(ap, str); + while (nullptr != (suf = va_arg(ap, const char*))) { + if (HasSuffix(str, suf)) { + matches = true; + break; + } + } + va_end(ap); + return matches; +} + +static bool IsLossyImageFile(const char *str) { + return HasOneOfTheseSuffixes(str, ".jpg", ".jpeg", nullptr); +} class App { public: @@ -233,7 +259,7 @@ public: BAIL_IF_ERR(err = AllocateBatchBuffer(&_src, _batchSize, src->width, src->height, NVCV_BGR, NVCV_F32, NVCV_PLANAR, NVCV_CUDA, 1)); BAIL_IF_ERR(err = AllocateBatchBuffer(&_dst, _batchSize, src->width, src->height, NVCV_BGR, NVCV_F32, NVCV_PLANAR, NVCV_CUDA, 1)); BAIL_IF_ERR(err = NvVFX_SetString(_eff, NVVFX_MODEL_DIRECTORY, FLAG_modelDir.c_str())); - BAIL_IF_ERR(err = NvVFX_SetU32(_eff, NVVFX_STRENGTH, (unsigned)FLAG_strength)); + BAIL_IF_ERR(err = NvVFX_SetU32(_eff, NVVFX_MODE, FLAG_mode)); } #endif // NVVFX_FX_ARTIFACT_REDUCTION #ifdef NVVFX_FX_SUPER_RES @@ -241,7 +267,7 @@ public: BAIL_IF_ERR(err = AllocateBatchBuffer(&_src, _batchSize, src->width, src->height, NVCV_BGR, NVCV_F32, NVCV_PLANAR, NVCV_CUDA, 1)); BAIL_IF_ERR(err = AllocateBatchBuffer(&_dst, _batchSize, dw, dh, NVCV_BGR, NVCV_F32, NVCV_PLANAR, NVCV_CUDA, 1)); BAIL_IF_ERR(err = NvVFX_SetString(_eff, NVVFX_MODEL_DIRECTORY, FLAG_modelDir.c_str())); - BAIL_IF_ERR(err = NvVFX_SetU32(_eff, NVVFX_STRENGTH, (unsigned)FLAG_strength)); + BAIL_IF_ERR(err = NvVFX_SetU32(_eff, NVVFX_MODE, FLAG_mode)); } #endif // NVVFX_FX_SUPER_RES else { @@ -339,6 +365,8 @@ NvCV_Status BatchProcessImages(const char* effectName, const std::vector _tmpVFX --> _srcGpuBuf BAIL_IF_ERR(vfxErr = NvVFX_SetImage(_eff, NVVFX_INPUT_IMAGE, &_srcGpuBuf)); - BAIL_IF_ERR(vfxErr = NvVFX_SetImage(_eff, NVVFX_OUTPUT_IMAGE, &_dstGpuBuf)); + BAIL_IF_ERR(vfxErr = NvVFX_SetImage(_eff, NVVFX_OUTPUT_IMAGE, &_dstGpuBuf)); BAIL_IF_ERR(vfxErr = NvVFX_SetF32(_eff, NVVFX_STRENGTH, FLAG_strength)); - + unsigned int stateSizeInBytes; BAIL_IF_ERR(vfxErr = NvVFX_GetU32(_eff, NVVFX_STATE_SIZE, &stateSizeInBytes)); cudaMalloc(&state, stateSizeInBytes); @@ -525,6 +529,8 @@ FXApp::Err FXApp::processImage(const char *inFile, const char *outFile) { BAIL_IF_ERR(vfxErr = NvCVImage_Transfer(&_dstGpuBuf, &_dstVFX, 255.f, stream, &_tmpVFX)); if (outFile && outFile[0]) { + if(IsLossyImageFile(outFile)) + fprintf(stderr, "WARNING: JPEG output file format will reduce image quality\n"); if (!cv::imwrite(outFile, _dstImg)) { printf("Error writing: \"%s\"\n", outFile); return errWrite; @@ -552,7 +558,7 @@ FXApp::Err FXApp::processMovie(const char *inFile, const char *outFile) { void* state = nullptr; void* stateArray[1]; - + if (inFile && !inFile[0]) inFile = nullptr; // Set file paths to NULL if zero length if (!FLAG_webcam && inFile) { @@ -574,7 +580,7 @@ FXApp::Err FXApp::processMovie(const char *inFile, const char *outFile) { printf("Filters only target H264 videos, not %.4s\n", (char*)&info.codec); BAIL_IF_ERR(vfxErr = allocBuffers(info.width, info.height)); - + if (outFile && !outFile[0]) outFile = nullptr; if (outFile) { ok = writer.open(outFile, StringToFourcc(FLAG_codec), info.frameRate, cv::Size(_dstVFX.width, _dstVFX.height)); @@ -586,7 +592,7 @@ FXApp::Err FXApp::processMovie(const char *inFile, const char *outFile) { } } - BAIL_IF_ERR(vfxErr = NvVFX_SetImage(_eff, NVVFX_INPUT_IMAGE, &_srcGpuBuf)); + BAIL_IF_ERR(vfxErr = NvVFX_SetImage(_eff, NVVFX_INPUT_IMAGE, &_srcGpuBuf)); BAIL_IF_ERR(vfxErr = NvVFX_SetImage(_eff, NVVFX_OUTPUT_IMAGE, &_dstGpuBuf)); BAIL_IF_ERR(vfxErr = NvVFX_SetF32(_eff, NVVFX_STRENGTH, FLAG_strength)); @@ -598,7 +604,7 @@ FXApp::Err FXApp::processMovie(const char *inFile, const char *outFile) { BAIL_IF_ERR(vfxErr = NvVFX_SetObject(_eff, NVVFX_STATE, (void*)stateArray)); BAIL_IF_ERR(vfxErr = NvVFX_Load(_eff)); - + for (frameNum = 0; reader.read(_srcImg); frameNum++) { if (_enableEffect) { BAIL_IF_ERR(vfxErr = NvCVImage_Transfer(&_srcVFX, &_srcGpuBuf, 1.f / 255.f, stream, &_tmpVFX)); @@ -613,7 +619,7 @@ FXApp::Err FXApp::processMovie(const char *inFile, const char *outFile) { writer.write(_dstImg); if (_show) { - if (_drawVisualization) drawEffectStatus(_dstImg); + if (_drawVisualization) drawEffectStatus(_dstImg); drawFrameRate(_dstImg); cv::imshow("Output", _dstImg); int key= cv::waitKey(1); @@ -630,7 +636,7 @@ FXApp::Err FXApp::processMovie(const char *inFile, const char *outFile) { if (_progress) fprintf(stderr, "\n"); reader.release(); if (outFile) - writer.release(); + writer.release(); bail: if (state) cudaFree(state); // release state memory return appErrFromVfxStatus(vfxErr); diff --git a/samples/DenoiseEffectApp/DenoiseEffectApp.exe b/samples/DenoiseEffectApp/DenoiseEffectApp.exe index 4f78833..bd49a7d 100644 Binary files a/samples/DenoiseEffectApp/DenoiseEffectApp.exe and b/samples/DenoiseEffectApp/DenoiseEffectApp.exe differ diff --git a/samples/UpscalePipelineApp/CMakeLists.txt b/samples/UpscalePipelineApp/CMakeLists.txt index 8ee5f68..7f441e8 100644 --- a/samples/UpscalePipelineApp/CMakeLists.txt +++ b/samples/UpscalePipelineApp/CMakeLists.txt @@ -12,29 +12,19 @@ target_include_directories(UpscalePipelineApp PUBLIC ${SDK_INCLUDES_PATH} ) -if(MSVC) - target_link_libraries(UpscalePipelineApp PUBLIC - opencv346 - NVVideoEffects - ${CMAKE_CURRENT_SOURCE_DIR}/../external/cuda/lib/x64/cudart.lib - ) +target_link_libraries(UpscalePipelineApp PUBLIC + opencv346 + NVVideoEffects + ${CMAKE_CURRENT_SOURCE_DIR}/../external/cuda/lib/x64/cudart.lib + ) - set(OPENCV_PATH_STR ${CMAKE_CURRENT_SOURCE_DIR}/../external/opencv/bin) - set(VFXSDK_PATH_STR ${CMAKE_CURRENT_SOURCE_DIR}/../../bin) # Also the location for CUDA/NVTRT/libcrypto - set(PATH_STR "PATH=%PATH%" ${VFXSDK_PATH_STR} ${OPENCV_PATH_STR}) - set(CMD_ARG_STR "--model_dir=\"${CMAKE_CURRENT_SOURCE_DIR}/../../bin/models\" --show --in_file=\"${CMAKE_CURRENT_SOURCE_DIR}/../input/input1.jpg\"") - set_target_properties(UpscalePipelineApp PROPERTIES - FOLDER SampleApps - VS_DEBUGGER_ENVIRONMENT "${PATH_STR}" - VS_DEBUGGER_COMMAND_ARGUMENTS "${CMD_ARG_STR}" - ) -else() +set(OPENCV_PATH_STR ${CMAKE_CURRENT_SOURCE_DIR}/../external/opencv/bin) +set(VFXSDK_PATH_STR ${CMAKE_CURRENT_SOURCE_DIR}/../../bin) # Also the location for CUDA/NVTRT/libcrypto +set(PATH_STR "PATH=%PATH%" ${VFXSDK_PATH_STR} ${OPENCV_PATH_STR}) +set(CMD_ARG_STR "--show --resolution=1440 --in_file=\"${CMAKE_CURRENT_SOURCE_DIR}/../input/input1.jpg\"") +set_target_properties(UpscalePipelineApp PROPERTIES + FOLDER SampleApps + VS_DEBUGGER_ENVIRONMENT "${PATH_STR}" + VS_DEBUGGER_COMMAND_ARGUMENTS "${CMD_ARG_STR}" + ) - target_link_libraries(UpscalePipelineApp PUBLIC - NVVideoEffects - NVCVImage - OpenCV - TensorRT - CUDA - ) -endif() diff --git a/samples/UpscalePipelineApp/UpscalePipeline.cpp b/samples/UpscalePipelineApp/UpscalePipeline.cpp index 2ef943d..3776ddd 100644 --- a/samples/UpscalePipelineApp/UpscalePipeline.cpp +++ b/samples/UpscalePipelineApp/UpscalePipeline.cpp @@ -66,7 +66,7 @@ bool FLAG_debug = false, FLAG_show = false, FLAG_progress = false; int FLAG_resolution = 0, - FLAG_arStrength = 0; + FLAG_arMode = 0; float FLAG_upscaleStrength = 0.2f; std::string FLAG_codec = DEFAULT_CODEC, FLAG_inFile, @@ -74,6 +74,11 @@ std::string FLAG_codec = DEFAULT_CODEC, FLAG_outDir, FLAG_modelDir; +// Set this when using OTA Updates +// This path is used by nvVideoEffectsProxy.cpp to load the SDK dll +// when using OTA Updates +char *g_nvVFXSDKPath = NULL; + static bool GetFlagArgVal(const char *flag, const char *arg, const char **val) { if (*arg != '-') return false; @@ -146,7 +151,7 @@ static void Usage() { " --in_file= input file to be processed\n" " --out_file= output file to be written\n" " --show display the results in a window\n" - " --ar_strength=(0|1) strength of artifact reduction filter (0: conservative, 1: aggressive, default 0)\n" + " --ar_mode=(0|1) mode of artifact reduction filter (0: conservative, 1: aggressive, default 0)\n" " --upscale_strength=(0 to 1) strength of upscale filter (float value between 0 to 1)\n" " --resolution= the desired height of the output\n" " --out_height= the desired height of the output\n" @@ -172,7 +177,7 @@ static int ParseMyArgs(int argc, char **argv) { GetFlagArgVal("out", arg, &FLAG_outFile) || GetFlagArgVal("out_file", arg, &FLAG_outFile) || GetFlagArgVal("show", arg, &FLAG_show) || - GetFlagArgVal("ar_strength", arg, &FLAG_arStrength) || + GetFlagArgVal("ar_mode", arg, &FLAG_arMode) || GetFlagArgVal("upscale_strength", arg, &FLAG_upscaleStrength) || GetFlagArgVal("resolution", arg, &FLAG_resolution) || GetFlagArgVal("model_dir", arg, &FLAG_modelDir) || @@ -226,6 +231,10 @@ static bool IsImageFile(const char *str) { return HasOneOfTheseSuffixes(str, ".bmp", ".jpg", ".jpeg", ".png", nullptr); } +static bool IsLossyImageFile(const char *str) { + return HasOneOfTheseSuffixes(str, ".jpg", ".jpeg", nullptr); +} + static const char* DurationString(double sc) { static char buf[16]; int hr, mn; @@ -488,7 +497,7 @@ FXApp::Err FXApp::processImage(const char *inFile, const char *outFile) { BAIL_IF_ERR(vfxErr = NvVFX_SetImage(_arEff, NVVFX_INPUT_IMAGE, &_srcGpuBuf)); BAIL_IF_ERR(vfxErr = NvVFX_SetImage(_arEff, NVVFX_OUTPUT_IMAGE, &_interGpuBGRf32pl)); BAIL_IF_ERR(vfxErr = NvVFX_SetCudaStream(_arEff, NVVFX_CUDA_STREAM, stream)); - BAIL_IF_ERR(vfxErr = NvVFX_SetU32(_arEff, NVVFX_STRENGTH, FLAG_arStrength)); + BAIL_IF_ERR(vfxErr = NvVFX_SetU32(_arEff, NVVFX_MODE, FLAG_arMode)); BAIL_IF_ERR(vfxErr = NvVFX_SetImage(_upscaleEff, NVVFX_INPUT_IMAGE, &_interGpuRGBAu8)); BAIL_IF_ERR(vfxErr = NvVFX_SetImage(_upscaleEff, NVVFX_OUTPUT_IMAGE, &_dstGpuBuf)); @@ -503,6 +512,8 @@ FXApp::Err FXApp::processImage(const char *inFile, const char *outFile) { BAIL_IF_ERR(vfxErr = NvCVImage_Transfer(&_dstGpuBuf, &_dstVFX, 1.f, stream, &_tmpVFX)); // _dstGpuBuf --> _dstTmpVFX --> _dstVFX if (outFile && outFile[0]) { + if(IsLossyImageFile(outFile)) + fprintf(stderr, "WARNING: JPEG output file format will reduce image quality\n"); if (!cv::imwrite(outFile, _dstImg)) { printf("Error writing: \"%s\"\n", outFile); return errWrite; @@ -552,7 +563,7 @@ FXApp::Err FXApp::processMovie(const char *inFile, const char *outFile) { BAIL_IF_ERR(vfxErr = NvVFX_SetImage(_arEff, NVVFX_INPUT_IMAGE, &_srcGpuBuf)); BAIL_IF_ERR(vfxErr = NvVFX_SetImage(_arEff, NVVFX_OUTPUT_IMAGE, &_interGpuBGRf32pl)); BAIL_IF_ERR(vfxErr = NvVFX_SetCudaStream(_arEff, NVVFX_CUDA_STREAM, stream)); - BAIL_IF_ERR(vfxErr = NvVFX_SetU32(_arEff, NVVFX_STRENGTH, FLAG_arStrength)); + BAIL_IF_ERR(vfxErr = NvVFX_SetU32(_arEff, NVVFX_MODE, FLAG_arMode)); BAIL_IF_ERR(vfxErr = NvVFX_Load(_arEff)); BAIL_IF_ERR(vfxErr = NvVFX_SetImage(_upscaleEff, NVVFX_INPUT_IMAGE, &_interGpuRGBAu8)); diff --git a/samples/UpscalePipelineApp/UpscalePipelineApp.exe b/samples/UpscalePipelineApp/UpscalePipelineApp.exe index 6121f65..ca2c5eb 100644 Binary files a/samples/UpscalePipelineApp/UpscalePipelineApp.exe and b/samples/UpscalePipelineApp/UpscalePipelineApp.exe differ diff --git a/samples/UpscalePipelineApp/run.bat b/samples/UpscalePipelineApp/run.bat index 2aa192f..bfd4f2d 100644 --- a/samples/UpscalePipelineApp/run.bat +++ b/samples/UpscalePipelineApp/run.bat @@ -1,6 +1,6 @@ SETLOCAL SET PATH=%PATH%;..\external\opencv\bin; REM Use --show to show the output in a window or use --out_file= to write output to file -UpscalePipelineApp.exe --in_file=..\input\input1.jpg --ar_strength=0 --upscale_strength=0 --resolution=1080 --show --out_file=ar_sr_0.png -UpscalePipelineApp.exe --in_file=..\input\input1.jpg --ar_strength=0 --upscale_strength=1 --resolution=1080 --show --out_file=ar_sr_1.png +UpscalePipelineApp.exe --in_file=..\input\input1.jpg --ar_mode=0 --upscale_strength=0 --resolution=1080 --show --out_file=ar_sr_0.png +UpscalePipelineApp.exe --in_file=..\input\input1.jpg --ar_mode=0 --upscale_strength=1 --resolution=1080 --show --out_file=ar_sr_1.png diff --git a/samples/VideoEffectsApp/CMakeLists.txt b/samples/VideoEffectsApp/CMakeLists.txt index c81d875..a9cd65a 100644 --- a/samples/VideoEffectsApp/CMakeLists.txt +++ b/samples/VideoEffectsApp/CMakeLists.txt @@ -7,29 +7,19 @@ add_executable(VideoEffectsApp ${SOURCE_FILES}) target_include_directories(VideoEffectsApp PRIVATE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/../utils) target_include_directories(VideoEffectsApp PUBLIC ${SDK_INCLUDES_PATH}) -if(MSVC) - target_link_libraries(VideoEffectsApp PUBLIC - opencv346 - NVVideoEffects - ${CMAKE_CURRENT_SOURCE_DIR}/../external/cuda/lib/x64/cudart.lib - ) +target_link_libraries(VideoEffectsApp PUBLIC + opencv346 + NVVideoEffects + ${CMAKE_CURRENT_SOURCE_DIR}/../external/cuda/lib/x64/cudart.lib + ) - set(OPENCV_PATH_STR ${CMAKE_CURRENT_SOURCE_DIR}/../external/opencv/bin) - set(VFXSDK_PATH_STR ${CMAKE_CURRENT_SOURCE_DIR}/../../bin) # Also the location for CUDA/NVTRT/libcrypto - set(PATH_STR "PATH=%PATH%" ${VFXSDK_PATH_STR} ${OPENCV_PATH_STR}) - set(CMD_ARG_STR "--model_dir=\"${CMAKE_CURRENT_SOURCE_DIR}/../../bin/models\" --show --effect=SuperRes --resolution=1080 --in_file=\"${CMAKE_CURRENT_SOURCE_DIR}/../input/input2.png\"") - set_target_properties(VideoEffectsApp PROPERTIES - FOLDER SampleApps - VS_DEBUGGER_ENVIRONMENT "${PATH_STR}" - VS_DEBUGGER_COMMAND_ARGUMENTS "${CMD_ARG_STR}" - ) -else() +set(OPENCV_PATH_STR ${CMAKE_CURRENT_SOURCE_DIR}/../external/opencv/bin) +set(VFXSDK_PATH_STR ${CMAKE_CURRENT_SOURCE_DIR}/../../bin) # Also the location for CUDA/NVTRT/libcrypto +set(PATH_STR "PATH=%PATH%" ${VFXSDK_PATH_STR} ${OPENCV_PATH_STR}) +set(CMD_ARG_STR "--show --effect=SuperRes --resolution=1080 --in_file=\"${CMAKE_CURRENT_SOURCE_DIR}/../input/input1.jpg\"") +set_target_properties(VideoEffectsApp PROPERTIES + FOLDER SampleApps + VS_DEBUGGER_ENVIRONMENT "${PATH_STR}" + VS_DEBUGGER_COMMAND_ARGUMENTS "${CMD_ARG_STR}" + ) - target_link_libraries(VideoEffectsApp PUBLIC - NVVideoEffects - NVCVImage - OpenCV - TensorRT - CUDA - ) -endif() diff --git a/samples/VideoEffectsApp/VideoEffectsApp.cpp b/samples/VideoEffectsApp/VideoEffectsApp.cpp index 00310dd..3576b0e 100644 --- a/samples/VideoEffectsApp/VideoEffectsApp.cpp +++ b/samples/VideoEffectsApp/VideoEffectsApp.cpp @@ -57,6 +57,7 @@ bool FLAG_debug = false, FLAG_progress = false, FLAG_webcam = false; float FLAG_strength = 0.f; +int FLAG_mode = 0; int FLAG_resolution = 0; std::string FLAG_codec = DEFAULT_CODEC, FLAG_camRes = "1280x720", @@ -66,6 +67,11 @@ std::string FLAG_codec = DEFAULT_CODEC, FLAG_modelDir, FLAG_effect; +// Set this when using OTA Updates +// This path is used by nvVideoEffectsProxy.cpp to load the SDK dll +// when using OTA Updates +char *g_nvVFXSDKPath = NULL; + static bool GetFlagArgVal(const char *flag, const char *arg, const char **val) { if (*arg != '-') return false; @@ -140,8 +146,9 @@ static void Usage() { " --out_file= output file to be written\n" " --effect= the effect to apply\n" " --show display the results in a window (for webcam, it is always true)\n" - " --strength= strength of an effect, 0 or 1 for super res and artifact reduction,\n" - " and [0.0, 1.0] for upscaling\n" + " --strength= strength of the upscaling effect, [0.0, 1.0]\n" + " --mode= mode of the super res or artifact reduction effect, 0 or 1, \n" + " where 0 - conservative and 1 - aggressive\n" " --cam_res=[WWWx]HHH specify camera resolution as height or width x height\n" " supports 720 and 1080 resolutions (default \"720\") \n" " --resolution= the desired height of the output\n" @@ -176,6 +183,7 @@ static int ParseMyArgs(int argc, char **argv) { GetFlagArgVal("webcam", arg, &FLAG_webcam) || GetFlagArgVal("cam_res", arg, &FLAG_camRes) || GetFlagArgVal("strength", arg, &FLAG_strength) || + GetFlagArgVal("mode", arg, &FLAG_mode) || GetFlagArgVal("resolution", arg, &FLAG_resolution) || GetFlagArgVal("model_dir", arg, &FLAG_modelDir) || GetFlagArgVal("codec", arg, &FLAG_codec) || @@ -228,6 +236,11 @@ static bool IsImageFile(const char *str) { return HasOneOfTheseSuffixes(str, ".bmp", ".jpg", ".jpeg", ".png", nullptr); } +static bool IsLossyImageFile(const char *str) { + return HasOneOfTheseSuffixes(str, ".jpg", ".jpeg", nullptr); +} + + static const char* DurationString(double sc) { static char buf[16]; int hr, mn; @@ -561,9 +574,9 @@ FXApp::Err FXApp::processImage(const char *inFile, const char *outFile) { BAIL_IF_ERR(vfxErr = NvVFX_SetImage(_eff, NVVFX_OUTPUT_IMAGE, &_dstGpuBuf)); BAIL_IF_ERR(vfxErr = NvVFX_SetCudaStream(_eff, NVVFX_CUDA_STREAM, stream)); if (!strcmp(_effectName, NVVFX_FX_ARTIFACT_REDUCTION)) { - BAIL_IF_ERR(vfxErr = NvVFX_SetU32(_eff, NVVFX_STRENGTH, (unsigned int)FLAG_strength)); + BAIL_IF_ERR(vfxErr = NvVFX_SetU32(_eff, NVVFX_MODE, (unsigned int)FLAG_mode)); } else if (!strcmp(_effectName, NVVFX_FX_SUPER_RES)) { - BAIL_IF_ERR(vfxErr = NvVFX_SetU32(_eff, NVVFX_STRENGTH, (unsigned int)FLAG_strength)); + BAIL_IF_ERR(vfxErr = NvVFX_SetU32(_eff, NVVFX_MODE, (unsigned int)FLAG_mode)); } BAIL_IF_ERR(vfxErr = NvVFX_Load(_eff)); @@ -571,6 +584,8 @@ FXApp::Err FXApp::processImage(const char *inFile, const char *outFile) { BAIL_IF_ERR(vfxErr = NvCVImage_Transfer(&_dstGpuBuf, &_dstVFX, 255.f, stream, &_tmpVFX)); // _dstGpuBuf --> _tmpVFX --> _dstVFX if (outFile && outFile[0]) { + if(IsLossyImageFile(outFile)) + fprintf(stderr, "WARNING: JPEG output file format will reduce image quality\n"); if (!cv::imwrite(outFile, _dstImg)) { printf("Error writing: \"%s\"\n", outFile); return errWrite; @@ -632,9 +647,9 @@ FXApp::Err FXApp::processMovie(const char *inFile, const char *outFile) { BAIL_IF_ERR(vfxErr = NvVFX_SetImage(_eff, NVVFX_OUTPUT_IMAGE, &_dstGpuBuf)); BAIL_IF_ERR(vfxErr = NvVFX_SetCudaStream(_eff, NVVFX_CUDA_STREAM, stream)); if (!strcmp(_effectName, NVVFX_FX_ARTIFACT_REDUCTION)) { - BAIL_IF_ERR(vfxErr = NvVFX_SetU32(_eff, NVVFX_STRENGTH, (unsigned int)FLAG_strength)); + BAIL_IF_ERR(vfxErr = NvVFX_SetU32(_eff, NVVFX_MODE, (unsigned int)FLAG_mode)); } else if (!strcmp(_effectName, NVVFX_FX_SUPER_RES)) { - BAIL_IF_ERR(vfxErr = NvVFX_SetU32(_eff, NVVFX_STRENGTH, (unsigned int)FLAG_strength)); + BAIL_IF_ERR(vfxErr = NvVFX_SetU32(_eff, NVVFX_MODE, (unsigned int)FLAG_mode)); } BAIL_IF_ERR(vfxErr = NvVFX_Load(_eff)); diff --git a/samples/VideoEffectsApp/VideoEffectsApp.exe b/samples/VideoEffectsApp/VideoEffectsApp.exe index 3464870..e91c458 100644 Binary files a/samples/VideoEffectsApp/VideoEffectsApp.exe and b/samples/VideoEffectsApp/VideoEffectsApp.exe differ diff --git a/samples/VideoEffectsApp/run.bat b/samples/VideoEffectsApp/run.bat index 9626a6b..7409e0e 100644 --- a/samples/VideoEffectsApp/run.bat +++ b/samples/VideoEffectsApp/run.bat @@ -1,7 +1,7 @@ SETLOCAL SET PATH=%PATH%;..\external\opencv\bin; REM Use --show to show the output in a window or use --out_file= to write output to file -VideoEffectsApp.exe --in_file=..\input\input1.jpg --out_file=ar_1.png --effect=ArtifactReduction --strength=1 --show -VideoEffectsApp.exe --in_file=..\input\input1.jpg --out_file=ar_0.png --effect=ArtifactReduction --strength=0 --show -VideoEffectsApp.exe --in_file=..\input\input2.jpg --out_file=sr_0.png --effect=SuperRes --resolution=2160 --strength=0 --show -VideoEffectsApp.exe --in_file=..\input\input2.jpg --out_file=sr_1.png --effect=SuperRes --resolution=2160 --strength=1 --show \ No newline at end of file +VideoEffectsApp.exe --in_file=..\input\input1.jpg --out_file=ar_1.png --effect=ArtifactReduction --mode=1 --show +VideoEffectsApp.exe --in_file=..\input\input1.jpg --out_file=ar_0.png --effect=ArtifactReduction --mode=0 --show +VideoEffectsApp.exe --in_file=..\input\input2.jpg --out_file=sr_0.png --effect=SuperRes --resolution=2160 --mode=0 --show +VideoEffectsApp.exe --in_file=..\input\input2.jpg --out_file=sr_1.png --effect=SuperRes --resolution=2160 --mode=1 --show \ No newline at end of file diff --git a/samples/external/CMakeLists.txt b/samples/external/CMakeLists.txt index 47d4497..0be0af0 100644 --- a/samples/external/CMakeLists.txt +++ b/samples/external/CMakeLists.txt @@ -40,7 +40,7 @@ else() message("OpenCV_LIBRARIES ${OpenCV_LIBRARIES}") message("OpenCV_LIBS ${OpenCV_LIBS}") - find_package(CUDA 11.1 REQUIRED) + find_package(CUDA 11.3 REQUIRED) add_library(CUDA INTERFACE) target_include_directories(CUDA INTERFACE ${CUDA_INCLUDE_DIRS}) target_link_libraries(CUDA INTERFACE "${CUDA_LIBRARIES};cuda") @@ -48,7 +48,7 @@ else() message("CUDA_INCLUDE_DIRS ${CUDA_INCLUDE_DIRS}") message("CUDA_LIBRARIES ${CUDA_LIBRARIES}") - find_package(TensorRT 7 REQUIRED) + find_package(TensorRT 8 REQUIRED) add_library(TensorRT INTERFACE) target_include_directories(TensorRT INTERFACE ${TensorRT_INCLUDE_DIRS}) target_link_libraries(TensorRT INTERFACE ${TensorRT_LIBRARIES}) diff --git a/samples/external/cuda/include/channel_descriptor.h b/samples/external/cuda/include/channel_descriptor.h index 77930f0..a375508 100644 --- a/samples/external/cuda/include/channel_descriptor.h +++ b/samples/external/cuda/include/channel_descriptor.h @@ -86,7 +86,7 @@ * \endcode * * where ::cudaChannelFormatKind is one of ::cudaChannelFormatKindSigned, - * ::cudaChannelFormatKindUnsigned, or ::cudaChannelFormatKindFloat. + * ::cudaChannelFormatKindUnsigned, cudaChannelFormatKindFloat or ::cudaChannelFormatKindNV12. * * \return * Channel descriptor with format \p f @@ -401,6 +401,12 @@ template<> __inline__ __host__ cudaChannelFormatDesc cudaCreateChannelDesc= 11) || (__clang_major__ < 3) || ((__clang_major__ == 3) && (__clang_minor__ < 3)) -#error -- unsupported clang version! clang version must be less than 11 and greater than 3.2 . The nvcc flag '-allow-unsupported-compiler' can be used to override this version check; however, using an unsupported host compiler may cause compilation failure or incorrect run time execution. Use at your own risk. +#if (__clang_major__ >= 12) || (__clang_major__ < 3) || ((__clang_major__ == 3) && (__clang_minor__ < 3)) +#error -- unsupported clang version! clang version must be less than 12 and greater than 3.2 . The nvcc flag '-allow-unsupported-compiler' can be used to override this version check; however, using an unsupported host compiler may cause compilation failure or incorrect run time execution. Use at your own risk. -#endif /* (__clang_major__ >= 11) || (__clang_major__ < 3) || ((__clang_major__ == 3) && (__clang_minor__ < 3)) */ +#endif /* (__clang_major__ >= 12) || (__clang_major__ < 3) || ((__clang_major__ == 3) && (__clang_minor__ < 3)) */ #endif /* defined(__clang__) && !defined(__ibmxl_vrm__) && !defined(__ICC) && !defined(__HORIZON__) && !defined(__APPLE__) */ @@ -155,15 +155,15 @@ #if defined(_WIN32) -#if _MSC_VER < 1700 || _MSC_VER >= 1930 +#if _MSC_VER < 1910 || _MSC_VER >= 1930 -#error -- unsupported Microsoft Visual Studio version! Only the versions between 2015 and 2019 (inclusive) are supported! The nvcc flag '-allow-unsupported-compiler' can be used to override this version check; however, using an unsupported host compiler may cause compilation failure or incorrect run time execution. Use at your own risk. +#error -- unsupported Microsoft Visual Studio version! Only the versions between 2017 and 2019 (inclusive) are supported! The nvcc flag '-allow-unsupported-compiler' can be used to override this version check; however, using an unsupported host compiler may cause compilation failure or incorrect run time execution. Use at your own risk. -#elif _MSC_VER >= 1700 && _MSC_VER < 1900 +#elif _MSC_VER >= 1910 && _MSC_VER < 1910 -#pragma message("support for this version of Microsoft Visual Studio has been deprecated! Only the versions between 2015 and 2019 (inclusive) are supported!") +#pragma message("support for this version of Microsoft Visual Studio has been deprecated! Only the versions between 2017 and 2019 (inclusive) are supported!") -#endif /* (_MSC_VER < 1700 || _MSC_VER >= 1930) || (_MSC_VER >= 1700 && _MSC_VER < 1900) */ +#endif /* (_MSC_VER < 1910 || _MSC_VER >= 1930) || (_MSC_VER >= 1910 && _MSC_VER < 1910) */ #endif /* _WIN32 */ #endif /* !__NV_NO_HOST_COMPILER_CHECK */ diff --git a/samples/external/cuda/include/crt/host_defines.h b/samples/external/cuda/include/crt/host_defines.h index b3c685c..ac38a21 100644 --- a/samples/external/cuda/include/crt/host_defines.h +++ b/samples/external/cuda/include/crt/host_defines.h @@ -97,6 +97,7 @@ #define __location__(a) \ __annotate__(a) #define CUDARTAPI +#define CUDARTAPI_CDECL #elif defined(_MSC_VER) @@ -133,6 +134,8 @@ __annotate__(__##a##__) #define CUDARTAPI \ __stdcall +#define CUDARTAPI_CDECL \ + __cdecl #else /* __GNUC__ || __CUDA_LIBDEVICE__ || __CUDACC_RTC__ */ diff --git a/samples/external/cuda/include/cuda_device_runtime_api.h b/samples/external/cuda/include/cuda_device_runtime_api.h index 6d131ce..83fbe8e 100644 --- a/samples/external/cuda/include/cuda_device_runtime_api.h +++ b/samples/external/cuda/include/cuda_device_runtime_api.h @@ -1,5 +1,5 @@ /* - * Copyright 1993-2018 NVIDIA Corporation. All rights reserved. + * Copyright 1993-2021 NVIDIA Corporation. All rights reserved. * * NOTICE TO LICENSEE: * @@ -66,43 +66,37 @@ extern "C" { struct cudaFuncAttributes; -#if defined(_WIN32) -#define __NV_WEAK__ __declspec(nv_weak) -#else -#define __NV_WEAK__ __attribute__((nv_weak)) -#endif -__device__ __NV_WEAK__ cudaError_t CUDARTAPI cudaMalloc(void **p, size_t s) +inline __device__ cudaError_t CUDARTAPI cudaMalloc(void **p, size_t s) { return cudaErrorUnknown; } -__device__ __NV_WEAK__ cudaError_t CUDARTAPI cudaFuncGetAttributes(struct cudaFuncAttributes *p, const void *c) +inline __device__ cudaError_t CUDARTAPI cudaFuncGetAttributes(struct cudaFuncAttributes *p, const void *c) { return cudaErrorUnknown; } -__device__ __NV_WEAK__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDeviceAttr attr, int device) +inline __device__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDeviceAttr attr, int device) { return cudaErrorUnknown; } -__device__ __NV_WEAK__ cudaError_t CUDARTAPI cudaGetDevice(int *device) +inline __device__ cudaError_t CUDARTAPI cudaGetDevice(int *device) { return cudaErrorUnknown; } -__device__ __NV_WEAK__ cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultiprocessor(int *numBlocks, const void *func, int blockSize, size_t dynamicSmemSize) +inline __device__ cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultiprocessor(int *numBlocks, const void *func, int blockSize, size_t dynamicSmemSize) { return cudaErrorUnknown; } -__device__ __NV_WEAK__ cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int *numBlocks, const void *func, int blockSize, size_t dynamicSmemSize, unsigned int flags) +inline __device__ cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int *numBlocks, const void *func, int blockSize, size_t dynamicSmemSize, unsigned int flags) { return cudaErrorUnknown; } -#undef __NV_WEAK__ #if defined(__cplusplus) } diff --git a/samples/external/cuda/include/cuda_runtime.h b/samples/external/cuda/include/cuda_runtime.h index 825b827..3d94c50 100644 --- a/samples/external/cuda/include/cuda_runtime.h +++ b/samples/external/cuda/include/cuda_runtime.h @@ -188,7 +188,8 @@ struct __device_builtin__ __nv_lambda_preheader_injection { }; * ::cudaErrorInvalidPtx, * ::cudaErrorUnsupportedPtxVersion, * ::cudaErrorNoKernelImageForDevice, - * ::cudaErrorJitCompilerNotFound + * ::cudaErrorJitCompilerNotFound, + * ::cudaErrorJitCompilationDisabled * \notefnerr * \note_async * \note_null_stream @@ -628,6 +629,57 @@ static __inline__ __host__ cudaError_t cudaMallocPitch( return ::cudaMallocPitch((void**)(void*)devPtr, pitch, width, height); } +/** + * \brief Allocate from a pool + * + * This is an alternate spelling for cudaMallocFromPoolAsync + * made available through operator overloading. + * + * \sa ::cudaMallocFromPoolAsync, + * \ref ::cudaMallocAsync(void** ptr, size_t size, cudaStream_t hStream) "cudaMallocAsync (C API)" + */ +static __inline__ __host__ cudaError_t cudaMallocAsync( + void **ptr, + size_t size, + cudaMemPool_t memPool, + cudaStream_t stream +) +{ + return ::cudaMallocFromPoolAsync(ptr, size, memPool, stream); +} + +template +static __inline__ __host__ cudaError_t cudaMallocAsync( + T **ptr, + size_t size, + cudaMemPool_t memPool, + cudaStream_t stream +) +{ + return ::cudaMallocFromPoolAsync((void**)(void*)ptr, size, memPool, stream); +} + +template +static __inline__ __host__ cudaError_t cudaMallocAsync( + T **ptr, + size_t size, + cudaStream_t stream +) +{ + return ::cudaMallocAsync((void**)(void*)ptr, size, stream); +} + +template +static __inline__ __host__ cudaError_t cudaMallocFromPoolAsync( + T **ptr, + size_t size, + cudaMemPool_t memPool, + cudaStream_t stream +) +{ + return ::cudaMallocFromPoolAsync((void**)(void*)ptr, size, memPool, stream); +} + #if defined(__CUDACC__) /** @@ -1190,6 +1242,59 @@ static __inline__ __host__ cudaError_t cudaGraphExecMemcpyNodeSetParamsFromSymbo return ::cudaGraphExecMemcpyNodeSetParamsFromSymbol(hGraphExec, node, dst, (const void*)&symbol, count, offset, kind); } +#if __cplusplus >= 201103 + +/** + * \brief Creates a user object by wrapping a C++ object + * + * TODO detail + * + * \param object_out - Location to return the user object handle + * \param objectToWrap - This becomes the \ptr argument to ::cudaUserObjectCreate. A + * lambda will be passed for the \p destroy argument, which calls + * delete on this object pointer. + * \param initialRefcount - The initial refcount to create the object with, typically 1. The + * initial references are owned by the calling thread. + * \param flags - Currently it is required to pass cudaUserObjectNoDestructorSync, + * which is the only defined flag. This indicates that the destroy + * callback cannot be waited on by any CUDA API. Users requiring + * synchronization of the callback should signal its completion + * manually. + * + * \return + * ::cudaSuccess, + * ::cudaErrorInvalidValue + * + * \sa + * ::cudaUserObjectCreate + */ +template +static __inline__ __host__ cudaError_t cudaUserObjectCreate( + cudaUserObject_t *object_out, + T *objectToWrap, + unsigned int initialRefcount, + unsigned int flags) +{ + return ::cudaUserObjectCreate( + object_out, + objectToWrap, + [](void *vpObj) { delete reinterpret_cast(vpObj); }, + initialRefcount, + flags); +} + +template +static __inline__ __host__ cudaError_t cudaUserObjectCreate( + cudaUserObject_t *object_out, + T *objectToWrap, + unsigned int initialRefcount, + cudaUserObjectFlags flags) +{ + return cudaUserObjectCreate(object_out, objectToWrap, initialRefcount, (unsigned int)flags); +} + +#endif + /** * \brief \hl Finds the address associated with a CUDA symbol * diff --git a/samples/external/cuda/include/cuda_runtime_api.h b/samples/external/cuda/include/cuda_runtime_api.h index 669fe6c..435e4a8 100644 --- a/samples/external/cuda/include/cuda_runtime_api.h +++ b/samples/external/cuda/include/cuda_runtime_api.h @@ -47,6 +47,8 @@ * Users Notice. */ + + #if !defined(__CUDA_RUNTIME_API_H__) #define __CUDA_RUNTIME_API_H__ @@ -133,7 +135,13 @@ */ /** CUDA Runtime API Version */ -#define CUDART_VERSION 11010 +#define CUDART_VERSION 11030 + +#if defined(__CUDA_API_VER_MAJOR__) && defined(__CUDA_API_VER_MINOR__) +# define __CUDART_API_VERSION ((__CUDA_API_VER_MAJOR__ * 1000) + (__CUDA_API_VER_MINOR__ * 10)) +#else +# define __CUDART_API_VERSION CUDART_VERSION +#endif #include "crt/host_defines.h" #include "builtin_types.h" @@ -149,6 +157,9 @@ #define __CUDART_API_PTSZ(api) api #endif +#define cudaSignalExternalSemaphoresAsync __CUDART_API_PTSZ(cudaSignalExternalSemaphoresAsync_v2) +#define cudaWaitExternalSemaphoresAsync __CUDART_API_PTSZ(cudaWaitExternalSemaphoresAsync_v2) + #if defined(__CUDART_API_PER_THREAD_DEFAULT_STREAM) #define cudaMemcpy __CUDART_API_PTDS(cudaMemcpy) #define cudaMemcpyToSymbol __CUDART_API_PTDS(cudaMemcpyToSymbol) @@ -169,8 +180,9 @@ #define cudaGraphLaunch __CUDART_API_PTSZ(cudaGraphLaunch) #define cudaStreamBeginCapture __CUDART_API_PTSZ(cudaStreamBeginCapture) #define cudaStreamEndCapture __CUDART_API_PTSZ(cudaStreamEndCapture) - #define cudaStreamIsCapturing __CUDART_API_PTSZ(cudaStreamIsCapturing) #define cudaStreamGetCaptureInfo __CUDART_API_PTSZ(cudaStreamGetCaptureInfo) + #define cudaStreamGetCaptureInfo_v2 __CUDART_API_PTSZ(cudaStreamGetCaptureInfo_v2) + #define cudaStreamIsCapturing __CUDART_API_PTSZ(cudaStreamIsCapturing) #define cudaMemcpyAsync __CUDART_API_PTSZ(cudaMemcpyAsync) #define cudaMemcpyToSymbolAsync __CUDART_API_PTSZ(cudaMemcpyToSymbolAsync) #define cudaMemcpyFromSymbolAsync __CUDART_API_PTSZ(cudaMemcpyFromSymbolAsync) @@ -197,11 +209,13 @@ #define cudaLaunchHostFunc __CUDART_API_PTSZ(cudaLaunchHostFunc) #define cudaMemPrefetchAsync __CUDART_API_PTSZ(cudaMemPrefetchAsync) #define cudaLaunchCooperativeKernel __CUDART_API_PTSZ(cudaLaunchCooperativeKernel) - #define cudaSignalExternalSemaphoresAsync __CUDART_API_PTSZ(cudaSignalExternalSemaphoresAsync) - #define cudaWaitExternalSemaphoresAsync __CUDART_API_PTSZ(cudaWaitExternalSemaphoresAsync) #define cudaStreamCopyAttributes __CUDART_API_PTSZ(cudaStreamCopyAttributes) #define cudaStreamGetAttribute __CUDART_API_PTSZ(cudaStreamGetAttribute) #define cudaStreamSetAttribute __CUDART_API_PTSZ(cudaStreamSetAttribute) + #define cudaMallocAsync __CUDART_API_PTSZ(cudaMallocAsync) + #define cudaFreeAsync __CUDART_API_PTSZ(cudaFreeAsync) + #define cudaMallocFromPoolAsync __CUDART_API_PTSZ(cudaMallocFromPoolAsync) + #define cudaGetDriverEntryPoint __CUDART_API_PTSZ(cudaGetDriverEntryPoint) #endif /** \cond impl_private */ @@ -370,8 +384,8 @@ extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaDeviceSynchronize(v * Values can range from 0B to 128B. This is purely a performance hint and * it can be ignored or clamped depending on the platform. * - * - cudaLimitPersistingL2CacheSize controls size of window in bytes available - * for ::cudaAccessPolicyWindow. This is purely a performance hint and it + * - ::cudaLimitPersistingL2CacheSize controls size in bytes available + * for persisting L2 cache. This is purely a performance hint and it * can be ignored or clamped depending on the platform. * * \param limit - Limit to set @@ -408,7 +422,7 @@ extern __host__ cudaError_t CUDARTAPI cudaDeviceSetLimit(enum cudaLimit limit, s * - ::cudaLimitDevRuntimePendingLaunchCount: maximum number of outstanding * device runtime launches. * - ::cudaLimitMaxL2FetchGranularity: L2 cache fetch granularity. - * - ::cudaLimitPersistingL2CacheSize: L2 cache persistings size in bytes + * - ::cudaLimitPersistingL2CacheSize: Persisting L2 cache size in bytes * * \param limit - Limit to query * \param pValue - Returned size of the limit @@ -447,7 +461,9 @@ extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaDeviceGetLimit(size * \sa * ::cuDeviceGetMaxTexture1DLinear, */ -extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaDeviceGetTexture1DLinearMaxWidth(size_t *maxWidthInElements, const struct cudaChannelFormatDesc *fmtDesc, int device); +#if __CUDART_API_VERSION >= 11010 + extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaDeviceGetTexture1DLinearMaxWidth(size_t *maxWidthInElements, const struct cudaChannelFormatDesc *fmtDesc, int device); +#endif /** * \brief Returns the preferred cache configuration for the current device. @@ -475,7 +491,7 @@ extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaDeviceGetTexture1DL * \note_init_rt * \note_callback * - * \sa cudaDeviceSetCacheConfig, + * \sa ::cudaDeviceSetCacheConfig, * \ref ::cudaFuncSetCacheConfig(const void*, enum cudaFuncCache) "cudaFuncSetCacheConfig (C API)", * \ref ::cudaFuncSetCacheConfig(T*, enum cudaFuncCache) "cudaFuncSetCacheConfig (C++ API)", * ::cuCtxGetCacheConfig @@ -927,8 +943,277 @@ extern __host__ cudaError_t CUDARTAPI cudaIpcOpenMemHandle(void **devPtr, cudaIp */ extern __host__ cudaError_t CUDARTAPI cudaIpcCloseMemHandle(void *devPtr); +/** + * \brief Blocks until remote writes are visible to the specified scope + * + * Blocks until remote writes to the target context via mappings created + * through GPUDirect RDMA APIs, like nvidia_p2p_get_pages (see + * https://docs.nvidia.com/cuda/gpudirect-rdma for more information), are + * visible to the specified scope. + * + * If the scope equals or lies within the scope indicated by + * ::cudaDevAttrGPUDirectRDMAWritesOrdering, the call will be a no-op and + * can be safely omitted for performance. This can be determined by + * comparing the numerical values between the two enums, with smaller + * scopes having smaller values. + * + * Users may query support for this API via ::cudaDevAttrGPUDirectRDMAFlushWritesOptions. + * + * \param target - The target of the operation, see cudaFlushGPUDirectRDMAWritesTarget + * \param scope - The scope of the operation, see cudaFlushGPUDirectRDMAWritesScope + * + * \return + * ::cudaSuccess, + * ::cudaErrorNotSupported, + * \notefnerr + * \note_init_rt + * \note_callback + * + * \sa + * ::cuFlushGPUDirectRDMAWrites + */ +#if __CUDART_API_VERSION >= 11030 +extern __host__ cudaError_t CUDARTAPI cudaDeviceFlushGPUDirectRDMAWrites(enum cudaFlushGPUDirectRDMAWritesTarget target, enum cudaFlushGPUDirectRDMAWritesScope scope); +#endif + /** @} */ /* END CUDART_DEVICE */ + +/** + * \defgroup CUDART_THREAD_DEPRECATED Thread Management [DEPRECATED] + * + * ___MANBRIEF___ deprecated thread management functions of the CUDA runtime + * API (___CURRENT_FILE___) ___ENDMANBRIEF___ + * + * This section describes deprecated thread management functions of the CUDA runtime + * application programming interface. + * + * @{ + */ + +/** + * \brief Exit and clean up from CUDA launches + * + * \deprecated + * + * Note that this function is deprecated because its name does not + * reflect its behavior. Its functionality is identical to the + * non-deprecated function ::cudaDeviceReset(), which should be used + * instead. + * + * Explicitly destroys all cleans up all resources associated with the current + * device in the current process. Any subsequent API call to this device will + * reinitialize the device. + * + * Note that this function will reset the device immediately. It is the caller's + * responsibility to ensure that the device is not being accessed by any + * other host threads from the process when this function is called. + * + * \return + * ::cudaSuccess + * \notefnerr + * \note_init_rt + * \note_callback + * + * \sa ::cudaDeviceReset + */ +extern __CUDA_DEPRECATED __host__ cudaError_t CUDARTAPI cudaThreadExit(void); + +/** + * \brief Wait for compute device to finish + * + * \deprecated + * + * Note that this function is deprecated because its name does not + * reflect its behavior. Its functionality is similar to the + * non-deprecated function ::cudaDeviceSynchronize(), which should be used + * instead. + * + * Blocks until the device has completed all preceding requested tasks. + * ::cudaThreadSynchronize() returns an error if one of the preceding tasks + * has failed. If the ::cudaDeviceScheduleBlockingSync flag was set for + * this device, the host thread will block until the device has finished + * its work. + * + * \return + * ::cudaSuccess + * \notefnerr + * \note_init_rt + * \note_callback + * + * \sa ::cudaDeviceSynchronize + */ +extern __CUDA_DEPRECATED __host__ cudaError_t CUDARTAPI cudaThreadSynchronize(void); + +/** + * \brief Set resource limits + * + * \deprecated + * + * Note that this function is deprecated because its name does not + * reflect its behavior. Its functionality is identical to the + * non-deprecated function ::cudaDeviceSetLimit(), which should be used + * instead. + * + * Setting \p limit to \p value is a request by the application to update + * the current limit maintained by the device. The driver is free to + * modify the requested value to meet h/w requirements (this could be + * clamping to minimum or maximum values, rounding up to nearest element + * size, etc). The application can use ::cudaThreadGetLimit() to find out + * exactly what the limit has been set to. + * + * Setting each ::cudaLimit has its own specific restrictions, so each is + * discussed here. + * + * - ::cudaLimitStackSize controls the stack size of each GPU thread. + * + * - ::cudaLimitPrintfFifoSize controls the size of the shared FIFO + * used by the ::printf() device system call. + * Setting ::cudaLimitPrintfFifoSize must be performed before + * launching any kernel that uses the ::printf() device + * system call, otherwise ::cudaErrorInvalidValue will be returned. + * + * - ::cudaLimitMallocHeapSize controls the size of the heap used + * by the ::malloc() and ::free() device system calls. Setting + * ::cudaLimitMallocHeapSize must be performed before launching + * any kernel that uses the ::malloc() or ::free() device system calls, + * otherwise ::cudaErrorInvalidValue will be returned. + * + * \param limit - Limit to set + * \param value - Size in bytes of limit + * + * \return + * ::cudaSuccess, + * ::cudaErrorUnsupportedLimit, + * ::cudaErrorInvalidValue + * \notefnerr + * \note_init_rt + * \note_callback + * + * \sa ::cudaDeviceSetLimit + */ +extern __CUDA_DEPRECATED __host__ cudaError_t CUDARTAPI cudaThreadSetLimit(enum cudaLimit limit, size_t value); + +/** + * \brief Returns resource limits + * + * \deprecated + * + * Note that this function is deprecated because its name does not + * reflect its behavior. Its functionality is identical to the + * non-deprecated function ::cudaDeviceGetLimit(), which should be used + * instead. + * + * Returns in \p *pValue the current size of \p limit. The supported + * ::cudaLimit values are: + * - ::cudaLimitStackSize: stack size of each GPU thread; + * - ::cudaLimitPrintfFifoSize: size of the shared FIFO used by the + * ::printf() device system call. + * - ::cudaLimitMallocHeapSize: size of the heap used by the + * ::malloc() and ::free() device system calls; + * + * \param limit - Limit to query + * \param pValue - Returned size in bytes of limit + * + * \return + * ::cudaSuccess, + * ::cudaErrorUnsupportedLimit, + * ::cudaErrorInvalidValue + * \notefnerr + * \note_init_rt + * \note_callback + * + * \sa ::cudaDeviceGetLimit + */ +extern __CUDA_DEPRECATED __host__ cudaError_t CUDARTAPI cudaThreadGetLimit(size_t *pValue, enum cudaLimit limit); + +/** + * \brief Returns the preferred cache configuration for the current device. + * + * \deprecated + * + * Note that this function is deprecated because its name does not + * reflect its behavior. Its functionality is identical to the + * non-deprecated function ::cudaDeviceGetCacheConfig(), which should be + * used instead. + * + * On devices where the L1 cache and shared memory use the same hardware + * resources, this returns through \p pCacheConfig the preferred cache + * configuration for the current device. This is only a preference. The + * runtime will use the requested configuration if possible, but it is free to + * choose a different configuration if required to execute functions. + * + * This will return a \p pCacheConfig of ::cudaFuncCachePreferNone on devices + * where the size of the L1 cache and shared memory are fixed. + * + * The supported cache configurations are: + * - ::cudaFuncCachePreferNone: no preference for shared memory or L1 (default) + * - ::cudaFuncCachePreferShared: prefer larger shared memory and smaller L1 cache + * - ::cudaFuncCachePreferL1: prefer larger L1 cache and smaller shared memory + * + * \param pCacheConfig - Returned cache configuration + * + * \return + * ::cudaSuccess + * \notefnerr + * \note_init_rt + * \note_callback + * + * \sa ::cudaDeviceGetCacheConfig + */ +extern __CUDA_DEPRECATED __host__ cudaError_t CUDARTAPI cudaThreadGetCacheConfig(enum cudaFuncCache *pCacheConfig); + +/** + * \brief Sets the preferred cache configuration for the current device. + * + * \deprecated + * + * Note that this function is deprecated because its name does not + * reflect its behavior. Its functionality is identical to the + * non-deprecated function ::cudaDeviceSetCacheConfig(), which should be + * used instead. + * + * On devices where the L1 cache and shared memory use the same hardware + * resources, this sets through \p cacheConfig the preferred cache + * configuration for the current device. This is only a preference. The + * runtime will use the requested configuration if possible, but it is free to + * choose a different configuration if required to execute the function. Any + * function preference set via + * \ref ::cudaFuncSetCacheConfig(const void*, enum cudaFuncCache) "cudaFuncSetCacheConfig (C API)" + * or + * \ref ::cudaFuncSetCacheConfig(T*, enum cudaFuncCache) "cudaFuncSetCacheConfig (C++ API)" + * will be preferred over this device-wide setting. Setting the device-wide + * cache configuration to ::cudaFuncCachePreferNone will cause subsequent + * kernel launches to prefer to not change the cache configuration unless + * required to launch the kernel. + * + * This setting does nothing on devices where the size of the L1 cache and + * shared memory are fixed. + * + * Launching a kernel with a different preference than the most recent + * preference setting may insert a device-side synchronization point. + * + * The supported cache configurations are: + * - ::cudaFuncCachePreferNone: no preference for shared memory or L1 (default) + * - ::cudaFuncCachePreferShared: prefer larger shared memory and smaller L1 cache + * - ::cudaFuncCachePreferL1: prefer larger L1 cache and smaller shared memory + * + * \param cacheConfig - Requested cache configuration + * + * \return + * ::cudaSuccess + * \notefnerr + * \note_init_rt + * \note_callback + * + * \sa ::cudaDeviceSetCacheConfig + */ +extern __CUDA_DEPRECATED __host__ cudaError_t CUDARTAPI cudaThreadSetCacheConfig(enum cudaFuncCache cacheConfig); + +/** @} */ /* END CUDART_THREAD_DEPRECATED */ + + + /** * \defgroup CUDART_ERROR Error Handling * @@ -978,7 +1263,8 @@ extern __host__ cudaError_t CUDARTAPI cudaIpcCloseMemHandle(void *devPtr); * ::cudaErrorInvalidPtx, * ::cudaErrorUnsupportedPtxVersion, * ::cudaErrorNoKernelImageForDevice, - * ::cudaErrorJitCompilerNotFound + * ::cudaErrorJitCompilerNotFound, + * ::cudaErrorJitCompilationDisabled * \notefnerr * \note_init_rt * \note_callback @@ -1025,7 +1311,8 @@ extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaGetLastError(void); * ::cudaErrorInvalidPtx, * ::cudaErrorUnsupportedPtxVersion, * ::cudaErrorNoKernelImageForDevice, - * ::cudaErrorJitCompilerNotFound + * ::cudaErrorJitCompilerNotFound, + * ::cudaErrorJitCompilationDisabled * \notefnerr * \note_init_rt * \note_callback @@ -1543,6 +1830,8 @@ extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaGetDeviceProperties * - ::cudaDevAttrMaxSharedMemoryPerBlockOptin: Maximum per block shared memory size on the device. This value can * be opted into when using ::cudaFuncSetAttribute * - ::cudaDevAttrMaxBlocksPerMultiprocessor: Maximum number of thread blocks that can reside on a multiprocessor. + * - ::cudaDevAttrMaxPersistingL2CacheSize: Maximum L2 persisting lines capacity setting in bytes. + * - ::cudaDevAttrMaxAccessPolicyWindowSize: Maximum value of cudaAccessPolicyWindow::num_bytes. * - ::cudaDevAttrHostRegisterReadOnly: Device supports using the ::cudaHostRegister flag cudaHostRegisterReadOnly * to register memory that must be mapped as read-only to the GPU * - ::cudaDevAttrSparseCudaArraySupported: 1 if the device supports sparse CUDA arrays and sparse CUDA mipmapped arrays. @@ -1565,6 +1854,67 @@ extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaGetDeviceProperties */ extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaDeviceGetAttribute(int *value, enum cudaDeviceAttr attr, int device); +/** + * \brief Returns the default mempool of a device + * + * The default mempool of a device contains device memory from that device. + * + * \return + * ::cudaSuccess, + * ::cudaErrorInvalidDevice, + * ::cudaErrorInvalidValue + * ::cudaErrorNotSupported + * \notefnerr + * \note_init_rt + * \note_callback + * + * \sa ::cuDeviceGetDefaultMemPool, ::cudaMallocAsync, ::cudaMemPoolTrimTo, ::cudaMemPoolGetAttribute, ::cudaDeviceSetMemPool, ::cudaMemPoolSetAttribute, ::cudaMemPoolSetAccess + */ +extern __host__ cudaError_t CUDARTAPI cudaDeviceGetDefaultMemPool(cudaMemPool_t *memPool, int device); + + +/** + * \brief Sets the current memory pool of a device + * + * The memory pool must be local to the specified device. + * Unless a mempool is specified in the ::cudaMallocAsync call, + * ::cudaMallocAsync allocates from the current mempool of the provided stream's device. + * By default, a device's current memory pool is its default memory pool. + * + * \note Use ::cudaMallocFromPoolAsync to specify asynchronous allocations from a device different + * than the one the stream runs on. + * + * \returns + * ::cudaSuccess, + * ::cudaErrorInvalidValue + * ::cudaErrorInvalidDevice + * ::cudaErrorNotSupported + * \notefnerr + * \note_callback + * + * \sa ::cuDeviceSetDefaultMemPool, ::cudaDeviceGetMemPool, ::cudaDeviceGetDefaultMemPool, ::cudaMemPoolCreate, ::cudaMemPoolDestroy, ::cudaMallocFromPoolAsync + */ +extern __host__ cudaError_t CUDARTAPI cudaDeviceSetMemPool(int device, cudaMemPool_t memPool); + +/** + * \brief Gets the current mempool for a device + * + * Returns the last pool provided to ::cudaDeviceSetMemPool for this device + * or the device's default memory pool if ::cudaDeviceSetMemPool has never been called. + * By default the current mempool is the default mempool for a device, + * otherwise the returned pool must have been set with ::cuDeviceSetMemPool or ::cudaDeviceSetMemPool. + * + * \returns + * ::cudaSuccess, + * ::cudaErrorInvalidValue + * ::cudaErrorNotSupported + * \notefnerr + * \note_init_rt + * \note_callback + * + * \sa ::cuDeviceGetMemPool, ::cudaDeviceGetDefaultMemPool, ::cudaDeviceSetMemPool + */ +extern __host__ cudaError_t CUDARTAPI cudaDeviceGetMemPool(cudaMemPool_t *memPool, int device); /** * \brief Return NvSciSync attributes that this device can support. @@ -2399,7 +2749,11 @@ extern __host__ cudaError_t CUDARTAPI cudaStreamQuery(cudaStream_t stream); * \sa ::cudaStreamCreate, ::cudaStreamCreateWithFlags, ::cudaStreamWaitEvent, ::cudaStreamSynchronize, ::cudaStreamAddCallback, ::cudaStreamDestroy, ::cudaMallocManaged, * ::cuStreamAttachMemAsync */ -extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaStreamAttachMemAsync(cudaStream_t stream, void *devPtr, size_t length __dv(0), unsigned int flags __dv(cudaMemAttachSingle)); +#if defined(__cplusplus) +extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaStreamAttachMemAsync(cudaStream_t stream, void *devPtr, size_t length __dv(0), unsigned int flags = cudaMemAttachSingle); +#else +extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaStreamAttachMemAsync(cudaStream_t stream, void *devPtr, size_t length __dv(0), unsigned int flags); +#endif /** * \brief Begins graph capture on a stream @@ -2557,10 +2911,13 @@ extern __host__ cudaError_t CUDARTAPI cudaStreamIsCapturing(cudaStream_t stream, /** * \brief Query capture status of a stream * + * Note there is a later version of this API, ::cudaStreamGetCaptureInfo_v2. It will + * supplant this version in 12.0, which is retained for minor version compatibility. + * * Query the capture status of a stream and get a unique id representing * the capture sequence over the lifetime of the process. * - * If called on ::cudaStreamLegacy (the "null stream") while a stream not created + * If called on ::cudaStreamLegacy (the "null stream") while a stream not created * with ::cudaStreamNonBlocking is capturing, returns ::cudaErrorStreamCaptureImplicit. * * A valid id is returned only if both of the following are true: @@ -2577,10 +2934,99 @@ extern __host__ cudaError_t CUDARTAPI cudaStreamIsCapturing(cudaStream_t stream, * \notefnerr * * \sa + * ::cudaStreamGetCaptureInfo_v2, * ::cudaStreamBeginCapture, * ::cudaStreamIsCapturing */ extern __host__ cudaError_t CUDARTAPI cudaStreamGetCaptureInfo(cudaStream_t stream, enum cudaStreamCaptureStatus *pCaptureStatus, unsigned long long *pId); + +/** + * \brief Query a stream's capture state (11.3+) + * + * Query stream state related to stream capture. + * + * If called on ::cudaStreamLegacy (the "null stream") while a stream not created + * with ::cudaStreamNonBlocking is capturing, returns ::cudaErrorStreamCaptureImplicit. + * + * Valid data (other than capture status) is returned only if both of the following are true: + * - the call returns cudaSuccess + * - the returned capture status is ::cudaStreamCaptureStatusActive + * + * This version of cudaStreamGetCaptureInfo is introduced in CUDA 11.3 and will supplant the + * previous version ::cudaStreamGetCaptureInfo in 12.0. Developers requiring compatibility + * across minor versions to CUDA 11.0 (driver version 445) can do one of the following: + * - Use the older version of the API, ::cudaStreamGetCaptureInfo + * - Pass null for all of \p graph_out, \p dependencies_out, and \p numDependencies_out. + * + * \param stream - The stream to query + * \param captureStatus_out - Location to return the capture status of the stream; required + * \param id_out - Optional location to return an id for the capture sequence, which is + * unique over the lifetime of the process + * \param graph_out - Optional location to return the graph being captured into. All + * operations other than destroy and node removal are permitted on the graph + * while the capture sequence is in progress. This API does not transfer + * ownership of the graph, which is transferred or destroyed at + * ::cudaStreamEndCapture. Note that the graph handle may be invalidated before + * end of capture for certain errors. Nodes that are or become + * unreachable from the original stream at ::cudaStreamEndCapture due to direct + * actions on the graph do not trigger ::cudaErrorStreamCaptureUnjoined. + * \param dependencies_out - Optional location to store a pointer to an array of nodes. + * The next node to be captured in the stream will depend on this set of nodes, + * absent operations such as event wait which modify this set. The array pointer + * is valid until the next API call which operates on the stream or until end of + * capture. The node handles may be copied out and are valid until they or the + * graph is destroyed. The driver-owned array may also be passed directly to + * APIs that operate on the graph (not the stream) without copying. + * \param numDependencies_out - Optional location to store the size of the array + * returned in dependencies_out. + * + * \return + * ::cudaSuccess, + * ::cudaErrorInvalidValue, + * ::cudaErrorStreamCaptureImplicit + * \note_graph_thread_safety + * \notefnerr + * + * \sa + * ::cudaStreamGetCaptureInfo, + * ::cudaStreamBeginCapture, + * ::cudaStreamIsCapturing, + * ::cudaStreamUpdateCaptureDependencies + */ +extern __host__ cudaError_t CUDARTAPI cudaStreamGetCaptureInfo_v2(cudaStream_t stream, enum cudaStreamCaptureStatus *captureStatus_out, unsigned long long *id_out __dv(0), cudaGraph_t *graph_out __dv(0), const cudaGraphNode_t **dependencies_out __dv(0), size_t *numDependencies_out __dv(0)); + +/** + * \brief Update the set of dependencies in a capturing stream (11.3+) + * + * Modifies the dependency set of a capturing stream. The dependency set is the set + * of nodes that the next captured node in the stream will depend on. + * + * Valid flags are ::cudaStreamAddCaptureDependencies and + * ::cudaStreamSetCaptureDependencies. These control whether the set passed to + * the API is added to the existing set or replaces it. A flags value of 0 defaults + * to ::cudaStreamAddCaptureDependencies. + * + * Nodes that are removed from the dependency set via this API do not result in + * ::cudaErrorStreamCaptureUnjoined if they are unreachable from the stream at + * ::cudaStreamEndCapture. + * + * Returns ::cudaErrorIllegalState if the stream is not capturing. + * + * This API is new in CUDA 11.3. Developers requiring compatibility across minor + * versions of the CUDA driver to 11.0 should not use this API or provide a fallback. + * + * \return + * ::cudaSuccess, + * ::cudaErrorInvalidValue, + * ::cudaErrorIllegalState + * \notefnerr + * + * \sa + * ::cudaStreamBeginCapture, + * ::cudaStreamGetCaptureInfo, + * ::cudaStreamGetCaptureInfo_v2 + */ +extern __host__ cudaError_t CUDARTAPI cudaStreamUpdateCaptureDependencies(cudaStream_t stream, cudaGraphNode_t *dependencies, size_t numDependencies, unsigned int flags __dv(0)); /** @} */ /* END CUDART_STREAM */ /** @@ -2737,10 +3183,12 @@ extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaEventRecord(cudaEve * ::cudaEventCreateWithFlags, ::cudaEventQuery, * ::cudaEventSynchronize, ::cudaEventDestroy, ::cudaEventElapsedTime, * ::cudaStreamWaitEvent, - * ::cudaEventRecord + * ::cudaEventRecord, * ::cuEventRecord, */ -extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaEventRecordWithFlags(cudaEvent_t event, cudaStream_t stream __dv(0), unsigned int flags __dv(0)); +#if __CUDART_API_VERSION >= 11010 + extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaEventRecordWithFlags(cudaEvent_t event, cudaStream_t stream __dv(0), unsigned int flags __dv(0)); +#endif /** * \brief Queries an event's status @@ -3100,7 +3548,7 @@ extern __host__ cudaError_t CUDARTAPI cudaImportExternalMemory(cudaExternalMemor * \note_init_rt * \note_callback * - * \sa ::cudaImportExternalMemory + * \sa ::cudaImportExternalMemory, * ::cudaDestroyExternalMemory, * ::cudaExternalMemoryGetMappedMipmappedArray */ @@ -3155,7 +3603,7 @@ extern __host__ cudaError_t CUDARTAPI cudaExternalMemoryGetMappedBuffer(void **d * \note_init_rt * \note_callback * - * \sa ::cudaImportExternalMemory + * \sa ::cudaImportExternalMemory, * ::cudaDestroyExternalMemory, * ::cudaExternalMemoryGetMappedBuffer * @@ -3183,7 +3631,7 @@ extern __host__ cudaError_t CUDARTAPI cudaExternalMemoryGetMappedMipmappedArray( * \note_callback * \note_destroy_ub * - * \sa ::cudaImportExternalMemory + * \sa ::cudaImportExternalMemory, * ::cudaExternalMemoryGetMappedBuffer, * ::cudaExternalMemoryGetMappedMipmappedArray */ @@ -3220,14 +3668,16 @@ extern __host__ cudaError_t CUDARTAPI cudaDestroyExternalMemory(cudaExternalMemo * * \code typedef enum cudaExternalSemaphoreHandleType_enum { - cudaExternalSemaphoreHandleTypeOpaqueFd = 1, - cudaExternalSemaphoreHandleTypeOpaqueWin32 = 2, - cudaExternalSemaphoreHandleTypeOpaqueWin32Kmt = 3, - cudaExternalSemaphoreHandleTypeD3D12Fence = 4, - cudaExternalSemaphoreHandleTypeD3D11Fence = 5, - cudaExternalSemaphoreHandleTypeNvSciSync = 6, - cudaExternalSemaphoreHandleTypeKeyedMutex = 7, - cudaExternalSemaphoreHandleTypeKeyedMutexKmt = 8 + cudaExternalSemaphoreHandleTypeOpaqueFd = 1, + cudaExternalSemaphoreHandleTypeOpaqueWin32 = 2, + cudaExternalSemaphoreHandleTypeOpaqueWin32Kmt = 3, + cudaExternalSemaphoreHandleTypeD3D12Fence = 4, + cudaExternalSemaphoreHandleTypeD3D11Fence = 5, + cudaExternalSemaphoreHandleTypeNvSciSync = 6, + cudaExternalSemaphoreHandleTypeKeyedMutex = 7, + cudaExternalSemaphoreHandleTypeKeyedMutexKmt = 8, + cudaExternalSemaphoreHandleTypeTimelineSemaphoreFd = 9, + cudaExternalSemaphoreHandleTypeTimelineSemaphoreWin32 = 10 } cudaExternalSemaphoreHandleType; * \endcode * @@ -3293,7 +3743,7 @@ extern __host__ cudaError_t CUDARTAPI cudaDestroyExternalMemory(cudaExternalMemo * ::cudaExternalSemaphoreHandleDesc::handle::win32::name must not be * NULL. If ::cudaExternalSemaphoreHandleDesc::handle::win32::handle * is not NULL, then it represent a valid shared NT handle that - * is returned by IDXGIResource1::CreateSharedHandle when referring to + * is returned by IDXGIResource1::CreateSharedHandle when referring to * a IDXGIKeyedMutex object. * * If ::cudaExternalSemaphoreHandleDesc::type is @@ -3304,6 +3754,26 @@ extern __host__ cudaError_t CUDARTAPI cudaDestroyExternalMemory(cudaExternalMemo * handle that is returned by IDXGIResource::GetSharedHandle when * referring to a IDXGIKeyedMutex object. * + * If ::cudaExternalSemaphoreHandleDesc::type is + * ::cudaExternalSemaphoreHandleTypeTimelineSemaphoreFd, then + * ::cudaExternalSemaphoreHandleDesc::handle::fd must be a valid file + * descriptor referencing a synchronization object. Ownership of the + * file descriptor is transferred to the CUDA driver when the handle + * is imported successfully. Performing any operations on the file + * descriptor after it is imported results in undefined behavior. + * + * If ::cudaExternalSemaphoreHandleDesc::type is + * ::cudaExternalSemaphoreHandleTypeTimelineSemaphoreWin32, then exactly one of + * ::cudaExternalSemaphoreHandleDesc::handle::win32::handle and + * ::cudaExternalSemaphoreHandleDesc::handle::win32::name must not be + * NULL. If ::cudaExternalSemaphoreHandleDesc::handle::win32::handle + * is not NULL, then it must represent a valid shared NT handle that + * references a synchronization object. Ownership of this handle is + * not transferred to CUDA after the import operation, so the + * application must release the handle using the appropriate system + * call. If ::cudaExternalSemaphoreHandleDesc::handle::win32::name is + * not NULL, then it must name a valid synchronization object. + * * \param extSem_out - Returned handle to an external semaphore * \param semHandleDesc - Semaphore import handle descriptor * @@ -3338,7 +3808,9 @@ extern __host__ cudaError_t CUDARTAPI cudaImportExternalSemaphore(cudaExternalSe * * If the semaphore object is any one of the following types: * ::cudaExternalSemaphoreHandleTypeD3D12Fence, - * ::cudaExternalSemaphoreHandleTypeD3D11Fence + * ::cudaExternalSemaphoreHandleTypeD3D11Fence, + * ::cudaExternalSemaphoreHandleTypeTimelineSemaphoreFd, + * ::cudaExternalSemaphoreHandleTypeTimelineSemaphoreWin32 * then the semaphore will be set to the value specified in * ::cudaExternalSemaphoreSignalParams::params::fence::value. * @@ -3406,7 +3878,9 @@ extern __host__ cudaError_t CUDARTAPI cudaSignalExternalSemaphoresAsync(const cu * * If the semaphore object is any one of the following types: * ::cudaExternalSemaphoreHandleTypeD3D12Fence, - * ::cudaExternalSemaphoreHandleTypeD3D11Fence + * ::cudaExternalSemaphoreHandleTypeD3D11Fence, + * ::cudaExternalSemaphoreHandleTypeTimelineSemaphoreFd, + * ::cudaExternalSemaphoreHandleTypeTimelineSemaphoreWin32 * then waiting on the semaphore will wait until the value of the * semaphore is greater than or equal to * ::cudaExternalSemaphoreWaitParams::params::fence::value. @@ -3536,7 +4010,8 @@ extern __host__ cudaError_t CUDARTAPI cudaDestroyExternalSemaphore(cudaExternalS * ::cudaErrorInvalidPtx, * ::cudaErrorUnsupportedPtxVersion, * ::cudaErrorNoKernelImageForDevice, - * ::cudaErrorJitCompilerNotFound + * ::cudaErrorJitCompilerNotFound, + * ::cudaErrorJitCompilationDisabled * \note_null_stream * \notefnerr * \note_init_rt @@ -3608,6 +4083,8 @@ extern __host__ cudaError_t CUDARTAPI cudaLaunchCooperativeKernel(const void *fu /** * \brief Launches device functions on multiple devices where thread blocks can cooperate and synchronize as they execute * + * \deprecated This function is deprecated as of CUDA 11.3. + * * Invokes kernels as specified in the \p launchParamsList array where each element * of the array specifies all the parameters required to perform a single kernel launch. * These kernels can cooperate and synchronize as they execute. The size of the array is @@ -3702,7 +4179,7 @@ extern __host__ cudaError_t CUDARTAPI cudaLaunchCooperativeKernel(const void *fu * ::cudaLaunchCooperativeKernel, * ::cuLaunchCooperativeKernelMultiDevice */ -extern __host__ cudaError_t CUDARTAPI cudaLaunchCooperativeKernelMultiDevice(struct cudaLaunchParams *launchParamsList, unsigned int numDevices, unsigned int flags __dv(0)); +extern __CUDA_DEPRECATED __host__ cudaError_t CUDARTAPI cudaLaunchCooperativeKernelMultiDevice(struct cudaLaunchParams *launchParamsList, unsigned int numDevices, unsigned int flags __dv(0)); /** * \brief Sets the preferred cache configuration for a device function @@ -3876,6 +4353,58 @@ extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaFuncGetAttributes(s */ extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaFuncSetAttribute(const void *func, enum cudaFuncAttribute attr, int value); + + +/** + * \brief Converts a double argument to be executed on a device + * + * \param d - Double to convert + * + * \deprecated This function is deprecated as of CUDA 7.5 + * + * Converts the double value of \p d to an internal float representation if + * the device does not support double arithmetic. If the device does natively + * support doubles, then this function does nothing. + * + * \return + * ::cudaSuccess + * \notefnerr + * \note_init_rt + * \note_callback + * + * \sa + * \ref ::cudaFuncSetCacheConfig(const void*, enum cudaFuncCache) "cudaFuncSetCacheConfig (C API)", + * \ref ::cudaFuncGetAttributes(struct cudaFuncAttributes*, const void*) "cudaFuncGetAttributes (C API)", + * ::cudaSetDoubleForHost + */ +extern __CUDA_DEPRECATED __host__ cudaError_t CUDARTAPI cudaSetDoubleForDevice(double *d); + +/** + * \brief Converts a double argument after execution on a device + * + * \deprecated This function is deprecated as of CUDA 7.5 + * + * Converts the double value of \p d from a potentially internal float + * representation if the device does not support double arithmetic. If the + * device does natively support doubles, then this function does nothing. + * + * \param d - Double to convert + * + * \return + * ::cudaSuccess + * \notefnerr + * \note_init_rt + * \note_callback + * + * \sa + * \ref ::cudaFuncSetCacheConfig(const void*, enum cudaFuncCache) "cudaFuncSetCacheConfig (C API)", + * \ref ::cudaFuncGetAttributes(struct cudaFuncAttributes*, const void*) "cudaFuncGetAttributes (C API)", + * ::cudaSetDoubleForDevice + */ +extern __CUDA_DEPRECATED __host__ cudaError_t CUDARTAPI cudaSetDoubleForHost(double *d); + + + /** * \brief Enqueues a host function call in a stream * @@ -4191,8 +4720,11 @@ extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaOccupancyMaxActiveB * ::cudaFreeHost, ::cudaHostAlloc, ::cudaDeviceGetAttribute, ::cudaStreamAttachMemAsync, * ::cuMemAllocManaged */ -extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMallocManaged(void **devPtr, size_t size, unsigned int flags __dv(cudaMemAttachGlobal)); - +#if defined(__cplusplus) +extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMallocManaged(void **devPtr, size_t size, unsigned int flags = cudaMemAttachGlobal); +#else +extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMallocManaged(void **devPtr, size_t size, unsigned int flags); +#endif /** * \brief Allocate memory on the device @@ -5380,6 +5912,35 @@ extern __host__ cudaError_t CUDARTAPI cudaMemGetInfo(size_t *free, size_t *total */ extern __host__ cudaError_t CUDARTAPI cudaArrayGetInfo(struct cudaChannelFormatDesc *desc, struct cudaExtent *extent, unsigned int *flags, cudaArray_t array); +/** + * \brief Gets a CUDA array plane from a CUDA array + * + * Returns in \p pPlaneArray a CUDA array that represents a single format plane + * of the CUDA array \p hArray. + * + * If \p planeIdx is greater than the maximum number of planes in this array or if the array does + * not have a multi-planar format e.g: ::cudaChannelFormatKindNV12, then ::cudaErrorInvalidValue is returned. + * + * Note that if the \p hArray has format ::cudaChannelFormatKindNV12, then passing in 0 for \p planeIdx returns + * a CUDA array of the same size as \p hArray but with one 8-bit channel and ::cudaChannelFormatKindUnsigned as its format kind. + * If 1 is passed for \p planeIdx, then the returned CUDA array has half the height and width + * of \p hArray with two 8-bit channels and ::cudaChannelFormatKindUnsigned as its format kind. + * + * \param pPlaneArray - Returned CUDA array referenced by the \p planeIdx + * \param hArray - CUDA array + * \param planeIdx - Plane index + * + * \return + * ::cudaSuccess, + * ::cudaErrorInvalidValue + * ::cudaErrorInvalidResourceHandle + * \notefnerr + * + * \sa + * ::cuArrayGetPlane + */ +extern __host__ cudaError_t CUDARTAPI cudaArrayGetPlane(cudaArray_t *pPlaneArray, cudaArray_t hArray, unsigned int planeIdx); + /** * \brief Returns the layout properties of a sparse CUDA array * @@ -5405,7 +5966,9 @@ extern __host__ cudaError_t CUDARTAPI cudaArrayGetInfo(struct cudaChannelFormatD * ::cudaMipmappedArrayGetSparseProperties, * ::cuMemMapArrayAsync */ -extern __host__ cudaError_t CUDARTAPI cudaArrayGetSparseProperties(struct cudaArraySparseProperties *sparseProperties, cudaArray_t array); +#if __CUDART_API_VERSION >= 11010 + extern __host__ cudaError_t CUDARTAPI cudaArrayGetSparseProperties(struct cudaArraySparseProperties *sparseProperties, cudaArray_t array); +#endif /** * \brief Returns the layout properties of a sparse CUDA mipmapped array @@ -5433,7 +5996,9 @@ extern __host__ cudaError_t CUDARTAPI cudaArrayGetSparseProperties(struct cudaAr * ::cudaArrayGetSparseProperties, * ::cuMemMapArrayAsync */ -extern __host__ cudaError_t CUDARTAPI cudaMipmappedArrayGetSparseProperties(struct cudaArraySparseProperties *sparseProperties, cudaMipmappedArray_t mipmap); +#if __CUDART_API_VERSION >= 11010 + extern __host__ cudaError_t CUDARTAPI cudaMipmappedArrayGetSparseProperties(struct cudaArraySparseProperties *sparseProperties, cudaMipmappedArray_t mipmap); +#endif /** * \brief Copies data between host and device @@ -5567,10 +6132,10 @@ extern __host__ cudaError_t CUDARTAPI cudaMemcpy2D(void *dst, size_t dpitch, con * \brief Copies data between host and device * * Copies a matrix (\p height rows of \p width bytes each) from the memory - * area pointed to by \p src to the CUDA array \p dst starting at the - * upper left corner (\p wOffset, \p hOffset) where \p kind specifies the - * direction of the copy, and must be one of ::cudaMemcpyHostToHost, - * ::cudaMemcpyHostToDevice, ::cudaMemcpyDeviceToHost, + * area pointed to by \p src to the CUDA array \p dst starting at + * \p hOffset rows and \p wOffset bytes from the upper left corner, + * where \p kind specifies the direction of the copy, and must be one + * of ::cudaMemcpyHostToHost, ::cudaMemcpyHostToDevice, ::cudaMemcpyDeviceToHost, * ::cudaMemcpyDeviceToDevice, or ::cudaMemcpyDefault. Passing * ::cudaMemcpyDefault is recommended, in which case the type of transfer is * inferred from the pointer values. However, ::cudaMemcpyDefault is only @@ -5582,8 +6147,8 @@ extern __host__ cudaError_t CUDARTAPI cudaMemcpy2D(void *dst, size_t dpitch, con * exceeds the maximum allowed. * * \param dst - Destination memory address - * \param wOffset - Destination starting X offset - * \param hOffset - Destination starting Y offset + * \param wOffset - Destination starting X offset (columns in bytes) + * \param hOffset - Destination starting Y offset (rows) * \param src - Source memory address * \param spitch - Pitch of source memory * \param width - Width of matrix transfer (columns in bytes) @@ -5617,8 +6182,8 @@ extern __host__ cudaError_t CUDARTAPI cudaMemcpy2DToArray(cudaArray_t dst, size_ * \brief Copies data between host and device * * Copies a matrix (\p height rows of \p width bytes each) from the CUDA - * array \p srcArray starting at the upper left corner - * (\p wOffset, \p hOffset) to the memory area pointed to by \p dst, where + * array \p src starting at \p hOffset rows and \p wOffset bytes from the + * upper left corner to the memory area pointed to by \p dst, where * \p kind specifies the direction of the copy, and must be one of * ::cudaMemcpyHostToHost, ::cudaMemcpyHostToDevice, ::cudaMemcpyDeviceToHost, * ::cudaMemcpyDeviceToDevice, or ::cudaMemcpyDefault. Passing @@ -5634,8 +6199,8 @@ extern __host__ cudaError_t CUDARTAPI cudaMemcpy2DToArray(cudaArray_t dst, size_ * \param dst - Destination memory address * \param dpitch - Pitch of destination memory * \param src - Source memory address - * \param wOffset - Source starting X offset - * \param hOffset - Source starting Y offset + * \param wOffset - Source starting X offset (columns in bytes) + * \param hOffset - Source starting Y offset (rows) * \param width - Width of matrix transfer (columns in bytes) * \param height - Height of matrix transfer (rows) * \param kind - Type of transfer @@ -5667,9 +6232,9 @@ extern __host__ cudaError_t CUDARTAPI cudaMemcpy2DFromArray(void *dst, size_t dp * \brief Copies data between host and device * * Copies a matrix (\p height rows of \p width bytes each) from the CUDA - * array \p srcArray starting at the upper left corner - * (\p wOffsetSrc, \p hOffsetSrc) to the CUDA array \p dst starting at - * the upper left corner (\p wOffsetDst, \p hOffsetDst), where \p kind + * array \p src starting at \p hOffsetSrc rows and \p wOffsetSrc bytes from the + * upper left corner to the CUDA array \p dst starting at \p hOffsetDst rows + * and \p wOffsetDst bytes from the upper left corner, where \p kind * specifies the direction of the copy, and must be one of * ::cudaMemcpyHostToHost, ::cudaMemcpyHostToDevice, ::cudaMemcpyDeviceToHost, * ::cudaMemcpyDeviceToDevice, or ::cudaMemcpyDefault. Passing @@ -5680,11 +6245,11 @@ extern __host__ cudaError_t CUDARTAPI cudaMemcpy2DFromArray(void *dst, size_t dp * \p wOffsetSrc + \p width must not exceed the width of the CUDA array \p src. * * \param dst - Destination memory address - * \param wOffsetDst - Destination starting X offset - * \param hOffsetDst - Destination starting Y offset + * \param wOffsetDst - Destination starting X offset (columns in bytes) + * \param hOffsetDst - Destination starting Y offset (rows) * \param src - Source memory address - * \param wOffsetSrc - Source starting X offset - * \param hOffsetSrc - Source starting Y offset + * \param wOffsetSrc - Source starting X offset (columns in bytes) + * \param hOffsetSrc - Source starting Y offset (rows) * \param width - Width of matrix transfer (columns in bytes) * \param height - Height of matrix transfer (rows) * \param kind - Type of transfer @@ -5845,7 +6410,7 @@ extern __host__ cudaError_t CUDARTAPI cudaMemcpyFromSymbol(void *dst, const void * ::cudaMemcpyFromSymbol, ::cudaMemcpy2DAsync, * ::cudaMemcpy2DToArrayAsync, * ::cudaMemcpy2DFromArrayAsync, - * ::cudaMemcpyToSymbolAsync, ::cudaMemcpyFromSymbolAsync + * ::cudaMemcpyToSymbolAsync, ::cudaMemcpyFromSymbolAsync, * ::cuMemcpyAsync, * ::cuMemcpyDtoHAsync, * ::cuMemcpyHtoDAsync, @@ -5955,9 +6520,9 @@ extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMemcpy2DAsync(void * \brief Copies data between host and device * * Copies a matrix (\p height rows of \p width bytes each) from the memory - * area pointed to by \p src to the CUDA array \p dst starting at the - * upper left corner (\p wOffset, \p hOffset) where \p kind specifies the - * direction of the copy, and must be one of ::cudaMemcpyHostToHost, + * area pointed to by \p src to the CUDA array \p dst starting at \p hOffset + * rows and \p wOffset bytes from the upper left corner, where \p kind specifies + * the direction of the copy, and must be one of ::cudaMemcpyHostToHost, * ::cudaMemcpyHostToDevice, ::cudaMemcpyDeviceToHost, * ::cudaMemcpyDeviceToDevice, or ::cudaMemcpyDefault. Passing * ::cudaMemcpyDefault is recommended, in which case the type of transfer is @@ -5977,8 +6542,8 @@ extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMemcpy2DAsync(void * streams. * * \param dst - Destination memory address - * \param wOffset - Destination starting X offset - * \param hOffset - Destination starting Y offset + * \param wOffset - Destination starting X offset (columns in bytes) + * \param hOffset - Destination starting Y offset (rows) * \param src - Source memory address * \param spitch - Pitch of source memory * \param width - Width of matrix transfer (columns in bytes) @@ -6013,9 +6578,9 @@ extern __host__ cudaError_t CUDARTAPI cudaMemcpy2DToArrayAsync(cudaArray_t dst, * \brief Copies data between host and device * * Copies a matrix (\p height rows of \p width bytes each) from the CUDA - * array \p srcArray starting at the upper left corner - * (\p wOffset, \p hOffset) to the memory area pointed to by \p dst, where - * \p kind specifies the direction of the copy, and must be one of + * array \p src starting at \p hOffset rows and \p wOffset bytes from the + * upper left corner to the memory area pointed to by \p dst, + * where \p kind specifies the direction of the copy, and must be one of * ::cudaMemcpyHostToHost, ::cudaMemcpyHostToDevice, ::cudaMemcpyDeviceToHost, * ::cudaMemcpyDeviceToDevice, or ::cudaMemcpyDefault. Passing * ::cudaMemcpyDefault is recommended, in which case the type of transfer is @@ -6036,8 +6601,8 @@ extern __host__ cudaError_t CUDARTAPI cudaMemcpy2DToArrayAsync(cudaArray_t dst, * \param dst - Destination memory address * \param dpitch - Pitch of destination memory * \param src - Source memory address - * \param wOffset - Source starting X offset - * \param hOffset - Source starting Y offset + * \param wOffset - Source starting X offset (columns in bytes) + * \param hOffset - Source starting Y offset (rows) * \param width - Width of matrix transfer (columns in bytes) * \param height - Height of matrix transfer (rows) * \param kind - Type of transfer @@ -6740,7 +7305,7 @@ extern __host__ cudaError_t CUDARTAPI cudaMemRangeGetAttribute(void *data, size_ * \note_init_rt * \note_callback * - * \sa ::cudaMemRangeGetAttribute, ::cudaMemAdvise + * \sa ::cudaMemRangeGetAttribute, ::cudaMemAdvise, * ::cudaMemPrefetchAsync, * ::cuMemRangeGetAttributes */ @@ -6769,8 +7334,8 @@ extern __host__ cudaError_t CUDARTAPI cudaMemRangeGetAttributes(void **data, siz * \deprecated * * Copies \p count bytes from the memory area pointed to by \p src to the - * CUDA array \p dst starting at the upper left corner - * (\p wOffset, \p hOffset), where \p kind specifies the direction + * CUDA array \p dst starting at \p hOffset rows and \p wOffset bytes from + * the upper left corner, where \p kind specifies the direction * of the copy, and must be one of ::cudaMemcpyHostToHost, * ::cudaMemcpyHostToDevice, ::cudaMemcpyDeviceToHost, * ::cudaMemcpyDeviceToDevice, or ::cudaMemcpyDefault. Passing @@ -6779,8 +7344,8 @@ extern __host__ cudaError_t CUDARTAPI cudaMemRangeGetAttributes(void **data, siz * allowed on systems that support unified virtual addressing. * * \param dst - Destination memory address - * \param wOffset - Destination starting X offset - * \param hOffset - Destination starting Y offset + * \param wOffset - Destination starting X offset (columns in bytes) + * \param hOffset - Destination starting Y offset (rows) * \param src - Source memory address * \param count - Size in bytes to copy * \param kind - Type of transfer @@ -6811,9 +7376,9 @@ extern __CUDA_DEPRECATED __host__ cudaError_t CUDARTAPI cudaMemcpyToArray(cudaAr * * \deprecated * - * Copies \p count bytes from the CUDA array \p src starting at the upper - * left corner (\p wOffset, hOffset) to the memory area pointed to by \p dst, - * where \p kind specifies the direction of the copy, and must be one of + * Copies \p count bytes from the CUDA array \p src starting at \p hOffset rows + * and \p wOffset bytes from the upper left corner to the memory area pointed to + * by \p dst, where \p kind specifies the direction of the copy, and must be one of * ::cudaMemcpyHostToHost, ::cudaMemcpyHostToDevice, ::cudaMemcpyDeviceToHost, * ::cudaMemcpyDeviceToDevice, or ::cudaMemcpyDefault. Passing * ::cudaMemcpyDefault is recommended, in which case the type of transfer is @@ -6822,8 +7387,8 @@ extern __CUDA_DEPRECATED __host__ cudaError_t CUDARTAPI cudaMemcpyToArray(cudaAr * * \param dst - Destination memory address * \param src - Source memory address - * \param wOffset - Source starting X offset - * \param hOffset - Source starting Y offset + * \param wOffset - Source starting X offset (columns in bytes) + * \param hOffset - Source starting Y offset (rows) * \param count - Size in bytes to copy * \param kind - Type of transfer * @@ -6853,10 +7418,10 @@ extern __CUDA_DEPRECATED __host__ cudaError_t CUDARTAPI cudaMemcpyFromArray(void * * \deprecated * - * Copies \p count bytes from the CUDA array \p src starting at the upper - * left corner (\p wOffsetSrc, \p hOffsetSrc) to the CUDA array \p dst - * starting at the upper left corner (\p wOffsetDst, \p hOffsetDst) where - * \p kind specifies the direction of the copy, and must be one of + * Copies \p count bytes from the CUDA array \p src starting at \p hOffsetSrc + * rows and \p wOffsetSrc bytes from the upper left corner to the CUDA array + * \p dst starting at \p hOffsetDst rows and \p wOffsetDst bytes from the upper + * left corner, where \p kind specifies the direction of the copy, and must be one of * ::cudaMemcpyHostToHost, ::cudaMemcpyHostToDevice, ::cudaMemcpyDeviceToHost, * ::cudaMemcpyDeviceToDevice, or ::cudaMemcpyDefault. Passing * ::cudaMemcpyDefault is recommended, in which case the type of transfer is @@ -6864,11 +7429,11 @@ extern __CUDA_DEPRECATED __host__ cudaError_t CUDARTAPI cudaMemcpyFromArray(void * allowed on systems that support unified virtual addressing. * * \param dst - Destination memory address - * \param wOffsetDst - Destination starting X offset - * \param hOffsetDst - Destination starting Y offset + * \param wOffsetDst - Destination starting X offset (columns in bytes) + * \param hOffsetDst - Destination starting Y offset (rows) * \param src - Source memory address - * \param wOffsetSrc - Source starting X offset - * \param hOffsetSrc - Source starting Y offset + * \param wOffsetSrc - Source starting X offset (columns in bytes) + * \param hOffsetSrc - Source starting Y offset (rows) * \param count - Size in bytes to copy * \param kind - Type of transfer * @@ -6897,8 +7462,8 @@ extern __CUDA_DEPRECATED __host__ cudaError_t CUDARTAPI cudaMemcpyArrayToArray(c * \deprecated * * Copies \p count bytes from the memory area pointed to by \p src to the - * CUDA array \p dst starting at the upper left corner - * (\p wOffset, \p hOffset), where \p kind specifies the + * CUDA array \p dst starting at \p hOffset rows and \p wOffset bytes from + * the upper left corner, where \p kind specifies the * direction of the copy, and must be one of ::cudaMemcpyHostToHost, * ::cudaMemcpyHostToDevice, ::cudaMemcpyDeviceToHost, * ::cudaMemcpyDeviceToDevice, or ::cudaMemcpyDefault. Passing @@ -6913,8 +7478,8 @@ extern __CUDA_DEPRECATED __host__ cudaError_t CUDARTAPI cudaMemcpyArrayToArray(c * is non-zero, the copy may overlap with operations in other streams. * * \param dst - Destination memory address - * \param wOffset - Destination starting X offset - * \param hOffset - Destination starting Y offset + * \param wOffset - Destination starting X offset (columns in bytes) + * \param hOffset - Destination starting Y offset (rows) * \param src - Source memory address * \param count - Size in bytes to copy * \param kind - Type of transfer @@ -6947,9 +7512,9 @@ extern __CUDA_DEPRECATED __host__ cudaError_t CUDARTAPI cudaMemcpyToArrayAsync(c * * \deprecated * - * Copies \p count bytes from the CUDA array \p src starting at the upper - * left corner (\p wOffset, hOffset) to the memory area pointed to by \p dst, - * where \p kind specifies the direction of the copy, and must be one of + * Copies \p count bytes from the CUDA array \p src starting at \p hOffset rows + * and \p wOffset bytes from the upper left corner to the memory area pointed to + * by \p dst, where \p kind specifies the direction of the copy, and must be one of * ::cudaMemcpyHostToHost, ::cudaMemcpyHostToDevice, ::cudaMemcpyDeviceToHost, * ::cudaMemcpyDeviceToDevice, or ::cudaMemcpyDefault. Passing * ::cudaMemcpyDefault is recommended, in which case the type of transfer is @@ -6964,8 +7529,8 @@ extern __CUDA_DEPRECATED __host__ cudaError_t CUDARTAPI cudaMemcpyToArrayAsync(c * * \param dst - Destination memory address * \param src - Source memory address - * \param wOffset - Source starting X offset - * \param hOffset - Source starting Y offset + * \param wOffset - Source starting X offset (columns in bytes) + * \param hOffset - Source starting Y offset (rows) * \param count - Size in bytes to copy * \param kind - Type of transfer * \param stream - Stream identifier @@ -6994,6 +7559,402 @@ extern __CUDA_DEPRECATED __host__ cudaError_t CUDARTAPI cudaMemcpyFromArrayAsync /** @} */ /* END CUDART_MEMORY_DEPRECATED */ +/** + * \defgroup CUDART_MEMORY_POOLS Stream Ordered Memory Allocator + * + * ___MANBRIEF___ Functions for performing allocation and free operations in stream order. + * Functions for controlling the behavior of the underlying allocator. + * (___CURRENT_FILE___) ___ENDMANBRIEF___ + * + * + * @{ + * + * \section CUDART_MEMORY_POOLS_overview overview + * + * The asynchronous allocator allows the user to allocate and free in stream order. + * All asynchronous accesses of the allocation must happen between + * the stream executions of the allocation and the free. If the memory is accessed + * outside of the promised stream order, a use before allocation / use after free error + * will cause undefined behavior. + * + * The allocator is free to reallocate the memory as long as it can guarantee + * that compliant memory accesses will not overlap temporally. + * The allocator may refer to internal stream ordering as well as inter-stream dependencies + * (such as CUDA events and null stream dependencies) when establishing the temporal guarantee. + * The allocator may also insert inter-stream dependencies to establish the temporal guarantee. + * + * \section CUDART_MEMORY_POOLS_support Supported Platforms + * + * Whether or not a device supports the integrated stream ordered memory allocator + * may be queried by calling ::cudaDeviceGetAttribute() with the device attribute + * ::cudaDevAttrMemoryPoolsSupported. + */ + +/** + * \brief Allocates memory with stream ordered semantics + * + * Inserts an allocation operation into \p hStream. + * A pointer to the allocated memory is returned immediately in *dptr. + * The allocation must not be accessed until the the allocation operation completes. + * The allocation comes from the memory pool associated with the stream's device. + * + * \note The default memory pool of a device contains device memory from that device. + * \note Basic stream ordering allows future work submitted into the same stream to use the allocation. + * Stream query, stream synchronize, and CUDA events can be used to guarantee that the allocation + * operation completes before work submitted in a separate stream runs. + * + * \param[out] devPtr - Returned device pointer + * \param[in] size - Number of bytes to allocate + * \param[in] hStream - The stream establishing the stream ordering contract and the memory pool to allocate from + * + * \return + * ::cudaSuccess, + * ::cudaErrorInvalidValue, + * ::cudaErrorNotSupported, + * ::cudaErrorOutOfMemory, + * \notefnerr + * \note_null_stream + * \note_init_rt + * \note_callback + * + * \sa ::cuMemAllocAsync, + * \ref ::cudaMallocAsync(void** ptr, size_t size, cudaMemPool_t memPool, cudaStream_t stream) "cudaMallocAsync (C++ API)", + * ::cudaMallocFromPoolAsync, ::cudaFreeAsync, ::cudaDeviceSetMemPool, ::cudaDeviceGetDefaultMemPool, ::cudaDeviceGetMemPool, ::cudaMemPoolSetAccess, ::cudaMemPoolSetAttribute, ::cudaMemPoolGetAttribute + */ +extern __host__ cudaError_t CUDARTAPI cudaMallocAsync(void **devPtr, size_t size, cudaStream_t hStream); + +/** + * \brief Frees memory with stream ordered semantics + * + * Inserts a free operation into \p hStream. + * The allocation must not be accessed after stream execution reaches the free. + * After this API returns, accessing the memory from any subsequent work launched on the GPU + * or querying its pointer attributes results in undefined behavior. + * + * \param dptr - memory to free + * \param hStream - The stream establishing the stream ordering promise + * \returns + * ::cudaSuccess, + * ::cudaErrorInvalidValue, + * ::cudaErrorNotSupported + * \notefnerr + * \note_null_stream + * \note_init_rt + * \note_callback + * + * \sa ::cuMemFreeAsync, ::cudaMallocAsync + */ +extern __host__ cudaError_t CUDARTAPI cudaFreeAsync(void *devPtr, cudaStream_t hStream); + +/** + * \brief Tries to release memory back to the OS + * + * Releases memory back to the OS until the pool contains fewer than minBytesToKeep + * reserved bytes, or there is no more memory that the allocator can safely release. + * The allocator cannot release OS allocations that back outstanding asynchronous allocations. + * The OS allocations may happen at different granularity from the user allocations. + * + * \note: Allocations that have not been freed count as outstanding. + * \note: Allocations that have been asynchronously freed but whose completion has + * not been observed on the host (eg. by a synchronize) can count as outstanding. + * + * \param[in] pool - The memory pool to trim + * \param[in] minBytesToKeep - If the pool has less than minBytesToKeep reserved, + * the TrimTo operation is a no-op. Otherwise the pool will be guaranteed to have + * at least minBytesToKeep bytes reserved after the operation. + * \returns + * ::cudaSuccess, + * ::cudaErrorInvalidValue + * \note_callback + * + * \sa ::cuMemPoolTrimTo, ::cudaMallocAsync, ::cudaFreeAsync, ::cudaDeviceGetDefaultMemPool, ::cudaDeviceGetMemPool, ::cudaMemPoolCreate + */ +extern __host__ cudaError_t CUDARTAPI cudaMemPoolTrimTo(cudaMemPool_t memPool, size_t minBytesToKeep); + +/** + * \brief Sets attributes of a memory pool + * + * Supported attributes are: + * - ::cudaMemPoolAttrReleaseThreshold: (value type = cuuint64_t) + * Amount of reserved memory in bytes to hold onto before trying + * to release memory back to the OS. When more than the release + * threshold bytes of memory are held by the memory pool, the + * allocator will try to release memory back to the OS on the + * next call to stream, event or context synchronize. (default 0) + * - ::cudaMemPoolReuseFollowEventDependencies: (value type = int) + * Allow ::cudaMallocAsync to use memory asynchronously freed + * in another stream as long as a stream ordering dependency + * of the allocating stream on the free action exists. + * Cuda events and null stream interactions can create the required + * stream ordered dependencies. (default enabled) + * - ::cudaMemPoolReuseAllowOpportunistic: (value type = int) + * Allow reuse of already completed frees when there is no dependency + * between the free and allocation. (default enabled) + * - ::cudaMemPoolReuseAllowInternalDependencies: (value type = int) + * Allow ::cudaMallocAsync to insert new stream dependencies + * in order to establish the stream ordering required to reuse + * a piece of memory released by ::cudaFreeAsync (default enabled). + * + * \param[in] pool - The memory pool to modify + * \param[in] attr - The attribute to modify + * \param[in] value - Pointer to the value to assign + * + * \returns + * ::cudaSuccess, + * ::cudaErrorInvalidValue + * \note_callback + * + * \sa ::cuMemPoolSetAttribute, ::cudaMallocAsync, ::cudaFreeAsync, ::cudaDeviceGetDefaultMemPool, ::cudaDeviceGetMemPool, ::cudaMemPoolCreate + + */ +extern __host__ cudaError_t CUDARTAPI cudaMemPoolSetAttribute(cudaMemPool_t memPool, enum cudaMemPoolAttr attr, void *value ); + +/** + * \brief Gets attributes of a memory pool + * + * Supported attributes are: + * - ::cudaMemPoolAttrReleaseThreshold: (value type = cuuint64_t) + * Amount of reserved memory in bytes to hold onto before trying + * to release memory back to the OS. When more than the release + * threshold bytes of memory are held by the memory pool, the + * allocator will try to release memory back to the OS on the + * next call to stream, event or context synchronize. (default 0) + * - ::cudaMemPoolReuseFollowEventDependencies: (value type = int) + * Allow ::cudaMallocAsync to use memory asynchronously freed + * in another stream as long as a stream ordering dependency + * of the allocating stream on the free action exists. + * Cuda events and null stream interactions can create the required + * stream ordered dependencies. (default enabled) + * - ::cudaMemPoolReuseAllowOpportunistic: (value type = int) + * Allow reuse of already completed frees when there is no dependency + * between the free and allocation. (default enabled) + * - ::cudaMemPoolReuseAllowInternalDependencies: (value type = int) + * Allow ::cudaMallocAsync to insert new stream dependencies + * in order to establish the stream ordering required to reuse + * a piece of memory released by ::cudaFreeAsync (default enabled). + * + * \param[in] pool - The memory pool to get attributes of + * \param[in] attr - The attribute to get + * \param[in] value - Retrieved value + * + * \returns + * ::cudaSuccess, + * ::cudaErrorInvalidValue + * \note_callback + * + * \sa ::cuMemPoolGetAttribute, ::cudaMallocAsync, ::cudaFreeAsync, ::cudaDeviceGetDefaultMemPool, ::cudaDeviceGetMemPool, ::cudaMemPoolCreate + + */ +extern __host__ cudaError_t CUDARTAPI cudaMemPoolGetAttribute(cudaMemPool_t memPool, enum cudaMemPoolAttr attr, void *value ); + +/** + * \brief Controls visibility of pools between devices + * + * \param[in] pool - The pool being modified + * \param[in] map - Array of access descriptors. Each descriptor instructs the access to enable for a single gpu + * \param[in] count - Number of descriptors in the map array. + * + * \returns + * ::cudaSuccess, + * ::cudaErrorInvalidValue + * + * \sa ::cuMemPoolSetAccess, ::cudaMemPoolGetAccess, ::cudaMallocAsync, cudaFreeAsync + */ +extern __host__ cudaError_t CUDARTAPI cudaMemPoolSetAccess(cudaMemPool_t memPool, const struct cudaMemAccessDesc *descList, size_t count); + +/** + * \brief Returns the accessibility of a pool from a device + * + * Returns the accessibility of the pool's memory from the specified location. + * + * \param[out] flags - the accessibility of the pool from the specified location + * \param[in] memPool - the pool being queried + * \param[in] location - the location accessing the pool + * + * \sa ::cuMemPoolGetAccess, ::cudaMemPoolSetAccess + */ +extern __host__ cudaError_t CUDARTAPI cudaMemPoolGetAccess(enum cudaMemAccessFlags *flags, cudaMemPool_t memPool, struct cudaMemLocation *location); + +/** + * \brief Creates a memory pool + * + * Creates a CUDA memory pool and returns the handle in \p pool. The \p poolProps determines + * the properties of the pool such as the backing device and IPC capabilities. + * + * By default, the pool's memory will be accessible from the device it is allocated on. + * + * \note Specifying cudaMemHandleTypeNone creates a memory pool that will not support IPC. + * + * \returns + * ::cudaSuccess, + * ::cudaErrorInvalidValue, + * ::cudaErrorNotSupported + * + * \sa ::cuMemPoolCreate, ::cudaDeviceSetMemPool, ::cudaMallocFromPoolAsync, ::cudaMemPoolExportToShareableHandle, ::cudaDeviceGetDefaultMemPool, ::cudaDeviceGetMemPool + + */ +extern __host__ cudaError_t CUDARTAPI cudaMemPoolCreate(cudaMemPool_t *memPool, const struct cudaMemPoolProps *poolProps); + +/** + * \brief Destroys the specified memory pool + * + * If any pointers obtained from this pool haven't been freed or + * the pool has free operations that haven't completed + * when ::cudaMemPoolDestroy is invoked, the function will return immediately and the + * resources associated with the pool will be released automatically + * once there are no more outstanding allocations. + * + * Destroying the current mempool of a device sets the default mempool of + * that device as the current mempool for that device. + * + * \note A device's default memory pool cannot be destroyed. + * + * \returns + * ::cudaSuccess, + * ::cudaErrorInvalidValue + * + * \sa cuMemPoolDestroy, ::cudaFreeAsync, ::cudaDeviceSetMemPool, ::cudaDeviceGetDefaultMemPool, ::cudaDeviceGetMemPool, ::cudaMemPoolCreate + */ +extern __host__ cudaError_t CUDARTAPI cudaMemPoolDestroy(cudaMemPool_t memPool); + +/** + * \brief Allocates memory from a specified pool with stream ordered semantics. + * + * Inserts an allocation operation into \p hStream. + * A pointer to the allocated memory is returned immediately in *dptr. + * The allocation must not be accessed until the the allocation operation completes. + * The allocation comes from the specified memory pool. + * + * \note + * - The specified memory pool may be from a device different than that of the specified \p hStream. + * + * - Basic stream ordering allows future work submitted into the same stream to use the allocation. + * Stream query, stream synchronize, and CUDA events can be used to guarantee that the allocation + * operation completes before work submitted in a separate stream runs. + * + * \param[out] ptr - Returned device pointer + * \param[in] bytesize - Number of bytes to allocate + * \param[in] memPool - The pool to allocate from + * \param[in] stream - The stream establishing the stream ordering semantic + * + * \returns + * ::cudaSuccess, + * ::cudaErrorInvalidValue, + * ::cudaErrorNotSupported, + * ::cudaErrorOutOfMemory + * + * \sa ::cuMemAllocFromPoolAsync, + * \ref ::cudaMallocAsync(void** ptr, size_t size, cudaMemPool_t memPool, cudaStream_t stream) "cudaMallocAsync (C++ API)", + * ::cudaMallocAsync, ::cudaFreeAsync, ::cudaDeviceGetDefaultMemPool, ::cudaMemPoolCreate, ::cudaMemPoolSetAccess, ::cudaMemPoolSetAttribute + */ +extern __host__ cudaError_t CUDARTAPI cudaMallocFromPoolAsync(void **ptr, size_t size, cudaMemPool_t memPool, cudaStream_t stream); + +/** + * \brief Exports a memory pool to the requested handle type. + * + * Given an IPC capable mempool, create an OS handle to share the pool with another process. + * A recipient process can convert the shareable handle into a mempool with ::cudaMemPoolImportFromShareableHandle. + * Individual pointers can then be shared with the ::cudaMemPoolExportPointer and ::cudaMemPoolImportPointer APIs. + * The implementation of what the shareable handle is and how it can be transferred is defined by the requested + * handle type. + * + * \note: To create an IPC capable mempool, create a mempool with a CUmemAllocationHandleType other than cudaMemHandleTypeNone. + * + * \param[out] handle_out - pointer to the location in which to store the requested handle + * \param[in] pool - pool to export + * \param[in] handleType - the type of handle to create + * \param[in] flags - must be 0 + * + * \returns + * ::cudaSuccess, + * ::cudaErrorInvalidValue, + * ::cudaErrorOutOfMemory + * + * \sa ::cuMemPoolExportToShareableHandle, ::cudaMemPoolImportFromShareableHandle, ::cudaMemPoolExportPointer, ::cudaMemPoolImportPointer + */ +extern __host__ cudaError_t CUDARTAPI cudaMemPoolExportToShareableHandle( + void *shareableHandle, + cudaMemPool_t memPool, + enum cudaMemAllocationHandleType handleType, + unsigned int flags); + +/** + * \brief imports a memory pool from a shared handle. + * + * Specific allocations can be imported from the imported pool with ::cudaMemPoolImportPointer. + * + * \note Imported memory pools do not support creating new allocations. + * As such imported memory pools may not be used in ::cudaDeviceSetMemPool + * or ::cudaMallocFromPoolAsync calls. + * + * \param[out] pool_out - Returned memory pool + * \param[in] handle - OS handle of the pool to open + * \param[in] handleType - The type of handle being imported + * \param[in] flags - must be 0 + * + * \returns + * ::cudaSuccess, + * ::cudaErrorInvalidValue, + * ::cudaErrorOutOfMemory + * + * \sa ::cuMemPoolImportFromShareableHandle, ::cudaMemPoolExportToShareableHandle, ::cudaMemPoolExportPointer, ::cudaMemPoolImportPointer + */ +extern __host__ cudaError_t CUDARTAPI cudaMemPoolImportFromShareableHandle( + cudaMemPool_t *memPool, + void *shareableHandle, + enum cudaMemAllocationHandleType handleType, + unsigned int flags); + +/** + * \brief Export data to share a memory pool allocation between processes. + * + * Constructs \p shareData_out for sharing a specific allocation from an already shared memory pool. + * The recipient process can import the allocation with the ::cudaMemPoolImportPointer api. + * The data is not a handle and may be shared through any IPC mechanism. + * + * \param[out] shareData_out - Returned export data + * \param[in] ptr - pointer to memory being exported + * + * \returns + * ::cudaSuccess, + * ::cudaErrorInvalidValue, + * ::cudaErrorOutOfMemory + * + * \sa ::cuMemPoolExportPointer, ::cudaMemPoolExportToShareableHandle, ::cudaMemPoolImportFromShareableHandle, ::cudaMemPoolImportPointer + */ +extern __host__ cudaError_t CUDARTAPI cudaMemPoolExportPointer(struct cudaMemPoolPtrExportData *exportData, void *ptr); + +/** + * \brief Import a memory pool allocation from another process. + * + * Returns in \p ptr_out a pointer to the imported memory. + * The imported memory must not be accessed before the allocation operation completes + * in the exporting process. The imported memory must be freed from all importing processes before + * being freed in the exporting process. The pointer may be freed with cudaFree + * or cudaFreeAsync. If ::cudaFreeAsync is used, the free must be completed + * on the importing process before the free operation on the exporting process. + * + * \note The ::cudaFreeAsync api may be used in the exporting process before + * the ::cudaFreeAsync operation completes in its stream as long as the + * ::cudaFreeAsync in the exporting process specifies a stream with + * a stream dependency on the importing process's ::cudaFreeAsync. + * + * \param[out] ptr_out - pointer to imported memory + * \param[in] pool - pool from which to import + * \param[in] shareData - data specifying the memory to import + * + * \returns + * ::CUDA_SUCCESS, + * ::CUDA_ERROR_INVALID_VALUE, + * ::CUDA_ERROR_NOT_INITIALIZED, + * ::CUDA_ERROR_OUT_OF_MEMORY + * + * \sa ::cuMemPoolImportPointer, ::cudaMemPoolExportToShareableHandle, ::cudaMemPoolImportFromShareableHandle, ::cudaMemPoolExportPointer + */ +extern __host__ cudaError_t CUDARTAPI cudaMemPoolImportPointer(void **ptr, cudaMemPool_t memPool, struct cudaMemPoolPtrExportData *exportData); + +/** @} */ /* END CUDART_MEMORY_POOLS */ + /** * \defgroup CUDART_UNIFIED Unified Addressing * @@ -7021,9 +7982,6 @@ extern __CUDA_DEPRECATED __host__ cudaError_t CUDARTAPI cudaMemcpyFromArrayAsync * * Unified addressing is automatically enabled in 64-bit processes . * - * Unified addressing is not yet supported on Windows Vista or - * Windows 7 for devices that do not use the TCC driver model. - * * \section CUDART_UNIFIED_lookup Looking Up Information from Pointer Values * * It is possible to look up information about the memory which backs a @@ -7721,7 +8679,7 @@ extern __CUDA_DEPRECATED __host__ cudaError_t CUDARTAPI cudaBindTextureToArray(c * \ref ::cudaUnbindTexture(const struct textureReference*) "cudaUnbindTexture (C API)", * \ref ::cudaGetTextureAlignmentOffset(size_t*, const struct textureReference*) "cudaGetTextureAlignmentOffset (C API)", * ::cuTexRefSetMipmappedArray, - * ::cuTexRefSetMipmapFilterMode + * ::cuTexRefSetMipmapFilterMode, * ::cuTexRefSetMipmapLevelClamp, * ::cuTexRefSetMipmapLevelBias, * ::cuTexRefSetFormat, @@ -8769,7 +9727,8 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphAddMemcpyNode(cudaGraphNode_t *pG * ::cudaGraphAddHostNode, * ::cudaGraphAddMemsetNode */ -extern __host__ cudaError_t CUDARTAPI cudaGraphAddMemcpyNodeToSymbol( +#if __CUDART_API_VERSION >= 11010 + extern __host__ cudaError_t CUDARTAPI cudaGraphAddMemcpyNodeToSymbol( cudaGraphNode_t *pGraphNode, cudaGraph_t graph, const cudaGraphNode_t *pDependencies, @@ -8779,6 +9738,7 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphAddMemcpyNodeToSymbol( size_t count, size_t offset, enum cudaMemcpyKind kind); +#endif /** * \brief Creates a memcpy node to copy from a symbol on the device and adds it to a graph @@ -8836,7 +9796,8 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphAddMemcpyNodeToSymbol( * ::cudaGraphAddHostNode, * ::cudaGraphAddMemsetNode */ -extern __host__ cudaError_t CUDARTAPI cudaGraphAddMemcpyNodeFromSymbol( +#if __CUDART_API_VERSION >= 11010 + extern __host__ cudaError_t CUDARTAPI cudaGraphAddMemcpyNodeFromSymbol( cudaGraphNode_t* pGraphNode, cudaGraph_t graph, const cudaGraphNode_t* pDependencies, @@ -8846,6 +9807,7 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphAddMemcpyNodeFromSymbol( size_t count, size_t offset, enum cudaMemcpyKind kind); +#endif /** * \brief Creates a 1D memcpy node and adds it to a graph @@ -8902,7 +9864,8 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphAddMemcpyNodeFromSymbol( * ::cudaGraphAddHostNode, * ::cudaGraphAddMemsetNode */ -extern __host__ cudaError_t CUDARTAPI cudaGraphAddMemcpyNode1D( +#if __CUDART_API_VERSION >= 11010 + extern __host__ cudaError_t CUDARTAPI cudaGraphAddMemcpyNode1D( cudaGraphNode_t *pGraphNode, cudaGraph_t graph, const cudaGraphNode_t *pDependencies, @@ -8911,6 +9874,7 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphAddMemcpyNode1D( const void* src, size_t count, enum cudaMemcpyKind kind); +#endif /** * \brief Returns a memcpy node's parameters @@ -8997,13 +9961,15 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphMemcpyNodeSetParams(cudaGraphNode * ::cudaGraphAddMemcpyNode, * ::cudaGraphMemcpyNodeGetParams */ -extern __host__ cudaError_t CUDARTAPI cudaGraphMemcpyNodeSetParamsToSymbol( +#if __CUDART_API_VERSION >= 11010 + extern __host__ cudaError_t CUDARTAPI cudaGraphMemcpyNodeSetParamsToSymbol( cudaGraphNode_t node, const void* symbol, const void* src, size_t count, size_t offset, enum cudaMemcpyKind kind); +#endif /** * \brief Sets a memcpy node's parameters to copy from a symbol on the device @@ -9041,13 +10007,15 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphMemcpyNodeSetParamsToSymbol( * ::cudaGraphAddMemcpyNode, * ::cudaGraphMemcpyNodeGetParams */ -extern __host__ cudaError_t CUDARTAPI cudaGraphMemcpyNodeSetParamsFromSymbol( +#if __CUDART_API_VERSION >= 11010 + extern __host__ cudaError_t CUDARTAPI cudaGraphMemcpyNodeSetParamsFromSymbol( cudaGraphNode_t node, void* dst, const void* symbol, size_t count, size_t offset, enum cudaMemcpyKind kind); +#endif /** * \brief Sets a memcpy node's parameters to perform a 1-dimensional copy @@ -9085,12 +10053,14 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphMemcpyNodeSetParamsFromSymbol( * ::cudaGraphAddMemcpyNode, * ::cudaGraphMemcpyNodeGetParams */ -extern __host__ cudaError_t CUDARTAPI cudaGraphMemcpyNodeSetParams1D( +#if __CUDART_API_VERSION >= 11010 + extern __host__ cudaError_t CUDARTAPI cudaGraphMemcpyNodeSetParams1D( cudaGraphNode_t node, void* dst, const void* src, size_t count, enum cudaMemcpyKind kind); +#endif /** * \brief Creates a memset node and adds it to a graph @@ -9369,7 +10339,7 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphAddEmptyNode(cudaGraphNode_t *pGr * \brief Creates an event record node and adds it to a graph * * Creates a new event record node and adds it to \p hGraph with \p numDependencies - * dependencies specified via \p dependencies and arguments specified in \p params. + * dependencies specified via \p dependencies and event specified in \p event. * It is possible for \p numDependencies to be 0, in which case the node will be placed * at the root of the graph. \p dependencies may not have any duplicate entries. * A handle to the new node will be returned in \p phGraphNode. @@ -9406,8 +10376,10 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphAddEmptyNode(cudaGraphNode_t *pGr * ::cudaGraphAddMemcpyNode, * ::cudaGraphAddMemsetNode, */ -extern __host__ cudaError_t CUDARTAPI cudaGraphAddEventRecordNode(cudaGraphNode_t *pGraphNode, cudaGraph_t graph, const cudaGraphNode_t *pDependencies, size_t numDependencies, cudaEvent_t event); - +#if __CUDART_API_VERSION >= 11010 + extern __host__ cudaError_t CUDARTAPI cudaGraphAddEventRecordNode(cudaGraphNode_t *pGraphNode, cudaGraph_t graph, const cudaGraphNode_t *pDependencies, size_t numDependencies, cudaEvent_t event); +#endif + /** * \brief Returns the event associated with an event record node * @@ -9428,10 +10400,12 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphAddEventRecordNode(cudaGraphNode_ * ::cudaGraphAddEventRecordNode, * ::cudaGraphEventRecordNodeSetEvent, * ::cudaGraphEventWaitNodeGetEvent, - * ::cudaEventRecord, + * ::cudaEventRecordWithFlags, * ::cudaStreamWaitEvent */ -extern __host__ cudaError_t CUDARTAPI cudaGraphEventRecordNodeGetEvent(cudaGraphNode_t node, cudaEvent_t *event_out); +#if __CUDART_API_VERSION >= 11010 + extern __host__ cudaError_t CUDARTAPI cudaGraphEventRecordNodeGetEvent(cudaGraphNode_t node, cudaEvent_t *event_out); +#endif /** * \brief Sets an event record node's event @@ -9453,16 +10427,18 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphEventRecordNodeGetEvent(cudaGraph * ::cudaGraphAddEventRecordNode, * ::cudaGraphEventRecordNodeGetEvent, * ::cudaGraphEventWaitNodeSetEvent, - * ::cudaEventRecord, + * ::cudaEventRecordWithFlags, * ::cudaStreamWaitEvent */ -extern __host__ cudaError_t CUDARTAPI cudaGraphEventRecordNodeSetEvent(cudaGraphNode_t node, cudaEvent_t event); +#if __CUDART_API_VERSION >= 11010 + extern __host__ cudaError_t CUDARTAPI cudaGraphEventRecordNodeSetEvent(cudaGraphNode_t node, cudaEvent_t event); +#endif /** * \brief Creates an event wait node and adds it to a graph * * Creates a new event wait node and adds it to \p hGraph with \p numDependencies - * dependencies specified via \p dependencies and arguments specified in \p params. + * dependencies specified via \p dependencies and event specified in \p event. * It is possible for \p numDependencies to be 0, in which case the node will be placed * at the root of the graph. \p dependencies may not have any duplicate entries. * A handle to the new node will be returned in \p phGraphNode. @@ -9501,7 +10477,9 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphEventRecordNodeSetEvent(cudaGraph * ::cudaGraphAddMemcpyNode, * ::cudaGraphAddMemsetNode, */ -extern __host__ cudaError_t CUDARTAPI cudaGraphAddEventWaitNode(cudaGraphNode_t *pGraphNode, cudaGraph_t graph, const cudaGraphNode_t *pDependencies, size_t numDependencies, cudaEvent_t event); +#if __CUDART_API_VERSION >= 11010 + extern __host__ cudaError_t CUDARTAPI cudaGraphAddEventWaitNode(cudaGraphNode_t *pGraphNode, cudaGraph_t graph, const cudaGraphNode_t *pDependencies, size_t numDependencies, cudaEvent_t event); +#endif /** * \brief Returns the event associated with an event wait node @@ -9523,10 +10501,12 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphAddEventWaitNode(cudaGraphNode_t * ::cudaGraphAddEventWaitNode, * ::cudaGraphEventWaitNodeSetEvent, * ::cudaGraphEventRecordNodeGetEvent, - * ::cudaEventRecord, + * ::cudaEventRecordWithFlags, * ::cudaStreamWaitEvent */ -extern __host__ cudaError_t CUDARTAPI cudaGraphEventWaitNodeGetEvent(cudaGraphNode_t node, cudaEvent_t *event_out); +#if __CUDART_API_VERSION >= 11010 + extern __host__ cudaError_t CUDARTAPI cudaGraphEventWaitNodeGetEvent(cudaGraphNode_t node, cudaEvent_t *event_out); +#endif /** * \brief Sets an event wait node's event @@ -9548,10 +10528,232 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphEventWaitNodeGetEvent(cudaGraphNo * ::cudaGraphAddEventWaitNode, * ::cudaGraphEventWaitNodeGetEvent, * ::cudaGraphEventRecordNodeSetEvent, - * ::cudaEventRecord, + * ::cudaEventRecordWithFlags, * ::cudaStreamWaitEvent */ -extern __host__ cudaError_t CUDARTAPI cudaGraphEventWaitNodeSetEvent(cudaGraphNode_t node, cudaEvent_t event); +#if __CUDART_API_VERSION >= 11010 + extern __host__ cudaError_t CUDARTAPI cudaGraphEventWaitNodeSetEvent(cudaGraphNode_t node, cudaEvent_t event); +#endif + +/** + * \brief Creates an external semaphore signal node and adds it to a graph + * + * Creates a new external semaphore signal node and adds it to \p graph with \p + * numDependencies dependencies specified via \p dependencies and arguments specified + * in \p nodeParams. It is possible for \p numDependencies to be 0, in which case the + * node will be placed at the root of the graph. \p dependencies may not have any + * duplicate entries. A handle to the new node will be returned in \p pGraphNode. + * + * Performs a signal operation on a set of externally allocated semaphore objects + * when the node is launched. The operation(s) will occur after all of the node's + * dependencies have completed. + * + * \param pGraphNode - Returns newly created node + * \param graph - Graph to which to add the node + * \param pDependencies - Dependencies of the node + * \param numDependencies - Number of dependencies + * \param nodeParams - Parameters for the node + * + * \return + * ::CUDA_SUCCESS, + * ::CUDA_ERROR_DEINITIALIZED, + * ::CUDA_ERROR_NOT_INITIALIZED, + * ::CUDA_ERROR_NOT_SUPPORTED, + * ::CUDA_ERROR_INVALID_VALUE + * \note_graph_thread_safety + * \notefnerr + * + * \sa + * ::cudaGraphExternalSemaphoresSignalNodeGetParams, + * ::cudaGraphExternalSemaphoresSignalNodeSetParams, + * ::cudaGraphExecExternalSemaphoresSignalNodeSetParams, + * ::cudaGraphAddExternalSemaphoresWaitNode, + * ::cudaImportExternalSemaphore, + * ::cudaSignalExternalSemaphoresAsync, + * ::cudaWaitExternalSemaphoresAsync, + * ::cudaGraphCreate, + * ::cudaGraphDestroyNode, + * ::cudaGraphAddEventRecordNode, + * ::cudaGraphAddEventWaitNode, + * ::cudaGraphAddChildGraphNode, + * ::cudaGraphAddEmptyNode, + * ::cudaGraphAddKernelNode, + * ::cudaGraphAddMemcpyNode, + * ::cudaGraphAddMemsetNode, + */ +#if __CUDART_API_VERSION >= 11020 +extern __host__ cudaError_t CUDARTAPI cudaGraphAddExternalSemaphoresSignalNode(cudaGraphNode_t *pGraphNode, cudaGraph_t graph, const cudaGraphNode_t *pDependencies, size_t numDependencies, const struct cudaExternalSemaphoreSignalNodeParams *nodeParams); +#endif + +/** + * \brief Returns an external semaphore signal node's parameters + * + * Returns the parameters of an external semaphore signal node \p hNode in \p params_out. + * The \p extSemArray and \p paramsArray returned in \p params_out, + * are owned by the node. This memory remains valid until the node is destroyed or its + * parameters are modified, and should not be modified + * directly. Use ::cudaGraphExternalSemaphoresSignalNodeSetParams to update the + * parameters of this node. + * + * \param hNode - Node to get the parameters for + * \param params_out - Pointer to return the parameters + * + * \return + * ::CUDA_SUCCESS, + * ::CUDA_ERROR_DEINITIALIZED, + * ::CUDA_ERROR_NOT_INITIALIZED, + * ::CUDA_ERROR_INVALID_VALUE + * \note_graph_thread_safety + * \notefnerr + * + * \sa + * ::cudaLaunchKernel, + * ::cudaGraphAddExternalSemaphoresSignalNode, + * ::cudaGraphExternalSemaphoresSignalNodeSetParams, + * ::cudaGraphAddExternalSemaphoresWaitNode, + * ::cudaSignalExternalSemaphoresAsync, + * ::cudaWaitExternalSemaphoresAsync + */ +#if __CUDART_API_VERSION >= 11020 +extern __host__ cudaError_t CUDARTAPI cudaGraphExternalSemaphoresSignalNodeGetParams(cudaGraphNode_t hNode, struct cudaExternalSemaphoreSignalNodeParams *params_out); +#endif + +/** + * \brief Sets an external semaphore signal node's parameters + * + * Sets the parameters of an external semaphore signal node \p hNode to \p nodeParams. + * + * \param hNode - Node to set the parameters for + * \param nodeParams - Parameters to copy + * + * \return + * ::CUDA_SUCCESS, + * ::CUDA_ERROR_INVALID_VALUE, + * ::CUDA_ERROR_INVALID_HANDLE, + * ::CUDA_ERROR_OUT_OF_MEMORY + * \note_graph_thread_safety + * \notefnerr + * + * \sa + * ::cudaGraphAddExternalSemaphoresSignalNode, + * ::cudaGraphExternalSemaphoresSignalNodeSetParams, + * ::cudaGraphAddExternalSemaphoresWaitNode, + * ::cudaSignalExternalSemaphoresAsync, + * ::cudaWaitExternalSemaphoresAsync + */ +#if __CUDART_API_VERSION >= 11020 +extern __host__ cudaError_t CUDARTAPI cudaGraphExternalSemaphoresSignalNodeSetParams(cudaGraphNode_t hNode, const struct cudaExternalSemaphoreSignalNodeParams *nodeParams); +#endif + +/** + * \brief Creates an external semaphore wait node and adds it to a graph + * + * Creates a new external semaphore wait node and adds it to \p graph with \p numDependencies + * dependencies specified via \p dependencies and arguments specified in \p nodeParams. + * It is possible for \p numDependencies to be 0, in which case the node will be placed + * at the root of the graph. \p dependencies may not have any duplicate entries. A handle + * to the new node will be returned in \p pGraphNode. + * + * Performs a wait operation on a set of externally allocated semaphore objects + * when the node is launched. The node's dependencies will not be launched until + * the wait operation has completed. + * + * \param pGraphNode - Returns newly created node + * \param graph - Graph to which to add the node + * \param pDependencies - Dependencies of the node + * \param numDependencies - Number of dependencies + * \param nodeParams - Parameters for the node + * + * \return + * ::CUDA_SUCCESS, + * ::CUDA_ERROR_DEINITIALIZED, + * ::CUDA_ERROR_NOT_INITIALIZED, + * ::CUDA_ERROR_NOT_SUPPORTED, + * ::CUDA_ERROR_INVALID_VALUE + * \note_graph_thread_safety + * \notefnerr + * + * \sa + * ::cudaGraphExternalSemaphoresWaitNodeGetParams, + * ::cudaGraphExternalSemaphoresWaitNodeSetParams, + * ::cudaGraphExecExternalSemaphoresWaitNodeSetParams, + * ::cudaGraphAddExternalSemaphoresSignalNode, + * ::cudaImportExternalSemaphore, + * ::cudaSignalExternalSemaphoresAsync, + * ::cudaWaitExternalSemaphoresAsync, + * ::cudaGraphCreate, + * ::cudaGraphDestroyNode, + * ::cudaGraphAddEventRecordNode, + * ::cudaGraphAddEventWaitNode, + * ::cudaGraphAddChildGraphNode, + * ::cudaGraphAddEmptyNode, + * ::cudaGraphAddKernelNode, + * ::cudaGraphAddMemcpyNode, + * ::cudaGraphAddMemsetNode, + */ +#if __CUDART_API_VERSION >= 11020 +extern __host__ cudaError_t CUDARTAPI cudaGraphAddExternalSemaphoresWaitNode(cudaGraphNode_t *pGraphNode, cudaGraph_t graph, const cudaGraphNode_t *pDependencies, size_t numDependencies, const struct cudaExternalSemaphoreWaitNodeParams *nodeParams); +#endif + +/** + * \brief Returns an external semaphore wait node's parameters + * + * Returns the parameters of an external semaphore wait node \p hNode in \p params_out. + * The \p extSemArray and \p paramsArray returned in \p params_out, + * are owned by the node. This memory remains valid until the node is destroyed or its + * parameters are modified, and should not be modified + * directly. Use ::cudaGraphExternalSemaphoresSignalNodeSetParams to update the + * parameters of this node. + * + * \param hNode - Node to get the parameters for + * \param params_out - Pointer to return the parameters + * + * \return + * ::CUDA_SUCCESS, + * ::CUDA_ERROR_DEINITIALIZED, + * ::CUDA_ERROR_NOT_INITIALIZED, + * ::CUDA_ERROR_INVALID_VALUE + * \note_graph_thread_safety + * \notefnerr + * + * \sa + * ::cudaLaunchKernel, + * ::cudaGraphAddExternalSemaphoresWaitNode, + * ::cudaGraphExternalSemaphoresWaitNodeSetParams, + * ::cudaGraphAddExternalSemaphoresWaitNode, + * ::cudaSignalExternalSemaphoresAsync, + * ::cudaWaitExternalSemaphoresAsync + */ +#if __CUDART_API_VERSION >= 11020 +extern __host__ cudaError_t CUDARTAPI cudaGraphExternalSemaphoresWaitNodeGetParams(cudaGraphNode_t hNode, struct cudaExternalSemaphoreWaitNodeParams *params_out); +#endif + +/** + * \brief Sets an external semaphore wait node's parameters + * + * Sets the parameters of an external semaphore wait node \p hNode to \p nodeParams. + * + * \param hNode - Node to set the parameters for + * \param nodeParams - Parameters to copy + * + * \return + * ::CUDA_SUCCESS, + * ::CUDA_ERROR_INVALID_VALUE, + * ::CUDA_ERROR_INVALID_HANDLE, + * ::CUDA_ERROR_OUT_OF_MEMORY + * \note_graph_thread_safety + * \notefnerr + * + * \sa + * ::cudaGraphAddExternalSemaphoresWaitNode, + * ::cudaGraphExternalSemaphoresWaitNodeSetParams, + * ::cudaGraphAddExternalSemaphoresWaitNode, + * ::cudaSignalExternalSemaphoresAsync, + * ::cudaWaitExternalSemaphoresAsync + */ +#if __CUDART_API_VERSION >= 11020 +extern __host__ cudaError_t CUDARTAPI cudaGraphExternalSemaphoresWaitNodeSetParams(cudaGraphNode_t hNode, const struct cudaExternalSemaphoreWaitNodeParams *nodeParams); +#endif /** * \brief Clones a graph @@ -9954,6 +11156,15 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphInstantiate(cudaGraphExec_t *pGra * \sa * ::cudaGraphAddKernelNode, * ::cudaGraphKernelNodeSetParams, + * ::cudaGraphExecMemcpyNodeSetParams, + * ::cudaGraphExecMemsetNodeSetParams, + * ::cudaGraphExecHostNodeSetParams, + * ::cudaGraphExecChildGraphNodeSetParams, + * ::cudaGraphExecEventRecordNodeSetEvent, + * ::cudaGraphExecEventWaitNodeSetEvent, + * ::cudaGraphExecExternalSemaphoresSignalNodeSetParams, + * ::cudaGraphExecExternalSemaphoresWaitNodeSetParams, + * ::cudaGraphExecUpdate, * ::cudaGraphInstantiate */ extern __host__ cudaError_t CUDARTAPI cudaGraphExecKernelNodeSetParams(cudaGraphExec_t hGraphExec, cudaGraphNode_t node, const struct cudaKernelNodeParams *pNodeParams); @@ -9992,13 +11203,19 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphExecKernelNodeSetParams(cudaGraph * \sa * ::cudaGraphAddMemcpyNode, * ::cudaGraphMemcpyNodeSetParams, - * ::cudaGraphInstantiate, * ::cudaGraphExecMemcpyNodeSetParamsToSymbol, * ::cudaGraphExecMemcpyNodeSetParamsFromSymbol, * ::cudaGraphExecMemcpyNodeSetParams1D, * ::cudaGraphExecKernelNodeSetParams, * ::cudaGraphExecMemsetNodeSetParams, - * ::cudaGraphExecHostNodeSetParams + * ::cudaGraphExecHostNodeSetParams, + * ::cudaGraphExecChildGraphNodeSetParams, + * ::cudaGraphExecEventRecordNodeSetEvent, + * ::cudaGraphExecEventWaitNodeSetEvent, + * ::cudaGraphExecExternalSemaphoresSignalNodeSetParams, + * ::cudaGraphExecExternalSemaphoresWaitNodeSetParams, + * ::cudaGraphExecUpdate, + * ::cudaGraphInstantiate */ extern __host__ cudaError_t CUDARTAPI cudaGraphExecMemcpyNodeSetParams(cudaGraphExec_t hGraphExec, cudaGraphNode_t node, const struct cudaMemcpy3DParms *pNodeParams); @@ -10041,14 +11258,21 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphExecMemcpyNodeSetParams(cudaGraph * ::cudaGraphAddMemcpyNodeToSymbol, * ::cudaGraphMemcpyNodeSetParams, * ::cudaGraphMemcpyNodeSetParamsToSymbol, - * ::cudaGraphInstantiate, * ::cudaGraphExecMemcpyNodeSetParams, * ::cudaGraphExecMemcpyNodeSetParamsFromSymbol, * ::cudaGraphExecKernelNodeSetParams, * ::cudaGraphExecMemsetNodeSetParams, - * ::cudaGraphExecHostNodeSetParams + * ::cudaGraphExecHostNodeSetParams, + * ::cudaGraphExecChildGraphNodeSetParams, + * ::cudaGraphExecEventRecordNodeSetEvent, + * ::cudaGraphExecEventWaitNodeSetEvent, + * ::cudaGraphExecExternalSemaphoresSignalNodeSetParams, + * ::cudaGraphExecExternalSemaphoresWaitNodeSetParams, + * ::cudaGraphExecUpdate, + * ::cudaGraphInstantiate */ -extern __host__ cudaError_t CUDARTAPI cudaGraphExecMemcpyNodeSetParamsToSymbol( +#if __CUDART_API_VERSION >= 11010 + extern __host__ cudaError_t CUDARTAPI cudaGraphExecMemcpyNodeSetParamsToSymbol( cudaGraphExec_t hGraphExec, cudaGraphNode_t node, const void* symbol, @@ -10056,6 +11280,7 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphExecMemcpyNodeSetParamsToSymbol( size_t count, size_t offset, enum cudaMemcpyKind kind); +#endif /** * \brief Sets the parameters for a memcpy node in the given graphExec to copy from a symbol on the device @@ -10096,14 +11321,21 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphExecMemcpyNodeSetParamsToSymbol( * ::cudaGraphAddMemcpyNodeFromSymbol, * ::cudaGraphMemcpyNodeSetParams, * ::cudaGraphMemcpyNodeSetParamsFromSymbol, - * ::cudaGraphInstantiate, * ::cudaGraphExecMemcpyNodeSetParams, * ::cudaGraphExecMemcpyNodeSetParamsToSymbol, * ::cudaGraphExecKernelNodeSetParams, * ::cudaGraphExecMemsetNodeSetParams, - * ::cudaGraphExecHostNodeSetParams + * ::cudaGraphExecHostNodeSetParams, + * ::cudaGraphExecChildGraphNodeSetParams, + * ::cudaGraphExecEventRecordNodeSetEvent, + * ::cudaGraphExecEventWaitNodeSetEvent, + * ::cudaGraphExecExternalSemaphoresSignalNodeSetParams, + * ::cudaGraphExecExternalSemaphoresWaitNodeSetParams, + * ::cudaGraphExecUpdate, + * ::cudaGraphInstantiate */ -extern __host__ cudaError_t CUDARTAPI cudaGraphExecMemcpyNodeSetParamsFromSymbol( +#if __CUDART_API_VERSION >= 11010 + extern __host__ cudaError_t CUDARTAPI cudaGraphExecMemcpyNodeSetParamsFromSymbol( cudaGraphExec_t hGraphExec, cudaGraphNode_t node, void* dst, @@ -10111,6 +11343,7 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphExecMemcpyNodeSetParamsFromSymbol size_t count, size_t offset, enum cudaMemcpyKind kind); +#endif /** * \brief Sets the parameters for a memcpy node in the given graphExec to perform a 1-dimensional copy @@ -10150,19 +11383,27 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphExecMemcpyNodeSetParamsFromSymbol * ::cudaGraphAddMemcpyNode1D, * ::cudaGraphMemcpyNodeSetParams, * ::cudaGraphMemcpyNodeSetParams1D, - * ::cudaGraphInstantiate, * ::cudaGraphExecMemcpyNodeSetParams, * ::cudaGraphExecKernelNodeSetParams, * ::cudaGraphExecMemsetNodeSetParams, - * ::cudaGraphExecHostNodeSetParams + * ::cudaGraphExecHostNodeSetParams, + * ::cudaGraphExecChildGraphNodeSetParams, + * ::cudaGraphExecEventRecordNodeSetEvent, + * ::cudaGraphExecEventWaitNodeSetEvent, + * ::cudaGraphExecExternalSemaphoresSignalNodeSetParams, + * ::cudaGraphExecExternalSemaphoresWaitNodeSetParams, + * ::cudaGraphExecUpdate, + * ::cudaGraphInstantiate */ -extern __host__ cudaError_t CUDARTAPI cudaGraphExecMemcpyNodeSetParams1D( +#if __CUDART_API_VERSION >= 11010 + extern __host__ cudaError_t CUDARTAPI cudaGraphExecMemcpyNodeSetParams1D( cudaGraphExec_t hGraphExec, cudaGraphNode_t node, void* dst, const void* src, size_t count, enum cudaMemcpyKind kind); +#endif /** * \brief Sets the parameters for a memset node in the given graphExec. @@ -10197,11 +11438,17 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphExecMemcpyNodeSetParams1D( * * \sa * ::cudaGraphAddMemsetNode, - * ::cudaGraphMemsetNodeSetParams + * ::cudaGraphMemsetNodeSetParams, + * ::cudaGraphExecKernelNodeSetParams, + * ::cudaGraphExecMemcpyNodeSetParams, + * ::cudaGraphExecHostNodeSetParams, + * ::cudaGraphExecChildGraphNodeSetParams, + * ::cudaGraphExecEventRecordNodeSetEvent, + * ::cudaGraphExecEventWaitNodeSetEvent, + * ::cudaGraphExecExternalSemaphoresSignalNodeSetParams, + * ::cudaGraphExecExternalSemaphoresWaitNodeSetParams, + * ::cudaGraphExecUpdate, * ::cudaGraphInstantiate - * ::cudaGraphExecKernelNodeSetParams - * ::cudaGraphExecMemcpyNodeSetParams - * ::cudaGraphExecHostNodeSetParams */ extern __host__ cudaError_t CUDARTAPI cudaGraphExecMemsetNodeSetParams(cudaGraphExec_t hGraphExec, cudaGraphNode_t node, const struct cudaMemsetParams *pNodeParams); @@ -10230,11 +11477,17 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphExecMemsetNodeSetParams(cudaGraph * * \sa * ::cudaGraphAddHostNode, - * ::cudaGraphHostNodeSetParams + * ::cudaGraphHostNodeSetParams, + * ::cudaGraphExecKernelNodeSetParams, + * ::cudaGraphExecMemcpyNodeSetParams, + * ::cudaGraphExecMemsetNodeSetParams, + * ::cudaGraphExecChildGraphNodeSetParams, + * ::cudaGraphExecEventRecordNodeSetEvent, + * ::cudaGraphExecEventWaitNodeSetEvent, + * ::cudaGraphExecExternalSemaphoresSignalNodeSetParams, + * ::cudaGraphExecExternalSemaphoresWaitNodeSetParams, + * ::cudaGraphExecUpdate, * ::cudaGraphInstantiate - * ::cudaGraphExecKernelNodeSetParams - * ::cudaGraphExecMemcpyNodeSetParams - * ::cudaGraphExecMemsetNodeSetParams */ extern __host__ cudaError_t CUDARTAPI cudaGraphExecHostNodeSetParams(cudaGraphExec_t hGraphExec, cudaGraphNode_t node, const struct cudaHostNodeParams *pNodeParams); @@ -10270,14 +11523,20 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphExecHostNodeSetParams(cudaGraphEx * \sa * ::cudaGraphAddChildGraphNode, * ::cudaGraphChildGraphNodeGetGraph, - * ::cudaGraphInstantiate, - * ::cudaGraphExecUpdate, * ::cudaGraphExecKernelNodeSetParams, * ::cudaGraphExecMemcpyNodeSetParams, * ::cudaGraphExecMemsetNodeSetParams, - * ::cudaGraphExecHostNodeSetParams + * ::cudaGraphExecHostNodeSetParams, + * ::cudaGraphExecEventRecordNodeSetEvent, + * ::cudaGraphExecEventWaitNodeSetEvent, + * ::cudaGraphExecExternalSemaphoresSignalNodeSetParams, + * ::cudaGraphExecExternalSemaphoresWaitNodeSetParams, + * ::cudaGraphExecUpdate, + * ::cudaGraphInstantiate */ -extern __host__ cudaError_t CUDARTAPI cudaGraphExecChildGraphNodeSetParams(cudaGraphExec_t hGraphExec, cudaGraphNode_t node, cudaGraph_t childGraph); +#if __CUDART_API_VERSION >= 11010 + extern __host__ cudaError_t CUDARTAPI cudaGraphExecChildGraphNodeSetParams(cudaGraphExec_t hGraphExec, cudaGraphNode_t node, cudaGraph_t childGraph); +#endif /** * \brief Sets the event for an event record node in the given graphExec @@ -10304,19 +11563,27 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphExecChildGraphNodeSetParams(cudaG * ::cudaGraphAddEventRecordNode, * ::cudaGraphEventRecordNodeGetEvent, * ::cudaGraphEventWaitNodeSetEvent, - * ::cudaEventRecord, - * ::cudaStreamWaitEvent - * ::cudaGraphCreate, - * ::cudaGraphDestroyNode, + * ::cudaEventRecordWithFlags, + * ::cudaStreamWaitEvent, + * ::cudaGraphExecKernelNodeSetParams, + * ::cudaGraphExecMemcpyNodeSetParams, + * ::cudaGraphExecMemsetNodeSetParams, + * ::cudaGraphExecHostNodeSetParams, + * ::cudaGraphExecChildGraphNodeSetParams, + * ::cudaGraphExecEventWaitNodeSetEvent, + * ::cudaGraphExecExternalSemaphoresSignalNodeSetParams, + * ::cudaGraphExecExternalSemaphoresWaitNodeSetParams, + * ::cudaGraphExecUpdate, * ::cudaGraphInstantiate */ -extern __host__ cudaError_t CUDARTAPI cudaGraphExecEventRecordNodeSetEvent(cudaGraphExec_t hGraphExec, cudaGraphNode_t hNode, cudaEvent_t event); - +#if __CUDART_API_VERSION >= 11010 + extern __host__ cudaError_t CUDARTAPI cudaGraphExecEventRecordNodeSetEvent(cudaGraphExec_t hGraphExec, cudaGraphNode_t hNode, cudaEvent_t event); +#endif /** - * \brief Sets the event for an event record node in the given graphExec + * \brief Sets the event for an event wait node in the given graphExec * - * Sets the event of an event record node in an executable graph \p hGraphExec. + * Sets the event of an event wait node in an executable graph \p hGraphExec. * The node is identified by the corresponding node \p hNode in the * non-executable graph, from which the executable graph was instantiated. * @@ -10338,13 +11605,112 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphExecEventRecordNodeSetEvent(cudaG * ::cudaGraphAddEventWaitNode, * ::cudaGraphEventWaitNodeGetEvent, * ::cudaGraphEventRecordNodeSetEvent, - * ::cudaEventRecord, - * ::cudaStreamWaitEvent - * ::cudaGraphCreate, - * ::cudaGraphDestroyNode, + * ::cudaEventRecordWithFlags, + * ::cudaStreamWaitEvent, + * ::cudaGraphExecKernelNodeSetParams, + * ::cudaGraphExecMemcpyNodeSetParams, + * ::cudaGraphExecMemsetNodeSetParams, + * ::cudaGraphExecHostNodeSetParams, + * ::cudaGraphExecChildGraphNodeSetParams, + * ::cudaGraphExecEventRecordNodeSetEvent, + * ::cudaGraphExecExternalSemaphoresSignalNodeSetParams, + * ::cudaGraphExecExternalSemaphoresWaitNodeSetParams, + * ::cudaGraphExecUpdate, * ::cudaGraphInstantiate */ -extern __host__ cudaError_t CUDARTAPI cudaGraphExecEventWaitNodeSetEvent(cudaGraphExec_t hGraphExec, cudaGraphNode_t hNode, cudaEvent_t event); +#if __CUDART_API_VERSION >= 11010 + extern __host__ cudaError_t CUDARTAPI cudaGraphExecEventWaitNodeSetEvent(cudaGraphExec_t hGraphExec, cudaGraphNode_t hNode, cudaEvent_t event); +#endif + +/** + * \brief Sets the parameters for an external semaphore signal node in the given graphExec + * + * Sets the parameters of an external semaphore signal node in an executable graph \p hGraphExec. + * The node is identified by the corresponding node \p hNode in the + * non-executable graph, from which the executable graph was instantiated. + * + * \p hNode must not have been removed from the original graph. + * + * The modifications only affect future launches of \p hGraphExec. Already + * enqueued or running launches of \p hGraphExec are not affected by this call. + * \p hNode is also not modified by this call. + * + * Changing \p nodeParams->numExtSems is not supported. + * + * \param hGraphExec - The executable graph in which to set the specified node + * \param hNode - semaphore signal node from the graph from which graphExec was instantiated + * \param nodeParams - Updated Parameters to set + * + * \return + * ::CUDA_SUCCESS, + * ::CUDA_ERROR_INVALID_VALUE, + * \note_graph_thread_safety + * \notefnerr + * + * \sa + * ::cudaGraphAddExternalSemaphoresSignalNode, + * ::cudaImportExternalSemaphore, + * ::cudaSignalExternalSemaphoresAsync, + * ::cudaWaitExternalSemaphoresAsync, + * ::cudaGraphExecKernelNodeSetParams, + * ::cudaGraphExecMemcpyNodeSetParams, + * ::cudaGraphExecMemsetNodeSetParams, + * ::cudaGraphExecHostNodeSetParams, + * ::cudaGraphExecChildGraphNodeSetParams, + * ::cudaGraphExecEventRecordNodeSetEvent, + * ::cudaGraphExecEventWaitNodeSetEvent, + * ::cudaGraphExecExternalSemaphoresWaitNodeSetParams, + * ::cudaGraphExecUpdate, + * ::cudaGraphInstantiate + */ +#if __CUDART_API_VERSION >= 11020 +extern __host__ cudaError_t CUDARTAPI cudaGraphExecExternalSemaphoresSignalNodeSetParams(cudaGraphExec_t hGraphExec, cudaGraphNode_t hNode, const struct cudaExternalSemaphoreSignalNodeParams *nodeParams); +#endif + +/** + * \brief Sets the parameters for an external semaphore wait node in the given graphExec + * + * Sets the parameters of an external semaphore wait node in an executable graph \p hGraphExec. + * The node is identified by the corresponding node \p hNode in the + * non-executable graph, from which the executable graph was instantiated. + * + * \p hNode must not have been removed from the original graph. + * + * The modifications only affect future launches of \p hGraphExec. Already + * enqueued or running launches of \p hGraphExec are not affected by this call. + * \p hNode is also not modified by this call. + * + * Changing \p nodeParams->numExtSems is not supported. + * + * \param hGraphExec - The executable graph in which to set the specified node + * \param hNode - semaphore wait node from the graph from which graphExec was instantiated + * \param nodeParams - Updated Parameters to set + * + * \return + * ::CUDA_SUCCESS, + * ::CUDA_ERROR_INVALID_VALUE, + * \note_graph_thread_safety + * \notefnerr + * + * \sa + * ::cudaGraphAddExternalSemaphoresWaitNode, + * ::cudaImportExternalSemaphore, + * ::cudaSignalExternalSemaphoresAsync, + * ::cudaWaitExternalSemaphoresAsync, + * ::cudaGraphExecKernelNodeSetParams, + * ::cudaGraphExecMemcpyNodeSetParams, + * ::cudaGraphExecMemsetNodeSetParams, + * ::cudaGraphExecHostNodeSetParams, + * ::cudaGraphExecChildGraphNodeSetParams, + * ::cudaGraphExecEventRecordNodeSetEvent, + * ::cudaGraphExecEventWaitNodeSetEvent, + * ::cudaGraphExecExternalSemaphoresSignalNodeSetParams, + * ::cudaGraphExecUpdate, + * ::cudaGraphInstantiate + */ +#if __CUDART_API_VERSION >= 11020 +extern __host__ cudaError_t CUDARTAPI cudaGraphExecExternalSemaphoresWaitNodeSetParams(cudaGraphExec_t hGraphExec, cudaGraphNode_t hNode, const struct cudaExternalSemaphoreWaitNodeParams *nodeParams); +#endif /** * \brief Check whether an executable graph can be updated with a graph and perform the update if possible @@ -10355,7 +11721,9 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphExecEventWaitNodeSetEvent(cudaGra * Limitations: * * - Kernel nodes: - * - The function must not change (same restriction as cudaGraphExecKernelNodeSetParams()) + * - The owning context of the function cannot change. + * - A node whose function originally did not use CUDA dynamic parallelism cannot be updated + * to a function which uses CDP * - Memset and memcpy nodes: * - The CUDA device(s) to which the operand(s) was allocated/mapped cannot change. * - The source/destination memory must be allocated from the same contexts as the original @@ -10367,10 +11735,6 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphExecEventWaitNodeSetEvent(cudaGra * * Note: The API may add further restrictions in future releases. The return code should always be checked. * - * Some node types are not currently supported: - * - Empty graph nodes(cudaGraphNodeTypeEmpty) - * - Child graphs(cudaGraphNodeTypeGraph). - * * cudaGraphExecUpdate sets \p updateResult_out to cudaGraphExecUpdateErrorTopologyChanged under * the following conditions: * @@ -10387,8 +11751,9 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphExecEventWaitNodeSetEvent(cudaGra * - cudaGraphExecUpdateErrorTopologyChanged if the graph topology changed * - cudaGraphExecUpdateErrorNodeTypeChanged if the type of a node changed, in which case * \p hErrorNode_out is set to the node from \p hGraph. - * - cudaGraphExecUpdateErrorFunctionChanged if the func field of a kernel changed, in which - * case \p hErrorNode_out is set to the node from \p hGraph + * - cudaGraphExecUpdateErrorFunctionChanged if the function of a kernel node changed (CUDA driver < 11.2) + * - cudaGraphExecUpdateErrorUnsupportedFunctionChange if the func field of a kernel changed in an + * unsupported way(see note above), in which case \p hErrorNode_out is set to the node from \p hGraph * - cudaGraphExecUpdateErrorParametersChanged if any parameters to a node changed in a way * that is not supported, in which case \p hErrorNode_out is set to the node from \p hGraph * - cudaGraphExecUpdateErrorNotSupported if something about a node is unsupported, like @@ -10442,7 +11807,9 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphExecUpdate(cudaGraphExec_t hGraph * ::cudaGraphLaunch, * ::cudaGraphExecDestroy */ -extern __host__ cudaError_t CUDARTAPI cudaGraphUpload(cudaGraphExec_t graphExec, cudaStream_t stream); +#if __CUDART_API_VERSION >= 11010 + extern __host__ cudaError_t CUDARTAPI cudaGraphUpload(cudaGraphExec_t graphExec, cudaStream_t stream); +#endif /** * \brief Launches an executable graph in a stream @@ -10514,8 +11881,234 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphExecDestroy(cudaGraphExec_t graph */ extern __host__ cudaError_t CUDARTAPI cudaGraphDestroy(cudaGraph_t graph); +/** + * \brief Write a DOT file describing graph structure + * + * Using the provided \p graph, write to \p path a DOT formatted description of the graph. + * By default this includes the graph topology, node types, node id, kernel names and memcpy direction. + * \p flags can be specified to write more detailed information about each node type such as + * parameter values, kernel attributes, node and function handles. + * + * \param graph - The graph to create a DOT file from + * \param path - The path to write the DOT file to + * \param flags - Flags from cudaGraphDebugDotFlags for specifying which additional node information to write + * + * \return + * ::cudaSuccess, + * ::cudaErrorInvalidValue, + * ::cudaErrorOperatingSystem + */ +extern __host__ cudaError_t CUDARTAPI cudaGraphDebugDotPrint(cudaGraph_t graph, const char *path, unsigned int flags); + +/** + * \brief Create a user object + * + * Create a user object with the specified destructor callback and initial reference count. The + * initial references are owned by the caller. + * + * Destructor callbacks cannot make CUDA API calls and should avoid blocking behavior, as they + * are executed by a shared internal thread. Another thread may be signaled to perform such + * actions, if it does not block forward progress of tasks scheduled through CUDA. + * + * See CUDA User Objects in the CUDA C++ Programming Guide for more information on user objects. + * + * \param object_out - Location to return the user object handle + * \param ptr - The pointer to pass to the destroy function + * \param destroy - Callback to free the user object when it is no longer in use + * \param initialRefcount - The initial refcount to create the object with, typically 1. The + * initial references are owned by the calling thread. + * \param flags - Currently it is required to pass ::cudaUserObjectNoDestructorSync, + * which is the only defined flag. This indicates that the destroy + * callback cannot be waited on by any CUDA API. Users requiring + * synchronization of the callback should signal its completion + * manually. + * + * \return + * ::cudaSuccess, + * ::cudaErrorInvalidValue + * + * \sa + * ::cudaUserObjectRetain, + * ::cudaUserObjectRelease, + * ::cudaGraphRetainUserObject, + * ::cudaGraphReleaseUserObject, + * ::cudaGraphCreate + */ +extern __host__ cudaError_t CUDARTAPI cudaUserObjectCreate(cudaUserObject_t *object_out, void *ptr, cudaHostFn_t destroy, unsigned int initialRefcount, unsigned int flags); + +/** + * \brief Retain a reference to a user object + * + * Retains new references to a user object. The new references are owned by the caller. + * + * See CUDA User Objects in the CUDA C++ Programming Guide for more information on user objects. + * + * \param object - The object to retain + * \param count - The number of references to retain, typically 1. Must be nonzero + * and not larger than INT_MAX. + * + * \return + * ::cudaSuccess, + * ::cudaErrorInvalidValue + * + * \sa + * ::cudaUserObjectCreate, + * ::cudaUserObjectRelease, + * ::cudaGraphRetainUserObject, + * ::cudaGraphReleaseUserObject, + * ::cudaGraphCreate + */ +extern __host__ cudaError_t CUDARTAPI cudaUserObjectRetain(cudaUserObject_t object, unsigned int count __dv(1)); + +/** + * \brief Release a reference to a user object + * + * Releases user object references owned by the caller. The object's destructor is invoked if + * the reference count reaches zero. + * + * It is undefined behavior to release references not owned by the caller, or to use a user + * object handle after all references are released. + * + * See CUDA User Objects in the CUDA C++ Programming Guide for more information on user objects. + * + * \param object - The object to release + * \param count - The number of references to release, typically 1. Must be nonzero + * and not larger than INT_MAX. + * + * \return + * ::cudaSuccess, + * ::cudaErrorInvalidValue + * + * \sa + * ::cudaUserObjectCreate, + * ::cudaUserObjectRetain, + * ::cudaGraphRetainUserObject, + * ::cudaGraphReleaseUserObject, + * ::cudaGraphCreate + */ +extern __host__ cudaError_t CUDARTAPI cudaUserObjectRelease(cudaUserObject_t object, unsigned int count __dv(1)); + +/** + * \brief Retain a reference to a user object from a graph + * + * Creates or moves user object references that will be owned by a CUDA graph. + * + * See CUDA User Objects in the CUDA C++ Programming Guide for more information on user objects. + * + * \param graph - The graph to associate the reference with + * \param object - The user object to retain a reference for + * \param count - The number of references to add to the graph, typically 1. Must be + * nonzero and not larger than INT_MAX. + * \param flags - The optional flag ::cudaGraphUserObjectMove transfers references + * from the calling thread, rather than create new references. Pass 0 + * to create new references. + * + * \return + * ::cudaSuccess, + * ::cudaErrorInvalidValue + * + * \sa + * ::cudaUserObjectCreate + * ::cudaUserObjectRetain, + * ::cudaUserObjectRelease, + * ::cudaGraphReleaseUserObject, + * ::cudaGraphCreate + */ +extern __host__ cudaError_t CUDARTAPI cudaGraphRetainUserObject(cudaGraph_t graph, cudaUserObject_t object, unsigned int count __dv(1), unsigned int flags __dv(0)); + +/** + * \brief Release a user object reference from a graph + * + * Releases user object references owned by a graph. + * + * See CUDA User Objects in the CUDA C++ Programming Guide for more information on user objects. + * + * \param graph - The graph that will release the reference + * \param object - The user object to release a reference for + * \param count - The number of references to release, typically 1. Must be nonzero + * and not larger than INT_MAX. + * + * \return + * ::cudaSuccess, + * ::cudaErrorInvalidValue + * + * \sa + * ::cudaUserObjectCreate + * ::cudaUserObjectRetain, + * ::cudaUserObjectRelease, + * ::cudaGraphRetainUserObject, + * ::cudaGraphCreate + */ +extern __host__ cudaError_t CUDARTAPI cudaGraphReleaseUserObject(cudaGraph_t graph, cudaUserObject_t object, unsigned int count __dv(1)); + /** @} */ /* END CUDART_GRAPH */ +/** + * \defgroup CUDART_DRIVER_ENTRY_POINT Driver Entry Point Access + * + * ___MANBRIEF___ driver entry point access functions of the CUDA runtime API + * (___CURRENT_FILE___) ___ENDMANBRIEF___ + * + * This section describes the driver entry point access functions of CUDA + * runtime application programming interface. + * + * @{ + */ + +/** + * \brief Returns the requested driver API function pointer + * + * Returns in \p **funcPtr the address of the CUDA driver function for the requested flags. + * + * For a requested driver symbol, if the CUDA version in which the driver symbol was + * introduced is less than or equal to the CUDA runtime version, the API will return + * the function pointer to the corresponding versioned driver function. + * + * The pointer returned by the API should be cast to a function pointer matching the + * requested driver function's definition in the API header file. The function pointer + * typedef can be picked up from the corresponding typedefs header file. For example, + * cudaTypedefs.h consists of function pointer typedefs for driver APIs defined in cuda.h. + * + * The API will return ::cudaErrorSymbolNotFound if the requested driver function is not + * supported on the platform, no ABI compatible driver function exists for the CUDA runtime + * version or if the driver symbol is invalid. + * + * The requested flags can be: + * - ::cudaEnableDefault: This is the default mode. This is equivalent to + * ::cudaEnablePerThreadDefaultStream if the code is compiled with + * --default-stream per-thread compilation flag or the macro CUDA_API_PER_THREAD_DEFAULT_STREAM + * is defined; ::cudaEnableLegacyStream otherwise. + * - ::cudaEnableLegacyStream: This will enable the search for all driver symbols + * that match the requested driver symbol name except the corresponding per-thread versions. + * - ::cudaEnablePerThreadDefaultStream: This will enable the search for all + * driver symbols that match the requested driver symbol name including the per-thread + * versions. If a per-thread version is not found, the API will return the legacy version + * of the driver function. + * + * \param symbol - The base name of the driver API function to look for. As an example, + * for the driver API ::cuMemAlloc_v2, \p symbol would be cuMemAlloc. + * Note that the API will use the CUDA runtime version to return the + * address to the most recent ABI compatible driver symbol, ::cuMemAlloc + * or ::cuMemAlloc_v2. + * \param funcPtr - Location to return the function pointer to the requested driver function + * \param flags - Flags to specify search options. + * + * \return + * ::cudaSuccess, + * ::cudaErrorInvalidValue, + * ::cudaErrorNotSupported, + * ::cudaErrorSymbolNotFound + * \note_version_mixing + * \note_init_rt + * \note_callback + * + * \sa + * ::cuGetProcAddress + */ +extern __host__ cudaError_t CUDARTAPI cudaGetDriverEntryPoint(const char *symbol, void **funcPtr, unsigned long long flags); + +/** @} */ /* END CUDART_DRIVER_ENTRY_POINT */ + /** \cond impl_private */ extern __host__ cudaError_t CUDARTAPI cudaGetExportTable(const void **ppExportTable, const cudaUUID_t *pExportTableId); /** \endcond impl_private */ @@ -10585,9 +12178,9 @@ extern __host__ cudaError_t CUDARTAPI cudaGetExportTable(const void **ppExportTa * Runtime API call on any thread which requires an active context will trigger the * reinitialization of that device's primary context. * - * Note that there is no reference counting of the primary context's lifetime. It is - * recommended that the primary context not be deinitialized except just before exit - * or to recover from an unspecified launch failure. + * Note that primary contexts are shared resources. It is recommended that + * the primary context not be reset except just before exit or to recover from an + * unspecified launch failure. * * \section CUDART_CUDA_context Context Interoperability * @@ -10693,7 +12286,7 @@ extern __host__ cudaError_t CUDARTAPI cudaGetExportTable(const void **ppExportTa * ::cudaSuccess * */ -extern __host__ cudaError_t cudaGetFuncBySymbol(cudaFunction_t* functionPtr, const void* symbolPtr); +extern __host__ cudaError_t CUDARTAPI_CDECL cudaGetFuncBySymbol(cudaFunction_t* functionPtr, const void* symbolPtr); /** @} */ /* END CUDART_DRIVER */ @@ -10747,9 +12340,15 @@ extern __host__ cudaError_t cudaGetFuncBySymbol(cudaFunction_t* functionPtr, con #undef cudaStreamEndCapture #undef cudaStreamIsCapturing #undef cudaStreamGetCaptureInfo + #undef cudaStreamGetCaptureInfo_v2 #undef cudaStreamCopyAttributes #undef cudaStreamGetAttribute #undef cudaStreamSetAttribute + #undef cudaMallocAsync + #undef cudaFreeAsync + #undef cudaMallocFromPoolAsync + #undef cudaGetDriverEntryPoint + extern __host__ cudaError_t CUDARTAPI cudaMemcpy(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind); extern __host__ cudaError_t CUDARTAPI cudaMemcpyToSymbol(const void *symbol, const void *src, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyHostToDevice)); extern __host__ cudaError_t CUDARTAPI cudaMemcpyFromSymbol(void *dst, const void *symbol, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToHost)); @@ -10791,18 +12390,29 @@ extern __host__ cudaError_t cudaGetFuncBySymbol(cudaFunction_t* functionPtr, con extern __host__ cudaError_t CUDARTAPI cudaLaunchCooperativeKernel(const void *func, dim3 gridDim, dim3 blockDim, void **args, size_t sharedMem, cudaStream_t stream); extern __host__ cudaError_t CUDARTAPI cudaLaunchHostFunc(cudaStream_t stream, cudaHostFn_t fn, void *userData); extern __host__ cudaError_t CUDARTAPI cudaMemPrefetchAsync(const void *devPtr, size_t count, int dstDevice, cudaStream_t stream); - extern __host__ cudaError_t CUDARTAPI cudaSignalExternalSemaphoresAsync(const cudaExternalSemaphore_t *extSemArray, const struct cudaExternalSemaphoreSignalParams *paramsArray, unsigned int numExtSems, cudaStream_t stream __dv(0)); - extern __host__ cudaError_t CUDARTAPI cudaWaitExternalSemaphoresAsync(const cudaExternalSemaphore_t *extSemArray, const struct cudaExternalSemaphoreWaitParams *paramsArray, unsigned int numExtSems, cudaStream_t stream __dv(0)); + extern __host__ cudaError_t CUDARTAPI cudaSignalExternalSemaphoresAsync(const cudaExternalSemaphore_t *extSemArray, const struct cudaExternalSemaphoreSignalParams_v1 *paramsArray, unsigned int numExtSems, cudaStream_t stream __dv(0)); + extern __host__ cudaError_t CUDARTAPI cudaSignalExternalSemaphoresAsync_ptsz(const cudaExternalSemaphore_t *extSemArray, const struct cudaExternalSemaphoreSignalParams_v1 *paramsArray, unsigned int numExtSems, cudaStream_t stream __dv(0)); + extern __host__ cudaError_t CUDARTAPI cudaSignalExternalSemaphoresAsync_v2(const cudaExternalSemaphore_t *extSemArray, const struct cudaExternalSemaphoreSignalParams *paramsArray, unsigned int numExtSems, cudaStream_t stream __dv(0)); + extern __host__ cudaError_t CUDARTAPI cudaWaitExternalSemaphoresAsync(const cudaExternalSemaphore_t *extSemArray, const struct cudaExternalSemaphoreWaitParams_v1 *paramsArray, unsigned int numExtSems, cudaStream_t stream __dv(0)); + extern __host__ cudaError_t CUDARTAPI cudaWaitExternalSemaphoresAsync_ptsz(const cudaExternalSemaphore_t *extSemArray, const struct cudaExternalSemaphoreWaitParams_v1 *paramsArray, unsigned int numExtSems, cudaStream_t stream __dv(0)); + extern __host__ cudaError_t CUDARTAPI cudaWaitExternalSemaphoresAsync_v2(const cudaExternalSemaphore_t *extSemArray, const struct cudaExternalSemaphoreWaitParams *paramsArray, unsigned int numExtSems, cudaStream_t stream __dv(0)); extern __host__ cudaError_t CUDARTAPI cudaGraphUpload(cudaGraphExec_t graphExec, cudaStream_t stream); extern __host__ cudaError_t CUDARTAPI cudaGraphLaunch(cudaGraphExec_t graphExec, cudaStream_t stream); extern __host__ cudaError_t CUDARTAPI cudaStreamBeginCapture(cudaStream_t stream, enum cudaStreamCaptureMode mode); extern __host__ cudaError_t CUDARTAPI cudaStreamEndCapture(cudaStream_t stream, cudaGraph_t *pGraph); extern __host__ cudaError_t CUDARTAPI cudaStreamIsCapturing(cudaStream_t stream, enum cudaStreamCaptureStatus *pCaptureStatus); - extern __host__ cudaError_t CUDARTAPI cudaStreamGetCaptureInfo(cudaStream_t hStream, enum cudaStreamCaptureStatus *pCaptureStatus, unsigned long long *pId); + extern __host__ cudaError_t CUDARTAPI cudaStreamGetCaptureInfo(cudaStream_t stream, enum cudaStreamCaptureStatus *captureStatus_out, unsigned long long *id_out); + extern __host__ cudaError_t CUDARTAPI cudaStreamGetCaptureInfo_v2(cudaStream_t stream, enum cudaStreamCaptureStatus *captureStatus_out, unsigned long long *id_out __dv(0), cudaGraph_t *graph_out __dv(0), const cudaGraphNode_t **dependencies_out __dv(0), size_t *numDependencies_out __dv(0)); + extern __host__ cudaError_t CUDARTAPI cudaStreamUpdateCaptureDependencies_ptsz(cudaStream_t stream, cudaGraphNode_t *dependencies, size_t numDependencies, unsigned int flags __dv(0)); extern __host__ cudaError_t CUDARTAPI cudaStreamCopyAttributes(cudaStream_t dstStream, cudaStream_t srcStream); extern __host__ cudaError_t CUDARTAPI cudaStreamGetAttribute(cudaStream_t stream, enum cudaStreamAttrID attr, union cudaStreamAttrValue *value); extern __host__ cudaError_t CUDARTAPI cudaStreamSetAttribute(cudaStream_t stream, enum cudaStreamAttrID attr, const union cudaStreamAttrValue *param); + extern __host__ cudaError_t CUDARTAPI cudaMallocAsync(void **devPtr, size_t size, cudaStream_t hStream); + extern __host__ cudaError_t CUDARTAPI cudaFreeAsync(void *devPtr, cudaStream_t hStream); + extern __host__ cudaError_t CUDARTAPI cudaMallocFromPoolAsync(void **ptr, size_t size, cudaMemPool_t memPool, cudaStream_t stream); + extern __host__ cudaError_t CUDARTAPI cudaGetDriverEntryPoint(const char *symbol, void **funcPtr, unsigned long long flags); + #elif defined(__CUDART_API_PER_THREAD_DEFAULT_STREAM) // nvcc stubs reference the 'cudaLaunch'/'cudaLaunchKernel' identifier even if it was defined // to 'cudaLaunch_ptsz'/'cudaLaunchKernel_ptsz'. Redirect through a static inline function. diff --git a/samples/external/cuda/include/driver_types.h b/samples/external/cuda/include/driver_types.h index ff006f8..c6fcbc2 100644 --- a/samples/external/cuda/include/driver_types.h +++ b/samples/external/cuda/include/driver_types.h @@ -393,6 +393,13 @@ enum __device_builtin__ cudaError */ cudaErrorMemoryValueTooLarge = 32, + /** + * This indicates that the CUDA driver that the application has loaded is a + * stub library. Applications that run with the stub rather than a real + * driver loaded will result in CUDA API returning this error. + */ + cudaErrorStubLibrary = 34, + /** * This indicates that the installed NVIDIA CUDA driver is older than the * CUDA runtime library. This is not a supported configuration. Users should @@ -538,10 +545,19 @@ enum __device_builtin__ cudaError cudaErrorInvalidDevice = 101, /** - * This indicates that the device doesn't have valid Grid License. + * This indicates that the device doesn't have a valid Grid License. */ cudaErrorDeviceNotLicensed = 102, + /** + * By default, the CUDA runtime may perform a minimal set of self-tests, + * as well as CUDA driver tests, to establish the validity of both. + * Introduced in CUDA 11.2, this error return indicates that at least one + * of these tests has failed and the validity of either the runtime + * or the driver could not be established. + */ + cudaErrorSoftwareValidityNotEstablished = 103, + /** * This indicates an internal startup failure in the CUDA runtime. */ @@ -668,6 +684,13 @@ enum __device_builtin__ cudaError */ cudaErrorUnsupportedPtxVersion = 222, + /** + * This indicates that the JIT compilation was disabled. The JIT compilation compiles + * PTX. The runtime may fall back to compiling PTX if an application does not contain + * a suitable binary for the current device. + */ + cudaErrorJitCompilationDisabled = 223, + /** * This indicates that the device kernel source is invalid. */ @@ -708,7 +731,8 @@ enum __device_builtin__ cudaError /** * This indicates that a named symbol was not found. Examples of symbols - * are global/constant variable names, texture names, and surface names. + * are global/constant variable names, driver function names, texture names, + * and surface names. */ cudaErrorSymbolNotFound = 500, @@ -1002,7 +1026,8 @@ enum __device_builtin__ cudaChannelFormatKind cudaChannelFormatKindSigned = 0, /**< Signed channel format */ cudaChannelFormatKindUnsigned = 1, /**< Unsigned channel format */ cudaChannelFormatKindFloat = 2, /**< Float channel format */ - cudaChannelFormatKindNone = 3 /**< No channel format */ + cudaChannelFormatKindNone = 3, /**< No channel format */ + cudaChannelFormatKindNV12 = 4 /**< Unsigned 8-bit integers, planar 4:2:0 YUV format */ }; /** @@ -1058,6 +1083,7 @@ struct __device_builtin__ cudaArraySparseProperties { unsigned int miptailFirstLevel; /**< First mip level at which the mip tail begins */ unsigned long long miptailSize; /**< Total size of the mip tail. */ unsigned int flags; /**< Flags will either be zero or ::cudaArraySparsePropertiesSingleMipTail */ + unsigned int reserved[4]; }; /** @@ -1163,7 +1189,7 @@ struct __device_builtin__ cudaMemsetParams { size_t pitch; /**< Pitch of destination device pointer. Unused if height is 1 */ unsigned int value; /**< Value to be set */ unsigned int elementSize; /**< Size of each element in bytes. Must be 1, 2, or 4. */ - size_t width; /**< Width in bytes, of the row */ + size_t width; /**< Width of the row in elements */ size_t height; /**< Number of rows */ }; @@ -1258,6 +1284,28 @@ union __device_builtin__ cudaStreamAttrValue { enum cudaSynchronizationPolicy syncPolicy; }; +/** + * Flags for ::cudaStreamUpdateCaptureDependencies + */ +enum __device_builtin__ cudaStreamUpdateCaptureDependenciesFlags { + cudaStreamAddCaptureDependencies = 0x0, /**< Add new nodes to the dependency set */ + cudaStreamSetCaptureDependencies = 0x1 /**< Replace the dependency set with the new nodes */ +}; + +/** + * Flags for user objects for graphs + */ +enum __device_builtin__ cudaUserObjectFlags { + cudaUserObjectNoDestructorSync = 0x1 /**< Indicates the destructor execution is not synchronized by any CUDA handle. */ +}; + +/** + * Flags for retaining user object references for graphs + */ +enum __device_builtin__ cudaUserObjectRetainFlags { + cudaGraphUserObjectMove = 0x1 /**< Transfer references from the caller rather than creating new references. */ +}; + /** * CUDA graphics interop resource */ @@ -1307,7 +1355,7 @@ enum __device_builtin__ cudaKernelNodeAttrID { }; /** - * Graph kernel node attributes union, used with ::cudaKernelNodeSetAttribute/::cudaKernelNodeGetAttribute + * Graph kernel node attributes union, used with ::cudaGraphKernelNodeSetAttribute/::cudaGraphKernelNodeGetAttribute */ union __device_builtin__ cudaKernelNodeAttrValue { struct cudaAccessPolicyWindow accessPolicyWindow; /**< Attribute ::CUaccessPolicyWindow. */ @@ -1619,6 +1667,39 @@ enum __device_builtin__ cudaOutputMode cudaCSV = 0x01 /**< Output mode Comma separated values format. */ }; +/** + * CUDA GPUDirect RDMA flush writes APIs supported on the device + */ +enum __device_builtin__ cudaFlushGPUDirectRDMAWritesOptions { + cudaFlushGPUDirectRDMAWritesOptionHost = 1<<0, /**< ::cudaDeviceFlushGPUDirectRDMAWrites() and its CUDA Driver API counterpart are supported on the device. */ + cudaFlushGPUDirectRDMAWritesOptionMemOps = 1<<1 /**< The ::CU_STREAM_WAIT_VALUE_FLUSH flag and the ::CU_STREAM_MEM_OP_FLUSH_REMOTE_WRITES MemOp are supported on the CUDA device. */ +}; + +/** + * CUDA GPUDirect RDMA flush writes ordering features of the device + */ +enum __device_builtin__ cudaGPUDirectRDMAWritesOrdering { + cudaGPUDirectRDMAWritesOrderingNone = 0, /**< The device does not natively support ordering of GPUDirect RDMA writes. ::cudaFlushGPUDirectRDMAWrites() can be leveraged if supported. */ + cudaGPUDirectRDMAWritesOrderingOwner = 100, /**< Natively, the device can consistently consume GPUDirect RDMA writes, although other CUDA devices may not. */ + cudaGPUDirectRDMAWritesOrderingAllDevices = 200 /**< Any CUDA device in the system can consistently consume GPUDirect RDMA writes to this device. */ +}; + +/** + * CUDA GPUDirect RDMA flush writes scopes + */ +enum __device_builtin__ cudaFlushGPUDirectRDMAWritesScope { + cudaFlushGPUDirectRDMAWritesToOwner = 100, /**< Blocks until remote writes are visible to the CUDA device context owning the data. */ + cudaFlushGPUDirectRDMAWritesToAllDevices = 200 /**< Blocks until remote writes are visible to all CUDA device contexts. */ +}; + +/** + * CUDA GPUDirect RDMA flush writes targets + */ +enum __device_builtin__ cudaFlushGPUDirectRDMAWritesTarget { + cudaFlushGPUDirectRDMAWritesTargetCurrentDevice /**< Sets the target for ::cudaDeviceFlushGPUDirectRDMAWrites() to the currently active CUDA device context. */ +}; + + /** * CUDA device attributes */ @@ -1725,9 +1806,166 @@ enum __device_builtin__ cudaDeviceAttr cudaDevAttrPageableMemoryAccessUsesHostPageTables = 100, /**< Device accesses pageable memory via the host's page tables. */ cudaDevAttrDirectManagedMemAccessFromHost = 101, /**< Host can directly access managed memory on the device without migration. */ cudaDevAttrMaxBlocksPerMultiprocessor = 106, /**< Maximum number of blocks per multiprocessor */ + cudaDevAttrMaxPersistingL2CacheSize = 108, /**< Maximum L2 persisting lines capacity setting in bytes. */ + cudaDevAttrMaxAccessPolicyWindowSize = 109, /**< Maximum value of cudaAccessPolicyWindow::num_bytes. */ cudaDevAttrReservedSharedMemoryPerBlock = 111, /**< Shared memory reserved by CUDA driver per block in bytes */ cudaDevAttrSparseCudaArraySupported = 112, /**< Device supports sparse CUDA arrays and sparse CUDA mipmapped arrays */ - cudaDevAttrHostRegisterReadOnlySupported = 113 /**< Device supports using the ::cuMemHostRegister flag CU_MEMHOSTERGISTER_READ_ONLY to register memory that must be mapped as read-only to the GPU */ + cudaDevAttrHostRegisterReadOnlySupported = 113, /**< Device supports using the ::cudaHostRegister flag cudaHostRegisterReadOnly to register memory that must be mapped as read-only to the GPU */ + cudaDevAttrMaxTimelineSemaphoreInteropSupported = 114, /**< External timeline semaphore interop is supported on the device */ + cudaDevAttrMemoryPoolsSupported = 115, /**< Device supports using the ::cudaMallocAsync and ::cudaMemPool family of APIs */ + cudaDevAttrGPUDirectRDMASupported = 116, /**< Device supports GPUDirect RDMA APIs, like nvidia_p2p_get_pages (see https://docs.nvidia.com/cuda/gpudirect-rdma for more information) */ + cudaDevAttrGPUDirectRDMAFlushWritesOptions = 117, /**< The returned attribute shall be interpreted as a bitmask, where the individual bits are listed in the ::cudaFlushGPUDirectRDMAWritesOptions enum */ + cudaDevAttrGPUDirectRDMAWritesOrdering = 118, /**< GPUDirect RDMA writes to the device do not need to be flushed for consumers within the scope indicated by the returned attribute. See ::cudaGPUDirectRDMAWritesOrdering for the numerical values returned here. */ + cudaDevAttrMemoryPoolSupportedHandleTypes = 119 /**< Handle types supported with mempool based IPC */ +}; + +/** + * CUDA memory pool attributes + */ +enum __device_builtin__ cudaMemPoolAttr +{ + /** + * (value type = int) + * Allow cuMemAllocAsync to use memory asynchronously freed + * in another streams as long as a stream ordering dependency + * of the allocating stream on the free action exists. + * Cuda events and null stream interactions can create the required + * stream ordered dependencies. (default enabled) + */ + cudaMemPoolReuseFollowEventDependencies = 0x1, + + /** + * (value type = int) + * Allow reuse of already completed frees when there is no dependency + * between the free and allocation. (default enabled) + */ + cudaMemPoolReuseAllowOpportunistic = 0x2, + + /** + * (value type = int) + * Allow cuMemAllocAsync to insert new stream dependencies + * in order to establish the stream ordering required to reuse + * a piece of memory released by cuFreeAsync (default enabled). + */ + cudaMemPoolReuseAllowInternalDependencies = 0x3, + + + /** + * (value type = cuuint64_t) + * Amount of reserved memory in bytes to hold onto before trying + * to release memory back to the OS. When more than the release + * threshold bytes of memory are held by the memory pool, the + * allocator will try to release memory back to the OS on the + * next call to stream, event or context synchronize. (default 0) + */ + cudaMemPoolAttrReleaseThreshold = 0x4, + + /** + * (value type = cuuint64_t) + * Amount of backing memory currently allocated for the mempool. + */ + cudaMemPoolAttrReservedMemCurrent = 0x5, + + /** + * (value type = cuuint64_t) + * High watermark of backing memory allocated for the mempool since the + * last time it was reset. High watermark can only be reset to zero. + */ + cudaMemPoolAttrReservedMemHigh = 0x6, + + /** + * (value type = cuuint64_t) + * Amount of memory from the pool that is currently in use by the application. + */ + cudaMemPoolAttrUsedMemCurrent = 0x7, + + /** + * (value type = cuuint64_t) + * High watermark of the amount of memory from the pool that was in use by the application since + * the last time it was reset. High watermark can only be reset to zero. + */ + cudaMemPoolAttrUsedMemHigh = 0x8 +}; + +/** + * Specifies the type of location + */ +enum __device_builtin__ cudaMemLocationType { + cudaMemLocationTypeInvalid = 0, + cudaMemLocationTypeDevice = 1 /**< Location is a device location, thus id is a device ordinal */ +}; + +/** + * Specifies a memory location. + * + * To specify a gpu, set type = ::cudaMemLocationTypeDevice and set id = the gpu's device ordinal. + */ +struct __device_builtin__ cudaMemLocation { + enum cudaMemLocationType type; /**< Specifies the location type, which modifies the meaning of id. */ + int id; /**< identifier for a given this location's ::CUmemLocationType. */ +}; + +/** + * Specifies the memory protection flags for mapping. + */ +enum __device_builtin__ cudaMemAccessFlags { + cudaMemAccessFlagsProtNone = 0, /**< Default, make the address range not accessible */ + cudaMemAccessFlagsProtRead = 1, /**< Make the address range read accessible */ + cudaMemAccessFlagsProtReadWrite = 3 /**< Make the address range read-write accessible */ +}; + +/** + * Memory access descriptor + */ +struct __device_builtin__ cudaMemAccessDesc { + struct cudaMemLocation location; /**< Location on which the request is to change it's accessibility */ + enum cudaMemAccessFlags flags; /**< ::CUmemProt accessibility flags to set on the request */ +}; + +/** + * Defines the allocation types available + */ +enum __device_builtin__ cudaMemAllocationType { + cudaMemAllocationTypeInvalid = 0x0, + /** This allocation type is 'pinned', i.e. cannot migrate from its current + * location while the application is actively using it + */ + cudaMemAllocationTypePinned = 0x1, + cudaMemAllocationTypeMax = 0x7FFFFFFF +}; + +/** + * Flags for specifying particular handle types + */ +enum __device_builtin__ cudaMemAllocationHandleType { + cudaMemHandleTypeNone = 0x0, /**< Does not allow any export mechanism. > */ + cudaMemHandleTypePosixFileDescriptor = 0x1, /**< Allows a file descriptor to be used for exporting. Permitted only on POSIX systems. (int) */ + cudaMemHandleTypeWin32 = 0x2, /**< Allows a Win32 NT handle to be used for exporting. (HANDLE) */ + cudaMemHandleTypeWin32Kmt = 0x4 /**< Allows a Win32 KMT handle to be used for exporting. (D3DKMT_HANDLE) */ +}; + +/** + * Specifies the properties of allocations made from the pool. + */ +struct __device_builtin__ cudaMemPoolProps { + enum cudaMemAllocationType allocType; /**< Allocation type. Currently must be specified as cudaMemAllocationTypePinned */ + enum cudaMemAllocationHandleType handleTypes; /**< Handle types that will be supported by allocations from the pool. */ + struct cudaMemLocation location; /**< Location allocations should reside. */ + /** + * Windows-specific LPSECURITYATTRIBUTES required when + * ::cudaMemHandleTypeWin32 is specified. This security attribute defines + * the scope of which exported allocations may be tranferred to other + * processes. In all other cases, this field is required to be zero. + */ + void *win32SecurityAttributes; + unsigned char reserved[64]; /**< reserved for future use, must be 0 */ +}; + +/** + * Opaque data for exporting a pool allocation + */ +struct __device_builtin__ cudaMemPoolPtrExportData { + unsigned char reserved[64]; }; /** @@ -1784,7 +2022,7 @@ struct __device_builtin__ cudaDeviceProp int computeMode; /**< Compute mode (See ::cudaComputeMode) */ int maxTexture1D; /**< Maximum 1D texture size */ int maxTexture1DMipmap; /**< Maximum 1D mipmapped texture size */ - int maxTexture1DLinear; /**< Maximum size for 1D textures bound to linear memory */ + int maxTexture1DLinear; /**< Deprecated, do not use. Use cudaDeviceGetTexture1DLinearMaxWidth() or cuDeviceGetTexture1DLinearMaxWidth() instead. */ int maxTexture2D[2]; /**< Maximum 2D texture dimensions */ int maxTexture2DMipmap[2]; /**< Maximum 2D mipmapped texture dimensions */ int maxTexture2DLinear[3]; /**< Maximum dimensions (width, height, pitch) for 2D textures bound to pitched memory */ @@ -1831,7 +2069,7 @@ struct __device_builtin__ cudaDeviceProp int computePreemptionSupported; /**< Device supports Compute Preemption */ int canUseHostPointerForRegisteredMem; /**< Device can access host registered memory at the same virtual address as the CPU */ int cooperativeLaunch; /**< Device supports launching cooperative kernels via ::cudaLaunchCooperativeKernel */ - int cooperativeMultiDeviceLaunch; /**< Device can participate in cooperative kernels launched via ::cudaLaunchCooperativeKernelMultiDevice */ + int cooperativeMultiDeviceLaunch; /**< Deprecated, cudaLaunchCooperativeKernelMultiDevice is deprecated. */ size_t sharedMemPerBlockOptin; /**< Per device maximum shared memory per block usable by special opt in */ int pageableMemoryAccessUsesHostPageTables; /**< Device accesses pageable memory via the host's page tables */ int directManagedMemAccessFromHost; /**< Host can directly access managed memory on the device without migration. */ @@ -2134,7 +2372,7 @@ enum __device_builtin__ cudaExternalSemaphoreHandleType { * Handle is an opaque shared NT handle */ cudaExternalSemaphoreHandleTypeOpaqueWin32 = 2, - /** + /** * Handle is an opaque, globally shared handle */ cudaExternalSemaphoreHandleTypeOpaqueWin32Kmt = 3, @@ -2157,7 +2395,15 @@ enum __device_builtin__ cudaExternalSemaphoreHandleType { /** * Handle is a shared KMT handle referencing a D3D11 keyed mutex object */ - cudaExternalSemaphoreHandleTypeKeyedMutexKmt = 8 + cudaExternalSemaphoreHandleTypeKeyedMutexKmt = 8, + /** + * Handle is an opaque handle file descriptor referencing a timeline semaphore + */ + cudaExternalSemaphoreHandleTypeTimelineSemaphoreFd = 9, + /** + * Handle is an opaque handle file descriptor referencing a timeline semaphore + */ + cudaExternalSemaphoreHandleTypeTimelineSemaphoreWin32 = 10 }; /** @@ -2170,8 +2416,10 @@ struct __device_builtin__ cudaExternalSemaphoreHandleDesc { enum cudaExternalSemaphoreHandleType type; union { /** - * File descriptor referencing the semaphore object. Valid - * when type is ::cudaExternalSemaphoreHandleTypeOpaqueFd + * File descriptor referencing the semaphore object. Valid when + * type is one of the following: + * - ::cudaExternalSemaphoreHandleTypeOpaqueFd + * - ::cudaExternalSemaphoreHandleTypeTimelineSemaphoreFd */ int fd; /** @@ -2182,6 +2430,7 @@ struct __device_builtin__ cudaExternalSemaphoreHandleDesc { * - ::cudaExternalSemaphoreHandleTypeD3D12Fence * - ::cudaExternalSemaphoreHandleTypeD3D11Fence * - ::cudaExternalSemaphoreHandleTypeKeyedMutex + * - ::cudaExternalSemaphoreHandleTypeTimelineSemaphoreWin32 * Exactly one of 'handle' and 'name' must be non-NULL. If * type is one of the following: * ::cudaExternalSemaphoreHandleTypeOpaqueWin32Kmt @@ -2210,10 +2459,11 @@ struct __device_builtin__ cudaExternalSemaphoreHandleDesc { unsigned int flags; }; +#if defined(__CUDA_API_VERSION_INTERNAL) /** - * External semaphore signal parameters + * External semaphore signal parameters(deprecated) */ -struct __device_builtin__ cudaExternalSemaphoreSignalParams { +struct __device_builtin__ cudaExternalSemaphoreSignalParams_v1 { struct { /** * Parameters for fence objects @@ -2256,9 +2506,9 @@ struct __device_builtin__ cudaExternalSemaphoreSignalParams { }; /** -* External semaphore wait parameters +* External semaphore wait parameters(deprecated) */ -struct __device_builtin__ cudaExternalSemaphoreWaitParams { +struct __device_builtin__ cudaExternalSemaphoreWaitParams_v1 { struct { /** * Parameters for fence objects @@ -2303,6 +2553,105 @@ struct __device_builtin__ cudaExternalSemaphoreWaitParams { */ unsigned int flags; }; +#endif + +/** + * External semaphore signal parameters, compatible with driver type + */ +struct __device_builtin__ cudaExternalSemaphoreSignalParams{ + struct { + /** + * Parameters for fence objects + */ + struct { + /** + * Value of fence to be signaled + */ + unsigned long long value; + } fence; + union { + /** + * Pointer to NvSciSyncFence. Valid if ::cudaExternalSemaphoreHandleType + * is of type ::cudaExternalSemaphoreHandleTypeNvSciSync. + */ + void *fence; + unsigned long long reserved; + } nvSciSync; + /** + * Parameters for keyed mutex objects + */ + struct { + /* + * Value of key to release the mutex with + */ + unsigned long long key; + } keyedMutex; + unsigned int reserved[12]; + } params; + /** + * Only when ::cudaExternalSemaphoreSignalParams is used to + * signal a ::cudaExternalSemaphore_t of type + * ::cudaExternalSemaphoreHandleTypeNvSciSync, the valid flag is + * ::cudaExternalSemaphoreSignalSkipNvSciBufMemSync: which indicates + * that while signaling the ::cudaExternalSemaphore_t, no memory + * synchronization operations should be performed for any external memory + * object imported as ::cudaExternalMemoryHandleTypeNvSciBuf. + * For all other types of ::cudaExternalSemaphore_t, flags must be zero. + */ + unsigned int flags; + unsigned int reserved[16]; +}; + +/** + * External semaphore wait parameters, compatible with driver type + */ +struct __device_builtin__ cudaExternalSemaphoreWaitParams { + struct { + /** + * Parameters for fence objects + */ + struct { + /** + * Value of fence to be waited on + */ + unsigned long long value; + } fence; + union { + /** + * Pointer to NvSciSyncFence. Valid if ::cudaExternalSemaphoreHandleType + * is of type ::cudaExternalSemaphoreHandleTypeNvSciSync. + */ + void *fence; + unsigned long long reserved; + } nvSciSync; + /** + * Parameters for keyed mutex objects + */ + struct { + /** + * Value of key to acquire the mutex with + */ + unsigned long long key; + /** + * Timeout in milliseconds to wait to acquire the mutex + */ + unsigned int timeoutMs; + } keyedMutex; + unsigned int reserved[10]; + } params; + /** + * Only when ::cudaExternalSemaphoreSignalParams is used to + * signal a ::cudaExternalSemaphore_t of type + * ::cudaExternalSemaphoreHandleTypeNvSciSync, the valid flag is + * ::cudaExternalSemaphoreSignalSkipNvSciBufMemSync: which indicates + * that while waiting for the ::cudaExternalSemaphore_t, no memory + * synchronization operations should be performed for any external memory + * object imported as ::cudaExternalMemoryHandleTypeNvSciBuf. + * For all other types of ::cudaExternalSemaphore_t, flags must be zero. + */ + unsigned int flags; + unsigned int reserved[16]; +}; /******************************************************************************* @@ -2356,11 +2705,21 @@ typedef __device_builtin__ struct CUgraph_st *cudaGraph_t; */ typedef __device_builtin__ struct CUgraphNode_st *cudaGraphNode_t; +/** + * CUDA user object for graphs + */ +typedef __device_builtin__ struct CUuserObject_st *cudaUserObject_t; + /** * CUDA function */ typedef __device_builtin__ struct CUfunc_st *cudaFunction_t; +/** + * CUDA memory pool + */ +typedef __device_builtin__ struct CUmemPoolHandle_st *cudaMemPool_t; + /** * CUDA cooperative group scope */ @@ -2395,16 +2754,36 @@ struct __device_builtin__ cudaKernelNodeParams { void **extra; /**< Pointer to kernel arguments in the "extra" format */ }; +/** + * External semaphore signal node parameters + */ +struct __device_builtin__ cudaExternalSemaphoreSignalNodeParams { + cudaExternalSemaphore_t* extSemArray; /**< Array of external semaphore handles. */ + const struct cudaExternalSemaphoreSignalParams* paramsArray; /**< Array of external semaphore signal parameters. */ + unsigned int numExtSems; /**< Number of handles and parameters supplied in extSemArray and paramsArray. */ +}; + +/** + * External semaphore wait node parameters + */ +struct __device_builtin__ cudaExternalSemaphoreWaitNodeParams { + cudaExternalSemaphore_t* extSemArray; /**< Array of external semaphore handles. */ + const struct cudaExternalSemaphoreWaitParams* paramsArray; /**< Array of external semaphore wait parameters. */ + unsigned int numExtSems; /**< Number of handles and parameters supplied in extSemArray and paramsArray. */ +}; + /** * CUDA Graph node types */ enum __device_builtin__ cudaGraphNodeType { - cudaGraphNodeTypeKernel = 0x00, /**< GPU kernel node */ - cudaGraphNodeTypeMemcpy = 0x01, /**< Memcpy node */ - cudaGraphNodeTypeMemset = 0x02, /**< Memset node */ - cudaGraphNodeTypeHost = 0x03, /**< Host (executable) node */ - cudaGraphNodeTypeGraph = 0x04, /**< Node which executes an embedded graph */ - cudaGraphNodeTypeEmpty = 0x05, /**< Empty (no-op) node */ + cudaGraphNodeTypeKernel = 0x00, /**< GPU kernel node */ + cudaGraphNodeTypeMemcpy = 0x01, /**< Memcpy node */ + cudaGraphNodeTypeMemset = 0x02, /**< Memset node */ + cudaGraphNodeTypeHost = 0x03, /**< Host (executable) node */ + cudaGraphNodeTypeGraph = 0x04, /**< Node which executes an embedded graph */ + cudaGraphNodeTypeEmpty = 0x05, /**< Empty (no-op) node */ + cudaGraphNodeTypeWaitEvent = 0x06, /**< External event wait node */ + cudaGraphNodeTypeEventRecord = 0x07, /**< External event record node */ cudaGraphNodeTypeCount }; @@ -2421,9 +2800,36 @@ enum __device_builtin__ cudaGraphExecUpdateResult { cudaGraphExecUpdateError = 0x1, /**< The update failed for an unexpected reason which is described in the return value of the function */ cudaGraphExecUpdateErrorTopologyChanged = 0x2, /**< The update failed because the topology changed */ cudaGraphExecUpdateErrorNodeTypeChanged = 0x3, /**< The update failed because a node type changed */ - cudaGraphExecUpdateErrorFunctionChanged = 0x4, /**< The update failed because the function of a kernel node changed */ + cudaGraphExecUpdateErrorFunctionChanged = 0x4, /**< The update failed because the function of a kernel node changed (CUDA driver < 11.2) */ cudaGraphExecUpdateErrorParametersChanged = 0x5, /**< The update failed because the parameters changed in a way that is not supported */ - cudaGraphExecUpdateErrorNotSupported = 0x6 /**< The update failed because something about the node is not supported */ + cudaGraphExecUpdateErrorNotSupported = 0x6, /**< The update failed because something about the node is not supported */ + cudaGraphExecUpdateErrorUnsupportedFunctionChange = 0x7 /**< The update failed because the function of a kernel node changed in an unsupported way */ +}; + +/** + * Flags to specify search options to be used with ::cudaGetDriverEntryPoint + * For more details see ::cuGetProcAddress + */ +enum __device_builtin__ cudaGetDriverEntryPointFlags { + cudaEnableDefault = 0x0, /**< Default search mode for driver symbols. */ + cudaEnableLegacyStream = 0x1, /**< Search for legacy versions of driver symbols. */ + cudaEnablePerThreadDefaultStream = 0x2 /**< Search for per-thread versions of driver symbols. */ +}; + +/** + * CUDA Graph debug write options + */ +enum __device_builtin__ cudaGraphDebugDotFlags { + cudaGraphDebugDotFlagsVerbose = 1<<0, /** Output all debug data as if every debug flag is enabled */ + cudaGraphDebugDotFlagsKernelNodeParams = 1<<2, /** Adds cudaKernelNodeParams to output */ + cudaGraphDebugDotFlagsMemcpyNodeParams = 1<<3, /** Adds cudaMemcpy3DParms to output */ + cudaGraphDebugDotFlagsMemsetNodeParams = 1<<4, /** Adds cudaMemsetParams to output */ + cudaGraphDebugDotFlagsHostNodeParams = 1<<5, /** Adds cudaHostNodeParams to output */ + cudaGraphDebugDotFlagsEventNodeParams = 1<<6, /** Adds cudaEvent_t handle from record and wait nodes to output */ + cudaGraphDebugDotFlagsExtSemasSignalNodeParams = 1<<7, /** Adds cudaExternalSemaphoreSignalNodeParams values to output */ + cudaGraphDebugDotFlagsExtSemasWaitNodeParams = 1<<8, /** Adds cudaExternalSemaphoreWaitNodeParams to output */ + cudaGraphDebugDotFlagsKernelNodeAttributes = 1<<9, /** Adds cudaKernelNodeAttrID values to output */ + cudaGraphDebugDotFlagsHandles = 1<<10 /** Adds node handles and every kernel function handle to output */ }; /** @} */ diff --git a/samples/external/cuda/lib/x64/cudart.lib b/samples/external/cuda/lib/x64/cudart.lib index b9c8ee5..19fe0c2 100644 Binary files a/samples/external/cuda/lib/x64/cudart.lib and b/samples/external/cuda/lib/x64/cudart.lib differ diff --git a/samples/utils/nvCVOpenCV.h b/samples/utils/nvCVOpenCV.h index ddb9274..271d59a 100644 --- a/samples/utils/nvCVOpenCV.h +++ b/samples/utils/nvCVOpenCV.h @@ -58,21 +58,22 @@ inline void CVWrapperForNvCVImage(const NvCVImage *nvcvIm, cv::Mat *cvIm) { // Wrap a cv::Mat in an NvCVImage. inline void NVWrapperForCVMat(const cv::Mat *cvIm, NvCVImage *nvcvIm) { static const NvCVImage_PixelFormat nvFormat[] = { NVCV_FORMAT_UNKNOWN, NVCV_Y, NVCV_YA, NVCV_BGR, NVCV_BGRA }; - static const NvCVImage_ComponentType nvType[] = { NVCV_U8, NVCV_TYPE_UNKNOWN, NVCV_U16, NVCV_S16, NVCV_S32, NVCV_F32, NVCV_F64 }; + static const NvCVImage_ComponentType nvType[] = {NVCV_U8, NVCV_TYPE_UNKNOWN, NVCV_U16, NVCV_S16, + NVCV_S32, NVCV_F32, NVCV_F64, NVCV_TYPE_UNKNOWN}; nvcvIm->pixels = cvIm->data; nvcvIm->width = cvIm->cols; nvcvIm->height = cvIm->rows; - nvcvIm->pitch = (unsigned)cvIm->step1(); - nvcvIm->pixelFormat = nvFormat[cvIm->channels()]; - nvcvIm->componentType = nvType[cvIm->depth()]; + nvcvIm->pitch = (int)cvIm->step[0]; + nvcvIm->pixelFormat = nvFormat[cvIm->channels() <= 4 ? cvIm->channels() : 0]; + nvcvIm->componentType = nvType[cvIm->depth() & 7]; nvcvIm->bufferBytes = 0; nvcvIm->deletePtr = nullptr; nvcvIm->deleteProc = nullptr; - nvcvIm->pixelBytes = (unsigned char)cvIm->elemSize(); + nvcvIm->pixelBytes = (unsigned char)cvIm->step[1]; nvcvIm->componentBytes = (unsigned char)cvIm->elemSize1(); nvcvIm->numComponents = (unsigned char)cvIm->channels(); - nvcvIm->planar = 0; - nvcvIm->gpuMem = 0; + nvcvIm->planar = NVCV_CHUNKY; + nvcvIm->gpuMem = NVCV_CPU; nvcvIm->reserved[0] = 0; nvcvIm->reserved[1] = 0; } diff --git a/version.h b/version.h index 4308db3..aba1505 100644 --- a/version.h +++ b/version.h @@ -1,10 +1,35 @@ +/*############################################################################### +# +# Copyright 2020 NVIDIA Corporation +# +# Permission is hereby granted, free of charge, to any person obtaining a copy of +# this software and associated documentation files (the "Software"), to deal in +# the Software without restriction, including without limitation the rights to +# use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies of +# the Software, and to permit persons to whom the Software is furnished to do so, +# subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in all +# copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS +# FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR +# COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER +# IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN +# CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. +# +###############################################################################*/ + + #define NVIDIA_VIDEOEFFECTS_SDK_VERSION_MAJOR 0 #define NVIDIA_VIDEOEFFECTS_SDK_VERSION_MINOR 6 -#define NVIDIA_VIDEOEFFECTS_SDK_VERSION_RELEASE 4 +#define NVIDIA_VIDEOEFFECTS_SDK_VERSION_RELEASE 5 +#define NVIDIA_VIDEOEFFECTS_SDK_VERSION_BUILD 2 -#define NVIDIA_VIDEOEFFECTS_SDK_VERSION 0,6,4,0 +#define NVIDIA_VIDEOEFFECTS_SDK_VERSION 0,6,5,2 #define NVIDIA_VIDEOEFFECTS_SDK_VERSION_MAJOR_MINOR 0,6 -#define NVIDIA_VIDEOEFFECTS_SDK_VERSION_STRING "0.6.4.0" -#define NVIDIA_VIDEOEFFECTS_SDK_VERSION_STRING_SHORT "0.6.4" +#define NVIDIA_VIDEOEFFECTS_SDK_VERSION_STRING "0.6.5.2" +#define NVIDIA_VIDEOEFFECTS_SDK_VERSION_STRING_SHORT "0.6.5" #define NVIDIA_VIDEOEFFECTS_SDK_VERSION_STRING_MAJOR_MINOR "0.6"