Update the CMakeList for remaining tegra samples and remove the old Makefile/NsightEclipse.xml

Samples/4_CUDA_Libraries/CMakeLists.txt

@@ -9,6 +9,7 @@ add_subdirectory(conjugateGradientMultiBlockCG)
 add_subdirectory(conjugateGradientMultiDeviceCG)
 add_subdirectory(conjugateGradientPrecond)
 add_subdirectory(conjugateGradientUM)
+add_subdirectory(cudaNvSci)
 add_subdirectory(cuSolverDn_LinearSolver)
 add_subdirectory(cuSolverRf)
 add_subdirectory(cuSolverSp_LinearSolver)

Samples/4_CUDA_Libraries/cudaNvSci/.vscode/c_cpp_properties.json

@@ -0,0 +1,18 @@
+{
+    "configurations": [
+        {
+            "name": "Linux",
+            "includePath": [
+                "${workspaceFolder}/**",
+                "${workspaceFolder}/../../../Common"
+            ],
+            "defines": [],
+            "compilerPath": "/usr/local/cuda/bin/nvcc",
+            "cStandard": "gnu17",
+            "cppStandard": "gnu++14",
+            "intelliSenseMode": "linux-gcc-x64",
+            "configurationProvider": "ms-vscode.makefile-tools"
+        }
+    ],
+    "version": 4
+}

Samples/4_CUDA_Libraries/cudaNvSci/.vscode/extensions.json

@@ -0,0 +1,7 @@
+{
+    "recommendations": [
+        "nvidia.nsight-vscode-edition",
+        "ms-vscode.cpptools",
+        "ms-vscode.makefile-tools"
+    ]
+}

Samples/4_CUDA_Libraries/cudaNvSci/.vscode/launch.json

@@ -0,0 +1,10 @@
+{
+    "configurations": [
+        {
+            "name": "CUDA C++: Launch",
+            "type": "cuda-gdb",
+            "request": "launch",
+            "program": "${workspaceFolder}/cudaNvSci"
+        }
+    ]
+}

Samples/4_CUDA_Libraries/cudaNvSci/.vscode/tasks.json

@@ -0,0 +1,15 @@
+{
+    "version": "2.0.0",
+    "tasks": [
+        {
+            "label": "sample",
+            "type": "shell",
+            "command": "make dbg=1",
+            "problemMatcher": ["$nvcc"],
+            "group": {
+                "kind": "build",
+                "isDefault": true
+            }
+        }
+    ]
+}

Samples/4_CUDA_Libraries/cudaNvSci/CMakeLists.txt

