diff --git a/CHANGELOG b/CHANGELOG index f046421..e717cda 100644 --- a/CHANGELOG +++ b/CHANGELOG @@ -1,8 +1,14 @@ -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 +Changelog (v0.7.1.0) +-------------------- + - NvCVImage_Transfer() now sets alpha to 255 or 1.0f when doing RGB -> RGBA + - NvCVImage_CompositeRect() has a premultiplied alpha mode added + - AI Green Screen filter + - Introduced NVVFX_MAX_INPUT_WIDTH, NVVFX_MAX_INPUT_HEIGHT parameter selectors and getters to set the maximum input resolution + for NVVFX_FX_GREEN_SCREEN pipeline to avoid lazy memory allocations + - Introduced the following parameter selectors and APIs for NVVFX_FX_GREEN_SCREEN to improve the quality of the model + - NVVFX_MAX_NUMBER_STREAMS parameter selector to set the maximum number of streams + - NVVFX_STATE, NVVFX_STATE_COUNT parameter selectors + - NvVFX_AllocateState, NvVFX_DeallocateState, NvVFX_ResetState, NvVFX_SetStateObjectHandleArray APIs + - BatchEffectApp does not support NVVFX_FX_GREEN_SCREEN anymore. Use BatchAigsEffectApp as the sample app for batch input + - Migrated to TensorRT 8.4.2.2 + - Migrated to CUDA 11.6u1 diff --git a/CMakeLists.txt b/CMakeLists.txt index efaa07f..1ebe6b8 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -21,9 +21,77 @@ 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}) -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}) +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}") + + + # Add target for NVCVImage + add_library(NVCVImage INTERFACE) + + # found in different locations depending on type of package + find_path(NVCVImage_INCLUDES + NAMES nvCVImage.h + PATHS + /usr/local/VideoFX/include + /usr/include/x86_64-linux-gnu + /usr/include + REQUIRED + ) + + target_include_directories(NVCVImage INTERFACE ${NVCVImage_INCLUDES}) + + + find_library(NVCVImage_LIB + NAMES libNVCVImage.so + PATHS + /usr/local/VideoFX/lib + /usr/lib/x86_64-linux-gnu + /usr/lib64 + /usr/lib + REQUIRED + NO_DEFAULT_PATH) + + target_link_libraries(NVCVImage INTERFACE "${NVCVImage_LIB}") + + message(STATUS "NVCVImage_LIB: ${NVCVImage_LIB}") + message(STATUS "NVCVImage_INCLUDES_PATH: ${NVCVImage_INCLUDES}") + +endif() add_subdirectory(samples) diff --git a/README.MD b/README.MD index 9140cfc..7a8ffce 100644 --- a/README.MD +++ b/README.MD @@ -1,40 +1,39 @@ # README ## NVIDIA MAXINE VideoEffects SDK: API Source Code and Sample Applications -NVIDIA MAXINE VideoEffects SDK is an SDK for enhancing and applying filters to videos at real-time. The SDK is powered by NVIDIA graphics processing units (GPUs) with Tensor Cores, and as a result, the algorithm throughput is greatly accelerated, and latency is reduced. +NVIDIA MAXINE Video Effects SDK enables AI-based visual effects that run with standard webcam input and can easily be integrated into video conference and content creation pipelines. The underlying deep learning models are optimized with NVIDIA AI using NVIDIA® TensorRT™ for high-performance inference, making it possible for developers to apply multiple effects in real-time applications. The SDK has the following AI features: -- **AI Green Screen**, which segments and masks the background areas in a video or image. -- **Background Blur**, which uses the segmentation mask from the AI Green Screen filter or other sources, and produces a blur effect over the background of a video or iamge. -- **Encoder Artifact Reduction**, which reduces the blocky and noisy artifacts from an encoded video while preserving the details of the original video. -- **Super Resolution**, which upscales a video while also reducing the blocky and noisy artifacts. It can enhance the details and sharpen the output while simultaneously preserving the content. This is suitable for upscaling lossy content. -- **Upscale**, which is a very fast and light-weight method for upscaling an input video. It also provides a sharpening parameter to sharpen the resulting output. This feature can be optionally pipelined with the encoder artifact reduction feature to enhance the scale while reducing the video artifacts. -- **Webcam Denoising**, which removes noise from a webcam video while preserving the texture details. +- **Virtual Background**,which segments and masks the background areas in a video or image to enable AI-powered background removal, replacement, or blur. +- **Artifact Reduction**, which reduces compression artifacts from an encoded video while preserving the details of the original video. +- **Super Resolution**, which generates a detail-enhanced video with up to 4X high-quality scaling, while also reducing blocky/noisy artifacts and preserving textures and content. It is suitable for upscaling lossy content. +- **Upscaler**, which is a very fast and light-weight method to deliver up to 4X high-quality scaled video with an adjustable sharpening parameter. This feature can be optionally pipelined with the Artifact Reduction feature to enhance the scale while reducing the video artifacts. +- **Video Noise Removal**, which removes low-light camera noise from a webcam video while preserving the texture details.

NVIDIA Super Resolution

-NVIDIA Webcam Denoising +NVIDIA Video Noise Removal