@@ -0,0 +1,75 @@
+cmake_minimum_required(VERSION 3.20)
+
+list(APPEND CMAKE_MODULE_PATH "${CMAKE_SOURCE_DIR}/cmake/Modules")
+
+project(cudaNvSci LANGUAGES C CXX CUDA)
+
+find_package(CUDAToolkit REQUIRED)
+
+set(CMAKE_POSITION_INDEPENDENT_CODE ON)
+
+set(CMAKE_CUDA_ARCHITECTURES 53 61 70 72 75 80 86 87 90)
+if(CMAKE_BUILD_TYPE STREQUAL "Debug")
+    # set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -G")  # enable cuda-gdb (expensive)
+endif()
+
+# Include directories and libraries
+include_directories(../../../Common)
+
+if(CMAKE_SYSTEM_NAME STREQUAL "Linux")
+    # Find the NVSCI libraries
+    # use CMAKE_LIBRARY_PATH so that users can also specify the NVSCI lib path in cmake command
+    set(CMAKE_LIBRARY_PATH "/usr/lib" ${CMAKE_LIBRARY_PATH})
+    file(GLOB_RECURSE NVSCIBUF_LIB
+         ${CMAKE_LIBRARY_PATH}/*/libnvscibuf.so
+    )
+    file(GLOB_RECURSE NVSCISYNC_LIB
+         ${CMAKE_LIBRARY_PATH}/*/libnvscisync.so
+    )
+
+    # Find the NVSCI header files
+    # use CMAKE_INCLUDE_PATH so that users can also specify the NVSCI include path in cmake command
+    set(CMAKE_INCLUDE_PATH "/usr/include" ${CMAKE_LIBRARY_PATH})
+    find_path(NVSCIBUF_INCLUDE_DIR nvscibuf.h PATHS ${CMAKE_INCLUDE_PATH})
+    find_path(NVSCISYNC_INCLUDE_DIR nvscisync.h PATHS ${CMAKE_INCLUDE_PATH})
+
+    if(NVSCIBUF_LIB AND NVSCISYNC_LIB AND NVSCIBUF_INCLUDE_DIR AND NVSCISYNC_INCLUDE_DIR)
+        message(STATUS "FOUND NVSCI libs: ${NVSCIBUF_LIB} ${NVSCISYNC_LIB}")
+        message(STATUS "Using NVSCI headers path: ${NVSCIBUF_INCLUDE_DIR} ${NVSCIBUF_INCLUDE_DIR}")
+        # Source file
+        # Add target for cudaNvSci
+        add_executable(cudaNvSci imageKernels.cu cudaNvSci.cpp main.cpp)
+
+        target_compile_options(cudaNvSci PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:--extended-lambda>)
+
+        target_compile_features(cudaNvSci PRIVATE cxx_std_17 cuda_std_17)
+
+        set_target_properties(cudaNvSci PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
+
+        target_include_directories(cudaNvSci PUBLIC
+            ${CUDAToolkit_INCLUDE_DIRS}
+            ${NVSCIBUF_INCLUDE_DIR}
+            ${NVSCISYNC_INCLUDE_DIR}
+        )
+
+        target_link_libraries(cudaNvSci
+            CUDA::cuda_driver
+            ${NVSCIBUF_LIB}
+            ${NVSCISYNC_LIB}
+        )
+        # Copy teapot1024.ppm to the output directory
+        add_custom_command(TARGET cudaNvSci POST_BUILD
+            COMMAND ${CMAKE_COMMAND} -E copy_if_different
+            ${CMAKE_CURRENT_SOURCE_DIR}/teapot1024.ppm ${CMAKE_CURRENT_BINARY_DIR}/teapot1024.ppm
+        )
+
+        # Specify additional clean files
+        set_target_properties(cudaNvSci PROPERTIES
+            ADDITIONAL_CLEAN_FILES "teapot1024_out.ppm"
+        )
+    else()
+        message(STATUS "NvSCI not found - will not build sample 'cudaNvSci'")
+    endif()
+else()
+    message(STATUS "Will not build sample cudaNvSci - requires Linux OS")
+endif()

Samples/4_CUDA_Libraries/cudaNvSci/README.md

@@ -0,0 +1,40 @@
+# cudaNvSci - CUDA NvSciBuf/NvSciSync Interop
+
+## Description
+
+This sample demonstrates CUDA-NvSciBuf/NvSciSync Interop. Two CPU threads import the NvSciBuf and NvSciSync into CUDA to perform two image processing algorithms on a ppm image - image rotation in 1st thread & rgba to grayscale conversion of rotated image in 2nd thread. Currently only supported on Ubuntu 18.04
+
+## Key Concepts
+
+CUDA NvSci Interop, Data Parallel Algorithms, Image Processing
+
+## Supported SM Architectures
+
+[SM 6.0 ](https://developer.nvidia.com/cuda-gpus)  [SM 6.1 ](https://developer.nvidia.com/cuda-gpus)  [SM 7.0 ](https://developer.nvidia.com/cuda-gpus)  [SM 7.2 ](https://developer.nvidia.com/cuda-gpus)  [SM 7.5 ](https://developer.nvidia.com/cuda-gpus)  [SM 8.0 ](https://developer.nvidia.com/cuda-gpus)  [SM 8.6 ](https://developer.nvidia.com/cuda-gpus)  [SM 8.7 ](https://developer.nvidia.com/cuda-gpus)  [SM 8.9 ](https://developer.nvidia.com/cuda-gpus)  [SM 9.0 ](https://developer.nvidia.com/cuda-gpus)
+
+## Supported OSes
+
+Linux
+
+## Supported CPU Architecture
+
+x86_64, aarch64
+
+## CUDA APIs involved
+
+### [CUDA Driver API](http://docs.nvidia.com/cuda/cuda-driver-api/index.html)
+cuDeviceGetUuid
+
+### [CUDA Runtime API](http://docs.nvidia.com/cuda/cuda-runtime-api/index.html)
+cudaExternalMemoryGetMappedBuffer, cudaImportExternalSemaphore, cudaDeviceGetAttribute, cudaNvSciSignal, cudaGetMipmappedArrayLevel, cudaImportNvSciRawBuf, cudaSetDevice, cudaImportNvSciImage, cudaNvSciApp, cudaDeviceId, cudaMallocHost, cudaSignalExternalSemaphoresAsync, cudaCreateTextureObject, cudaFreeHost, cudaNvSci, cudaNvSciWait, cudaGetDeviceCount, cudaMemcpyAsync, cudaStreamCreateWithFlags, cudaExternalMemoryGetMappedMipmappedArray, cudaStreamDestroy, cudaDeviceGetNvSciSyncAttributes, cudaDestroyTextureObject, cudaDestroyExternalMemory, cudaImportExternalMemory, cudaDestroyExternalSemaphore, cudaFreeMipmappedArray, cudaFree, cudaStreamSynchronize, cudaWaitExternalSemaphoresAsync, cudaImportNvSciSemaphore
+
+## Dependencies needed to build/run
+[NVSCI](../../../README.md#nvsci)
+
+## Prerequisites
+
+Download and install the [CUDA Toolkit 12.5](https://developer.nvidia.com/cuda-downloads) for your corresponding platform.
+Make sure the dependencies mentioned in [Dependencies]() section above are installed.
+
+## References (for more details)
+

Samples/4_CUDA_Libraries/cudaNvSci/cudaNvSci.cpp

@@ -0,0 +1,610 @@
+/* Copyright (c) 2022, NVIDIA CORPORATION. All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions
+ * are met:
+ *  * Redistributions of source code must retain the above copyright
+ *    notice, this list of conditions and the following disclaimer.
+ *  * Redistributions in binary form must reproduce the above copyright
+ *    notice, this list of conditions and the following disclaimer in the
+ *    documentation and/or other materials provided with the distribution.
+ *  * Neither the name of NVIDIA CORPORATION nor the names of its
+ *    contributors may be used to endorse or promote products derived
+ *    from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
+ * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ * PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
+ * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+
+#include "cudaNvSci.h"
+#include <cuda.h>
+#include <condition_variable>
+#include <iostream>
+#include <thread>
+
+std::mutex m_mutex;
+std::condition_variable m_condVar;
+bool workSubmitted = false;
+
+class cudaNvSciSignal {
+ private:
+  NvSciSyncModule m_syncModule;
+  NvSciBufModule m_bufModule;
+
+  NvSciSyncAttrList m_syncAttrList;
+  NvSciSyncFence *m_fence;
+
+  NvSciBufAttrList m_rawBufAttrList;
+  NvSciBufAttrList m_imageBufAttrList;
+  NvSciBufAttrList m_buffAttrListOut[2];
+  NvSciBufAttrKeyValuePair pairArrayOut[10];
+
+  cudaExternalMemory_t extMemRawBuf, extMemImageBuf;
+  cudaMipmappedArray_t d_mipmapArray;
+  cudaArray_t d_mipLevelArray;
+  cudaTextureObject_t texObject;
+  cudaExternalSemaphore_t signalSem;
+
+  cudaStream_t streamToRun;
+  int m_cudaDeviceId;
+  CUuuid m_devUUID;
+  uint64_t m_imageWidth;
+  uint64_t m_imageHeight;
+  void *d_outputBuf;
+  size_t m_bufSize;
+
+ public:
+  cudaNvSciSignal(NvSciBufModule bufModule, NvSciSyncModule syncModule,
+                  int cudaDeviceId, int bufSize, uint64_t imageWidth,
+                  uint64_t imageHeight, NvSciSyncFence *fence)
+      : m_syncModule(syncModule),
+        m_bufModule(bufModule),
+        m_cudaDeviceId(cudaDeviceId),
+        d_outputBuf(NULL),
+        m_bufSize(bufSize),
+        m_imageWidth(imageWidth),
+        m_imageHeight(imageHeight),
+        m_fence(fence) {
+    initCuda();
+
+    checkNvSciErrors(NvSciSyncAttrListCreate(m_syncModule, &m_syncAttrList));
+    checkNvSciErrors(NvSciBufAttrListCreate(m_bufModule, &m_rawBufAttrList));
+    checkNvSciErrors(NvSciBufAttrListCreate(m_bufModule, &m_imageBufAttrList));
+
+    setRawBufAttrList(m_bufSize);
+    setImageBufAttrList(m_imageWidth, m_imageHeight);
+
+    checkCudaErrors(cudaDeviceGetNvSciSyncAttributes(
+        m_syncAttrList, m_cudaDeviceId, cudaNvSciSyncAttrSignal));
+  }
+
+  ~cudaNvSciSignal() {
+    checkCudaErrors(cudaSetDevice(m_cudaDeviceId));
+    checkCudaErrors(cudaFreeMipmappedArray(d_mipmapArray));
+    checkCudaErrors(cudaFree(d_outputBuf));
+    checkCudaErrors(cudaDestroyExternalSemaphore(signalSem));
+    checkCudaErrors(cudaDestroyExternalMemory(extMemRawBuf));
+    checkCudaErrors(cudaDestroyExternalMemory(extMemImageBuf));
+    checkCudaErrors(cudaDestroyTextureObject(texObject));
+    checkCudaErrors(cudaStreamDestroy(streamToRun));
+  }
+
+  void initCuda() {
+    checkCudaErrors(cudaSetDevice(m_cudaDeviceId));
+    checkCudaErrors(
+        cudaStreamCreateWithFlags(&streamToRun, cudaStreamNonBlocking));
+
+    int major = 0, minor = 0;
+    checkCudaErrors(cudaDeviceGetAttribute(
+        &major, cudaDevAttrComputeCapabilityMajor, m_cudaDeviceId));
+    checkCudaErrors(cudaDeviceGetAttribute(
+        &minor, cudaDevAttrComputeCapabilityMinor, m_cudaDeviceId));
+    printf(
+        "[cudaNvSciSignal] GPU Device %d: \"%s\" with compute capability "
+        "%d.%d\n\n",
+        m_cudaDeviceId, _ConvertSMVer2ArchName(major, minor), major, minor);
+
+#ifdef cuDeviceGetUuid_v2
+    CUresult res = cuDeviceGetUuid_v2(&m_devUUID, m_cudaDeviceId);
+#else
+    CUresult res = cuDeviceGetUuid(&m_devUUID, m_cudaDeviceId);
+#endif
+
+    if (res != CUDA_SUCCESS) {
+      fprintf(stderr, "Driver API error = %04d \n", res);
+      exit(EXIT_FAILURE);
+    }
+  }
+
+  void setRawBufAttrList(uint64_t size) {
+    NvSciBufType bufType = NvSciBufType_RawBuffer;
+    bool cpuAccess = false;
+    NvSciBufAttrValAccessPerm perm = NvSciBufAccessPerm_ReadWrite;
+    NvSciBufAttrKeyValuePair rawBufAttrs[] = {
+        {NvSciBufRawBufferAttrKey_Size, &size, sizeof(size)},
+        {NvSciBufGeneralAttrKey_Types, &bufType, sizeof(bufType)},
+        {NvSciBufGeneralAttrKey_NeedCpuAccess, &cpuAccess, sizeof(cpuAccess)},
+        {NvSciBufGeneralAttrKey_RequiredPerm, &perm, sizeof(perm)},
+        {NvSciBufGeneralAttrKey_GpuId, &m_devUUID, sizeof(m_devUUID)},
+    };
+
+    checkNvSciErrors(NvSciBufAttrListSetAttrs(
+        m_rawBufAttrList, rawBufAttrs,
+        sizeof(rawBufAttrs) / sizeof(NvSciBufAttrKeyValuePair)));
+  }
+
+  void setImageBufAttrList(uint32_t width, uint32_t height) {
+    NvSciBufType bufType = NvSciBufType_Image;
+    NvSciBufAttrValImageLayoutType layout = NvSciBufImage_BlockLinearType;
+    NvSciBufAttrValAccessPerm perm = NvSciBufAccessPerm_ReadWrite;
+
+    uint32_t planeCount = 1;
+    uint32_t planeWidths[] = {width};
+    uint32_t planeHeights[] = {height};
+    uint64_t lrpad = 0, tbpad = 100;
+
+    bool cpuAccessFlag = false;
+
+    NvSciBufAttrValColorFmt planecolorfmts[] = {NvSciColor_B8G8R8A8};
+    NvSciBufAttrValColorStd planecolorstds[] = {NvSciColorStd_SRGB};
+    NvSciBufAttrValImageScanType planescantype[] = {NvSciBufScan_InterlaceType};
+
+    NvSciBufAttrKeyValuePair imgBufAttrs[] = {
+        {NvSciBufGeneralAttrKey_Types, &bufType, sizeof(bufType)},
+        {NvSciBufImageAttrKey_PlaneCount, &planeCount, sizeof(planeCount)},
+        {NvSciBufImageAttrKey_Layout, &layout, sizeof(layout)},
+        {NvSciBufImageAttrKey_TopPadding, &tbpad, sizeof(tbpad)},
+        {NvSciBufImageAttrKey_BottomPadding, &tbpad, sizeof(tbpad)},
+        {NvSciBufImageAttrKey_LeftPadding, &lrpad, sizeof(lrpad)},
+        {NvSciBufImageAttrKey_RightPadding, &lrpad, sizeof(lrpad)},
+        {NvSciBufImageAttrKey_PlaneColorFormat, planecolorfmts,
+         sizeof(planecolorfmts)},
+        {NvSciBufImageAttrKey_PlaneColorStd, planecolorstds,
+         sizeof(planecolorstds)},
+        {NvSciBufImageAttrKey_PlaneWidth, planeWidths, sizeof(planeWidths)},
+        {NvSciBufImageAttrKey_PlaneHeight, planeHeights, sizeof(planeHeights)},
+        {NvSciBufGeneralAttrKey_NeedCpuAccess, &cpuAccessFlag,
+         sizeof(cpuAccessFlag)},
+        {NvSciBufGeneralAttrKey_RequiredPerm, &perm, sizeof(perm)},
+        {NvSciBufImageAttrKey_PlaneScanType, planescantype,
+         sizeof(planescantype)},
+        {NvSciBufGeneralAttrKey_GpuId, &m_devUUID, sizeof(m_devUUID)},
+    };
+
+    checkNvSciErrors(NvSciBufAttrListSetAttrs(
+        m_imageBufAttrList, imgBufAttrs,
+        sizeof(imgBufAttrs) / sizeof(NvSciBufAttrKeyValuePair)));
+  }
+
+  NvSciSyncAttrList getNvSciSyncAttrList() { return m_syncAttrList; }
+
+  NvSciBufAttrList getNvSciRawBufAttrList() { return m_rawBufAttrList; }
+
+  NvSciBufAttrList getNvSciImageBufAttrList() { return m_imageBufAttrList; }
+
+  void runRotateImageAndSignal(unsigned char *imageData) {
+    int numOfGPUs = 0;
+    checkCudaErrors(cudaGetDeviceCount(&numOfGPUs));  // For cuda init purpose
+    checkCudaErrors(cudaSetDevice(m_cudaDeviceId));
+
+    copyDataToImageArray(imageData);
+    createTexture();
+
+    float angle = 0.5f;  // angle to rotate image by (in radians)
+    rotateKernel(texObject, angle, (unsigned int *)d_outputBuf, m_imageWidth,
+                 m_imageHeight, streamToRun);
+
+    signalExternalSemaphore();
+  }
+
+  void cudaImportNvSciSemaphore(NvSciSyncObj syncObj) {
+    checkCudaErrors(cudaSetDevice(m_cudaDeviceId));
+
+    cudaExternalSemaphoreHandleDesc extSemDesc;
+    memset(&extSemDesc, 0, sizeof(extSemDesc));
+    extSemDesc.type = cudaExternalSemaphoreHandleTypeNvSciSync;
+    extSemDesc.handle.nvSciSyncObj = (void *)syncObj;
+
+    checkCudaErrors(cudaImportExternalSemaphore(&signalSem, &extSemDesc));
+  }
+
+  void signalExternalSemaphore() {
+    cudaExternalSemaphoreSignalParams signalParams;
+    memset(&signalParams, 0, sizeof(signalParams));
+    // For cross-process signaler-waiter applications need to use NvSciIpc
+    // and NvSciSync[Export|Import] utilities to share the NvSciSyncFence
+    // across process. This step is optional in single-process.
+    signalParams.params.nvSciSync.fence = (void *)m_fence;
+    signalParams.flags = 0;
+
+    checkCudaErrors(cudaSignalExternalSemaphoresAsync(&signalSem, &signalParams,
+                                                      1, streamToRun));
+  }
+
+  void cudaImportNvSciRawBuf(NvSciBufObj inputBufObj) {
+    checkCudaErrors(cudaSetDevice(m_cudaDeviceId));
+    checkNvSciErrors(
+        NvSciBufObjGetAttrList(inputBufObj, &m_buffAttrListOut[0]));
+
+    memset(pairArrayOut, 0, sizeof(NvSciBufAttrKeyValuePair) * 10);
+    pairArrayOut[0].key = NvSciBufRawBufferAttrKey_Size;
+
+    checkNvSciErrors(
+        NvSciBufAttrListGetAttrs(m_buffAttrListOut[0], pairArrayOut, 1));
+
+    uint64_t size = *(uint64_t *)pairArrayOut[0].value;
+
+    cudaExternalMemoryHandleDesc memHandleDesc;
+    memset(&memHandleDesc, 0, sizeof(memHandleDesc));
+    memHandleDesc.type = cudaExternalMemoryHandleTypeNvSciBuf;
+    memHandleDesc.handle.nvSciBufObject = inputBufObj;
+    memHandleDesc.size = size;
+    checkCudaErrors(cudaImportExternalMemory(&extMemRawBuf, &memHandleDesc));
+
+    cudaExternalMemoryBufferDesc bufferDesc;
+    memset(&bufferDesc, 0, sizeof(bufferDesc));
+    bufferDesc.offset = 0;
+    bufferDesc.size = size;
+    m_bufSize = size;
+    checkCudaErrors(cudaExternalMemoryGetMappedBuffer(
+        &d_outputBuf, extMemRawBuf, &bufferDesc));
+  }
+
+  void cudaImportNvSciImage(NvSciBufObj inputBufObj) {
+    checkCudaErrors(cudaSetDevice(m_cudaDeviceId));
+    checkNvSciErrors(
+        NvSciBufObjGetAttrList(inputBufObj, &m_buffAttrListOut[1]));
+
+    memset(pairArrayOut, 0, sizeof(NvSciBufAttrKeyValuePair) * 10);
+    pairArrayOut[0].key = NvSciBufImageAttrKey_Size;
+    pairArrayOut[1].key = NvSciBufImageAttrKey_Alignment;
+    pairArrayOut[2].key = NvSciBufImageAttrKey_PlaneCount;
+    pairArrayOut[3].key = NvSciBufImageAttrKey_PlaneWidth;
+    pairArrayOut[4].key = NvSciBufImageAttrKey_PlaneHeight;
+
+    checkNvSciErrors(
+        NvSciBufAttrListGetAttrs(m_buffAttrListOut[1], pairArrayOut, 5));
+
+    uint64_t size = *(uint64_t *)pairArrayOut[0].value;
+    uint64_t alignment = *(uint64_t *)pairArrayOut[1].value;
+    uint64_t planeCount = *(uint64_t *)pairArrayOut[2].value;
+    uint64_t imageWidth = *(uint64_t *)pairArrayOut[3].value;
+    uint64_t imageHeight = *(uint64_t *)pairArrayOut[4].value;
+
+    cudaExternalMemoryHandleDesc memHandleDesc;
+    memset(&memHandleDesc, 0, sizeof(memHandleDesc));
+    memHandleDesc.type = cudaExternalMemoryHandleTypeNvSciBuf;
+    memHandleDesc.handle.nvSciBufObject = inputBufObj;
+    memHandleDesc.size = size;
+    checkCudaErrors(cudaImportExternalMemory(&extMemImageBuf, &memHandleDesc));
+
+    cudaExtent extent = {};
+    memset(&extent, 0, sizeof(extent));
+    extent.width = imageWidth;
+    extent.height = imageHeight;
+    extent.depth = 0;
+
+    cudaChannelFormatDesc desc;
+    desc.x = 8;
+    desc.y = 8;
+    desc.z = 8;
+    desc.w = 8;
+    desc.f = cudaChannelFormatKindUnsigned;
+
+    cudaExternalMemoryMipmappedArrayDesc mipmapDesc = {0};
+    mipmapDesc.offset = 0;
+    mipmapDesc.formatDesc = desc;
+    mipmapDesc.extent = extent;
+    mipmapDesc.flags = 0;
+
+    mipmapDesc.numLevels = 1;
+    checkCudaErrors(cudaExternalMemoryGetMappedMipmappedArray(
+        &d_mipmapArray, extMemImageBuf, &mipmapDesc));
+  }
+
+  void copyDataToImageArray(unsigned char *imageData) {
+    uint32_t mipLevelId = 0;
+    checkCudaErrors(cudaGetMipmappedArrayLevel(&d_mipLevelArray, d_mipmapArray,
+                                               mipLevelId));
+
+    checkCudaErrors(cudaMemcpy2DToArrayAsync(
+        d_mipLevelArray, 0, 0, imageData, m_imageWidth * sizeof(unsigned int),
+        m_imageWidth * sizeof(unsigned int), m_imageHeight,
+        cudaMemcpyHostToDevice, streamToRun));
+  }
+
+  void createTexture() {
+    cudaResourceDesc texRes;
+    memset(&texRes, 0, sizeof(cudaResourceDesc));
+
+    texRes.resType = cudaResourceTypeArray;
+    texRes.res.array.array = d_mipLevelArray;
+
+    cudaTextureDesc texDescr;
+    memset(&texDescr, 0, sizeof(cudaTextureDesc));
+
+    texDescr.normalizedCoords = true;
+    texDescr.filterMode = cudaFilterModeLinear;
+    texDescr.addressMode[0] = cudaAddressModeWrap;
+    texDescr.addressMode[1] = cudaAddressModeWrap;
+    texDescr.readMode = cudaReadModeNormalizedFloat;
+
+    checkCudaErrors(
+        cudaCreateTextureObject(&texObject, &texRes, &texDescr, NULL));
+  }
+};
+
+class cudaNvSciWait {
+ private:
+  NvSciSyncModule m_syncModule;
+  NvSciBufModule m_bufModule;
+
+  NvSciSyncAttrList m_syncAttrList;
+  NvSciBufAttrList m_rawBufAttrList;
+  NvSciBufAttrList m_buffAttrListOut;
+  NvSciBufAttrKeyValuePair pairArrayOut[10];
+  NvSciSyncFence *m_fence;
+
+  cudaExternalMemory_t extMemRawBuf;
+  cudaExternalSemaphore_t waitSem;
+  cudaStream_t streamToRun;
+  int m_cudaDeviceId;
+  CUuuid m_devUUID;
+  void *d_outputBuf;
+  size_t m_bufSize;
+  size_t imageWidth;
+  size_t imageHeight;
+
+ public:
+  cudaNvSciWait(NvSciBufModule bufModule, NvSciSyncModule syncModule,
+                int cudaDeviceId, int bufSize, NvSciSyncFence *fence)
+      : m_bufModule(bufModule),
+        m_syncModule(syncModule),
+        m_cudaDeviceId(cudaDeviceId),
+        m_bufSize(bufSize),
+        m_fence(fence) {
+    initCuda();
+    checkNvSciErrors(NvSciSyncAttrListCreate(m_syncModule, &m_syncAttrList));
+    checkNvSciErrors(NvSciBufAttrListCreate(m_bufModule, &m_rawBufAttrList));
+
+    setRawBufAttrList(m_bufSize);
+    checkCudaErrors(cudaDeviceGetNvSciSyncAttributes(
+        m_syncAttrList, m_cudaDeviceId, cudaNvSciSyncAttrWait));
+  }
+
+  ~cudaNvSciWait() {
+    checkCudaErrors(cudaStreamDestroy(streamToRun));
+    checkCudaErrors(cudaDestroyExternalSemaphore(waitSem));
+    checkCudaErrors(cudaDestroyExternalMemory(extMemRawBuf));
+    checkCudaErrors(cudaFree(d_outputBuf));
+  }
+
+  void initCuda() {
+    checkCudaErrors(cudaSetDevice(m_cudaDeviceId));
+    checkCudaErrors(
+        cudaStreamCreateWithFlags(&streamToRun, cudaStreamNonBlocking));
+#ifdef cuDeviceGetUuid_v2
+    CUresult res = cuDeviceGetUuid_v2(&m_devUUID, m_cudaDeviceId);
+#else
+    CUresult res = cuDeviceGetUuid(&m_devUUID, m_cudaDeviceId);
+#endif
+    if (res != CUDA_SUCCESS) {
+      fprintf(stderr, "Driver API error = %04d \n", res);
+      exit(EXIT_FAILURE);
+    }
+
+    int major = 0, minor = 0;
+    checkCudaErrors(cudaDeviceGetAttribute(
+        &major, cudaDevAttrComputeCapabilityMajor, m_cudaDeviceId));
+    checkCudaErrors(cudaDeviceGetAttribute(
+        &minor, cudaDevAttrComputeCapabilityMinor, m_cudaDeviceId));
+    printf(
+        "[cudaNvSciWait] GPU Device %d: \"%s\" with compute capability "
+        "%d.%d\n\n",
+        m_cudaDeviceId, _ConvertSMVer2ArchName(major, minor), major, minor);
+  }
+
+  void setRawBufAttrList(uint64_t size) {
+    NvSciBufType bufType = NvSciBufType_RawBuffer;
+    bool cpuAccess = false;
+    NvSciBufAttrValAccessPerm perm = NvSciBufAccessPerm_ReadWrite;
+    NvSciBufAttrKeyValuePair rawBufAttrs[] = {
+        {NvSciBufRawBufferAttrKey_Size, &size, sizeof(size)},
+        {NvSciBufGeneralAttrKey_Types, &bufType, sizeof(bufType)},
+        {NvSciBufGeneralAttrKey_NeedCpuAccess, &cpuAccess, sizeof(cpuAccess)},
+        {NvSciBufGeneralAttrKey_RequiredPerm, &perm, sizeof(perm)},
+        {NvSciBufGeneralAttrKey_GpuId, &m_devUUID, sizeof(m_devUUID)},
+    };
+
+    checkNvSciErrors(NvSciBufAttrListSetAttrs(
+        m_rawBufAttrList, rawBufAttrs,
+        sizeof(rawBufAttrs) / sizeof(NvSciBufAttrKeyValuePair)));
+  }
+
+  NvSciSyncAttrList getNvSciSyncAttrList() { return m_syncAttrList; }
+
+  NvSciBufAttrList getNvSciRawBufAttrList() { return m_rawBufAttrList; }
+
+  void runImageGrayscale(std::string image_filename, size_t imageWidth,
+                         size_t imageHeight) {
+    int numOfGPUs = 0;
+    checkCudaErrors(cudaGetDeviceCount(&numOfGPUs));  // For cuda init purpose
+    checkCudaErrors(cudaSetDevice(m_cudaDeviceId));
+
+    waitExternalSemaphore();
+    launchGrayScaleKernel((unsigned int *)d_outputBuf, image_filename,
+                          imageWidth, imageHeight, streamToRun);
+  }
+
+  void cudaImportNvSciSemaphore(NvSciSyncObj syncObj) {
+    checkCudaErrors(cudaSetDevice(m_cudaDeviceId));
+
+    cudaExternalSemaphoreHandleDesc extSemDesc;
+    memset(&extSemDesc, 0, sizeof(extSemDesc));
+    extSemDesc.type = cudaExternalSemaphoreHandleTypeNvSciSync;
+    extSemDesc.handle.nvSciSyncObj = (void *)syncObj;
+
+    checkCudaErrors(cudaImportExternalSemaphore(&waitSem, &extSemDesc));
+  }
+
+  void waitExternalSemaphore() {
+    checkCudaErrors(cudaSetDevice(m_cudaDeviceId));
+
+    cudaExternalSemaphoreWaitParams waitParams;
+    memset(&waitParams, 0, sizeof(waitParams));
+    // For cross-process signaler-waiter applications need to use NvSciIpc
+    // and NvSciSync[Export|Import] utilities to share the NvSciSyncFence
+    // across process. This step is optional in single-process.
+    waitParams.params.nvSciSync.fence = (void *)m_fence;
+    waitParams.flags = 0;
+
+    checkCudaErrors(
+        cudaWaitExternalSemaphoresAsync(&waitSem, &waitParams, 1, streamToRun));
+  }
+
+  void cudaImportNvSciRawBuf(NvSciBufObj inputBufObj) {
+    checkCudaErrors(cudaSetDevice(m_cudaDeviceId));
+
+    checkNvSciErrors(NvSciBufObjGetAttrList(inputBufObj, &m_buffAttrListOut));
+
+    memset(pairArrayOut, 0, sizeof(NvSciBufAttrKeyValuePair) * 10);
+    pairArrayOut[0].key = NvSciBufRawBufferAttrKey_Size;
+
+    checkNvSciErrors(
+        NvSciBufAttrListGetAttrs(m_buffAttrListOut, pairArrayOut, 1));
+
+    uint64_t size = *(uint64_t *)pairArrayOut[0].value;
+
+    cudaExternalMemoryHandleDesc memHandleDesc;
+    memset(&memHandleDesc, 0, sizeof(memHandleDesc));
+    memHandleDesc.type = cudaExternalMemoryHandleTypeNvSciBuf;
+    memHandleDesc.handle.nvSciBufObject = inputBufObj;
+    memHandleDesc.size = size;
+    checkCudaErrors(cudaImportExternalMemory(&extMemRawBuf, &memHandleDesc));
+
+    cudaExternalMemoryBufferDesc bufferDesc;
+    memset(&bufferDesc, 0, sizeof(bufferDesc));
+    bufferDesc.offset = 0;
+    bufferDesc.size = size;
+    m_bufSize = size;
+
+    checkCudaErrors(cudaExternalMemoryGetMappedBuffer(
+        &d_outputBuf, extMemRawBuf, &bufferDesc));
+  }
+};
+
+void thread_rotateAndSignal(cudaNvSciSignal *cudaNvSciSignalObj,
+                            unsigned char *imageData) {
+  std::lock_guard<std::mutex> guard(m_mutex);
+  cudaNvSciSignalObj->runRotateImageAndSignal(imageData);
+  workSubmitted = true;
+  m_condVar.notify_one();
+}
+
+void thread_waitAndGrayscale(cudaNvSciWait *cudaNvSciWaitObj,
+                             std::string image_filename, size_t imageWidth,
+                             size_t imageHeight) {
+  // Acquire the lock
+  std::unique_lock<std::mutex> mlock(m_mutex);
+  m_condVar.wait(mlock, [] { return workSubmitted; });
+  cudaNvSciWaitObj->runImageGrayscale(image_filename, imageWidth, imageHeight);
+}
+
+cudaNvSci::cudaNvSci(int isMultiGPU, std::vector<int> &deviceIds,
+                     unsigned char *imageData, size_t width, size_t height)
+    : m_isMultiGPU(isMultiGPU),
+      image_data(imageData),
+      imageWidth(width),
+      imageHeight(height) {
+  if (isMultiGPU) {
+    m_cudaNvSciSignalDeviceId = deviceIds[0];
+    m_cudaNvSciWaitDeviceId = deviceIds[1];
+  } else {
+    m_cudaNvSciSignalDeviceId = m_cudaNvSciWaitDeviceId = deviceIds[0];
+  }
+
+  m_bufSize = imageWidth * imageHeight * sizeof(unsigned int);
+}
+
+void cudaNvSci::initNvSci() {
+  checkNvSciErrors(NvSciSyncModuleOpen(&syncModule));
+  checkNvSciErrors(NvSciBufModuleOpen(&buffModule));
+  fence = (NvSciSyncFence *)calloc(1, sizeof(NvSciSyncFence));
+}
+
+void cudaNvSci::runCudaNvSci(std::string &image_filename) {
+  initNvSci();
+
+  cudaNvSciSignal rotateAndSignal(buffModule, syncModule,
+                                  m_cudaNvSciSignalDeviceId, m_bufSize,
+                                  imageWidth, imageHeight, fence);
+  cudaNvSciWait waitAndGrayscale(buffModule, syncModule,
+                                 m_cudaNvSciWaitDeviceId, m_bufSize, fence);
+
+  rawBufUnreconciledList[0] = rotateAndSignal.getNvSciRawBufAttrList();
+  rawBufUnreconciledList[1] = waitAndGrayscale.getNvSciRawBufAttrList();
+
+  createNvSciRawBufObj();
+
+  imageBufUnreconciledList[0] = rotateAndSignal.getNvSciImageBufAttrList();
+
+  createNvSciBufImageObj();
+
+  rotateAndSignal.cudaImportNvSciRawBuf(rawBufObj);
+  rotateAndSignal.cudaImportNvSciImage(imageBufObj);
+
+  waitAndGrayscale.cudaImportNvSciRawBuf(rawBufObj);
+
+  syncUnreconciledList[0] = rotateAndSignal.getNvSciSyncAttrList();
+  syncUnreconciledList[1] = waitAndGrayscale.getNvSciSyncAttrList();
+
+  createNvSciSyncObj();
+
+  rotateAndSignal.cudaImportNvSciSemaphore(syncObj);
+  waitAndGrayscale.cudaImportNvSciSemaphore(syncObj);
+
+  std::thread rotateThread(&thread_rotateAndSignal, &rotateAndSignal,
+                           image_data);
+
+  std::thread grayscaleThread(&thread_waitAndGrayscale, &waitAndGrayscale,
+                              image_filename, imageWidth, imageHeight);
+
+  rotateThread.join();
+  grayscaleThread.join();
+}
+
+void cudaNvSci::createNvSciRawBufObj() {
+  int numAttrList = 2;
+  checkNvSciErrors(NvSciBufAttrListReconcile(rawBufUnreconciledList,
+                                             numAttrList, &rawBufReconciledList,
+                                             &buffConflictList));
+  checkNvSciErrors(NvSciBufObjAlloc(rawBufReconciledList, &rawBufObj));
+  printf("created NvSciBufObj\n");
+}
+
+void cudaNvSci::createNvSciBufImageObj() {
+  int numAttrList = 1;
+  checkNvSciErrors(NvSciBufAttrListReconcile(
+      imageBufUnreconciledList, numAttrList, &imageBufReconciledList,
+      &imageBufConflictList));
+  checkNvSciErrors(NvSciBufObjAlloc(imageBufReconciledList, &imageBufObj));
+  printf("created NvSciBufImageObj\n");
+}
+
+void cudaNvSci::createNvSciSyncObj() {
+  int numAttrList = 2;
+  checkNvSciErrors(NvSciSyncAttrListReconcile(syncUnreconciledList, numAttrList,
+                                              &syncReconciledList,
+                                              &syncConflictList));
+  checkNvSciErrors(NvSciSyncObjAlloc(syncReconciledList, &syncObj));
+  printf("created NvSciSyncObj\n");
+}

Samples/4_CUDA_Libraries/cudaNvSci/cudaNvSci.h

@@ -0,0 +1,101 @@
+/* Copyright (c) 2022, NVIDIA CORPORATION. All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions
+ * are met:
+ *  * Redistributions of source code must retain the above copyright
+ *    notice, this list of conditions and the following disclaimer.
+ *  * Redistributions in binary form must reproduce the above copyright
+ *    notice, this list of conditions and the following disclaimer in the
+ *    documentation and/or other materials provided with the distribution.
+ *  * Neither the name of NVIDIA CORPORATION nor the names of its
+ *    contributors may be used to endorse or promote products derived
+ *    from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
+ * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ * PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
+ * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+
+#ifndef CUDANVSCI_H
+#define CUDANVSCI_H
+
+#include <cuda_runtime.h>
+#include <helper_cuda.h>
+#include <nvscibuf.h>
+#include <nvscisync.h>
+#include <vector>
+
+#define checkNvSciErrors(call)                              \
+  do {                                                      \
+    NvSciError _status = call;                              \
+    if (NvSciError_Success != _status) {                    \
+      printf(                                               \
+          "NVSCI call in file '%s' in line %i returned"     \
+          " %d, expected %d\n",                             \
+          __FILE__, __LINE__, _status, NvSciError_Success); \
+      fflush(stdout);                                       \
+      exit(EXIT_FAILURE);                                   \
+    }                                                       \
+  } while (0)
+
+extern void rotateKernel(cudaTextureObject_t &texObj, const float angle,
+                         unsigned int *d_outputData, const int imageWidth,
+                         const int imageHeight, cudaStream_t stream);
+extern void launchGrayScaleKernel(unsigned int *d_rgbaImage,
+                                  std::string image_filename, size_t imageWidth,
+                                  size_t imageHeight, cudaStream_t stream);
+
+class cudaNvSci {
+ private:
+  int m_isMultiGPU;
+  int m_cudaNvSciSignalDeviceId;
+  int m_cudaNvSciWaitDeviceId;
+  unsigned char *image_data;
+  size_t m_bufSize;
+  size_t imageWidth;
+  size_t imageHeight;
+
+ public:
+  NvSciSyncModule syncModule;
+  NvSciBufModule buffModule;
+  NvSciSyncAttrList syncUnreconciledList[2];
+  NvSciSyncAttrList syncReconciledList;
+  NvSciSyncAttrList syncConflictList;
+
+  NvSciBufAttrList rawBufUnreconciledList[2];
+  NvSciBufAttrList imageBufUnreconciledList[2];
+  NvSciBufAttrList rawBufReconciledList;
+  NvSciBufAttrList buffConflictList;
+  NvSciBufAttrList imageBufReconciledList;
+  NvSciBufAttrList imageBufConflictList;
+  NvSciBufAttrList buffAttrListOut;
+
+  NvSciSyncObj syncObj;
+  NvSciBufObj rawBufObj;
+  NvSciBufObj imageBufObj;
+  NvSciSyncFence *fence;
+
+  cudaNvSci(int isMultiGPU, std::vector<int> &deviceIds,
+            unsigned char *image_data, size_t imageWidth, size_t imageHeight);
+
+  void initNvSci();
+
+  void runCudaNvSci(std::string &image_filename);
+
+  void createNvSciRawBufObj();
+
+  void createNvSciSyncObj();
+
+  void createNvSciBufImageObj();
+};
+
+#endif  // CUDANVSCI_H

Samples/4_CUDA_Libraries/cudaNvSci/imageKernels.cu

@@ -0,0 +1,120 @@
+/* Copyright (c) 2022, NVIDIA CORPORATION. All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions
+ * are met:
+ *  * Redistributions of source code must retain the above copyright
+ *    notice, this list of conditions and the following disclaimer.
+ *  * Redistributions in binary form must reproduce the above copyright
+ *    notice, this list of conditions and the following disclaimer in the
+ *    documentation and/or other materials provided with the distribution.
+ *  * Neither the name of NVIDIA CORPORATION nor the names of its
+ *    contributors may be used to endorse or promote products derived
+ *    from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
+ * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ * PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
+ * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+
+#include <cuda.h>
+#include <helper_cuda.h>
+#include <helper_image.h>
+
+// convert floating point rgba color to 32-bit integer
+__device__ unsigned int rgbaFloatToInt(float4 rgba) {
+  rgba.x = __saturatef(rgba.x);  // clamp to [0.0, 1.0]
+  rgba.y = __saturatef(rgba.y);
+  rgba.z = __saturatef(rgba.z);
+  rgba.w = __saturatef(rgba.w);
+  return ((unsigned int)(rgba.w * 255.0f) << 24) |
+         ((unsigned int)(rgba.z * 255.0f) << 16) |
+         ((unsigned int)(rgba.y * 255.0f) << 8) |
+         ((unsigned int)(rgba.x * 255.0f));
+}
+
+////////////////////////////////////////////////////////////////////////////////
+//! Rotate an image using texture lookups
+//! @param outputData  output data in global memory
+////////////////////////////////////////////////////////////////////////////////
+static __global__ void transformKernel(unsigned int *outputData, int width,
+                                       int height, float theta,
+                                       cudaTextureObject_t tex) {
+  // calculate normalized texture coordinates
+  unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
+  unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
+
+  float u = (float)x - (float)width / 2;
+  float v = (float)y - (float)height / 2;
+  float tu = u * cosf(theta) - v * sinf(theta);
+  float tv = v * cosf(theta) + u * sinf(theta);
+
+  tu /= (float)width;
+  tv /= (float)height;
+
+  // read from texture and write to global memory
+  float4 pix = tex2D<float4>(tex, tu + 0.5f, tv + 0.5f);
+  unsigned int pixelInt = rgbaFloatToInt(pix);
+  outputData[y * width + x] = pixelInt;
+}
+
+static __global__ void rgbToGrayscaleKernel(unsigned int *rgbaImage,
+                                            size_t imageWidth,
+                                            size_t imageHeight) {
+  size_t gidX = blockDim.x * blockIdx.x + threadIdx.x;
+
+  uchar4 *pixArray = (uchar4 *)rgbaImage;
+
+  for (int pixId = gidX; pixId < imageWidth * imageHeight;
+       pixId += gridDim.x * blockDim.x) {
+    uchar4 dataA = pixArray[pixId];
+    unsigned char grayscale =
+        (unsigned char)(dataA.x * 0.3 + dataA.y * 0.59 + dataA.z * 0.11);
+    uchar4 dataB = make_uchar4(grayscale, grayscale, grayscale, 0);
+    pixArray[pixId] = dataB;
+  }
+}
+
+void launchGrayScaleKernel(unsigned int *d_rgbaImage,
+                           std::string image_filename, size_t imageWidth,
+                           size_t imageHeight, cudaStream_t stream) {
+  int numThreadsPerBlock = 1024;
+  int numOfBlocks = (imageWidth * imageHeight) / numThreadsPerBlock;
+
+  rgbToGrayscaleKernel<<<numOfBlocks, numThreadsPerBlock, 0, stream>>>(
+      d_rgbaImage, imageWidth, imageHeight);
+
+  unsigned int *outputData;
+  checkCudaErrors(cudaMallocHost((void**)&outputData, sizeof(unsigned int) * imageWidth * imageHeight));
+  checkCudaErrors(cudaMemcpyAsync(
+      outputData, d_rgbaImage, sizeof(unsigned int) * imageWidth * imageHeight,
+      cudaMemcpyDeviceToHost, stream));
+  checkCudaErrors(cudaStreamSynchronize(stream));
+
+  char outputFilename[1024];
+  strcpy(outputFilename, image_filename.c_str());
+  strcpy(outputFilename + image_filename.length() - 4, "_out.ppm");
+  sdkSavePPM4ub(outputFilename, (unsigned char *)outputData, imageWidth,
+                imageHeight);
+  printf("Wrote '%s'\n", outputFilename);
+
+  checkCudaErrors(cudaFreeHost(outputData));
+}
+
+void rotateKernel(cudaTextureObject_t &texObj, const float angle,
+                  unsigned int *d_outputData, const int imageWidth,
+                  const int imageHeight, cudaStream_t stream) {
+  dim3 dimBlock(8, 8, 1);
+  dim3 dimGrid(imageWidth / dimBlock.x, imageHeight / dimBlock.y, 1);
+
+  transformKernel<<<dimGrid, dimBlock, 0, stream>>>(d_outputData, imageWidth,
+                                                    imageHeight, angle, texObj);
+}

Samples/4_CUDA_Libraries/cudaNvSci/main.cpp

@@ -0,0 +1,101 @@
+/* Copyright (c) 2022, NVIDIA CORPORATION. All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions
+ * are met:
+ *  * Redistributions of source code must retain the above copyright
+ *    notice, this list of conditions and the following disclaimer.
+ *  * Redistributions in binary form must reproduce the above copyright
+ *    notice, this list of conditions and the following disclaimer in the
+ *    documentation and/or other materials provided with the distribution.
+ *  * Neither the name of NVIDIA CORPORATION nor the names of its
+ *    contributors may be used to endorse or promote products derived
+ *    from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
+ * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ * PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
+ * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+
+#include <cuda.h>
+#include <vector>
+#include "cudaNvSci.h"
+#include <helper_cuda.h>
+#include <helper_image.h>
+
+void loadImageData(const std::string &filename, const char **argv,
+                   unsigned char **image_data, uint32_t &imageWidth,
+                   uint32_t &imageHeight) {
+  // load image (needed so we can get the width and height before we create
+  // the window
+  char *image_path = sdkFindFilePath(filename.c_str(), argv[0]);
+
+  if (image_path == 0) {
+    printf("Error finding image file '%s'\n", filename.c_str());
+    exit(EXIT_FAILURE);
+  }
+
+  sdkLoadPPM4(image_path, image_data, &imageWidth, &imageHeight);
+
+  if (!image_data) {
+    printf("Error opening file '%s'\n", image_path);
+    exit(EXIT_FAILURE);
+  }
+
+  printf("Loaded '%s', %d x %d pixels\n", image_path, imageWidth, imageHeight);
+}
+
+int main(int argc, const char **argv) {
+  int numOfGPUs = 0;
+  std::vector<int> deviceIds;
+  checkCudaErrors(cudaGetDeviceCount(&numOfGPUs));
+
+  printf("%d GPUs found\n", numOfGPUs);
+  if (!numOfGPUs) {
+    exit(EXIT_WAIVED);
+  } else {
+    for (int devID = 0; devID < numOfGPUs; devID++) {
+      int major = 0, minor = 0;
+      checkCudaErrors(cudaDeviceGetAttribute(
+          &major, cudaDevAttrComputeCapabilityMajor, devID));
+      checkCudaErrors(cudaDeviceGetAttribute(
+          &minor, cudaDevAttrComputeCapabilityMinor, devID));
+      if (major >= 6) {
+        deviceIds.push_back(devID);
+      }
+    }
+    if (deviceIds.size() == 0) {
+      printf(
+          "cudaNvSci requires one or more GPUs of Pascal(SM 6.0) or higher "
+          "archs\nWaiving..\n");
+      exit(EXIT_WAIVED);
+    }
+  }
+
+  std::string image_filename = "teapot1024.ppm";
+
+  if (checkCmdLineFlag(argc, (const char **)argv, "file")) {
+    getCmdLineArgumentString(argc, (const char **)argv, "file",
+                             (char **)&image_filename);
+  }
+
+  uint32_t imageWidth = 0;
+  uint32_t imageHeight = 0;
+  unsigned char *image_data = NULL;
+
+  loadImageData(image_filename, argv, &image_data, imageWidth, imageHeight);
+
+  cudaNvSci cudaNvSciApp(deviceIds.size() > 1, deviceIds, image_data,
+                         imageWidth, imageHeight);
+  cudaNvSciApp.runCudaNvSci(image_filename);
+
+  return EXIT_SUCCESS;
+}