The SDK provides several sample applications that demonstrate the features listed above in real time by using offline videos. -- **AI Green Screen App**, which is a sample app that demonstrates the background segmentation feature. -- **VideoEffects App**, which is a sample app that can invoke each of Encoder Artifact Reduction, Super Resolution or Upscale features individually. -- **UpscalePipeline App**, which is a sample app that pipelines the Encoder Artifact Reduction feature with the Upscale feature. -- **DenoiseEffect App**, which is a sample app that demonstrates the webcam denoising feature. +- **AI Green Screen App**, which is a sample app that demonstrates the Virtual background feature. +- **VideoEffects App**, which is a sample app that can invoke each of Artifact Reduction, Super Resolution or Upscaler features individually. +- **UpscalePipeline App**, which is a sample app that pipelines the Artifact Reduction feature with the Upscaler feature. +- **DenoiseEffect App**, which is a sample app that demonstrates the Video Noise Removal feature. The input and output resolutions supported by the features of the SDK are listed below. -- The Encoder Artifact Reduction feature supports between 90p to 1080p as input resolutions. +- The 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. - 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. +- The Upscaler feature supports any input resolution, and the following scaling factors: 4/3x (~1.33x), 1.5x, 2x, 3x and 4x. +- The Video Noise Removal feature supports between 80p to 1080p as input resolutions. +- The Virtual Background and Background Blur features require that an input image/video be at least 288 pixels high. NVIDIA MAXINE VideoEffects SDK is distributed in the following parts: @@ -44,12 +43,12 @@ NVIDIA MAXINE VideoEffects SDK is distributed in the following parts: 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. +The SDK is supported on NVIDIA GPUs that are based on the NVIDIA® Turing™, Ampere™ or Ada™ architecture and have Tensor Cores. * 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: 465.89 or later +* NVIDIA Graphics Driver for Windows: 511.65 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/) diff --git a/nvvfx/include/nvCVImage.h b/nvvfx/include/nvCVImage.h index 014cd94..543df6e 100644 --- a/nvvfx/include/nvCVImage.h +++ b/nvvfx/include/nvCVImage.h @@ -323,6 +323,12 @@ NvCV_Status NvCV_API NvCVImage_Realloc(NvCVImage *im, unsigned width, unsigned h void NvCV_API NvCVImage_Dealloc(NvCVImage *im); +//! Deallocate the image buffer from the image asynchronously on the specified stream. The image is not deallocated. +//! param[in,out] im the image whose buffer is to be deallocated. +//! param[int] stream the CUDA stream on which the image buffer is to be deallocated.. +void NvCV_API NvCVImage_DeallocAsync(NvCVImage *im, struct CUstream_st *stream); + + //! Allocate a new image, with storage (C-style constructor). //! \param[in] width the desired width of the image, in pixels. //! \param[in] height the desired height of the image, in pixels. @@ -378,7 +384,7 @@ void NvCV_API NvCVImage_ComponentOffsets(NvCVImage_PixelFormat format, int *rOff //! | RGB --> RGB | X | X | X | X | //! | RGB --> RGBA | X | X | X | X | //! | RGBA --> Y | X | X | | | -//! | RGBA --> A | | X | | | +//! | RGBA --> A | X | | | | //! | RGBA --> RGB | X | X | X | X | //! | RGBA --> RGBA | X | X | X | X | //! | RGB --> YUV420 | X | | X | | @@ -559,7 +565,7 @@ NvCV_Status NvCV_API NvCVImage_Composite(const NvCVImage *fg, const NvCVImage *b //! \param[in] mat the matte image, indicating where the src should come through. //! This determines the size of the rectangle to be composited. //! If this is multi-channel, the alpha channel is used as the matte. -//! \param[in] mode the composition mode. Only 0 (straight alpha over) is implemented at this time. +//! \param[in] mode the composition mode: 0 (straight alpha over) or 1 (premultiplied alpha over). //! \param[out] dst the destination image. This can be the same as fg or bg. //! \param[in] dstOrg the upper-left corner of the dst image to be updated (NULL implies (0,0)). //! \param[in] stream the CUDA stream on which the composition is to be performed. @@ -631,6 +637,23 @@ NvCV_Status NvCV_API NvCVImage_GetYUVPointers(NvCVImage *im, int *yPixBytes, int *cPixBytes, int *yRowBytes, int *cRowBytes); +//! Sharpen an image. +//! The src and dst should be the same type - conversions are not performed. +//! This function is only implemented for NVCV_CHUNKY NVCV_U8 pixels, of format NVCV_RGB or NVCV_BGR. +//! \param[in] sharpness the sharpness strength, calibrated so that 1 and 2 yields Adobe's Sharpen and Sharpen More. +//! \param[in] src the source image to be sharpened. +//! \param[out] dst the resultant image (may be the same as the src). +//! \param[in] stream the CUDA stream on which to perform the computations. +//! \param[in] tmp a temporary working image. This can be NULL, but may result in lower performance. +//! It is best if it resides on the same processor (CPU or GPU) as the destination. +//! @return NVCV_SUCCESS if the operation completed successfully. +//! NVCV_ERR_MISMATCH if the source and destination formats are different. +//! NVCV_ERR_PIXELFORMAT if the function has not been implemented for the chosen pixel type. + +NvCV_Status NvCV_API NvCVImage_Sharpen(float sharpness, const NvCVImage *src, NvCVImage *dst, + struct CUstream_st *stream, NvCVImage *tmp); + + #ifdef __cplusplus } // extern "C" diff --git a/nvvfx/include/nvCVStatus.h b/nvvfx/include/nvCVStatus.h index dd47ba3..2931a9a 100644 --- a/nvvfx/include/nvCVStatus.h +++ b/nvvfx/include/nvCVStatus.h @@ -77,7 +77,14 @@ typedef enum NvCV_Status { NVCV_ERR_TRT_ENGINE = -30, ///< There was a problem deserializing the inference runtime engine. NVCV_ERR_NPP = -31, //!< An error has occurred in the NPP library. NVCV_ERR_CONFIG = -32, //!< No suitable model exists for the specified parameter configuration. + NVCV_ERR_TOOSMALL = -33, //!< A supplied parameter or buffer is not large enough. + NVCV_ERR_TOOBIG = -34, //!< A supplied parameter is too big. + NVCV_ERR_WRONGSIZE = -35, //!< A supplied parameter is not the expected size. + NVCV_ERR_OBJECTNOTFOUND = -36, //!< The specified object was not found. + NVCV_ERR_SINGULAR = -37, //!< A mathematical singularity has been encountered. + NVCV_ERR_NOTHINGRENDERED = -38, //!< Nothing was rendered in the specified region. + NVCV_ERR_OPENGL = -98, //!< An OpenGL error has occurred. NVCV_ERR_DIRECT3D = -99, //!< A Direct3D error has occurred. NVCV_ERR_CUDA_BASE = -100, //!< CUDA errors are offset from this value. diff --git a/nvvfx/include/nvVideoEffects.h b/nvvfx/include/nvVideoEffects.h index 7434175..df60a27 100644 --- a/nvvfx/include/nvVideoEffects.h +++ b/nvvfx/include/nvVideoEffects.h @@ -55,6 +55,12 @@ typedef const char* NvVFX_ParameterSelector; struct NvVFX_Object; typedef struct NvVFX_Object NvVFX_Object, *NvVFX_Handle; +///! +///! Effect may use this handle to manage state objects. +///! +struct NvVFX_StateObjectHandleBase; +typedef struct NvVFX_StateObjectHandleBase* NvVFX_StateObjectHandle; + //! Get the SDK version //! \param[in,out] version Pointer to an unsigned int set to //! (major << 24) | (minor << 16) | (build << 8) | 0 @@ -88,6 +94,7 @@ NvCV_Status NvVFX_API NvVFX_SetF32(NvVFX_Handle effect, NvVFX_ParameterSelector NvCV_Status NvVFX_API NvVFX_SetF64(NvVFX_Handle effect, NvVFX_ParameterSelector paramName, double val); NvCV_Status NvVFX_API NvVFX_SetU64(NvVFX_Handle effect, NvVFX_ParameterSelector paramName, unsigned long long val); NvCV_Status NvVFX_API NvVFX_SetObject(NvVFX_Handle effect, NvVFX_ParameterSelector paramName, void *ptr); +NvCV_Status NvVFX_API NvVFX_SetStateObjectHandleArray(NvVFX_Handle effect, NvVFX_ParameterSelector paramName, NvVFX_StateObjectHandle* handle); NvCV_Status NvVFX_API NvVFX_SetCudaStream(NvVFX_Handle effect, NvVFX_ParameterSelector paramName, CUstream stream); //! Set the selected image descriptor. @@ -167,28 +174,50 @@ NvCV_Status NvVFX_API NvVFX_GetString(NvVFX_Handle effect, NvVFX_ParameterSelect //! \param[in] async run the effect asynchronously if nonzero; otherwise run synchronously. //! \todo Should async instead be a pointer to a place to store a token that can be useful //! for synchronizing two streams alter? -//! \return NVFVX_SUCCESS if the operation was successful. +//! \return NVCV_SUCCESS if the operation was successful. //! \return NVCV_ERR_EFFECT if an invalid effect handle was supplied. NvCV_Status NvVFX_API NvVFX_Run(NvVFX_Handle effect, int async); //! Load the model based on the set params. //! \param[in] effect the effect object handle. -//! \return NVFVX_SUCCESS if the operation was successful. +//! \return NVCV_SUCCESS if the operation was successful. //! \return NVCV_ERR_EFFECT if an invalid effect handle was supplied. NvCV_Status NvVFX_API NvVFX_Load(NvVFX_Handle effect); //! Wrapper for cudaStreamCreate(), if it is desired to avoid linking with the cuda lib. //! \param[out] stream A place to store the newly allocated stream. -//! \return NVFVX_SUCCESS if the operation was successful, +//! \return NVCV_SUCCESS if the operation was successful, //! NVCV_ERR_CUDA_VALUE if not. NvCV_Status NvVFX_API NvVFX_CudaStreamCreate(CUstream *stream); //! Wrapper for cudaStreamDestroy(), if it is desired to avoid linking with the cuda lib. //! \param[in] stream The stream to destroy. -//! \return NVFVX_SUCCESS if the operation was successful, +//! \return NVCV_SUCCESS if the operation was successful, //! NVCV_ERR_CUDA_VALUE if not. NvCV_Status NvVFX_API NvVFX_CudaStreamDestroy(CUstream stream); +//! Allocate the state object handle for a feature. +//! \param[in] effect the effect object handle. +//! \param[in] handle handle to the state object +//! \return NVCV_SUCCESS if the operation was successful. +//! \return NVCV_ERR_EFFECT if an invalid effect handle was supplied. +//! \note This may depend on prior settings of parameters. +NvCV_Status NvVFX_API NvVFX_AllocateState(NvVFX_Handle effect, NvVFX_StateObjectHandle* handle); + +//! Deallocate the state object handle for stateful feature. +//! \param[in] effect the effect object handle. +//! \param[in] handle handle to the state object +//! \return NVCV_SUCCESS if the operation was successful. +//! \return NVCV_ERR_EFFECT if an invalid effect handle was supplied. +NvCV_Status NvVFX_API NvVFX_DeallocateState(NvVFX_Handle effect, NvVFX_StateObjectHandle handle); + +//! Reset the state object handle for stateful feature. +//! \param[in] effect the effect object handle. +//! \param[in] handle handle to the state object +//! \return NVCV_SUCCESS if the operation was successful. +//! \return NVCV_ERR_EFFECT if an invalid effect handle was supplied. +NvCV_Status NvVFX_API NvVFX_ResetState(NvVFX_Handle effect, NvVFX_StateObjectHandle handle); + // Filter selectors #define NVVFX_FX_TRANSFER "Transfer" @@ -209,6 +238,9 @@ NvCV_Status NvVFX_API NvVFX_CudaStreamDestroy(CUstream stream); #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_MAX_INPUT_WIDTH "MaxInputWidth" //!< Maximum width of the input supported +#define NVVFX_MAX_INPUT_HEIGHT "MaxInputHeight" //!< Maximum height of the input supported +#define NVVFX_MAX_NUMBER_STREAMS "MaxNumberStreams" //!< Maximum number of concurrent input streams #define NVVFX_SCALE "Scale" //!< Scale factor #define NVVFX_STRENGTH "Strength" //!< Strength for different filters #define NVVFX_STRENGTH_LEVELS "StrengthLevels" //!< Number of strength levels @@ -219,7 +251,7 @@ NvCV_Status NvVFX_API NvVFX_CudaStreamDestroy(CUstream stream); #define NVVFX_MODEL_BATCH "ModelBatch" //!< The preferred batching model to use (default 1) #define NVVFX_STATE "State" //!< State variable #define NVVFX_STATE_SIZE "StateSize" //!< Number of bytes needed to store state - +#define NVVFX_STATE_COUNT "NumStateObjects" //!< Number of active state object handles #ifdef __cplusplus diff --git a/nvvfx/src/NVVideoEffectsProxy.cpp b/nvvfx/src/NVVideoEffectsProxy.cpp index 631e81e..4166d1e 100644 --- a/nvvfx/src/NVVideoEffectsProxy.cpp +++ b/nvvfx/src/NVVideoEffectsProxy.cpp @@ -169,6 +169,13 @@ NvCV_Status NvVFX_API NvVFX_SetObject(NvVFX_Handle obj, NvVFX_ParameterSelector return funcPtr(obj, paramName, ptr); } +NvCV_Status NvVFX_API NvVFX_SetStateObjectHandleArray(NvVFX_Handle obj, NvVFX_ParameterSelector paramName, NvVFX_StateObjectHandle* handle) { + static const auto funcPtr = (decltype(NvVFX_SetStateObjectHandleArray)*)nvGetProcAddress(getNvVfxLib(), "NvVFX_SetStateObjectHandleArray"); + + if (nullptr == funcPtr) return NVCV_ERR_LIBRARY; + return funcPtr(obj, paramName, handle); +} + NvCV_Status NvVFX_API NvVFX_SetString(NvVFX_Handle obj, NvVFX_ParameterSelector paramName, const char* str) { static const auto funcPtr = (decltype(NvVFX_SetString)*)nvGetProcAddress(getNvVfxLib(), "NvVFX_SetString"); @@ -276,4 +283,25 @@ NvCV_Status NvVFX_API NvVFX_CudaStreamDestroy(CUstream stream) { return funcPtr(stream); } +NvCV_Status NvVFX_API NvVFX_AllocateState(NvVFX_Handle obj, NvVFX_StateObjectHandle* handle) { + static const auto funcPtr = (decltype(NvVFX_AllocateState)*)nvGetProcAddress(getNvVfxLib(), "NvVFX_AllocateState"); + + if (nullptr == funcPtr) return NVCV_ERR_LIBRARY; + return funcPtr(obj, handle); +} + +NvCV_Status NvVFX_API NvVFX_DeallocateState(NvVFX_Handle obj, NvVFX_StateObjectHandle handle) { + static const auto funcPtr = (decltype(NvVFX_DeallocateState)*)nvGetProcAddress(getNvVfxLib(), "NvVFX_DeallocateState"); + + if (nullptr == funcPtr) return NVCV_ERR_LIBRARY; + return funcPtr(obj, handle); +} + +NvCV_Status NvVFX_API NvVFX_ResetState(NvVFX_Handle obj, NvVFX_StateObjectHandle handle) { + static const auto funcPtr = (decltype(NvVFX_ResetState)*)nvGetProcAddress(getNvVfxLib(), "NvVFX_ResetState"); + + if (nullptr == funcPtr) return NVCV_ERR_LIBRARY; + return funcPtr(obj, handle); +} + #endif // enabling for this file diff --git a/nvvfx/src/nvCVImageProxy.cpp b/nvvfx/src/nvCVImageProxy.cpp index 4f5199a..4657939 100644 --- a/nvvfx/src/nvCVImageProxy.cpp +++ b/nvvfx/src/nvCVImageProxy.cpp @@ -132,10 +132,16 @@ NvCV_Status NvCV_API NvCVImage_Realloc(NvCVImage* im, unsigned width, unsigned h void NvCV_API NvCVImage_Dealloc(NvCVImage* im) { static const auto funcPtr = (decltype(NvCVImage_Dealloc)*)nvGetProcAddress(getNvCVImageLib(), "NvCVImage_Dealloc"); - + if (nullptr != funcPtr) funcPtr(im); } +void NvCV_API NvCVImage_DeallocAsync(NvCVImage* im, CUstream_st* stream) { + static const auto funcPtr = (decltype(NvCVImage_DeallocAsync)*)nvGetProcAddress(getNvCVImageLib(), "NvCVImage_DeallocAsync"); + + if (nullptr != funcPtr) funcPtr(im, stream); +} + NvCV_Status NvCV_API NvCVImage_Create(unsigned width, unsigned height, NvCVImage_PixelFormat format, NvCVImage_ComponentType type, unsigned isPlanar, unsigned onGPU, unsigned alignment, NvCVImage** out) { @@ -266,6 +272,14 @@ NvCV_Status NvCV_API NvCVImage_FlipY(const NvCVImage *src, NvCVImage *dst) { return funcPtr(src, dst); } +NvCV_Status NvCV_API NvCVImage_Sharpen(float sharpness, const NvCVImage *src, NvCVImage *dst, + struct CUstream_st *stream, NvCVImage *tmp) { + static const auto funcPtr = (decltype(NvCVImage_Sharpen)*)nvGetProcAddress(getNvCVImageLib(), "NvCVImage_Sharpen"); + + if (nullptr == funcPtr) return NVCV_ERR_LIBRARY; + return funcPtr(sharpness, src, dst, stream, tmp); +} + #ifdef _WIN32 __declspec(dllexport) const char* __cdecl #else diff --git a/samples/AigsEffectApp/AigsEffectApp.cpp b/samples/AigsEffectApp/AigsEffectApp.cpp index 981a44d..d914467 100644 --- a/samples/AigsEffectApp/AigsEffectApp.cpp +++ b/samples/AigsEffectApp/AigsEffectApp.cpp @@ -28,6 +28,7 @@ #include #include #include +#include #include "nvCVOpenCV.h" #include "nvVideoEffects.h" @@ -154,6 +155,7 @@ static void Usage() { " 4 (show input - compNone),\n" " 5 (composite over a specified background image - compBG),\n" " 6 (blur the background of the image - compBlur) }\n" + " --blur_strength=[0-1] strength of the background blur, when applicable\n" " --cuda_graph Enable cuda graph.\n" ); } @@ -330,14 +332,13 @@ struct FXApp { _framePeriod = 0.f; _lastTime = std::chrono::high_resolution_clock::time_point::min(); _blurStrength = 0.5f; + _maxInputWidth = 3840u; + _maxInputHeight = 2160u; + _maxNumberStreams = 1u; + _batchOfStates = nullptr; } ~FXApp() { - NvVFX_DestroyEffect(_eff); - NvVFX_DestroyEffect(_bgblurEff); - - if (_stream) { - NvVFX_CudaStreamDestroy(_stream); - } + destroyEffect(); } void setShow(bool show) { _show = show; } @@ -375,6 +376,11 @@ struct FXApp { NvCVImage _dstNvVFXImage; NvCVImage _blurNvVFXImage; float _blurStrength; + unsigned int _maxInputWidth; + unsigned int _maxInputHeight; + unsigned int _maxNumberStreams; + std::vector _stateArray; + NvVFX_StateObjectHandle* _batchOfStates; }; const char *FXApp::errorStringFromCode(Err code) { @@ -530,12 +536,41 @@ NvCV_Status FXApp::createAigsEffect() { return vfxErr; } + // Set maximum width, height and number of streams and then call Load() again + vfxErr = NvVFX_SetU32(_eff, NVVFX_MAX_INPUT_WIDTH, _maxInputWidth); + if (vfxErr != NVCV_SUCCESS) { + std::cerr << "Error setting the mode \n"; + return vfxErr; + } + + vfxErr = NvVFX_SetU32(_eff, NVVFX_MAX_INPUT_HEIGHT, _maxInputHeight); + if (vfxErr != NVCV_SUCCESS) { + std::cerr << "Error setting the mode \n"; + return vfxErr; + } + + vfxErr = NvVFX_SetU32(_eff, NVVFX_MAX_NUMBER_STREAMS, _maxNumberStreams); + if (vfxErr != NVCV_SUCCESS) { + std::cerr << "Error setting the mode \n"; + return vfxErr; + } + vfxErr = NvVFX_Load(_eff); if (vfxErr != NVCV_SUCCESS) { std::cerr << "Error loading the model \n"; return vfxErr; } + for (unsigned int i = 0; i < _maxNumberStreams; i++) { + NvVFX_StateObjectHandle state; + vfxErr = NvVFX_AllocateState(_eff, &state); + if (NVCV_SUCCESS != vfxErr) { + std::cerr << "Error allocate state variable for effect \"" << NVVFX_FX_GREEN_SCREEN << "\"\n"; + return vfxErr; + } + _stateArray.push_back(state); + } + // ------------------ create Background blur effect ------------------ // vfxErr = NvVFX_CreateEffect(NVVFX_FX_BGBLUR, &_bgblurEff); if (NVCV_SUCCESS != vfxErr) { @@ -559,8 +594,26 @@ NvCV_Status FXApp::createAigsEffect() { } void FXApp::destroyEffect() { + // If DeallocateState fails, all memory allocated in the SDK returns to the heap when the effect handle is destroyed. + for (unsigned int i = 0; i < _stateArray.size(); i++) { + NvVFX_DeallocateState(_eff, _stateArray[i]); + } + _stateArray.clear(); + + if (_batchOfStates != nullptr) { + free(_batchOfStates); + _batchOfStates = nullptr; + } + NvVFX_DestroyEffect(_eff); _eff = nullptr; + + NvVFX_DestroyEffect(_bgblurEff); + _bgblurEff = nullptr; + + if (_stream) { + NvVFX_CudaStreamDestroy(_stream); + } } static void overlay(const cv::Mat &image, const cv::Mat &mask, float alpha, cv::Mat &result) { @@ -573,6 +626,17 @@ FXApp::Err FXApp::processImage(const char *inFile, const char *outFile) { NvCV_Status vfxErr; bool ok; cv::Mat result; + NvCVImage fxSrcChunkyGPU, fxDstChunkyGPU; + + // Allocate space for batchOfStates to hold state variable addresses + // Assume that MODEL_BATCH Size is enough for this scenario + unsigned int modelBatch = 1; + BAIL_IF_ERR(vfxErr = NvVFX_GetU32(_eff, NVVFX_MODEL_BATCH, &modelBatch)); + _batchOfStates = (NvVFX_StateObjectHandle*) malloc(sizeof(NvVFX_StateObjectHandle) * modelBatch); + if (_batchOfStates == nullptr) { + vfxErr = NVCV_ERR_MEMORY; + goto bail; + } if (!_eff) return errEffect; _srcImg = cv::imread(inFile); @@ -584,12 +648,26 @@ FXApp::Err FXApp::processImage(const char *inFile, const char *outFile) { (void)NVWrapperForCVMat(&_srcImg, &_srcVFX); (void)NVWrapperForCVMat(&_dstImg, &_dstVFX); - NvCVImage fxSrcChunkyGPU(_srcImg.cols, _srcImg.rows, NVCV_BGR, NVCV_U8, NVCV_CHUNKY, NVCV_GPU, 1); - NvCVImage fxDstChunkyGPU(_srcImg.cols, _srcImg.rows, NVCV_A, NVCV_U8, NVCV_CHUNKY, NVCV_GPU, 1); + if (!fxSrcChunkyGPU.pixels) + { + BAIL_IF_ERR(vfxErr = + NvCVImage_Alloc(&fxSrcChunkyGPU, _srcImg.cols, _srcImg.rows, NVCV_BGR, NVCV_U8, NVCV_CHUNKY, NVCV_GPU, 1)); + } + + if (!fxDstChunkyGPU.pixels) { + BAIL_IF_ERR(vfxErr = + NvCVImage_Alloc(&fxDstChunkyGPU, _srcImg.cols, _srcImg.rows, NVCV_A, NVCV_U8, NVCV_CHUNKY, NVCV_GPU, 1)); + } BAIL_IF_ERR(vfxErr = NvVFX_SetImage(_eff, NVVFX_INPUT_IMAGE, &fxSrcChunkyGPU)); BAIL_IF_ERR(vfxErr = NvVFX_SetImage(_eff, NVVFX_OUTPUT_IMAGE, &fxDstChunkyGPU)); BAIL_IF_ERR(vfxErr = NvCVImage_Transfer(&_srcVFX, &fxSrcChunkyGPU, 1.0f, _stream, NULL)); + + // Assign states from stateArray in batchOfStates + // There is only one stream in this app + _batchOfStates[0] = _stateArray[0]; + BAIL_IF_ERR(vfxErr = NvVFX_SetStateObjectHandleArray(_eff, NVVFX_STATE, _batchOfStates)); + BAIL_IF_ERR(vfxErr = NvVFX_Run(_eff, 0)); BAIL_IF_ERR(vfxErr = NvCVImage_Transfer(&fxDstChunkyGPU, &_dstVFX, 1.0f, _stream, NULL)); @@ -627,6 +705,7 @@ FXApp::Err FXApp::processMovie(const char *inFile, const char *outFile) { cv::VideoWriter writer; unsigned frameNum; VideoInfo info; + unsigned int modelBatch = 1; if (inFile && !inFile[0]) inFile = nullptr; // Set file paths to NULL if zero length if (outFile && !outFile[0]) outFile = nullptr; @@ -698,6 +777,15 @@ FXApp::Err FXApp::processMovie(const char *inFile, const char *outFile) { } } + // Allocate space for batchOfStates to hold state variable addresses + // Assume that MODEL_BATCH Size is enough for this scenario + BAIL_IF_ERR(vfxErr = NvVFX_GetU32(_eff, NVVFX_MODEL_BATCH, &modelBatch)); + _batchOfStates = (NvVFX_StateObjectHandle*) malloc(sizeof(NvVFX_StateObjectHandle) * modelBatch); + if (_batchOfStates == nullptr) { + vfxErr = NVCV_ERR_MEMORY; + goto bail; + } + // allocate src for GPU if (!_srcNvVFXImage.pixels) BAIL_IF_ERR(vfxErr = @@ -726,6 +814,11 @@ FXApp::Err FXApp::processMovie(const char *inFile, const char *outFile) { BAIL_IF_ERR(vfxErr = NvVFX_SetImage(_eff, NVVFX_OUTPUT_IMAGE, &_dstNvVFXImage)); BAIL_IF_ERR(vfxErr = NvCVImage_Transfer(&_srcVFX, &_srcNvVFXImage, 1.0f, _stream, NULL)); + // Assign states from stateArray in batchOfStates + // There is only one stream in this app + _batchOfStates[0] = _stateArray[0]; + BAIL_IF_ERR(vfxErr = NvVFX_SetStateObjectHandleArray(_eff, NVVFX_STATE, _batchOfStates)); + auto startTime = std::chrono::high_resolution_clock::now(); BAIL_IF_ERR(vfxErr = NvVFX_Run(_eff, 0)); auto endTime = std::chrono::high_resolution_clock::now(); diff --git a/samples/AigsEffectApp/AigsEffectApp.exe b/samples/AigsEffectApp/AigsEffectApp.exe index 02274be..48bc6ab 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 fcb804d..1ad9c9a 100644 --- a/samples/AigsEffectApp/CMakeLists.txt +++ b/samples/AigsEffectApp/CMakeLists.txt @@ -15,17 +15,29 @@ target_include_directories(AigsEffectApp PUBLIC ${SDK_INCLUDES_PATH} ) -target_link_libraries(AigsEffectApp PUBLIC - opencv346 - NVVideoEffects - ${CMAKE_CURRENT_SOURCE_DIR}/../external/cuda/lib/x64/cudart.lib - ) +if(MSVC) -set(OPENCV_PATH_STR ${CMAKE_CURRENT_SOURCE_DIR}/../external/opencv/bin) -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 + target_link_libraries(AigsEffectApp 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(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}" + ) +else() + + target_link_libraries(AigsEffectApp PUBLIC + NVVideoEffects + NVCVImage + OpenCV + TensorRT + CUDA + ) +endif() diff --git a/samples/BatchEffectApp/BatchAigsEffectApp.cpp b/samples/BatchEffectApp/BatchAigsEffectApp.cpp new file mode 100644 index 0000000..2e9e8bd --- /dev/null +++ b/samples/BatchEffectApp/BatchAigsEffectApp.cpp @@ -0,0 +1,370 @@ +/*############################################################################### +# +# Copyright (c) 2022 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. +# +###############################################################################*/ + +#include +#include + +#include +#include "BatchUtilities.h" +#include "nvCVOpenCV.h" +#include "nvVideoEffects.h" +#include "opencv2/opencv.hpp" + +#ifdef _MSC_VER + #define strcasecmp _stricmp +#endif // _MSC_VER + +#define BAIL_IF_ERR(err) do { if (0 != (err)) { goto bail; } } while(0) +#define BAIL_IF_NULL(x, err, code) do { if ((void*)(x) == NULL) { err = code; goto bail; } } while(0) +#define BAIL_IF_FALSE(x, err, code) do { if (!(x)) { err = code; goto bail; } } while(0) +#define BAIL(err, code) do { err = code; goto bail; } while(0) + +#ifdef _WIN32 + #define DEFAULT_CODEC "avc1" +#else // !_WIN32 + #define DEFAULT_CODEC "H264" +#endif // _WIN32 + +bool FLAG_verbose = false; +int FLAG_mode = 0; +std::string FLAG_outFile, + FLAG_modelDir, + FLAG_codec = DEFAULT_CODEC; +std::vector FLAG_inFiles; + +// 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; + while (*++arg == '-') + continue; + const char *s = strchr(arg, '='); + if (s == NULL) { + if (strcmp(flag, arg) != 0) + return false; + *val = NULL; + return true; + } + size_t n = s - arg; + if ((strlen(flag) != n) || (strncmp(flag, arg, n) != 0)) + return false; + *val = s + 1; + return true; +} + +static bool GetFlagArgVal(const char *flag, const char *arg, std::string *val) { + const char *valStr; + if (!GetFlagArgVal(flag, arg, &valStr)) + return false; + val->assign(valStr ? valStr : ""); + return true; +} + +static bool GetFlagArgVal(const char *flag, const char *arg, bool *val) { + const char *valStr; + bool success = GetFlagArgVal(flag, arg, &valStr); + if (success) { + *val = (valStr == NULL || + strcasecmp(valStr, "true") == 0 || + strcasecmp(valStr, "on") == 0 || + strcasecmp(valStr, "yes") == 0 || + strcasecmp(valStr, "1") == 0 + ); + } + return success; +} + +static bool GetFlagArgVal(const char *flag, const char *arg, long *val) { + const char *valStr; + bool success = GetFlagArgVal(flag, arg, &valStr); + if (success) + *val = strtol(valStr, NULL, 10); + return success; +} + +static bool GetFlagArgVal(const char *flag, const char *arg, int *val) { + long longVal; + bool success = GetFlagArgVal(flag, arg, &longVal); + if (success) + *val = (int)longVal; + return success; +} + +static int StringToFourcc(const std::string &str) { + union chint { + int i; + char c[4]; + }; + chint x = {0}; + for (int n = (str.size() < 4) ? (int)str.size() : 4; n--;) x.c[n] = str[n]; + return x.i; +} + +static void Usage() { + printf( + "BatchAigsEffectApp [flags ...] inFile1 [ inFileN ...]\n" + " where flags is:\n" + " --out_file= output video files to be written (a pattern with one %%u or %%d), default \"BatchOut_%%02u.mp4\"\n" + " --model_dir= the path to the directory that contains the models\n" + " --mode= which model to pick for processing (default: 0)\n" + " --verbose verbose output\n" + " --codec= the fourcc code for the desired codec (default " DEFAULT_CODEC ")\n" + " and inFile1 ... are identically sized video files\n" + ); +} + +static int ParseMyArgs(int argc, char **argv) { + int errs = 0; + for (--argc, ++argv; argc--; ++argv) { + bool help; + const char *arg = *argv; + if (arg[0] == '-') { + if (arg[1] == '-') { // double-dash + if (GetFlagArgVal("verbose", arg, &FLAG_verbose) || + GetFlagArgVal("mode", arg, &FLAG_mode) || + GetFlagArgVal("model_dir", arg, &FLAG_modelDir) || + GetFlagArgVal("out_file", arg, &FLAG_outFile) || + GetFlagArgVal("codec", arg, &FLAG_codec) + ) { + continue; + } else if (GetFlagArgVal("help", arg, &help)) { // --help + Usage(); + errs = 1; + } + } + else { // single dash + for (++arg; *arg; ++arg) { + if (*arg == 'v') { + FLAG_verbose = true; + } else { + printf("Unknown flag ignored: \"-%c\"\n", *arg); + } + } + continue; + } + } + else { // no dash + FLAG_inFiles.push_back(arg); + } + } + return errs; +} + + +class App { +public: + NvVFX_Handle _eff; + NvCVImage _src, _stg, _dst; + CUstream _stream; + unsigned _batchSize; + + + App() : _eff(nullptr), _stream(0), _batchSize(0) {} + ~App() { + NvVFX_DestroyEffect(_eff); if (_stream) NvVFX_CudaStreamDestroy(_stream); + } + + NvCV_Status init(const char* effectName, unsigned batchSize, unsigned int mode, const NvCVImage *srcImg) { + NvCV_Status err = NVCV_ERR_UNIMPLEMENTED; + + _batchSize = batchSize; + BAIL_IF_ERR(err = NvVFX_CreateEffect(effectName, &_eff)); + + BAIL_IF_ERR(err = AllocateBatchBuffer(&_src, _batchSize, srcImg->width, srcImg->height, NVCV_BGR, NVCV_U8, NVCV_CHUNKY, NVCV_GPU, 1)); + BAIL_IF_ERR(err = AllocateBatchBuffer(&_dst, _batchSize, srcImg->width, srcImg->height, NVCV_A, NVCV_U8, NVCV_CHUNKY, NVCV_GPU, 1)); + BAIL_IF_ERR(err = NvVFX_SetString(_eff, NVVFX_MODEL_DIRECTORY, FLAG_modelDir.c_str())); + + + { // Set parameters. + NvCVImage nth; + BAIL_IF_ERR(err = NvVFX_SetImage(_eff, NVVFX_INPUT_IMAGE, NthImage(0, srcImg->height, &_src, &nth))); // Set the first of the batched images in ... + BAIL_IF_ERR(err = NvVFX_SetImage(_eff, NVVFX_OUTPUT_IMAGE, NthImage(0, _dst.height / _batchSize, &_dst, &nth))); // ... and out + BAIL_IF_ERR(err = NvVFX_CudaStreamCreate(&_stream)); + BAIL_IF_ERR(err = NvVFX_SetCudaStream(_eff, NVVFX_CUDA_STREAM, _stream)); + BAIL_IF_ERR(err = NvVFX_SetU32(_eff, NVVFX_MODE, mode)); + } + + bail: + return err; + } +}; + + +NvCV_Status BatchProcess(const char* effectName, unsigned int mode, + const std::vector& srcVideos, const char *outfilePattern, std::string codec) { + NvCV_Status err = NVCV_SUCCESS; + App app; + cv::Mat ocv1, ocv2; + NvCVImage nvx1, nvx2; + unsigned srcWidth, srcHeight, dstHeight; + + std::vector arrayOfStates; + NvVFX_StateObjectHandle* batchOfStates = nullptr; + + unsigned int numOfVideoStreams = static_cast(srcVideos.size()); + + // If valid states are passed for inference, then - + // 1. Effect can only process a batch which is equal to maximum number of video streams + // 2. Multiple frames from the same video stream should not be present in the same batch + unsigned batchSize = numOfVideoStreams; + + std::vector srcCaptures(numOfVideoStreams); + std::vector dstWriters(numOfVideoStreams); + for (unsigned int i = 0; i < numOfVideoStreams; i++) { + srcCaptures[i].open(srcVideos[i]); + if (srcCaptures[i].isOpened()==false) BAIL(err, NVCV_ERR_READ); + + int width, height; + double fps; + width = (int)srcCaptures[i].get(cv::CAP_PROP_FRAME_WIDTH); + height = (int)srcCaptures[i].get(cv::CAP_PROP_FRAME_HEIGHT); + fps = srcCaptures[i].get(cv::CAP_PROP_FPS); + + const int fourcc = StringToFourcc(codec); + char fileName[1024]; + snprintf(fileName, sizeof(fileName), outfilePattern, i); + dstWriters[i].open(fileName, fourcc, fps, cv::Size2i(width,height), false); + if (dstWriters[i].isOpened() == false) BAIL(err, NVCV_ERR_WRITE); + } + + // Read in the first image, to determine the resolution for init() + BAIL_IF_FALSE(srcVideos.size() > 0, err, NVCV_ERR_MISSINGINPUT); + srcCaptures[0] >> ocv1; + srcCaptures[0].set(cv::CAP_PROP_POS_FRAMES, 0); //resetting to first frame + if (!ocv1.data) { + printf("Cannot read video file \"%s\"\n", srcVideos[0]); + BAIL(err, NVCV_ERR_READ); + } + NVWrapperForCVMat(&ocv1, &nvx1); + srcWidth = nvx1.width; + srcHeight = nvx1.height; + + BAIL_IF_ERR(err = app.init(effectName, batchSize, mode, &nvx1)); // Init effect and buffers + BAIL_IF_ERR(err = NvVFX_SetU32(app._eff, NVVFX_MAX_NUMBER_STREAMS, numOfVideoStreams)); + BAIL_IF_ERR(err = NvVFX_SetU32(app._eff, NVVFX_MODEL_BATCH, numOfVideoStreams>1?8:1)); + BAIL_IF_ERR(err = NvVFX_Load(app._eff)); + + // Creating state objects, one per stream. + for (unsigned int i = 0; i < numOfVideoStreams; i++) { + NvVFX_StateObjectHandle state; + BAIL_IF_ERR(err = NvVFX_AllocateState(app._eff, &state)); + arrayOfStates.push_back(state); + } + + //Creating batch array to hold states + batchOfStates = (NvVFX_StateObjectHandle*)malloc(sizeof(NvVFX_StateObjectHandle) * batchSize); + if (batchOfStates == nullptr) { + err = NVCV_ERR_MEMORY; + goto bail; + } + + dstHeight = app._dst.height / batchSize; + BAIL_IF_ERR(err = NvCVImage_Alloc(&nvx2, app._dst.width, dstHeight, NVCV_A, NVCV_U8, NVCV_CHUNKY, NVCV_CPU, 0)); + CVWrapperForNvCVImage(&nvx2, &ocv2); + for(int j=0;;j++) + { + for (unsigned int i = 0; i < batchSize; i++) { + int capIdx = i%numOfVideoStreams; // interlacing frames from different video stream, but can in any order + srcCaptures[capIdx] >> ocv1; + if (ocv1.empty()) goto bail; + batchOfStates[i] = arrayOfStates[capIdx]; + + NVWrapperForCVMat(&ocv1, &nvx1); + if (!(nvx1.width == srcWidth && nvx1.height == srcHeight)) { + printf("Input video file \"%s\" %ux%u does not match %ux%u\n" + "Batching requires all video frames to be of the same size\n", srcVideos[i], nvx1.width, nvx1.height, srcWidth, srcHeight); + BAIL(err, NVCV_ERR_MISMATCH); + } + BAIL_IF_ERR(err = TransferToNthImage(i, &nvx1, &app._src, 1.f, app._stream, NULL)); + ocv1.release(); + } + + // Run batch + BAIL_IF_ERR(err = NvVFX_SetU32(app._eff, NVVFX_BATCH_SIZE, (unsigned)batchSize)); // The batchSize can change every Run + BAIL_IF_ERR(err = NvVFX_SetStateObjectHandleArray(app._eff, NVVFX_STATE, batchOfStates)); // The batch of states can change every Run + BAIL_IF_ERR(err = NvVFX_Run(app._eff, 0)); + + + for (unsigned int i = 0; i < batchSize; ++i) { + int writerIdx = i % numOfVideoStreams; + BAIL_IF_ERR(err = TransferFromNthImage(i, &app._dst, &nvx2, 1.0f, app._stream, NULL)); + dstWriters[writerIdx] << ocv2; + } + // NvCVImage_Dealloc() is called in the destructors + } +bail: + // If DeallocateState fails, all memory allocated in the SDK returns to the heap when the effect handle is destroyed. + for (unsigned int i = 0; i < arrayOfStates.size(); i++) { + NvVFX_DeallocateState(app._eff, arrayOfStates[i]); + } + arrayOfStates.clear(); + + if (batchOfStates) { + free(batchOfStates); + batchOfStates = nullptr; + } + + for (auto& cap : srcCaptures) { + if (cap.isOpened()) cap.release(); + } + + for (auto& writer : dstWriters) { + if (writer.isOpened()) writer.release(); + } + + return err; +} + + +int main(int argc, char** argv) { + int nErrs; + NvCV_Status vfxErr; + + nErrs = ParseMyArgs(argc, argv); + if (nErrs) + return nErrs; + + // If the outFile is missing a stream index + // insert one, assuming a period followed by a three-character extension + if (FLAG_outFile.empty()) + FLAG_outFile = "BatchOut_%02u.mp4"; + else if (std::string::npos == FLAG_outFile.find_first_of('%')) + FLAG_outFile.insert(FLAG_outFile.size() - 4, "_%02u"); + +#ifdef NVVFX_FX_GREEN_SCREEN + vfxErr = BatchProcess(NVVFX_FX_GREEN_SCREEN, FLAG_mode, FLAG_inFiles, FLAG_outFile.c_str(), FLAG_codec); +#elif defined(NVVFX_FX_GREEN_SCREEN_I) + vfxErr = BatchProcess(NVVFX_FX_GREEN_SCREEN_I, FLAG_mode, FLAG_inFiles, FLAG_outFile.c_str(), FLAG_codec); +#endif + if (NVCV_SUCCESS != vfxErr) { + Usage(); + printf("Error: %s\n", NvCV_GetErrorStringFromCode(vfxErr)); + nErrs = (int)vfxErr; + } + + return nErrs; +} diff --git a/samples/BatchEffectApp/BatchAigsEffectApp.exe b/samples/BatchEffectApp/BatchAigsEffectApp.exe new file mode 100644 index 0000000..cb18344 Binary files /dev/null and b/samples/BatchEffectApp/BatchAigsEffectApp.exe differ diff --git a/samples/BatchEffectApp/BatchDenoiseEffectApp.cpp b/samples/BatchEffectApp/BatchDenoiseEffectApp.cpp index 58dd01e..61854dc 100644 --- a/samples/BatchEffectApp/BatchDenoiseEffectApp.cpp +++ b/samples/BatchEffectApp/BatchDenoiseEffectApp.cpp @@ -217,16 +217,16 @@ NvCV_Status BatchProcess(const char* effectName, const std::vector& App app; cv::Mat ocv1, ocv2; NvCVImage nvx1, nvx2; - unsigned srcWidth, srcHeight, dstHeight, i; + unsigned srcWidth, srcHeight, dstHeight; void** arrayOfStates = nullptr; void** batchOfStates = nullptr; unsigned int stateSizeInBytes; - int numOfVideoStreams = srcVideos.size(); + unsigned int numOfVideoStreams = static_cast(srcVideos.size()); std::vector srcCaptures(numOfVideoStreams); std::vector dstWriters(numOfVideoStreams); - for (int i = 0; i < numOfVideoStreams; i++) { + for (unsigned int i = 0; i < numOfVideoStreams; i++) { srcCaptures[i].open(srcVideos[i]); if (srcCaptures[i].isOpened()==false) BAIL(err, NVCV_ERR_READ); @@ -260,7 +260,7 @@ NvCV_Status BatchProcess(const char* effectName, const std::vector& // Creating state objects, one per stream. BAIL_IF_ERR(err = NvVFX_GetU32(app._eff, NVVFX_STATE_SIZE, &stateSizeInBytes)); arrayOfStates = (void**)calloc(numOfVideoStreams, sizeof(void*)); // allocating void* array of numOfVideoStreams elements - for (int i = 0; i < numOfVideoStreams; i++) { + for (unsigned int i = 0; i < numOfVideoStreams; i++) { cudaMalloc(&arrayOfStates[i], stateSizeInBytes); cudaMemsetAsync(arrayOfStates[i], 0, stateSizeInBytes,app._stream); } @@ -273,7 +273,7 @@ NvCV_Status BatchProcess(const char* effectName, const std::vector& CVWrapperForNvCVImage(&nvx2, &ocv2); for(int j=0;;j++) { - for (int i = 0; i < batchSize; i++) { + for (unsigned int i = 0; i < batchSize; i++) { int capIdx = i%numOfVideoStreams; // interlacing frames from different video stream, but can in any order srcCaptures[capIdx] >> ocv1; if (ocv1.empty()) goto bail; @@ -295,7 +295,7 @@ NvCV_Status BatchProcess(const char* effectName, const std::vector& BAIL_IF_ERR(err = NvVFX_Run(app._eff, 0)); - for (i = 0; i < batchSize; ++i) { + for (unsigned int i = 0; i < batchSize; ++i) { int writerIdx = i % numOfVideoStreams; BAIL_IF_ERR(err = TransferFromNthImage(i, &app._dst, &nvx2, 255.f, app._stream, &app._stg)); dstWriters[writerIdx] << ocv2; @@ -304,7 +304,7 @@ NvCV_Status BatchProcess(const char* effectName, const std::vector& } bail: if (arrayOfStates) { - for (unsigned i = 0; i < numOfVideoStreams; i++) { + for (unsigned int i = 0; i < numOfVideoStreams; i++) { if (arrayOfStates[i]) cudaFree(arrayOfStates[i]); } free(arrayOfStates); diff --git a/samples/BatchEffectApp/BatchDenoiseEffectApp.exe b/samples/BatchEffectApp/BatchDenoiseEffectApp.exe index 7a056fc..9af42f9 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 2e67dcc..ae0380a 100644 --- a/samples/BatchEffectApp/BatchEffectApp.cpp +++ b/samples/BatchEffectApp/BatchEffectApp.cpp @@ -244,16 +244,9 @@ public: else if (!strcmp(effectName, NVVFX_FX_SR_UPSCALE)) { BAIL_IF_ERR(err = AllocateBatchBuffer(&_src, _batchSize, src->width, src->height, NVCV_RGBA, NVCV_U8, NVCV_CHUNKY, NVCV_CUDA, 32)); // n*32, n>=0 BAIL_IF_ERR(err = AllocateBatchBuffer(&_dst, _batchSize, dw, dh, NVCV_RGBA, NVCV_U8, NVCV_CHUNKY, NVCV_CUDA, 32)); + BAIL_IF_ERR(err = NvVFX_SetF32(_eff, NVVFX_STRENGTH, FLAG_strength)); } #endif // NVVFX_FX_SR_UPSCALE -#ifdef NVVFX_FX_GREEN_SCREEN - else if (!strcmp(effectName, NVVFX_FX_GREEN_SCREEN)) { - BAIL_IF_ERR(err = AllocateBatchBuffer(&_src, _batchSize, src->width, src->height, NVCV_BGR, NVCV_U8, NVCV_CHUNKY, NVCV_CUDA, 1)); - BAIL_IF_ERR(err = AllocateBatchBuffer(&_dst, _batchSize, src->width, src->height, NVCV_Y, NVCV_U8, NVCV_CHUNKY, 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_MODE, FLAG_mode)); - } -#endif // NVVFX_FX_GREEN_SCREEN #ifdef NVVFX_FX_ARTIFACT_REDUCTION else if (!strcmp(effectName, NVVFX_FX_ARTIFACT_REDUCTION)) { BAIL_IF_ERR(err = AllocateBatchBuffer(&_src, _batchSize, src->width, src->height, NVCV_BGR, NVCV_F32, NVCV_PLANAR, NVCV_CUDA, 1)); @@ -268,6 +261,7 @@ public: 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_MODE, FLAG_mode)); + BAIL_IF_ERR(err = NvVFX_SetF32(_eff, NVVFX_STRENGTH, FLAG_strength)); } #endif // NVVFX_FX_SUPER_RES else { diff --git a/samples/BatchEffectApp/BatchEffectApp.exe b/samples/BatchEffectApp/BatchEffectApp.exe index 58144a7..0582115 100644 Binary files a/samples/BatchEffectApp/BatchEffectApp.exe and b/samples/BatchEffectApp/BatchEffectApp.exe differ diff --git a/samples/BatchEffectApp/CMakeLists.txt b/samples/BatchEffectApp/CMakeLists.txt index dee8de8..83db6f9 100644 --- a/samples/BatchEffectApp/CMakeLists.txt +++ b/samples/BatchEffectApp/CMakeLists.txt @@ -17,20 +17,32 @@ target_include_directories(BatchEffectApp PUBLIC ${SDK_INCLUDES_PATH} ) -target_link_libraries(BatchEffectApp PUBLIC - opencv346 - NVVideoEffects - ${CMAKE_CURRENT_SOURCE_DIR}/../external/cuda/lib/x64/cudart.lib - ) +if(MSVC) -set(OPENCV_PATH_STR ${CMAKE_CURRENT_SOURCE_DIR}/../external/opencv/bin) -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(BatchEffectApp PROPERTIES - FOLDER SampleApps - VS_DEBUGGER_ENVIRONMENT "${PATH_STR}" - VS_DEBUGGER_COMMAND_ARGUMENTS "${CMD_ARG_STR}" - ) + target_link_libraries(BatchEffectApp 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(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(BatchEffectApp PROPERTIES + FOLDER SampleApps + VS_DEBUGGER_ENVIRONMENT "${PATH_STR}" + VS_DEBUGGER_COMMAND_ARGUMENTS "${CMD_ARG_STR}" + ) +else() + + target_link_libraries(BatchEffectApp PUBLIC + NVVideoEffects + NVCVImage + OpenCV + TensorRT + CUDA + ) +endif() #Batch denoise effect set(SOURCE_FILES @@ -51,19 +63,75 @@ target_include_directories(BatchDenoiseEffectApp PUBLIC ${SDK_INCLUDES_PATH} ) - -target_link_libraries(BatchDenoiseEffectApp PUBLIC - opencv346 - NVVideoEffects - ${CMAKE_CURRENT_SOURCE_DIR}/../external/cuda/lib/x64/cudart.lib - ) -target_include_directories(BatchDenoiseEffectApp PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}/../external/cuda/include) +if(MSVC) + target_link_libraries(BatchDenoiseEffectApp PUBLIC + opencv346 + NVVideoEffects + ${CMAKE_CURRENT_SOURCE_DIR}/../external/cuda/lib/x64/cudart.lib + ) + target_include_directories(BatchDenoiseEffectApp PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}/../external/cuda/include) -set(OPENCV_PATH_STR ${CMAKE_CURRENT_SOURCE_DIR}/../external/opencv/bin) -set(PATH_STR "PATH=%PATH%" ${OPENCV_PATH_STR}) -set(CMD_ARG_STR "video1.mp4 video2.mp4 ") -set_target_properties(BatchDenoiseEffectApp PROPERTIES - FOLDER SampleApps - VS_DEBUGGER_ENVIRONMENT "${PATH_STR}" - VS_DEBUGGER_COMMAND_ARGUMENTS "${CMD_ARG_STR}" - ) + set(OPENCV_PATH_STR ${CMAKE_CURRENT_SOURCE_DIR}/../external/opencv/bin) + set(PATH_STR "PATH=%PATH%" ${OPENCV_PATH_STR}) + set(CMD_ARG_STR "video1.mp4 video2.mp4 ") + set_target_properties(BatchDenoiseEffectApp PROPERTIES + FOLDER SampleApps + VS_DEBUGGER_ENVIRONMENT "${PATH_STR}" + VS_DEBUGGER_COMMAND_ARGUMENTS "${CMD_ARG_STR}" + ) +else() + + target_link_libraries(BatchDenoiseEffectApp PUBLIC + NVVideoEffects + NVCVImage + OpenCV + TensorRT + CUDA + ) +endif() + +#Batch aigs effect +set(SOURCE_FILES + BatchAigsEffectApp.cpp + BatchUtilities.cpp + ../../nvvfx/src/nvVideoEffectsProxy.cpp + ../../nvvfx/src/nvCVImageProxy.cpp) + +# Set Visual Studio source filters +source_group("Source Files" FILES ${SOURCE_FILES}) + +add_executable(BatchAigsEffectApp ${SOURCE_FILES}) +target_include_directories(BatchAigsEffectApp PRIVATE + ${CMAKE_CURRENT_SOURCE_DIR} + ${CMAKE_CURRENT_SOURCE_DIR}/../utils + ) +target_include_directories(BatchAigsEffectApp PUBLIC + ${SDK_INCLUDES_PATH} + ) + +if(MSVC) + target_link_libraries(BatchAigsEffectApp PUBLIC + opencv346 + NVVideoEffects + ${CMAKE_CURRENT_SOURCE_DIR}/../external/cuda/lib/x64/cudart.lib + ) + target_include_directories(BatchAigsEffectApp PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}/../external/cuda/include) + + set(OPENCV_PATH_STR ${CMAKE_CURRENT_SOURCE_DIR}/../external/opencv/bin) + set(PATH_STR "PATH=%PATH%" ${OPENCV_PATH_STR}) + set(CMD_ARG_STR "video1.mp4 video2.mp4 ") + set_target_properties(BatchAigsEffectApp PROPERTIES + FOLDER SampleApps + VS_DEBUGGER_ENVIRONMENT "${PATH_STR}" + VS_DEBUGGER_COMMAND_ARGUMENTS "${CMD_ARG_STR}" + ) +else() + + target_link_libraries(BatchAigsEffectApp PUBLIC + NVVideoEffects + NVCVImage + OpenCV + TensorRT + CUDA + ) +endif() diff --git a/samples/BatchEffectApp/run.bat b/samples/BatchEffectApp/run.bat index cc804b6..934633a 100644 --- a/samples/BatchEffectApp/run.bat +++ b/samples/BatchEffectApp/run.bat @@ -1,7 +1,9 @@ SETLOCAL SET PATH=%PATH%;..\external\opencv\bin; SET IMAGE_LIST=..\input\LeFret_000900.jpg ..\input\LeFret_001400.jpg ..\input\LeFret_003400.jpg ..\input\LeFret_012300.jpg -BatchEffectApp.exe --effect=GreenScreen --out_file=GreenScreen_%%04u.png %IMAGE_LIST% BatchEffectApp.exe --effect=ArtifactReduction --out_file=ArtifactReduction_%%04u.png %IMAGE_LIST% BatchEffectApp.exe --effect=SuperRes --out_file=SuperRes_%%04u.png --scale=1.5 %IMAGE_LIST% -BatchEffectApp.exe --effect=Upscale --out_file=Upscale_%%04u.png --scale=1.5 %IMAGE_LIST% \ No newline at end of file +BatchEffectApp.exe --effect=Upscale --out_file=Upscale_%%04u.png --scale=1.5 %IMAGE_LIST% + +SET VIDEO_LIST=..\input\input_0_100_frames.mp4 ..\input\input_100_200_frames.mp4 +BatchAigsEffectApp.exe --out_file=GreenScreen_%04u.mp4 %VIDEO_LIST% \ No newline at end of file diff --git a/samples/DenoiseEffectApp/CMakeLists.txt b/samples/DenoiseEffectApp/CMakeLists.txt index ba86e9f..c0d8b95 100644 --- a/samples/DenoiseEffectApp/CMakeLists.txt +++ b/samples/DenoiseEffectApp/CMakeLists.txt @@ -7,20 +7,29 @@ add_executable(DenoiseEffectApp ${SOURCE_FILES}) target_include_directories(DenoiseEffectApp PRIVATE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/../utils) target_include_directories(DenoiseEffectApp PUBLIC ${SDK_INCLUDES_PATH}) +if(MSVC) + target_link_libraries(DenoiseEffectApp PUBLIC + opencv346 + NVVideoEffects + ${CMAKE_CURRENT_SOURCE_DIR}/../external/cuda/lib/x64/cudart.lib + ) + target_include_directories(DenoiseEffectApp PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}/../external/cuda/include) + 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 --webcam") + set_target_properties(DenoiseEffectApp PROPERTIES + FOLDER SampleApps + VS_DEBUGGER_ENVIRONMENT "${PATH_STR}" + VS_DEBUGGER_COMMAND_ARGUMENTS "${CMD_ARG_STR}" + ) +else() -target_link_libraries(DenoiseEffectApp PUBLIC - opencv346 - NVVideoEffects - ${CMAKE_CURRENT_SOURCE_DIR}/../external/cuda/lib/x64/cudart.lib - ) -target_include_directories(DenoiseEffectApp PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}/../external/cuda/include) -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 --webcam") -set_target_properties(DenoiseEffectApp PROPERTIES - FOLDER SampleApps - VS_DEBUGGER_ENVIRONMENT "${PATH_STR}" - VS_DEBUGGER_COMMAND_ARGUMENTS "${CMD_ARG_STR}" - ) - + target_link_libraries(DenoiseEffectApp PUBLIC + NVVideoEffects + NVCVImage + OpenCV + TensorRT + CUDA + ) +endif() diff --git a/samples/DenoiseEffectApp/DenoiseEffectApp.exe b/samples/DenoiseEffectApp/DenoiseEffectApp.exe index bd49a7d..9994382 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 7f441e8..70fd3df 100644 --- a/samples/UpscalePipelineApp/CMakeLists.txt +++ b/samples/UpscalePipelineApp/CMakeLists.txt @@ -12,19 +12,29 @@ target_include_directories(UpscalePipelineApp PUBLIC ${SDK_INCLUDES_PATH} ) -target_link_libraries(UpscalePipelineApp PUBLIC - opencv346 - NVVideoEffects - ${CMAKE_CURRENT_SOURCE_DIR}/../external/cuda/lib/x64/cudart.lib - ) +if(MSVC) + 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 "--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}" - ) + 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 --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}" + ) +else() + target_link_libraries(UpscalePipelineApp PUBLIC + NVVideoEffects + NVCVImage + OpenCV + TensorRT + CUDA + ) +endif() diff --git a/samples/UpscalePipelineApp/UpscalePipelineApp.exe b/samples/UpscalePipelineApp/UpscalePipelineApp.exe index ca2c5eb..1d25885 100644 Binary files a/samples/UpscalePipelineApp/UpscalePipelineApp.exe and b/samples/UpscalePipelineApp/UpscalePipelineApp.exe differ diff --git a/samples/VideoEffectsApp/CMakeLists.txt b/samples/VideoEffectsApp/CMakeLists.txt index a9cd65a..6601fbc 100644 --- a/samples/VideoEffectsApp/CMakeLists.txt +++ b/samples/VideoEffectsApp/CMakeLists.txt @@ -7,19 +7,29 @@ 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}) -target_link_libraries(VideoEffectsApp PUBLIC - opencv346 - NVVideoEffects - ${CMAKE_CURRENT_SOURCE_DIR}/../external/cuda/lib/x64/cudart.lib - ) +if(MSVC) + 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 "--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}" - ) + 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/input1.jpg\"") + set_target_properties(VideoEffectsApp PROPERTIES + FOLDER SampleApps + VS_DEBUGGER_ENVIRONMENT "${PATH_STR}" + VS_DEBUGGER_COMMAND_ARGUMENTS "${CMD_ARG_STR}" + ) +else() + target_link_libraries(VideoEffectsApp PUBLIC + NVVideoEffects + NVCVImage + OpenCV + TensorRT + CUDA + ) +endif() diff --git a/samples/VideoEffectsApp/VideoEffectsApp.cpp b/samples/VideoEffectsApp/VideoEffectsApp.cpp index 3576b0e..2954193 100644 --- a/samples/VideoEffectsApp/VideoEffectsApp.cpp +++ b/samples/VideoEffectsApp/VideoEffectsApp.cpp @@ -57,7 +57,7 @@ bool FLAG_debug = false, FLAG_progress = false, FLAG_webcam = false; float FLAG_strength = 0.f; -int FLAG_mode = 0; +int FLAG_mode = 0; int FLAG_resolution = 0; std::string FLAG_codec = DEFAULT_CODEC, FLAG_camRes = "1280x720", @@ -519,6 +519,7 @@ NvCV_Status FXApp::allocBuffers(unsigned width, unsigned height) { printf("--resolution has not been specified\n"); return NVCV_ERR_PARAMETER; } + BAIL_IF_ERR(vfxErr = NvVFX_SetF32(_eff, NVVFX_STRENGTH, FLAG_strength)); int dstWidth = _srcImg.cols * FLAG_resolution / _srcImg.rows; _dstImg.create(FLAG_resolution, dstWidth, _srcImg.type()); // dst CPU BAIL_IF_NULL(_dstImg.data, vfxErr, NVCV_ERR_MEMORY); diff --git a/samples/VideoEffectsApp/VideoEffectsApp.exe b/samples/VideoEffectsApp/VideoEffectsApp.exe index e91c458..5974b7c 100644 Binary files a/samples/VideoEffectsApp/VideoEffectsApp.exe and b/samples/VideoEffectsApp/VideoEffectsApp.exe differ diff --git a/samples/external/cuda/include/channel_descriptor.h b/samples/external/cuda/include/channel_descriptor.h index a375508..1d61d29 100644 --- a/samples/external/cuda/include/channel_descriptor.h +++ b/samples/external/cuda/include/channel_descriptor.h @@ -86,7 +86,27 @@ * \endcode * * where ::cudaChannelFormatKind is one of ::cudaChannelFormatKindSigned, - * ::cudaChannelFormatKindUnsigned, cudaChannelFormatKindFloat or ::cudaChannelFormatKindNV12. + * ::cudaChannelFormatKindUnsigned, cudaChannelFormatKindFloat, + * ::cudaChannelFormatKindSignedNormalized8X1, ::cudaChannelFormatKindSignedNormalized8X2, + * ::cudaChannelFormatKindSignedNormalized8X4, + * ::cudaChannelFormatKindUnsignedNormalized8X1, ::cudaChannelFormatKindUnsignedNormalized8X2, + * ::cudaChannelFormatKindUnsignedNormalized8X4, + * ::cudaChannelFormatKindSignedNormalized16X1, ::cudaChannelFormatKindSignedNormalized16X2, + * ::cudaChannelFormatKindSignedNormalized16X4, + * ::cudaChannelFormatKindUnsignedNormalized16X1, ::cudaChannelFormatKindUnsignedNormalized16X2, + * ::cudaChannelFormatKindUnsignedNormalized16X4 + * or ::cudaChannelFormatKindNV12. + * + * The format is specified by the template specialization. + * + * The template function specializes for the following scalar types: + * char, signed char, unsigned char, short, unsigned short, int, unsigned int, long, unsigned long, and float. + * The template function specializes for the following vector types: + * char{1|2|4}, uchar{1|2|4}, short{1|2|4}, ushort{1|2|4}, int{1|2|4}, uint{1|2|4}, long{1|2|4}, ulong{1|2|4}, float{1|2|4}. + * The template function specializes for following cudaChannelFormatKind enum values: + * ::cudaChannelFormatKind{Uns|S}ignedNormalized{8|16}X{1|2|4}, and ::cudaChannelFormatKindNV12. + * + * Invoking the function on a type without a specialization defaults to creating a channel format of kind ::cudaChannelFormatKindNone * * \return * Channel descriptor with format \p f @@ -407,6 +427,166 @@ static __inline__ __host__ cudaChannelFormatDesc cudaCreateChannelDescNV12(void) return cudaCreateChannelDesc(e, e, e, 0, cudaChannelFormatKindNV12); } + +template __inline__ __host__ cudaChannelFormatDesc cudaCreateChannelDesc(void) +{ + return cudaCreateChannelDesc(0, 0, 0, 0, cudaChannelFormatKindNone); +} + +/* Signed 8-bit normalized integer formats */ +template<> __inline__ __host__ cudaChannelFormatDesc cudaCreateChannelDesc(void) +{ + return cudaCreateChannelDesc(8, 0, 0, 0, cudaChannelFormatKindSignedNormalized8X1); +} + +template<> __inline__ __host__ cudaChannelFormatDesc cudaCreateChannelDesc(void) +{ + return cudaCreateChannelDesc(8, 8, 0, 0, cudaChannelFormatKindSignedNormalized8X2); +} + +template<> __inline__ __host__ cudaChannelFormatDesc cudaCreateChannelDesc(void) +{ + return cudaCreateChannelDesc(8, 8, 8, 8, cudaChannelFormatKindSignedNormalized8X4); +} + +/* Unsigned 8-bit normalized integer formats */ +template<> __inline__ __host__ cudaChannelFormatDesc cudaCreateChannelDesc(void) +{ + return cudaCreateChannelDesc(8, 0, 0, 0, cudaChannelFormatKindUnsignedNormalized8X1); +} + +template<> __inline__ __host__ cudaChannelFormatDesc cudaCreateChannelDesc(void) +{ + return cudaCreateChannelDesc(8, 8, 0, 0, cudaChannelFormatKindUnsignedNormalized8X2); +} + +template<> __inline__ __host__ cudaChannelFormatDesc cudaCreateChannelDesc(void) +{ + return cudaCreateChannelDesc(8, 8, 8, 8, cudaChannelFormatKindUnsignedNormalized8X4); +} + +/* Signed 16-bit normalized integer formats */ +template<> __inline__ __host__ cudaChannelFormatDesc cudaCreateChannelDesc(void) +{ + return cudaCreateChannelDesc(16, 0, 0, 0, cudaChannelFormatKindSignedNormalized16X1); +} + +template<> __inline__ __host__ cudaChannelFormatDesc cudaCreateChannelDesc(void) +{ + return cudaCreateChannelDesc(16, 16, 0, 0, cudaChannelFormatKindSignedNormalized16X2); +} + +template<> __inline__ __host__ cudaChannelFormatDesc cudaCreateChannelDesc(void) +{ + return cudaCreateChannelDesc(16, 16, 16, 16, cudaChannelFormatKindSignedNormalized16X4); +} + +/* Unsigned 16-bit normalized integer formats */ +template<> __inline__ __host__ cudaChannelFormatDesc cudaCreateChannelDesc(void) +{ + return cudaCreateChannelDesc(16, 0, 0, 0, cudaChannelFormatKindUnsignedNormalized16X1); +} + +template<> __inline__ __host__ cudaChannelFormatDesc cudaCreateChannelDesc(void) +{ + return cudaCreateChannelDesc(16, 16, 0, 0, cudaChannelFormatKindUnsignedNormalized16X2); +} + +template<> __inline__ __host__ cudaChannelFormatDesc cudaCreateChannelDesc(void) +{ + return cudaCreateChannelDesc(16, 16, 16, 16, cudaChannelFormatKindUnsignedNormalized16X4); +} + +/* NV12 format */ +template<> __inline__ __host__ cudaChannelFormatDesc cudaCreateChannelDesc(void) +{ + return cudaCreateChannelDesc(8, 8, 8, 0, cudaChannelFormatKindNV12); +} + +/* BC1 format */ +template<> __inline__ __host__ cudaChannelFormatDesc cudaCreateChannelDesc(void) +{ + return cudaCreateChannelDesc(8, 8, 8, 8, cudaChannelFormatKindUnsignedBlockCompressed1); +} + +/* BC1sRGB format */ +template<> __inline__ __host__ cudaChannelFormatDesc cudaCreateChannelDesc(void) +{ + return cudaCreateChannelDesc(8, 8, 8, 8, cudaChannelFormatKindUnsignedBlockCompressed1SRGB); +} + +/* BC2 format */ +template<> __inline__ __host__ cudaChannelFormatDesc cudaCreateChannelDesc(void) +{ + return cudaCreateChannelDesc(8, 8, 8, 8, cudaChannelFormatKindUnsignedBlockCompressed2); +} + +/* BC2sRGB format */ +template<> __inline__ __host__ cudaChannelFormatDesc cudaCreateChannelDesc(void) +{ + return cudaCreateChannelDesc(8, 8, 8, 8, cudaChannelFormatKindUnsignedBlockCompressed2SRGB); +} + +/* BC3 format */ +template<> __inline__ __host__ cudaChannelFormatDesc cudaCreateChannelDesc(void) +{ + return cudaCreateChannelDesc(8, 8, 8, 8, cudaChannelFormatKindUnsignedBlockCompressed3); +} + +/* BC3sRGB format */ +template<> __inline__ __host__ cudaChannelFormatDesc cudaCreateChannelDesc(void) +{ + return cudaCreateChannelDesc(8, 8, 8, 8, cudaChannelFormatKindUnsignedBlockCompressed3SRGB); +} + +/* BC4 unsigned format */ +template<> __inline__ __host__ cudaChannelFormatDesc cudaCreateChannelDesc(void) +{ + return cudaCreateChannelDesc(8, 0, 0, 0, cudaChannelFormatKindUnsignedBlockCompressed4); +} + +/* BC4 signed format */ +template<> __inline__ __host__ cudaChannelFormatDesc cudaCreateChannelDesc(void) +{ + return cudaCreateChannelDesc(8, 0, 0, 0, cudaChannelFormatKindSignedBlockCompressed4); +} + +/* BC5 unsigned format */ +template<> __inline__ __host__ cudaChannelFormatDesc cudaCreateChannelDesc(void) +{ + return cudaCreateChannelDesc(8, 8, 0, 0, cudaChannelFormatKindUnsignedBlockCompressed5); +} + +/* BC5 signed format */ +template<> __inline__ __host__ cudaChannelFormatDesc cudaCreateChannelDesc(void) +{ + return cudaCreateChannelDesc(8, 8, 0, 0, cudaChannelFormatKindSignedBlockCompressed5); +} + +/* BC6H unsigned format */ +template<> __inline__ __host__ cudaChannelFormatDesc cudaCreateChannelDesc(void) +{ + return cudaCreateChannelDesc(16, 16, 16, 0, cudaChannelFormatKindUnsignedBlockCompressed6H); +} + +/* BC6H signed format */ +template<> __inline__ __host__ cudaChannelFormatDesc cudaCreateChannelDesc(void) +{ + return cudaCreateChannelDesc(16, 16, 16, 0, cudaChannelFormatKindSignedBlockCompressed6H); +} + +/* BC7 format */ +template<> __inline__ __host__ cudaChannelFormatDesc cudaCreateChannelDesc(void) +{ + return cudaCreateChannelDesc(8, 8, 8, 8, cudaChannelFormatKindUnsignedBlockCompressed7); +} + +/* BC7sRGB format */ +template<> __inline__ __host__ cudaChannelFormatDesc cudaCreateChannelDesc(void) +{ + return cudaCreateChannelDesc(8, 8, 8, 8, cudaChannelFormatKindUnsignedBlockCompressed7SRGB); +} + #endif /* __cplusplus */ /** @} */ diff --git a/samples/external/cuda/include/crt/host_config.h b/samples/external/cuda/include/crt/host_config.h index 13b6d09..5d56290 100644 --- a/samples/external/cuda/include/crt/host_config.h +++ b/samples/external/cuda/include/crt/host_config.h @@ -1,5 +1,5 @@ /* - * Copyright 1993-2020 NVIDIA Corporation. All rights reserved. + * Copyright 1993-2022 NVIDIA Corporation. All rights reserved. * * NOTICE TO LICENSEE: * @@ -105,21 +105,14 @@ #if defined(__ICC) -#if (__ICC != 1500 && __ICC != 1600 && __ICC != 1700 && __ICC != 1800 && !(__ICC >= 1900 && __ICC <= 1999)) || !defined(__GNUC__) || !defined(__LP64__) +#if (__ICC != 1500 && __ICC != 1600 && __ICC != 1700 && __ICC != 1800 && !(__ICC >= 1900 && __ICC <= 2021)) || !defined(__GNUC__) || !defined(__LP64__) -#error -- unsupported ICC configuration! Only ICC 15.0, ICC 16.0, ICC 17.0, ICC 18.0 and ICC 19.x on Linux x86_64 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 ICC configuration! Only ICC 15.0, ICC 16.0, ICC 17.0, ICC 18.0, ICC 19.x and 20.x on Linux x86_64 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. #endif /* (__ICC != 1500 && __ICC != 1600 && __ICC != 1700 && __ICC != 1800 && __ICC != 1900) || !__GNUC__ || !__LP64__ */ #endif /* __ICC */ -#if defined(__PGIC__) -#if ((__PGIC__ != 18) && (__PGIC__ != 19) && (__PGIC__ != 20) && (__PGIC__ != 21) && !(__PGIC__ == 99 && __PGIC_MINOR__ == 99)) -#error -- unsupported pgc++ configuration! Only pgc++ 18, 19, 20 and 21 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. - -#endif -#endif /* __PGIC__ */ - #if defined(__powerpc__) #if defined(__ibmxl_vrm__) && !(__ibmxl_vrm__ >= 0x0d010000 && __ibmxl_vrm__ < 0x0d020000) && \ @@ -134,19 +127,19 @@ #if defined(__GNUC__) -#if __GNUC__ > 10 +#if __GNUC__ > 11 -#error -- unsupported GNU version! gcc versions later than 10 are not 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 GNU version! gcc versions later than 11 are not 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. -#endif /* __GNUC__ > 10 */ +#endif /* __GNUC__ > 11 */ #if defined(__clang__) && !defined(__ibmxl_vrm__) && !defined(__ICC) && !defined(__HORIZON__) && !defined(__APPLE__) -#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. +#if (__clang_major__ >= 14) || (__clang_major__ < 3) || ((__clang_major__ == 3) && (__clang_minor__ < 3)) +#error -- unsupported clang version! clang version must be less than 14 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__ >= 12) || (__clang_major__ < 3) || ((__clang_major__ == 3) && (__clang_minor__ < 3)) */ +#endif /* (__clang_major__ >= 14) || (__clang_major__ < 3) || ((__clang_major__ == 3) && (__clang_minor__ < 3)) */ #endif /* defined(__clang__) && !defined(__ibmxl_vrm__) && !defined(__ICC) && !defined(__HORIZON__) && !defined(__APPLE__) */ @@ -155,15 +148,15 @@ #if defined(_WIN32) -#if _MSC_VER < 1910 || _MSC_VER >= 1930 +#if _MSC_VER < 1910 || _MSC_VER >= 1940 -#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. +#error -- unsupported Microsoft Visual Studio version! Only the versions between 2017 and 2022 (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 >= 1910 && _MSC_VER < 1910 -#pragma message("support for this version of Microsoft Visual Studio has been deprecated! Only the versions between 2017 and 2019 (inclusive) are supported!") +#pragma message("support for this version of Microsoft Visual Studio has been deprecated! Only the versions between 2017 and 2022 (inclusive) are supported!") -#endif /* (_MSC_VER < 1910 || _MSC_VER >= 1930) || (_MSC_VER >= 1910 && _MSC_VER < 1910) */ +#endif /* (_MSC_VER < 1910 || _MSC_VER >= 1940) || (_MSC_VER >= 1910 && _MSC_VER < 1910) */ #endif /* _WIN32 */ #endif /* !__NV_NO_HOST_COMPILER_CHECK */ diff --git a/samples/external/cuda/include/cuda_device_runtime_api.h b/samples/external/cuda/include/cuda_device_runtime_api.h index 83fbe8e..193c1c9 100644 --- a/samples/external/cuda/include/cuda_device_runtime_api.h +++ b/samples/external/cuda/include/cuda_device_runtime_api.h @@ -106,6 +106,22 @@ inline __device__ cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultipro #endif /* !defined(__CUDACC_RTC__) */ +#if defined(__DOXYGEN_ONLY__) || defined(CUDA_ENABLE_DEPRECATED) +# define __DEPRECATED__(msg) +#elif defined(_WIN32) +# define __DEPRECATED__(msg) __declspec(deprecated(msg)) +#elif (defined(__GNUC__) && (__GNUC__ < 4 || (__GNUC__ == 4 && __GNUC_MINOR__ < 5 && !defined(__clang__)))) +# define __DEPRECATED__(msg) __attribute__((deprecated)) +#else +# define __DEPRECATED__(msg) __attribute__((deprecated(msg))) +#endif + +#if defined(__CUDA_ARCH__) && !defined(__CDPRT_SUPPRESS_SYNC_DEPRECATION_WARNING) +# define __CDPRT_DEPRECATED(func_name) __DEPRECATED__("Use of "#func_name" from device code is deprecated and will not be supported in a future release. Disable this warning with -D__CDPRT_SUPPRESS_SYNC_DEPRECATION_WARNING.") +#else +# define __CDPRT_DEPRECATED(func_name) +#endif + #if defined(__cplusplus) && defined(__CUDACC__) /* Visible to nvcc front-end only */ #if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 350) // Visible to SM>=3.5 and "__host__ __device__" only @@ -118,7 +134,8 @@ extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaDeviceGetAttribut extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaDeviceGetLimit(size_t *pValue, enum cudaLimit limit); extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaDeviceGetCacheConfig(enum cudaFuncCache *pCacheConfig); extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaDeviceGetSharedMemConfig(enum cudaSharedMemConfig *pConfig); -extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaDeviceSynchronize(void); +extern __device__ __cudart_builtin__ __CDPRT_DEPRECATED(cudaDeviceSynchronize) cudaError_t CUDARTAPI cudaDeviceSynchronize(void); +extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaDeviceSynchronizeDeprecationAvoidance(void); extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaGetLastError(void); extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI cudaPeekAtLastError(void); extern __device__ __cudart_builtin__ const char* CUDARTAPI cudaGetErrorString(cudaError_t error); @@ -242,4 +259,7 @@ template static __inline__ __device__ __cudart_builtin__ cudaError_ #endif // !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 350) #endif /* defined(__cplusplus) && defined(__CUDACC__) */ +#undef __DEPRECATED__ +#undef __CDPRT_DEPRECATED + #endif /* !__CUDA_DEVICE_RUNTIME_API_H__ */ diff --git a/samples/external/cuda/include/cuda_runtime_api.h b/samples/external/cuda/include/cuda_runtime_api.h index 435e4a8..ef1a7e3 100644 --- a/samples/external/cuda/include/cuda_runtime_api.h +++ b/samples/external/cuda/include/cuda_runtime_api.h @@ -135,7 +135,7 @@ */ /** CUDA Runtime API Version */ -#define CUDART_VERSION 11030 +#define CUDART_VERSION 11060 #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)) @@ -143,7 +143,9 @@ # define __CUDART_API_VERSION CUDART_VERSION #endif +#ifndef __DOXYGEN_ONLY__ #include "crt/host_defines.h" +#endif #include "builtin_types.h" #include "cuda_device_runtime_api.h" @@ -281,8 +283,13 @@ extern "C" { * in the current process. * * Explicitly destroys and cleans up all resources associated with the current - * device in the current process. Any subsequent API call to this device will - * reinitialize the device. + * device in the current process. It is the caller's responsibility to ensure + * that the resources are not accessed or passed in subsequent API calls and + * doing so will result in undefined behavior. These resources include CUDA types + * such as ::cudaStream_t, ::cudaEvent_t, ::cudaArray_t, ::cudaMipmappedArray_t, + * ::cudaTextureObject_t, ::cudaSurfaceObject_t, ::textureReference, ::surfaceReference, + * ::cudaExternalMemory_t, ::cudaExternalSemaphore_t and ::cudaGraphicsResource_t. + * 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 @@ -309,6 +316,7 @@ extern __host__ cudaError_t CUDARTAPI cudaDeviceReset(void); * * \return * ::cudaSuccess + * \note_device_sync_deprecated * \notefnerr * \note_init_rt * \note_callback @@ -1664,93 +1672,93 @@ extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaGetDeviceProperties * * Returns in \p *value the integer value of the attribute \p attr on device * \p device. The supported attributes are: - * - ::cudaDevAttrMaxThreadsPerBlock: Maximum number of threads per block; - * - ::cudaDevAttrMaxBlockDimX: Maximum x-dimension of a block; - * - ::cudaDevAttrMaxBlockDimY: Maximum y-dimension of a block; - * - ::cudaDevAttrMaxBlockDimZ: Maximum z-dimension of a block; - * - ::cudaDevAttrMaxGridDimX: Maximum x-dimension of a grid; - * - ::cudaDevAttrMaxGridDimY: Maximum y-dimension of a grid; - * - ::cudaDevAttrMaxGridDimZ: Maximum z-dimension of a grid; + * - ::cudaDevAttrMaxThreadsPerBlock: Maximum number of threads per block + * - ::cudaDevAttrMaxBlockDimX: Maximum x-dimension of a block + * - ::cudaDevAttrMaxBlockDimY: Maximum y-dimension of a block + * - ::cudaDevAttrMaxBlockDimZ: Maximum z-dimension of a block + * - ::cudaDevAttrMaxGridDimX: Maximum x-dimension of a grid + * - ::cudaDevAttrMaxGridDimY: Maximum y-dimension of a grid + * - ::cudaDevAttrMaxGridDimZ: Maximum z-dimension of a grid * - ::cudaDevAttrMaxSharedMemoryPerBlock: Maximum amount of shared memory - * available to a thread block in bytes; + * available to a thread block in bytes * - ::cudaDevAttrTotalConstantMemory: Memory available on device for - * __constant__ variables in a CUDA C kernel in bytes; - * - ::cudaDevAttrWarpSize: Warp size in threads; + * __constant__ variables in a CUDA C kernel in bytes + * - ::cudaDevAttrWarpSize: Warp size in threads * - ::cudaDevAttrMaxPitch: Maximum pitch in bytes allowed by the memory copy - * functions that involve memory regions allocated through ::cudaMallocPitch(); - * - ::cudaDevAttrMaxTexture1DWidth: Maximum 1D texture width; + * functions that involve memory regions allocated through ::cudaMallocPitch() + * - ::cudaDevAttrMaxTexture1DWidth: Maximum 1D texture width * - ::cudaDevAttrMaxTexture1DLinearWidth: Maximum width for a 1D texture bound - * to linear memory; - * - ::cudaDevAttrMaxTexture1DMipmappedWidth: Maximum mipmapped 1D texture width; - * - ::cudaDevAttrMaxTexture2DWidth: Maximum 2D texture width; - * - ::cudaDevAttrMaxTexture2DHeight: Maximum 2D texture height; + * to linear memory + * - ::cudaDevAttrMaxTexture1DMipmappedWidth: Maximum mipmapped 1D texture width + * - ::cudaDevAttrMaxTexture2DWidth: Maximum 2D texture width + * - ::cudaDevAttrMaxTexture2DHeight: Maximum 2D texture height * - ::cudaDevAttrMaxTexture2DLinearWidth: Maximum width for a 2D texture - * bound to linear memory; + * bound to linear memory * - ::cudaDevAttrMaxTexture2DLinearHeight: Maximum height for a 2D texture - * bound to linear memory; + * bound to linear memory * - ::cudaDevAttrMaxTexture2DLinearPitch: Maximum pitch in bytes for a 2D - * texture bound to linear memory; + * texture bound to linear memory * - ::cudaDevAttrMaxTexture2DMipmappedWidth: Maximum mipmapped 2D texture - * width; + * width * - ::cudaDevAttrMaxTexture2DMipmappedHeight: Maximum mipmapped 2D texture - * height; - * - ::cudaDevAttrMaxTexture3DWidth: Maximum 3D texture width; - * - ::cudaDevAttrMaxTexture3DHeight: Maximum 3D texture height; - * - ::cudaDevAttrMaxTexture3DDepth: Maximum 3D texture depth; + * height + * - ::cudaDevAttrMaxTexture3DWidth: Maximum 3D texture width + * - ::cudaDevAttrMaxTexture3DHeight: Maximum 3D texture height + * - ::cudaDevAttrMaxTexture3DDepth: Maximum 3D texture depth * - ::cudaDevAttrMaxTexture3DWidthAlt: Alternate maximum 3D texture width, - * 0 if no alternate maximum 3D texture size is supported; + * 0 if no alternate maximum 3D texture size is supported * - ::cudaDevAttrMaxTexture3DHeightAlt: Alternate maximum 3D texture height, - * 0 if no alternate maximum 3D texture size is supported; + * 0 if no alternate maximum 3D texture size is supported * - ::cudaDevAttrMaxTexture3DDepthAlt: Alternate maximum 3D texture depth, - * 0 if no alternate maximum 3D texture size is supported; + * 0 if no alternate maximum 3D texture size is supported * - ::cudaDevAttrMaxTextureCubemapWidth: Maximum cubemap texture width or - * height; - * - ::cudaDevAttrMaxTexture1DLayeredWidth: Maximum 1D layered texture width; + * height + * - ::cudaDevAttrMaxTexture1DLayeredWidth: Maximum 1D layered texture width * - ::cudaDevAttrMaxTexture1DLayeredLayers: Maximum layers in a 1D layered - * texture; - * - ::cudaDevAttrMaxTexture2DLayeredWidth: Maximum 2D layered texture width; - * - ::cudaDevAttrMaxTexture2DLayeredHeight: Maximum 2D layered texture height; + * texture + * - ::cudaDevAttrMaxTexture2DLayeredWidth: Maximum 2D layered texture width + * - ::cudaDevAttrMaxTexture2DLayeredHeight: Maximum 2D layered texture height * - ::cudaDevAttrMaxTexture2DLayeredLayers: Maximum layers in a 2D layered - * texture; + * texture * - ::cudaDevAttrMaxTextureCubemapLayeredWidth: Maximum cubemap layered - * texture width or height; + * texture width or height * - ::cudaDevAttrMaxTextureCubemapLayeredLayers: Maximum layers in a cubemap - * layered texture; - * - ::cudaDevAttrMaxSurface1DWidth: Maximum 1D surface width; - * - ::cudaDevAttrMaxSurface2DWidth: Maximum 2D surface width; - * - ::cudaDevAttrMaxSurface2DHeight: Maximum 2D surface height; - * - ::cudaDevAttrMaxSurface3DWidth: Maximum 3D surface width; - * - ::cudaDevAttrMaxSurface3DHeight: Maximum 3D surface height; - * - ::cudaDevAttrMaxSurface3DDepth: Maximum 3D surface depth; - * - ::cudaDevAttrMaxSurface1DLayeredWidth: Maximum 1D layered surface width; + * layered texture + * - ::cudaDevAttrMaxSurface1DWidth: Maximum 1D surface width + * - ::cudaDevAttrMaxSurface2DWidth: Maximum 2D surface width + * - ::cudaDevAttrMaxSurface2DHeight: Maximum 2D surface height + * - ::cudaDevAttrMaxSurface3DWidth: Maximum 3D surface width + * - ::cudaDevAttrMaxSurface3DHeight: Maximum 3D surface height + * - ::cudaDevAttrMaxSurface3DDepth: Maximum 3D surface depth + * - ::cudaDevAttrMaxSurface1DLayeredWidth: Maximum 1D layered surface width * - ::cudaDevAttrMaxSurface1DLayeredLayers: Maximum layers in a 1D layered - * surface; - * - ::cudaDevAttrMaxSurface2DLayeredWidth: Maximum 2D layered surface width; - * - ::cudaDevAttrMaxSurface2DLayeredHeight: Maximum 2D layered surface height; + * surface + * - ::cudaDevAttrMaxSurface2DLayeredWidth: Maximum 2D layered surface width + * - ::cudaDevAttrMaxSurface2DLayeredHeight: Maximum 2D layered surface height * - ::cudaDevAttrMaxSurface2DLayeredLayers: Maximum layers in a 2D layered - * surface; - * - ::cudaDevAttrMaxSurfaceCubemapWidth: Maximum cubemap surface width; + * surface + * - ::cudaDevAttrMaxSurfaceCubemapWidth: Maximum cubemap surface width * - ::cudaDevAttrMaxSurfaceCubemapLayeredWidth: Maximum cubemap layered - * surface width; + * surface width * - ::cudaDevAttrMaxSurfaceCubemapLayeredLayers: Maximum layers in a cubemap - * layered surface; + * layered surface * - ::cudaDevAttrMaxRegistersPerBlock: Maximum number of 32-bit registers - * available to a thread block; - * - ::cudaDevAttrClockRate: Peak clock frequency in kilohertz; + * available to a thread block + * - ::cudaDevAttrClockRate: Peak clock frequency in kilohertz * - ::cudaDevAttrTextureAlignment: Alignment requirement; texture base * addresses aligned to ::textureAlign bytes do not need an offset applied - * to texture fetches; + * to texture fetches * - ::cudaDevAttrTexturePitchAlignment: Pitch alignment requirement for 2D - * texture references bound to pitched memory; + * texture references bound to pitched memory * - ::cudaDevAttrGpuOverlap: 1 if the device can concurrently copy memory - * between host and device while executing a kernel, or 0 if not; - * - ::cudaDevAttrMultiProcessorCount: Number of multiprocessors on the device; + * between host and device while executing a kernel, or 0 if not + * - ::cudaDevAttrMultiProcessorCount: Number of multiprocessors on the device * - ::cudaDevAttrKernelExecTimeout: 1 if there is a run time limit for kernels - * executed on the device, or 0 if not; + * executed on the device, or 0 if not * - ::cudaDevAttrIntegrated: 1 if the device is integrated with the memory - * subsystem, or 0 if not; + * subsystem, or 0 if not * - ::cudaDevAttrCanMapHostMemory: 1 if the device can map host memory into - * the CUDA address space, or 0 if not; + * the CUDA address space, or 0 if not * - ::cudaDevAttrComputeMode: Compute mode is the compute mode that the device * is currently in. Available modes are as follows: * - ::cudaComputeModeDefault: Default mode - Device is not restricted and @@ -1766,75 +1774,84 @@ extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaGetDeviceProperties * multiple kernels within the same context simultaneously, or 0 if * not. It is not guaranteed that multiple kernels will be resident on the * device concurrently so this feature should not be relied upon for - * correctness; + * correctness. * - ::cudaDevAttrEccEnabled: 1 if error correction is enabled on the device, - * 0 if error correction is disabled or not supported by the device; - * - ::cudaDevAttrPciBusId: PCI bus identifier of the device; + * 0 if error correction is disabled or not supported by the device + * - ::cudaDevAttrPciBusId: PCI bus identifier of the device * - ::cudaDevAttrPciDeviceId: PCI device (also known as slot) identifier of - * the device; + * the device * - ::cudaDevAttrTccDriver: 1 if the device is using a TCC driver. TCC is only - * available on Tesla hardware running Windows Vista or later; - * - ::cudaDevAttrMemoryClockRate: Peak memory clock frequency in kilohertz; - * - ::cudaDevAttrGlobalMemoryBusWidth: Global memory bus width in bits; + * available on Tesla hardware running Windows Vista or later. + * - ::cudaDevAttrMemoryClockRate: Peak memory clock frequency in kilohertz + * - ::cudaDevAttrGlobalMemoryBusWidth: Global memory bus width in bits * - ::cudaDevAttrL2CacheSize: Size of L2 cache in bytes. 0 if the device - * doesn't have L2 cache; + * doesn't have L2 cache. * - ::cudaDevAttrMaxThreadsPerMultiProcessor: Maximum resident threads per - * multiprocessor; + * multiprocessor * - ::cudaDevAttrUnifiedAddressing: 1 if the device shares a unified address - * space with the host, or 0 if not; + * space with the host, or 0 if not * - ::cudaDevAttrComputeCapabilityMajor: Major compute capability version - * number; + * number * - ::cudaDevAttrComputeCapabilityMinor: Minor compute capability version - * number; + * number * - ::cudaDevAttrStreamPrioritiesSupported: 1 if the device supports stream - * priorities, or 0 if not; + * priorities, or 0 if not * - ::cudaDevAttrGlobalL1CacheSupported: 1 if device supports caching globals - * in L1 cache, 0 if not; + * in L1 cache, 0 if not * - ::cudaDevAttrLocalL1CacheSupported: 1 if device supports caching locals - * in L1 cache, 0 if not; + * in L1 cache, 0 if not * - ::cudaDevAttrMaxSharedMemoryPerMultiprocessor: Maximum amount of shared memory * available to a multiprocessor in bytes; this amount is shared by all - * thread blocks simultaneously resident on a multiprocessor; + * thread blocks simultaneously resident on a multiprocessor * - ::cudaDevAttrMaxRegistersPerMultiprocessor: Maximum number of 32-bit registers * available to a multiprocessor; this number is shared by all thread blocks - * simultaneously resident on a multiprocessor; + * simultaneously resident on a multiprocessor * - ::cudaDevAttrManagedMemory: 1 if device supports allocating - * managed memory, 0 if not; - * - ::cudaDevAttrIsMultiGpuBoard: 1 if device is on a multi-GPU board, 0 if not; + * managed memory, 0 if not + * - ::cudaDevAttrIsMultiGpuBoard: 1 if device is on a multi-GPU board, 0 if not * - ::cudaDevAttrMultiGpuBoardGroupID: Unique identifier for a group of devices on the - * same multi-GPU board; + * same multi-GPU board * - ::cudaDevAttrHostNativeAtomicSupported: 1 if the link between the device and the - * host supports native atomic operations; + * host supports native atomic operations * - ::cudaDevAttrSingleToDoublePrecisionPerfRatio: Ratio of single precision performance - * (in floating-point operations per second) to double precision performance; + * (in floating-point operations per second) to double precision performance * - ::cudaDevAttrPageableMemoryAccess: 1 if the device supports coherently accessing - * pageable memory without calling cudaHostRegister on it, and 0 otherwise. + * pageable memory without calling cudaHostRegister on it, and 0 otherwise * - ::cudaDevAttrConcurrentManagedAccess: 1 if the device can coherently access managed - * memory concurrently with the CPU, and 0 otherwise. + * memory concurrently with the CPU, and 0 otherwise * - ::cudaDevAttrComputePreemptionSupported: 1 if the device supports - * Compute Preemption, 0 if not. + * Compute Preemption, 0 if not * - ::cudaDevAttrCanUseHostPointerForRegisteredMem: 1 if the device can access host - * registered memory at the same virtual address as the CPU, and 0 otherwise. + * registered memory at the same virtual address as the CPU, and 0 otherwise * - ::cudaDevAttrCooperativeLaunch: 1 if the device supports launching cooperative kernels - * via ::cudaLaunchCooperativeKernel, and 0 otherwise. + * via ::cudaLaunchCooperativeKernel, and 0 otherwise * - ::cudaDevAttrCooperativeMultiDeviceLaunch: 1 if the device supports launching cooperative - * kernels via ::cudaLaunchCooperativeKernelMultiDevice, and 0 otherwise. + * kernels via ::cudaLaunchCooperativeKernelMultiDevice, and 0 otherwise * - ::cudaDevAttrCanFlushRemoteWrites: 1 if the device supports flushing of outstanding - * remote writes, and 0 otherwise. + * remote writes, and 0 otherwise * - ::cudaDevAttrHostRegisterSupported: 1 if the device supports host memory registration - * via ::cudaHostRegister, and 0 otherwise. + * via ::cudaHostRegister, and 0 otherwise * - ::cudaDevAttrPageableMemoryAccessUsesHostPageTables: 1 if the device accesses pageable memory via the - * host's page tables, and 0 otherwise. + * host's page tables, and 0 otherwise * - ::cudaDevAttrDirectManagedMemAccessFromHost: 1 if the host can directly access managed memory on the device - * without migration, and 0 otherwise. + * without migration, and 0 otherwise * - ::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 + * - ::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 + * - ::cudaDevAttrReservedSharedMemoryPerBlock: Shared memory reserved by CUDA driver per block in bytes * - ::cudaDevAttrSparseCudaArraySupported: 1 if the device supports sparse CUDA arrays and sparse CUDA mipmapped arrays. + * - ::cudaDevAttrHostRegisterReadOnlySupported: Device supports using the ::cudaHostRegister flag cudaHostRegisterReadOnly + * to register memory that must be mapped as read-only to the GPU + * - ::cudaDevAttrMemoryPoolsSupported: 1 if the device supports using the cudaMallocAsync and cudaMemPool family of APIs, and 0 otherwise + * - ::cudaDevAttrGPUDirectRDMASupported: 1 if the device supports GPUDirect RDMA APIs, and 0 otherwise + * - ::cudaDevAttrGPUDirectRDMAFlushWritesOptions: bitmask to be interpreted according to the ::cudaFlushGPUDirectRDMAWritesOptions enum + * - ::cudaDevAttrGPUDirectRDMAWritesOrdering: see the ::cudaGPUDirectRDMAWritesOrdering enum for numerical values + * - ::cudaDevAttrMemoryPoolSupportedHandleTypes: Bitmask of handle types supported with mempool based IPC + + * - ::cudaDevAttrDeferredMappingCudaArraySupported : 1 if the device supports deferred mapping CUDA arrays and CUDA mipmapped arrays. + * * \param value - Returned device attribute value * \param attr - Device attribute to query @@ -2044,6 +2061,10 @@ extern __host__ cudaError_t CUDARTAPI cudaChooseDevice(int *device, const struct * This call may be made from any host thread, to any device, and at * any time. This function will do no synchronization with the previous * or new device, and should be considered a very low overhead call. + * If the current context bound to the calling thread is not the primary context, + * this call will bind the primary context to the calling thread and all the + * subsequent memory allocations, stream and event creations, and kernel launches + * will be associated with the primary context. * * \param device - Device on which the active host thread should execute the * device code. @@ -2116,19 +2137,14 @@ extern __host__ cudaError_t CUDARTAPI cudaSetValidDevices(int *device_arr, int l /** * \brief Sets flags to be used for device executions - * - * Records \p flags as the flags to use when initializing the current - * device. If no device has been made current to the calling thread, - * then \p flags will be applied to the initialization of any device - * initialized by the calling host thread, unless that device has had - * its initialization flags set explicitly by this or any host thread. * - * If the current device has been set and that device has already been - * initialized then this call will fail with the error - * ::cudaErrorSetOnActiveProcess. In this case it is necessary - * to reset \p device using ::cudaDeviceReset() before the device's - * initialization flags may be set. - * + * Records \p flags as the flags for the current device. If the current device + * has been set and that device has already been initialized, the previous flags + * are overwritten. If the current device has not been initialized, it is + * initialized with the provided flags. If no device has been made current to + * the calling thread, a default device is selected and initialized with the + * provided flags. + * * The two LSBs of the \p flags parameter can be used to control how the CPU * thread interacts with the OS scheduler when waiting for results from the * device. @@ -2164,14 +2180,15 @@ extern __host__ cudaError_t CUDARTAPI cudaSetValidDevices(int *device_arr, int l * - ::cudaDeviceLmemResizeToMax: Instruct CUDA to not reduce local memory * after resizing local memory for a kernel. This can prevent thrashing by * local memory allocations when launching many kernels with high local - * memory usage at the cost of potentially increased memory usage. + * memory usage at the cost of potentially increased memory usage.
+ * \ref deprecated "Deprecated:" This flag is deprecated and the behavior enabled + * by this flag is now the default and cannot be disabled. * * \param flags - Parameters for device operation * * \return * ::cudaSuccess, * ::cudaErrorInvalidValue, - * ::cudaErrorSetOnActiveProcess * \notefnerr * \note_init_rt * \note_callback @@ -2186,14 +2203,12 @@ extern __host__ cudaError_t CUDARTAPI cudaSetDeviceFlags( unsigned int flags ); /** * \brief Gets the flags for the current device * - * Returns in \p flags the flags for the current device. If there is a - * current device for the calling thread, and the device has been initialized - * or flags have been set on that device specifically, the flags for the - * device are returned. If there is no current device, but flags have been - * set for the thread with ::cudaSetDeviceFlags, the thread flags are returned. - * Finally, if there is no current device and no thread flags, the flags for - * the first device are returned, which may be the default flags. Compare - * to the behavior of ::cudaSetDeviceFlags. + * + * Returns in \p flags the flags for the current device. If there is a current + * device for the calling thread, the flags for the device are returned. If + * there is no current device, the flags for the first device are returned, + * which may be the default flags. Compare to the behavior of + * ::cudaSetDeviceFlags. * * Typically, the flags returned should match the behavior that will be seen * if the calling thread uses a device after this call, without any change to @@ -3106,7 +3121,7 @@ extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaEventCreateWithFlag * \brief Records an event * * Captures in \p event the contents of \p stream at the time of this call. - * \p event and \p stream must be on the same device. + * \p event and \p stream must be on the same CUDA context. * Calls such as ::cudaEventQuery() or ::cudaStreamWaitEvent() will then * examine or wait for completion of the work that was captured. Uses of * \p stream after this call do not modify \p event. See note on default @@ -3146,7 +3161,7 @@ extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaEventRecord(cudaEve * \brief Records an event * * Captures in \p event the contents of \p stream at the time of this call. - * \p event and \p stream must be on the same device. + * \p event and \p stream must be on the same CUDA context. * Calls such as ::cudaEventQuery() or ::cudaStreamWaitEvent() will then * examine or wait for completion of the work that was captured. Uses of * \p stream after this call do not modify \p event. See note on default @@ -3463,6 +3478,8 @@ extern __host__ cudaError_t CUDARTAPI cudaEventElapsedTime(float *ms, cudaEvent_ * If the NvSciBuf object imported into CUDA is also mapped by other drivers, then the * application must use ::cudaWaitExternalSemaphoresAsync or ::cudaSignalExternalSemaphoresAsync * as approprriate barriers to maintain coherence between CUDA and the other drivers. + * See ::cudaExternalSemaphoreWaitSkipNvSciBufMemSync and ::cudaExternalSemaphoreSignalSkipNvSciBufMemSync + * for memory synchronization. * * The size of the memory object must be specified in * ::cudaExternalMemoryHandleDesc::size. @@ -3482,6 +3499,7 @@ extern __host__ cudaError_t CUDARTAPI cudaEventElapsedTime(float *ms, cudaEvent_ * * \return * ::cudaSuccess, + * ::cudaErrorInvalidValue, * ::cudaErrorInvalidResourceHandle * \notefnerr * \note_init_rt @@ -3543,6 +3561,7 @@ extern __host__ cudaError_t CUDARTAPI cudaImportExternalMemory(cudaExternalMemor * * \return * ::cudaSuccess, + * ::cudaErrorInvalidValue, * ::cudaErrorInvalidResourceHandle * \notefnerr * \note_init_rt @@ -3598,6 +3617,7 @@ extern __host__ cudaError_t CUDARTAPI cudaExternalMemoryGetMappedBuffer(void **d * * \return * ::cudaSuccess, + * ::cudaErrorInvalidValue, * ::cudaErrorInvalidResourceHandle * \notefnerr * \note_init_rt @@ -4854,8 +4874,13 @@ extern __host__ cudaError_t CUDARTAPI cudaMallocPitch(void **devPtr, size_t *pit * - ::cudaArraySurfaceLoadStore: Allocates an array that can be read from or written to using a surface reference * - ::cudaArrayTextureGather: This flag indicates that texture gather operations will be performed on the array. * - ::cudaArraySparse: Allocates a CUDA array without physical backing memory. The subregions within this sparse array - * can later be mapped to physical memory by calling ::cuMemMapArrayAsync. The physical backing memory must be allocated - * via ::cuMemCreate. + * can later be mapped onto a physical memory allocation by calling ::cuMemMapArrayAsync. + * The physical backing memory must be allocated via ::cuMemCreate. + + * - ::cudaArrayDeferredMapping: Allocates a CUDA array without physical backing memory. The entire array can + * later be mapped onto a physical memory allocation by calling ::cuMemMapArrayAsync. + * The physical backing memory must be allocated via ::cuMemCreate. + * * \p width and \p height must meet certain size requirements. See ::cudaMalloc3DArray() for more details. * @@ -5316,8 +5341,13 @@ extern __host__ cudaError_t CUDARTAPI cudaMalloc3D(struct cudaPitchedPtr* pitche * - ::cudaArrayTextureGather: This flag indicates that texture gather operations will be performed on the CUDA * array. Texture gather can only be performed on 2D CUDA arrays. * - ::cudaArraySparse: Allocates a CUDA array without physical backing memory. The subregions within this sparse array - * can later be mapped to physical memory by calling ::cuMemMapArrayAsync. This flag can only be used for - * creating 2D, 3D or 2D layered sparse CUDA arrays. The physical backing memory must be allocated via ::cuMemCreate. + * can later be mapped onto a physical memory allocation by calling ::cuMemMapArrayAsync. This flag can only be used for + * creating 2D, 3D or 2D layered sparse CUDA arrays. The physical backing memory must be allocated via ::cuMemCreate. + + * - ::cudaArrayDeferredMapping: Allocates a CUDA array without physical backing memory. The entire array can + * later be mapped onto a physical memory allocation by calling ::cuMemMapArrayAsync. The physical backing memory must be allocated + * via ::cuMemCreate. + * * The width, height and depth extents must meet certain size requirements as listed in the following table. * All values are specified in elements. @@ -5459,9 +5489,14 @@ extern __host__ cudaError_t CUDARTAPI cudaMalloc3DArray(cudaArray_t *array, cons * - ::cudaArrayTextureGather: This flag indicates that texture gather operations will be performed on the CUDA * array. Texture gather can only be performed on 2D CUDA mipmapped arrays, and the gather operations are * performed only on the most detailed mipmap level. - * - ::cudaArraySparse: Allocates a CUDA array without physical backing memory. The subregions within this sparse array - * can later be mapped to physical memory by calling ::cuMemMapArrayAsync. This flag can only be used for creating + * - ::cudaArraySparse: Allocates a CUDA mipmapped array without physical backing memory. The subregions within this sparse array + * can later be mapped onto a physical memory allocation by calling ::cuMemMapArrayAsync. This flag can only be used for creating * 2D, 3D or 2D layered sparse CUDA mipmapped arrays. The physical backing memory must be allocated via ::cuMemCreate. + + * - ::cudaArrayDeferredMapping: Allocates a CUDA mipmapped array without physical backing memory. The entire array can + * later be mapped onto a physical memory allocation by calling ::cuMemMapArrayAsync. The physical backing memory must be allocated + * via ::cuMemCreate. + * * The width, height and depth extents must meet certain size requirements as listed in the following table. * All values are specified in elements. @@ -5867,8 +5902,20 @@ extern __host__ cudaError_t CUDARTAPI cudaMemcpy3DPeerAsync(const struct cudaMem /** * \brief Gets free and total device memory * - * Returns in \p *free and \p *total respectively, the free and total amount of - * memory available for allocation by the device in bytes. + * Returns in \p *total the total amount of memory available to the the current context. + * Returns in \p *free the amount of memory on the device that is free according to the OS. + * CUDA is not guaranteed to be able to allocate all of the memory that the OS reports as free. + * In a multi-tenet situation, free estimate returned is prone to race condition where + * a new allocation/free done by a different process or a different thread in the same + * process between the time when free memory was estimated and reported, will result in + * deviation in free value reported and actual free memory. + * + * The integrated GPU on Tegra shares memory with CPU and other component + * of the SoC. The free and total values returned by the API excludes + * the SWAP memory space maintained by the OS on some platforms. + * The OS may move some of the memory pages into swap area as the GPU or + * CPU allocate or access memory. See Tegra app note on how to calculate + * total and free memory on Tegra. * * \param free - Returned free memory in bytes * \param total - Returned total memory in bytes @@ -5941,6 +5988,55 @@ extern __host__ cudaError_t CUDARTAPI cudaArrayGetInfo(struct cudaChannelFormatD */ extern __host__ cudaError_t CUDARTAPI cudaArrayGetPlane(cudaArray_t *pPlaneArray, cudaArray_t hArray, unsigned int planeIdx); + +/** + * \brief Returns the memory requirements of a CUDA array + * + * Returns the memory requirements of a CUDA array in \p memoryRequirements + * If the CUDA array is not allocated with flag ::cudaArrayDeferredMapping + * ::cudaErrorInvalidValue will be returned. + * + * The returned value in ::cudaArrayMemoryRequirements::size + * represents the total size of the CUDA array. + * The returned value in ::cudaArrayMemoryRequirements::alignment + * represents the alignment necessary for mapping the CUDA array. + * + * \return + * ::cudaSuccess + * ::cudaErrorInvalidValue + * + * \param[out] memoryRequirements - Pointer to ::cudaArrayMemoryRequirements + * \param[in] array - CUDA array to get the memory requirements of + * \param[in] device - Device to get the memory requirements for + * \sa ::cudaMipmappedArrayGetMemoryRequirements + */ +extern __host__ cudaError_t CUDARTAPI cudaArrayGetMemoryRequirements(struct cudaArrayMemoryRequirements *memoryRequirements, cudaArray_t array, int device); + +/** + * \brief Returns the memory requirements of a CUDA mipmapped array + * + * Returns the memory requirements of a CUDA mipmapped array in \p memoryRequirements + * If the CUDA mipmapped array is not allocated with flag ::cudaArrayDeferredMapping + * ::cudaErrorInvalidValue will be returned. + * + * The returned value in ::cudaArrayMemoryRequirements::size + * represents the total size of the CUDA mipmapped array. + * The returned value in ::cudaArrayMemoryRequirements::alignment + * represents the alignment necessary for mapping the CUDA mipmapped + * array. + * + * \return + * ::cudaSuccess + * ::cudaErrorInvalidValue + * + * \param[out] memoryRequirements - Pointer to ::cudaArrayMemoryRequirements + * \param[in] mipmap - CUDA mipmapped array to get the memory requirements of + * \param[in] device - Device to get the memory requirements for + * \sa ::cudaArrayGetMemoryRequirements + */ +extern __host__ cudaError_t CUDARTAPI cudaMipmappedArrayGetMemoryRequirements(struct cudaArrayMemoryRequirements *memoryRequirements, cudaMipmappedArray_t mipmap, int device); + + /** * \brief Returns the layout properties of a sparse CUDA array * @@ -7602,6 +7698,9 @@ extern __CUDA_DEPRECATED __host__ cudaError_t CUDARTAPI cudaMemcpyFromArrayAsync * \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. + * \note During stream capture, this function results in the creation of an allocation node. In this case, + * the allocation is owned by the graph instead of the memory pool. The memory pool's properties + * are used to set the node's creation parameters. * * \param[out] devPtr - Returned device pointer * \param[in] size - Number of bytes to allocate @@ -7631,6 +7730,9 @@ extern __host__ cudaError_t CUDARTAPI cudaMallocAsync(void **devPtr, size_t size * After this API returns, accessing the memory from any subsequent work launched on the GPU * or querying its pointer attributes results in undefined behavior. * + * \note During stream capture, this function results in the creation of a free node and + * must therefore be passed the address of a graph allocation. + * * \param dptr - memory to free * \param hStream - The stream establishing the stream ordering promise * \returns @@ -7694,6 +7796,12 @@ extern __host__ cudaError_t CUDARTAPI cudaMemPoolTrimTo(cudaMemPool_t memPool, s * 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). + * - ::cudaMemPoolAttrReservedMemHigh: (value type = cuuint64_t) + * Reset the high watermark that tracks the amount of backing memory that was + * allocated for the memory pool. It is illegal to set this attribute to a non-zero value. + * - ::cudaMemPoolAttrUsedMemHigh: (value type = cuuint64_t) + * Reset the high watermark that tracks the amount of used memory that was + * allocated for the memory pool. It is illegal to set this attribute to a non-zero value. * * \param[in] pool - The memory pool to modify * \param[in] attr - The attribute to modify @@ -7732,6 +7840,16 @@ extern __host__ cudaError_t CUDARTAPI cudaMemPoolSetAttribute(cudaMemPool_t memP * 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). + * - ::cudaMemPoolAttrReservedMemCurrent: (value type = cuuint64_t) + * Amount of backing memory currently allocated for the mempool. + * - ::cudaMemPoolAttrReservedMemHigh: (value type = cuuint64_t) + * High watermark of backing memory allocated for the mempool since + * the last time it was reset. + * - ::cudaMemPoolAttrUsedMemCurrent: (value type = cuuint64_t) + * Amount of memory from the pool that is currently in use by the application. + * - ::cudaMemPoolAttrUsedMemHigh: (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. * * \param[in] pool - The memory pool to get attributes of * \param[in] attr - The attribute to get @@ -7832,6 +7950,10 @@ extern __host__ cudaError_t CUDARTAPI cudaMemPoolDestroy(cudaMemPool_t memPool); * 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. * + * \note During stream capture, this function results in the creation of an allocation node. In this case, + * the allocation is owned by the graph instead of the memory pool. The memory pool's properties + * are used to set the node's creation parameters. + * * \param[out] ptr - Returned device pointer * \param[in] bytesize - Number of bytes to allocate * \param[in] memPool - The pool to allocate from @@ -9003,6 +9125,7 @@ extern __host__ struct cudaChannelFormatDesc CUDARTAPI cudaCreateChannelDesc(int float minMipmapLevelClamp; float maxMipmapLevelClamp; int disableTrilinearOptimization; + int seamlessCubemap; }; * \endcode * where @@ -9062,6 +9185,11 @@ extern __host__ struct cudaChannelFormatDesc CUDARTAPI cudaCreateChannelDesc(int * * - ::cudaTextureDesc::disableTrilinearOptimization specifies whether the trilinear filtering optimizations will be disabled. * + * - ::cudaTextureDesc::seamlessCubemap specifies whether seamless cube map filtering is enabled. This flag can only be specified if the + * underlying resource is a CUDA array or a CUDA mipmapped array that was created with the flag ::cudaArrayCubemap. + * When seamless cube map filtering is enabled, texture address modes specified by ::cudaTextureDesc::addressMode are ignored. + * Instead, if the ::cudaTextureDesc::filterMode is set to ::cudaFilterModePoint the address mode ::cudaAddressModeClamp will be applied for all dimensions. + * If the ::cudaTextureDesc::filterMode is set to ::cudaFilterModeLinear seamless cube map filtering will be performed when sampling along the cube face borders. * * The ::cudaResourceViewDesc struct is defined as * \code @@ -10245,6 +10373,8 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphHostNodeSetParams(cudaGraphNode_t * at the root of the graph. \p pDependencies may not have any duplicate entries. * A handle to the new node will be returned in \p pGraphNode. * + * If \p hGraph contains allocation or free nodes, this call will return an error. + * * The node executes an embedded child graph. The child graph is cloned in this call. * * \param pGraphNode - Returns newly created node @@ -10281,6 +10411,9 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphAddChildGraphNode(cudaGraphNode_t * does not clone the graph. Changes to the graph will be reflected in * the node, and the node retains ownership of the graph. * + * Allocation and free nodes cannot be added to the returned graph. + * Attempting to do so will return an error. + * * \param node - Node to get the embedded graph for * \param pGraph - Location to store a handle to the graph * @@ -10356,13 +10489,12 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphAddEmptyNode(cudaGraphNode_t *pGr * \param event - Event for the node * * \return - * ::CUDA_SUCCESS, - * ::CUDA_ERROR_DEINITIALIZED, - * ::CUDA_ERROR_NOT_INITIALIZED, - * ::CUDA_ERROR_NOT_SUPPORTED, - * ::CUDA_ERROR_INVALID_VALUE + * ::cudaSuccess, + * ::cudaErrorInvalidValue * \note_graph_thread_safety * \notefnerr + * \note_init_rt + * \note_callback * * \sa * ::cudaGraphAddEventWaitNode, @@ -10389,12 +10521,12 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphAddEmptyNode(cudaGraphNode_t *pGr * \param event_out - Pointer to return the event * * \return - * ::CUDA_SUCCESS, - * ::CUDA_ERROR_DEINITIALIZED, - * ::CUDA_ERROR_NOT_INITIALIZED, - * ::CUDA_ERROR_INVALID_VALUE + * ::cudaSuccess, + * ::cudaErrorInvalidValue * \note_graph_thread_safety * \notefnerr + * \note_init_rt + * \note_callback * * \sa * ::cudaGraphAddEventRecordNode, @@ -10416,12 +10548,12 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphAddEmptyNode(cudaGraphNode_t *pGr * \param event - Event to use * * \return - * ::CUDA_SUCCESS, - * ::CUDA_ERROR_INVALID_VALUE, - * ::CUDA_ERROR_INVALID_HANDLE, - * ::CUDA_ERROR_OUT_OF_MEMORY + * ::cudaSuccess, + * ::cudaErrorInvalidValue * \note_graph_thread_safety * \notefnerr + * \note_init_rt + * \note_callback * * \sa * ::cudaGraphAddEventRecordNode, @@ -10457,13 +10589,12 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphAddEmptyNode(cudaGraphNode_t *pGr * \param event - Event for the node * * \return - * ::CUDA_SUCCESS, - * ::CUDA_ERROR_DEINITIALIZED, - * ::CUDA_ERROR_NOT_INITIALIZED, - * ::CUDA_ERROR_NOT_SUPPORTED, - * ::CUDA_ERROR_INVALID_VALUE + * ::cudaSuccess, + * ::cudaErrorInvalidValue * \note_graph_thread_safety * \notefnerr + * \note_init_rt + * \note_callback * * \sa * ::cudaGraphAddEventRecordNode, @@ -10490,12 +10621,12 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphAddEmptyNode(cudaGraphNode_t *pGr * \param event_out - Pointer to return the event * * \return - * ::CUDA_SUCCESS, - * ::CUDA_ERROR_DEINITIALIZED, - * ::CUDA_ERROR_NOT_INITIALIZED, - * ::CUDA_ERROR_INVALID_VALUE + * ::cudaSuccess, + * ::cudaErrorInvalidValue * \note_graph_thread_safety * \notefnerr + * \note_init_rt + * \note_callback * * \sa * ::cudaGraphAddEventWaitNode, @@ -10517,12 +10648,12 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphAddEmptyNode(cudaGraphNode_t *pGr * \param event - Event to use * * \return - * ::CUDA_SUCCESS, - * ::CUDA_ERROR_INVALID_VALUE, - * ::CUDA_ERROR_INVALID_HANDLE, - * ::CUDA_ERROR_OUT_OF_MEMORY + * ::cudaSuccess, + * ::cudaErrorInvalidValue * \note_graph_thread_safety * \notefnerr + * \note_init_rt + * \note_callback * * \sa * ::cudaGraphAddEventWaitNode, @@ -10555,13 +10686,12 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphAddEmptyNode(cudaGraphNode_t *pGr * \param nodeParams - Parameters for the node * * \return - * ::CUDA_SUCCESS, - * ::CUDA_ERROR_DEINITIALIZED, - * ::CUDA_ERROR_NOT_INITIALIZED, - * ::CUDA_ERROR_NOT_SUPPORTED, - * ::CUDA_ERROR_INVALID_VALUE + * ::cudaSuccess, + * ::cudaErrorInvalidValue * \note_graph_thread_safety * \notefnerr + * \note_init_rt + * \note_callback * * \sa * ::cudaGraphExternalSemaphoresSignalNodeGetParams, @@ -10599,12 +10729,12 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphAddExternalSemaphoresSignalNode(c * \param params_out - Pointer to return the parameters * * \return - * ::CUDA_SUCCESS, - * ::CUDA_ERROR_DEINITIALIZED, - * ::CUDA_ERROR_NOT_INITIALIZED, - * ::CUDA_ERROR_INVALID_VALUE + * ::cudaSuccess, + * ::cudaErrorInvalidValue * \note_graph_thread_safety * \notefnerr + * \note_init_rt + * \note_callback * * \sa * ::cudaLaunchKernel, @@ -10627,12 +10757,12 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphExternalSemaphoresSignalNodeGetPa * \param nodeParams - Parameters to copy * * \return - * ::CUDA_SUCCESS, - * ::CUDA_ERROR_INVALID_VALUE, - * ::CUDA_ERROR_INVALID_HANDLE, - * ::CUDA_ERROR_OUT_OF_MEMORY + * ::cudaSuccess, + * ::cudaErrorInvalidValue * \note_graph_thread_safety * \notefnerr + * \note_init_rt + * \note_callback * * \sa * ::cudaGraphAddExternalSemaphoresSignalNode, @@ -10665,13 +10795,12 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphExternalSemaphoresSignalNodeSetPa * \param nodeParams - Parameters for the node * * \return - * ::CUDA_SUCCESS, - * ::CUDA_ERROR_DEINITIALIZED, - * ::CUDA_ERROR_NOT_INITIALIZED, - * ::CUDA_ERROR_NOT_SUPPORTED, - * ::CUDA_ERROR_INVALID_VALUE + * ::cudaSuccess, + * ::cudaErrorInvalidValue * \note_graph_thread_safety * \notefnerr + * \note_init_rt + * \note_callback * * \sa * ::cudaGraphExternalSemaphoresWaitNodeGetParams, @@ -10709,12 +10838,12 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphAddExternalSemaphoresWaitNode(cud * \param params_out - Pointer to return the parameters * * \return - * ::CUDA_SUCCESS, - * ::CUDA_ERROR_DEINITIALIZED, - * ::CUDA_ERROR_NOT_INITIALIZED, - * ::CUDA_ERROR_INVALID_VALUE + * ::cudaSuccess, + * ::cudaErrorInvalidValue * \note_graph_thread_safety * \notefnerr + * \note_init_rt + * \note_callback * * \sa * ::cudaLaunchKernel, @@ -10737,12 +10866,12 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphExternalSemaphoresWaitNodeGetPara * \param nodeParams - Parameters to copy * * \return - * ::CUDA_SUCCESS, - * ::CUDA_ERROR_INVALID_VALUE, - * ::CUDA_ERROR_INVALID_HANDLE, - * ::CUDA_ERROR_OUT_OF_MEMORY + * ::cudaSuccess, + * ::cudaErrorInvalidValue * \note_graph_thread_safety * \notefnerr + * \note_init_rt + * \note_callback * * \sa * ::cudaGraphAddExternalSemaphoresWaitNode, @@ -10755,6 +10884,293 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphExternalSemaphoresWaitNodeGetPara extern __host__ cudaError_t CUDARTAPI cudaGraphExternalSemaphoresWaitNodeSetParams(cudaGraphNode_t hNode, const struct cudaExternalSemaphoreWaitNodeParams *nodeParams); #endif +/** + * \brief Creates an allocation node and adds it to a graph + * + * Creates a new allocation node and adds it to \p graph with \p numDependencies + * dependencies specified via \p pDependencies 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 pDependencies may not have any duplicate entries. A handle + * to the new node will be returned in \p pGraphNode. + * + * \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 + * + * When ::cudaGraphAddMemAllocNode creates an allocation node, it returns the address of the allocation in + * \p nodeParams.dptr. The allocation's address remains fixed across instantiations and launches. + * + * If the allocation is freed in the same graph, by creating a free node using ::cudaGraphAddMemFreeNode, + * the allocation can be accessed by nodes ordered after the allocation node but before the free node. + * These allocations cannot be freed outside the owning graph, and they can only be freed once in the + * owning graph. + * + * If the allocation is not freed in the same graph, then it can be accessed not only by nodes in the + * graph which are ordered after the allocation node, but also by stream operations ordered after the + * graph's execution but before the allocation is freed. + * + * Allocations which are not freed in the same graph can be freed by: + * - passing the allocation to ::cudaMemFreeAsync or ::cudaMemFree; + * - launching a graph with a free node for that allocation; or + * - specifying ::cudaGraphInstantiateFlagAutoFreeOnLaunch during instantiation, which makes + * each launch behave as though it called ::cudaMemFreeAsync for every unfreed allocation. + * + * It is not possible to free an allocation in both the owning graph and another graph. If the allocation + * is freed in the same graph, a free node cannot be added to another graph. If the allocation is freed + * in another graph, a free node can no longer be added to the owning graph. + * + * The following restrictions apply to graphs which contain allocation and/or memory free nodes: + * - Nodes and edges of the graph cannot be deleted. + * - The graph cannot be used in a child node. + * - Only one instantiation of the graph may exist at any point in time. + * - The graph cannot be cloned. + * + * \return + * ::cudaSuccess, + * ::cudaErrorCudartUnloading, + * ::cudaErrorInitializationError, + * ::cudaErrorNotSupported, + * ::cudaErrorInvalidValue, + * ::cudaErrorOutOfMemory + * \note_graph_thread_safety + * \notefnerr + * + * \sa + * ::cudaGraphAddMemFreeNode, + * ::cudaGraphMemAllocNodeGetParams, + * ::cudaDeviceGraphMemTrim, + * ::cudaDeviceGetGraphMemAttribute, + * ::cudaDeviceSetGraphMemAttribute, + * ::cudaMallocAsync, + * ::cudaFreeAsync, + * ::cudaGraphCreate, + * ::cudaGraphDestroyNode, + * ::cudaGraphAddChildGraphNode, + * ::cudaGraphAddEmptyNode, + * ::cudaGraphAddEventRecordNode, + * ::cudaGraphAddEventWaitNode, + * ::cudaGraphAddExternalSemaphoresSignalNode, + * ::cudaGraphAddExternalSemaphoresWaitNode, + * ::cudaGraphAddKernelNode, + * ::cudaGraphAddMemcpyNode, + * ::cudaGraphAddMemsetNode + */ +#if __CUDART_API_VERSION >= 11040 +extern __host__ cudaError_t CUDARTAPI cudaGraphAddMemAllocNode(cudaGraphNode_t *pGraphNode, cudaGraph_t graph, const cudaGraphNode_t *pDependencies, size_t numDependencies, struct cudaMemAllocNodeParams *nodeParams); +#endif + +/** + * \brief Returns a memory alloc node's parameters + * + * Returns the parameters of a memory alloc node \p hNode in \p params_out. + * The \p poolProps and \p accessDescs returned in \p params_out, are owned by the + * node. This memory remains valid until the node is destroyed. The returned + * parameters must not be modified. + * + * \param node - Node to get the parameters for + * \param params_out - Pointer to return the parameters + * + * \return + * ::cudaSuccess, + * ::cudaErrorInvalidValue + * \note_graph_thread_safety + * \notefnerr + * \note_init_rt + * \note_callback + * + * \sa + * ::cudaGraphAddMemAllocNode, + * ::cudaGraphMemFreeNodeGetParams + */ +#if __CUDART_API_VERSION >= 11040 +extern __host__ cudaError_t CUDARTAPI cudaGraphMemAllocNodeGetParams(cudaGraphNode_t node, struct cudaMemAllocNodeParams *params_out); +#endif + +/** + * \brief Creates a memory free node and adds it to a graph + * + * Creates a new memory free node and adds it to \p graph with \p numDependencies + * dependencies specified via \p pDependencies and address specified in \p dptr. + * It is possible for \p numDependencies to be 0, in which case the node will be placed + * at the root of the graph. \p pDependencies may not have any duplicate entries. A handle + * to the new node will be returned in \p pGraphNode. + * + * \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 dptr - Address of memory to free + * + * ::cudaGraphAddMemFreeNode will return ::cudaErrorInvalidValue if the user attempts to free: + * - an allocation twice in the same graph. + * - an address that was not returned by an allocation node. + * - an invalid address. + * + * The following restrictions apply to graphs which contain allocation and/or memory free nodes: + * - Nodes and edges of the graph cannot be deleted. + * - The graph cannot be used in a child node. + * - Only one instantiation of the graph may exist at any point in time. + * - The graph cannot be cloned. + * + * \return + * ::cudaSuccess, + * ::cudaErrorCudartUnloading, + * ::cudaErrorInitializationError, + * ::cudaErrorNotSupported, + * ::cudaErrorInvalidValue, + * ::cudaErrorOutOfMemory + * \note_graph_thread_safety + * \notefnerr + * + * \sa + * ::cudaGraphAddMemAllocNode, + * ::cudaGraphMemFreeNodeGetParams, + * ::cudaDeviceGraphMemTrim, + * ::cudaDeviceGetGraphMemAttribute, + * ::cudaDeviceSetGraphMemAttribute, + * ::cudaMallocAsync, + * ::cudaFreeAsync, + * ::cudaGraphCreate, + * ::cudaGraphDestroyNode, + * ::cudaGraphAddChildGraphNode, + * ::cudaGraphAddEmptyNode, + * ::cudaGraphAddEventRecordNode, + * ::cudaGraphAddEventWaitNode, + * ::cudaGraphAddExternalSemaphoresSignalNode, + * ::cudaGraphAddExternalSemaphoresWaitNode, + * ::cudaGraphAddKernelNode, + * ::cudaGraphAddMemcpyNode, + * ::cudaGraphAddMemsetNode + */ +#if __CUDART_API_VERSION >= 11040 +extern __host__ cudaError_t CUDARTAPI cudaGraphAddMemFreeNode(cudaGraphNode_t *pGraphNode, cudaGraph_t graph, const cudaGraphNode_t *pDependencies, size_t numDependencies, void *dptr); +#endif + +/** + * \brief Returns a memory free node's parameters + * + * Returns the address of a memory free node \p hNode in \p dptr_out. + * + * \param node - Node to get the parameters for + * \param dptr_out - Pointer to return the device address + * + * \return + * ::cudaSuccess, + * ::cudaErrorInvalidValue + * \note_graph_thread_safety + * \notefnerr + * \note_init_rt + * \note_callback + * + * \sa + * ::cudaGraphAddMemFreeNode, + * ::cudaGraphMemFreeNodeGetParams + */ +#if __CUDART_API_VERSION >= 11040 +extern __host__ cudaError_t CUDARTAPI cudaGraphMemFreeNodeGetParams(cudaGraphNode_t node, void *dptr_out); +#endif + +/** + * \brief Free unused memory that was cached on the specified device for use with graphs back to the OS. + * + * Blocks which are not in use by a graph that is either currently executing or scheduled to execute are + * freed back to the operating system. + * + * \param device - The device for which cached memory should be freed. + * + * \return + * ::cudaSuccess, + * ::cudaErrorInvalidValue + * \note_graph_thread_safety + * \notefnerr + * \note_init_rt + * \note_callback + * + * \sa + * ::cudaGraphAddMemAllocNode, + * ::cudaGraphAddMemFreeNode, + * ::cudaDeviceGetGraphMemAttribute, + * ::cudaDeviceSetGraphMemAttribute, + * ::cudaMallocAsync, + * ::cudaFreeAsync, + */ +#if __CUDART_API_VERSION >= 11040 +extern __host__ cudaError_t CUDARTAPI cudaDeviceGraphMemTrim(int device); +#endif + +/** + * \brief Query asynchronous allocation attributes related to graphs + * + * Valid attributes are: + * + * - ::cudaGraphMemAttrUsedMemCurrent: Amount of memory, in bytes, currently associated with graphs + * - ::cudaGraphMemAttrUsedMemHigh: High watermark of memory, in bytes, associated with graphs since the + * last time it was reset. High watermark can only be reset to zero. + * - ::cudaGraphMemAttrReservedMemCurrent: Amount of memory, in bytes, currently allocated for use by + * the CUDA graphs asynchronous allocator. + * - ::cudaGraphMemAttrReservedMemHigh: High watermark of memory, in bytes, currently allocated for use by + * the CUDA graphs asynchronous allocator. + * + * \param device - Specifies the scope of the query + * \param attr - attribute to get + * \param value - retrieved value + * + * \return + * ::cudaSuccess, + * ::cudaErrorInvalidDevice + * \note_graph_thread_safety + * \notefnerr + * \note_init_rt + * \note_callback + * + * \sa + * ::cudaDeviceSetGraphMemAttribute, + * ::cudaGraphAddMemAllocNode, + * ::cudaGraphAddMemFreeNode, + * ::cudaDeviceGraphMemTrim, + * ::cudaMallocAsync, + * ::cudaFreeAsync, + */ +#if __CUDART_API_VERSION >= 11040 +extern __host__ cudaError_t CUDARTAPI cudaDeviceGetGraphMemAttribute(int device, enum cudaGraphMemAttributeType attr, void* value); +#endif + +/** + * \brief Set asynchronous allocation attributes related to graphs + * + * Valid attributes are: + * + * - ::cudaGraphMemAttrUsedMemHigh: High watermark of memory, in bytes, associated with graphs since the + * last time it was reset. High watermark can only be reset to zero. + * - ::cudaGraphMemAttrReservedMemHigh: High watermark of memory, in bytes, currently allocated for use by + * the CUDA graphs asynchronous allocator. + * + * \param device - Specifies the scope of the query + * \param attr - attribute to get + * \param value - pointer to value to set + * + * \return + * ::cudaSuccess, + * ::cudaErrorInvalidDevice + * \note_graph_thread_safety + * \notefnerr + * \note_init_rt + * \note_callback + * + * \sa + * ::cudaDeviceGetGraphMemAttribute, + * ::cudaGraphAddMemAllocNode, + * ::cudaGraphAddMemFreeNode, + * ::cudaDeviceGraphMemTrim, + * ::cudaMallocAsync, + * ::cudaFreeAsync, + */ +#if __CUDART_API_VERSION >= 11040 +extern __host__ cudaError_t CUDARTAPI cudaDeviceSetGraphMemAttribute(int device, enum cudaGraphMemAttributeType attr, void* value); +#endif + /** * \brief Clones a graph * @@ -11068,6 +11484,9 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphRemoveDependencies(cudaGraph_t gr * Removes \p node from its graph. This operation also severs any dependencies of other nodes * on \p node and vice versa. * + * Dependencies cannot be removed from graphs which contain allocation or free nodes. + * Any attempt to do so will return an error. + * * \param node - Node to remove * * \return @@ -11119,6 +11538,7 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphDestroyNode(cudaGraphNode_t node) * \note_callback * * \sa + * ::cudaGraphInstantiateWithFlags, * ::cudaGraphCreate, * ::cudaGraphUpload, * ::cudaGraphLaunch, @@ -11126,6 +11546,50 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphDestroyNode(cudaGraphNode_t node) */ extern __host__ cudaError_t CUDARTAPI cudaGraphInstantiate(cudaGraphExec_t *pGraphExec, cudaGraph_t graph, cudaGraphNode_t *pErrorNode, char *pLogBuffer, size_t bufferSize); +/** + * \brief Creates an executable graph from a graph + * + * Instantiates \p graph as an executable graph. The graph is validated for any + * structural constraints or intra-node constraints which were not previously + * validated. If instantiation is successful, a handle to the instantiated graph + * is returned in \p pGraphExec. + * + * The \p flags parameter controls the behavior of instantiation and subsequent + * graph launches. Valid flags are: + * + * - ::cudaGraphInstantiateFlagAutoFreeOnLaunch, which configures a + * graph containing memory allocation nodes to automatically free any + * unfreed memory allocations before the graph is relaunched. + * + * If \p graph contains any allocation or free nodes, there can be at most one + * executable graph in existence for that graph at a time. + * + * An attempt to instantiate a second executable graph before destroying the first + * with ::cudaGraphExecDestroy will result in an error. + * + * \param pGraphExec - Returns instantiated graph + * \param graph - Graph to instantiate + * \param flags - Flags to control instantiation. See ::CUgraphInstantiate_flags. + * + * \return + * ::cudaSuccess, + * ::cudaErrorInvalidValue + * \note_graph_thread_safety + * \notefnerr + * \note_init_rt + * \note_callback + * + * \sa + * ::cudaGraphInstantiate, + * ::cudaGraphCreate, + * ::cudaGraphUpload, + * ::cudaGraphLaunch, + * ::cudaGraphExecDestroy + */ +#if __CUDART_API_VERSION >= 11040 +extern __host__ cudaError_t CUDARTAPI cudaGraphInstantiateWithFlags(cudaGraphExec_t *pGraphExec, cudaGraph_t graph, unsigned long long flags); +#endif + /** * \brief Sets the parameters for a kernel node in the given graphExec * @@ -11554,10 +12018,12 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphExecHostNodeSetParams(cudaGraphEx * \param event - Updated event to use * * \return - * ::CUDA_SUCCESS, - * ::CUDA_ERROR_INVALID_VALUE, + * ::cudaSuccess, + * ::cudaErrorInvalidValue, * \note_graph_thread_safety * \notefnerr + * \note_init_rt + * \note_callback * * \sa * ::cudaGraphAddEventRecordNode, @@ -11596,10 +12062,12 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphExecHostNodeSetParams(cudaGraphEx * \param event - Updated event to use * * \return - * ::CUDA_SUCCESS, - * ::CUDA_ERROR_INVALID_VALUE, + * ::cudaSuccess, + * ::cudaErrorInvalidValue, * \note_graph_thread_safety * \notefnerr + * \note_init_rt + * \note_callback * * \sa * ::cudaGraphAddEventWaitNode, @@ -11642,10 +12110,12 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphExecHostNodeSetParams(cudaGraphEx * \param nodeParams - Updated Parameters to set * * \return - * ::CUDA_SUCCESS, - * ::CUDA_ERROR_INVALID_VALUE, + * ::cudaSuccess, + * ::cudaErrorInvalidValue, * \note_graph_thread_safety * \notefnerr + * \note_init_rt + * \note_callback * * \sa * ::cudaGraphAddExternalSemaphoresSignalNode, @@ -11687,10 +12157,12 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphExecExternalSemaphoresSignalNodeS * \param nodeParams - Updated Parameters to set * * \return - * ::CUDA_SUCCESS, - * ::CUDA_ERROR_INVALID_VALUE, + * ::cudaSuccess, + * ::cudaErrorInvalidValue, * \note_graph_thread_safety * \notefnerr + * \note_init_rt + * \note_callback * * \sa * ::cudaGraphAddExternalSemaphoresWaitNode, @@ -11712,6 +12184,80 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphExecExternalSemaphoresSignalNodeS extern __host__ cudaError_t CUDARTAPI cudaGraphExecExternalSemaphoresWaitNodeSetParams(cudaGraphExec_t hGraphExec, cudaGraphNode_t hNode, const struct cudaExternalSemaphoreWaitNodeParams *nodeParams); #endif +/** + * \brief Enables or disables the specified node in the given graphExec + * + * Sets \p hNode to be either enabled or disabled. Disabled nodes are functionally equivalent + * to empty nodes until they are reenabled. Existing node parameters are not affected by + * disabling/enabling the node. + * + * 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. + * + * \note Currently only kernel nodes are supported. + * + * \param hGraphExec - The executable graph in which to set the specified node + * \param hNode - Node from the graph from which graphExec was instantiated + * \param isEnabled - Node is enabled if != 0, otherwise the node is disabled + * + * \return + * ::cudaSuccess, + * ::cudaErrorInvalidValue, + * \note_graph_thread_safety + * \notefnerr + * \note_init_rt + * \note_callback + * + * \sa + * ::cudaGraphNodeGetEnabled, + * ::cudaGraphExecUpdate, + * ::cudaGraphInstantiate + * ::cudaGraphLaunch + */ +#if __CUDART_API_VERSION >= 11060 +extern __host__ cudaError_t CUDARTAPI cudaGraphNodeSetEnabled(cudaGraphExec_t hGraphExec, cudaGraphNode_t hNode, unsigned int isEnabled); +#endif + +/** + * \brief Query whether a node in the given graphExec is enabled + * + * Sets isEnabled to 1 if \p hNode is enabled, or 0 if \p hNode is disabled. + * + * 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. + * + * \note Currently only kernel nodes are supported. + * + * \param hGraphExec - The executable graph in which to set the specified node + * \param hNode - Node from the graph from which graphExec was instantiated + * \param isEnabled - Location to return the enabled status of the node + * + * \return + * ::cudaSuccess, + * ::cudaErrorInvalidValue, + * \note_graph_thread_safety + * \notefnerr + * \note_init_rt + * \note_callback + * + * \sa + * ::cudaGraphNodeSetEnabled, + * ::cudaGraphExecUpdate, + * ::cudaGraphInstantiate + * ::cudaGraphLaunch + */ +#if __CUDART_API_VERSION >= 11060 +extern __host__ cudaError_t CUDARTAPI cudaGraphNodeGetEnabled(cudaGraphExec_t hGraphExec, cudaGraphNode_t hNode, unsigned int *isEnabled); +#endif + /** * \brief Check whether an executable graph can be updated with a graph and perform the update if possible * @@ -11723,7 +12269,8 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphExecExternalSemaphoresWaitNodeSet * - Kernel nodes: * - 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 + * to a function which uses CDP. + * - A cooperative node cannot be updated to a non-cooperative node, and vice-versa. * - 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 @@ -11756,6 +12303,8 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphExecExternalSemaphoresWaitNodeSet * 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 + * - cudaGraphExecUpdateErrorAttributesChanged if any attributes of 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 * the node's type or configuration, in which case \p hErrorNode_out is set to the node from \p hGraph * @@ -11792,6 +12341,7 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphExecUpdate(cudaGraphExec_t hGraph * Uploads \p hGraphExec to the device in \p hStream without executing it. Uploads of * the same \p hGraphExec will be serialized. Each upload is ordered behind both any * previous work in \p hStream and any previous launches of \p hGraphExec. + * Uses memory cached by \p stream to back the allocations owned by \p graphExec. * * \param hGraphExec - Executable graph to upload * \param hStream - Stream in which to upload the graph @@ -11819,6 +12369,10 @@ extern __host__ cudaError_t CUDARTAPI cudaGraphExecUpdate(cudaGraphExec_t hGraph * and any previous launches of \p graphExec. To execute a graph concurrently, it must be * instantiated multiple times into multiple executable graphs. * + * If any allocations created by \p graphExec remain unfreed (from a previous launch) and + * \p graphExec was not instantiated with ::cudaGraphInstantiateFlagAutoFreeOnLaunch, + * the launch will fail with ::cudaErrorInvalidValue. + * * \param graphExec - Executable graph to launch * \param stream - Stream in which to launch the graph * diff --git a/samples/external/cuda/include/device_types.h b/samples/external/cuda/include/device_types.h index 29a261f..4b575a1 100644 --- a/samples/external/cuda/include/device_types.h +++ b/samples/external/cuda/include/device_types.h @@ -55,7 +55,9 @@ #define __UNDEF_CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS_DEVICE_TYPES_H__ #endif +#ifndef __DOXYGEN_ONLY__ #include "crt/host_defines.h" +#endif /******************************************************************************* * * diff --git a/samples/external/cuda/include/driver_types.h b/samples/external/cuda/include/driver_types.h index c6fcbc2..cede349 100644 --- a/samples/external/cuda/include/driver_types.h +++ b/samples/external/cuda/include/driver_types.h @@ -55,9 +55,13 @@ #define __UNDEF_CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS_DRIVER_TYPES_H__ #endif +#ifndef __DOXYGEN_ONLY__ #include "crt/host_defines.h" +#endif #include "vector_types.h" + + /** * \defgroup CUDART_TYPES Data types used by CUDA Runtime * \ingroup CUDART @@ -145,6 +149,9 @@ #define cudaArrayColorAttachment 0x20 /**< Must be set in cudaExternalMemoryGetMappedMipmappedArray if the mipmapped array is used as a color target in a graphics API */ #define cudaArraySparse 0x40 /**< Must be set in cudaMallocArray, cudaMalloc3DArray or cudaMallocMipmappedArray in order to create a sparse CUDA array or CUDA mipmapped array */ +#define cudaArrayDeferredMapping 0x80 /**< Must be set in cudaMallocArray, cudaMalloc3DArray or cudaMallocMipmappedArray in order to create a deferred mapping CUDA array or CUDA mipmapped array */ + + #define cudaIpcMemLazyEnablePeerAccess 0x01 /**< Automatically enable peer access between remote devices as needed */ #define cudaMemAttachGlobal 0x01 /**< Memory can be accessed by any stream on any device*/ @@ -540,7 +547,8 @@ enum __device_builtin__ cudaError /** * This indicates that the device ordinal supplied by the user does not - * correspond to a valid CUDA device. + * correspond to a valid CUDA device or that the action requested is + * invalid for the specified device. */ cudaErrorInvalidDevice = 101, @@ -691,6 +699,11 @@ enum __device_builtin__ cudaError */ cudaErrorJitCompilationDisabled = 223, + /** + * This indicates that the provided execution affinity is not supported by the device. + */ + cudaErrorUnsupportedExecAffinity = 224, + /** * This indicates that the device kernel source is invalid. */ @@ -939,6 +952,32 @@ enum __device_builtin__ cudaError */ cudaErrorCompatNotSupportedOnDevice = 804, + /** + * This error indicates that the MPS client failed to connect to the MPS control daemon or the MPS server. + */ + cudaErrorMpsConnectionFailed = 805, + + /** + * This error indicates that the remote procedural call between the MPS server and the MPS client failed. + */ + cudaErrorMpsRpcFailure = 806, + + /** + * This error indicates that the MPS server is not ready to accept new MPS client requests. + * This error can be returned when the MPS server is in the process of recovering from a fatal failure. + */ + cudaErrorMpsServerNotReady = 807, + + /** + * This error indicates that the hardware resources required to create MPS client have been exhausted. + */ + cudaErrorMpsMaxClientsReached = 808, + + /** + * This error indicates the the hardware resources required to device connections have been exhausted. + */ + cudaErrorMpsMaxConnectionsReached = 809, + /** * The operation is not permitted when the stream is capturing. */ @@ -1004,6 +1043,24 @@ enum __device_builtin__ cudaError */ cudaErrorGraphExecUpdateFailure = 910, + /** + * This indicates that an async error has occurred in a device outside of CUDA. + * If CUDA was waiting for an external device's signal before consuming shared data, + * the external device signaled an error indicating that the data is not valid for + * consumption. This leaves the process in an inconsistent state and any further CUDA + * work will return the same error. To continue using CUDA, the process must be + * terminated and relaunched. + */ + cudaErrorExternalDevice = 911, + + + + + + + + + /** * This indicates that an unknown internal error has occurred. */ @@ -1023,11 +1080,37 @@ enum __device_builtin__ cudaError */ enum __device_builtin__ cudaChannelFormatKind { - cudaChannelFormatKindSigned = 0, /**< Signed channel format */ - cudaChannelFormatKindUnsigned = 1, /**< Unsigned channel format */ - cudaChannelFormatKindFloat = 2, /**< Float channel format */ - cudaChannelFormatKindNone = 3, /**< No channel format */ - cudaChannelFormatKindNV12 = 4 /**< Unsigned 8-bit integers, planar 4:2:0 YUV format */ + cudaChannelFormatKindSigned = 0, /**< Signed channel format */ + cudaChannelFormatKindUnsigned = 1, /**< Unsigned channel format */ + cudaChannelFormatKindFloat = 2, /**< Float channel format */ + cudaChannelFormatKindNone = 3, /**< No channel format */ + cudaChannelFormatKindNV12 = 4, /**< Unsigned 8-bit integers, planar 4:2:0 YUV format */ + cudaChannelFormatKindUnsignedNormalized8X1 = 5, /**< 1 channel unsigned 8-bit normalized integer */ + cudaChannelFormatKindUnsignedNormalized8X2 = 6, /**< 2 channel unsigned 8-bit normalized integer */ + cudaChannelFormatKindUnsignedNormalized8X4 = 7, /**< 4 channel unsigned 8-bit normalized integer */ + cudaChannelFormatKindUnsignedNormalized16X1 = 8, /**< 1 channel unsigned 16-bit normalized integer */ + cudaChannelFormatKindUnsignedNormalized16X2 = 9, /**< 2 channel unsigned 16-bit normalized integer */ + cudaChannelFormatKindUnsignedNormalized16X4 = 10, /**< 4 channel unsigned 16-bit normalized integer */ + cudaChannelFormatKindSignedNormalized8X1 = 11, /**< 1 channel signed 8-bit normalized integer */ + cudaChannelFormatKindSignedNormalized8X2 = 12, /**< 2 channel signed 8-bit normalized integer */ + cudaChannelFormatKindSignedNormalized8X4 = 13, /**< 4 channel signed 8-bit normalized integer */ + cudaChannelFormatKindSignedNormalized16X1 = 14, /**< 1 channel signed 16-bit normalized integer */ + cudaChannelFormatKindSignedNormalized16X2 = 15, /**< 2 channel signed 16-bit normalized integer */ + cudaChannelFormatKindSignedNormalized16X4 = 16, /**< 4 channel signed 16-bit normalized integer */ + cudaChannelFormatKindUnsignedBlockCompressed1 = 17, /**< 4 channel unsigned normalized block-compressed (BC1 compression) format */ + cudaChannelFormatKindUnsignedBlockCompressed1SRGB = 18, /**< 4 channel unsigned normalized block-compressed (BC1 compression) format with sRGB encoding*/ + cudaChannelFormatKindUnsignedBlockCompressed2 = 19, /**< 4 channel unsigned normalized block-compressed (BC2 compression) format */ + cudaChannelFormatKindUnsignedBlockCompressed2SRGB = 20, /**< 4 channel unsigned normalized block-compressed (BC2 compression) format with sRGB encoding */ + cudaChannelFormatKindUnsignedBlockCompressed3 = 21, /**< 4 channel unsigned normalized block-compressed (BC3 compression) format */ + cudaChannelFormatKindUnsignedBlockCompressed3SRGB = 22, /**< 4 channel unsigned normalized block-compressed (BC3 compression) format with sRGB encoding */ + cudaChannelFormatKindUnsignedBlockCompressed4 = 23, /**< 1 channel unsigned normalized block-compressed (BC4 compression) format */ + cudaChannelFormatKindSignedBlockCompressed4 = 24, /**< 1 channel signed normalized block-compressed (BC4 compression) format */ + cudaChannelFormatKindUnsignedBlockCompressed5 = 25, /**< 2 channel unsigned normalized block-compressed (BC5 compression) format */ + cudaChannelFormatKindSignedBlockCompressed5 = 26, /**< 2 channel signed normalized block-compressed (BC5 compression) format */ + cudaChannelFormatKindUnsignedBlockCompressed6H = 27, /**< 3 channel unsigned half-float block-compressed (BC6H compression) format */ + cudaChannelFormatKindSignedBlockCompressed6H = 28, /**< 3 channel signed half-float block-compressed (BC6H compression) format */ + cudaChannelFormatKindUnsignedBlockCompressed7 = 29, /**< 4 channel unsigned normalized block-compressed (BC7 compression) format */ + cudaChannelFormatKindUnsignedBlockCompressed7SRGB = 30 /**< 4 channel unsigned normalized block-compressed (BC7 compression) format with sRGB encoding */ }; /** @@ -1086,6 +1169,17 @@ struct __device_builtin__ cudaArraySparseProperties { unsigned int reserved[4]; }; + +/** + * CUDA array and CUDA mipmapped array memory requirements + */ +struct __device_builtin__ cudaArrayMemoryRequirements { + size_t size; /**< Total size of the array. */ + size_t alignment; /**< Alignment necessary for mapping the array. */ + unsigned int reserved[4]; +}; + + /** * CUDA memory types */ @@ -1566,6 +1660,55 @@ struct __device_builtin__ cudaFuncAttributes * See ::cudaFuncSetAttribute */ int preferredShmemCarveout; + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + }; /** @@ -1575,6 +1718,14 @@ enum __device_builtin__ cudaFuncAttribute { cudaFuncAttributeMaxDynamicSharedMemorySize = 8, /**< Maximum dynamic shared memory size */ cudaFuncAttributePreferredSharedMemoryCarveout = 9, /**< Preferred shared memory-L1 cache split */ + + + + + + + + cudaFuncAttributeMax }; @@ -1609,6 +1760,17 @@ enum __device_builtin__ cudaSharedCarveout { cudaSharedmemCarveoutMaxL1 = 0 /**< Prefer maximum available L1 cache, minimum shared memory */ }; + + + + + + + + + + + /** * CUDA device compute modes */ @@ -1799,7 +1961,7 @@ enum __device_builtin__ cudaDeviceAttr cudaDevAttrReserved93 = 93, cudaDevAttrReserved94 = 94, cudaDevAttrCooperativeLaunch = 95, /**< Device supports launching cooperative kernels via ::cudaLaunchCooperativeKernel*/ - cudaDevAttrCooperativeMultiDeviceLaunch = 96, /**< Device can participate in cooperative kernels launched via ::cudaLaunchCooperativeKernelMultiDevice */ + cudaDevAttrCooperativeMultiDeviceLaunch = 96, /**< Deprecated, cudaLaunchCooperativeKernelMultiDevice is deprecated. */ cudaDevAttrMaxSharedMemoryPerBlockOptin = 97, /**< The maximum optin shared memory per block. This value may vary by chip. See ::cudaFuncSetAttribute */ cudaDevAttrCanFlushRemoteWrites = 98, /**< Device supports flushing of outstanding remote writes. */ cudaDevAttrHostRegisterSupported = 99, /**< Device supports host memory registration via ::cudaHostRegister. */ @@ -1811,12 +1973,20 @@ enum __device_builtin__ cudaDeviceAttr 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 ::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 */ + cudaDevAttrTimelineSemaphoreInteropSupported = 114, /**< External timeline semaphore interop is supported on the device */ + cudaDevAttrMaxTimelineSemaphoreInteropSupported = 114, /**< Deprecated, 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 */ + cudaDevAttrMemoryPoolSupportedHandleTypes = 119, /**< Handle types supported with mempool based IPC */ + + + + + cudaDevAttrDeferredMappingCudaArraySupported = 121, /**< Device supports deferred mapping CUDA arrays and CUDA mipmapped arrays */ + + cudaDevAttrMax }; /** @@ -1968,6 +2138,53 @@ struct __device_builtin__ cudaMemPoolPtrExportData { unsigned char reserved[64]; }; +/** + * Memory allocation node parameters + */ +struct __device_builtin__ cudaMemAllocNodeParams { + /** + * in: location where the allocation should reside (specified in ::location). + * ::handleTypes must be ::cudaMemHandleTypeNone. IPC is not supported. + */ + struct cudaMemPoolProps poolProps; /**< in: array of memory access descriptors. Used to describe peer GPU access */ + const struct cudaMemAccessDesc *accessDescs; /**< in: number of memory access descriptors. Must not exceed the number of GPUs. */ + size_t accessDescCount; /**< in: Number of `accessDescs`s */ + size_t bytesize; /**< in: size in bytes of the requested allocation */ + void *dptr; /**< out: address of the allocation returned by CUDA */ +}; + +/** + * Graph memory attributes + */ +enum __device_builtin__ cudaGraphMemAttributeType { + /** + * (value type = cuuint64_t) + * Amount of memory, in bytes, currently associated with graphs. + */ + cudaGraphMemAttrUsedMemCurrent = 0x0, + + /** + * (value type = cuuint64_t) + * High watermark of memory, in bytes, associated with graphs since the + * last time it was reset. High watermark can only be reset to zero. + */ + cudaGraphMemAttrUsedMemHigh = 0x1, + + /** + * (value type = cuuint64_t) + * Amount of memory, in bytes, currently allocated for use by + * the CUDA graphs asynchronous allocator. + */ + cudaGraphMemAttrReservedMemCurrent = 0x2, + + /** + * (value type = cuuint64_t) + * High watermark of memory, in bytes, currently allocated for use by + * the CUDA graphs asynchronous allocator. + */ + cudaGraphMemAttrReservedMemHigh = 0x3 +}; + /** * CUDA device P2P attributes */ @@ -2161,6 +2378,11 @@ struct __device_builtin__ cudaDeviceProp 0, /* size_t reservedSharedMemPerBlock */ \ } /**< Empty device properties */ + + + + + /** * CUDA IPC Handle Size */ @@ -2459,7 +2681,6 @@ struct __device_builtin__ cudaExternalSemaphoreHandleDesc { unsigned int flags; }; -#if defined(__CUDA_API_VERSION_INTERNAL) /** * External semaphore signal parameters(deprecated) */ @@ -2553,7 +2774,6 @@ struct __device_builtin__ cudaExternalSemaphoreWaitParams_v1 { */ unsigned int flags; }; -#endif /** * External semaphore signal parameters, compatible with driver type @@ -2784,6 +3004,10 @@ enum __device_builtin__ cudaGraphNodeType { cudaGraphNodeTypeEmpty = 0x05, /**< Empty (no-op) node */ cudaGraphNodeTypeWaitEvent = 0x06, /**< External event wait node */ cudaGraphNodeTypeEventRecord = 0x07, /**< External event record node */ + cudaGraphNodeTypeExtSemaphoreSignal = 0x08, /**< External semaphore signal node */ + cudaGraphNodeTypeExtSemaphoreWait = 0x09, /**< External semaphore wait node */ + cudaGraphNodeTypeMemAlloc = 0x0a, /**< Memory allocation node */ + cudaGraphNodeTypeMemFree = 0x0b, /**< Memory free node */ cudaGraphNodeTypeCount }; @@ -2803,7 +3027,8 @@ enum __device_builtin__ cudaGraphExecUpdateResult { 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 */ - cudaGraphExecUpdateErrorUnsupportedFunctionChange = 0x7 /**< The update failed because the function of a kernel node changed in an unsupported way */ + cudaGraphExecUpdateErrorUnsupportedFunctionChange = 0x7, /**< The update failed because the function of a kernel node changed in an unsupported way */ + cudaGraphExecUpdateErrorAttributesChanged = 0x8 /**< The update failed because the node attributes changed in a way that is not supported */ }; /** @@ -2820,16 +3045,23 @@ enum __device_builtin__ cudaGetDriverEntryPointFlags { * 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 */ + 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 */ +}; + +/** + * Flags for instantiating a graph + */ +enum __device_builtin__ cudaGraphInstantiateFlags { + cudaGraphInstantiateFlagAutoFreeOnLaunch = 1 /**< Automatically free memory allocated in a graph before relaunching. */ }; /** @} */ diff --git a/samples/external/cuda/include/texture_types.h b/samples/external/cuda/include/texture_types.h index bfaf692..6c41742 100644 --- a/samples/external/cuda/include/texture_types.h +++ b/samples/external/cuda/include/texture_types.h @@ -212,6 +212,10 @@ struct __device_builtin__ cudaTextureDesc * Disable any trilinear filtering optimizations. */ int disableTrilinearOptimization; + /** + * Enable seamless cube map filtering. + */ + int seamlessCubemap; }; /** diff --git a/samples/external/cuda/include/vector_types.h b/samples/external/cuda/include/vector_types.h index ea8e6ce..4cfabcf 100644 --- a/samples/external/cuda/include/vector_types.h +++ b/samples/external/cuda/include/vector_types.h @@ -61,7 +61,9 @@ * * *******************************************************************************/ +#ifndef __DOXYGEN_ONLY__ #include "crt/host_defines.h" +#endif /******************************************************************************* * * diff --git a/samples/external/cuda/lib/x64/cudart.lib b/samples/external/cuda/lib/x64/cudart.lib index 19fe0c2..efe71be 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/input/input_0_100_frames.mp4 b/samples/input/input_0_100_frames.mp4 new file mode 100644 index 0000000..0180a8b Binary files /dev/null and b/samples/input/input_0_100_frames.mp4 differ diff --git a/samples/input/input_100_200_frames.mp4 b/samples/input/input_100_200_frames.mp4 new file mode 100644 index 0000000..a9ff2ee Binary files /dev/null and b/samples/input/input_100_200_frames.mp4 differ diff --git a/samples/utils/nvCVOpenCV.h b/samples/utils/nvCVOpenCV.h index 271d59a..9b5ebbb 100644 --- a/samples/utils/nvCVOpenCV.h +++ b/samples/utils/nvCVOpenCV.h @@ -58,13 +58,13 @@ 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, NVCV_TYPE_UNKNOWN}; + 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 = (int)cvIm->step[0]; - nvcvIm->pixelFormat = nvFormat[cvIm->channels() <= 4 ? cvIm->channels() : 0]; + nvcvIm->pixelFormat = nvFormat[cvIm->channels() <= 4 ? cvIm->channels() : 0]; nvcvIm->componentType = nvType[cvIm->depth() & 7]; nvcvIm->bufferBytes = 0; nvcvIm->deletePtr = nullptr; diff --git a/version.h b/version.h index aba1505..19928fa 100644 --- a/version.h +++ b/version.h @@ -23,13 +23,13 @@ #define NVIDIA_VIDEOEFFECTS_SDK_VERSION_MAJOR 0 -#define NVIDIA_VIDEOEFFECTS_SDK_VERSION_MINOR 6 -#define NVIDIA_VIDEOEFFECTS_SDK_VERSION_RELEASE 5 -#define NVIDIA_VIDEOEFFECTS_SDK_VERSION_BUILD 2 +#define NVIDIA_VIDEOEFFECTS_SDK_VERSION_MINOR 7 +#define NVIDIA_VIDEOEFFECTS_SDK_VERSION_RELEASE 1 +#define NVIDIA_VIDEOEFFECTS_SDK_VERSION_BUILD 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.5.2" -#define NVIDIA_VIDEOEFFECTS_SDK_VERSION_STRING_SHORT "0.6.5" -#define NVIDIA_VIDEOEFFECTS_SDK_VERSION_STRING_MAJOR_MINOR "0.6" +#define NVIDIA_VIDEOEFFECTS_SDK_VERSION 0,7,1,0 +#define NVIDIA_VIDEOEFFECTS_SDK_VERSION_MAJOR_MINOR 0,7 +#define NVIDIA_VIDEOEFFECTS_SDK_VERSION_STRING "0.7.1.0" +#define NVIDIA_VIDEOEFFECTS_SDK_VERSION_STRING_SHORT "0.7.1" +#define NVIDIA_VIDEOEFFECTS_SDK_VERSION_STRING_MAJOR_MINOR "0.7"