CUDA Interoperability

Complex image/data processing pipelines consist of multiple computation stages. Some pipeline stages may be mapped to the GPU along with the PVA processing stages to fully utilize device resources. The cuPVA runtime provides CUDA Interoperability APIs to conveniently realize such applications.

CUDA Interoperability APIs provide functionality to:

• Use CUDA device pointers directly in PVA programs to pass data between PVA and CUDA processing stages.

• Submit PVA commands directly to CUDA stream handles so PVA programs and CUDA kernels can be scheduled consecutively on the same stream.

In this tutorial, we extend the contrast stretching example used in the previous tutorials to demonstrate CUDA Interoperability APIs. The first stage, which computes an image histogram to determine the input pixel dynamic range, is mapped to the PVA. The subsequent contrast stretching stage is performed on the GPU using CUDA kernels.

Note

The current PVA context used by the CUDA interoperability APIs in this tutorial must have been created after a CUDA context was loaded. Any CUDA Runtime API call causes a CUDA context to be loaded, as does a CUDA Driver API call to cuCtxCreate(). The CUDA context must also be current on the host thread that calls the interoperability API. If the API is called from a worker thread, including a task launched with std::async or a thread-pool task, first call a CUDA Runtime or Driver API on that thread, for example cudaFree(nullptr) in C++ or cudaFree(NULL) in C, to make the CUDA context current. The CUDA allocation and stream-creation calls below establish this ordering on the main thread.

Host Code

C++

1. The main function starts by initializing the image buffer and parameter pointers to nullptr.

   int main(int argc, char **argv)
   {
       int err = 0;
       if (GetAssetsDirectory(argc, argv, assetsDirectory, MAX_IMAGE_PATH_LENGTH) != 0)
       {
           return 1;
       }
   
       uint8_t *image_cpu                        = nullptr;
       uint8_t *image_cuda                       = nullptr;
       ContrastStretchParams *stretchParams_cpu  = nullptr;
       ContrastStretchParams *stretchParams_cuda = nullptr;
       cudaStream_t cudaStream                   = nullptr;
   
       try
       {
           // Submitting PVA commands directly to a CUDA stream requires PVA driver version 2007 or later.
           if (cupva::GetDriverVersion() < 2007U)
           {
               printf("Unsupported platform, skipping execution...\n");
               return 0;
           }
   

2. The second stage, i.e., contrast stretching stage, of the algorithm is executed on the GPU. Contrast stretching stage uses the algorithm parameters computed on PVA in the first stage.

   Memory for algorithm parameters structure and image buffer are allocated using the cudaHostAlloc() CUDA API call. cudaHostAlloc() allocates a page-locked memory on the host and maps the allocation into device space when the cudaHostAllocMapped option is selected.

   On supported platforms, cuPVA can access CUDA device pointers directly, so they can be passed to cuPVA APIs without explicit import. The PVA program that executes in the first stage loads the image and stores the computed image dynamic range at those memory locations.

           if (cudaSuccess != cudaHostAlloc(&stretchParams_cpu, sizeof(ContrastStretchParams), cudaHostAllocMapped))
           {
               throw std::runtime_error("cudaHostAlloc failed");
           }
           if (cudaSuccess != cudaHostGetDevicePointer(&stretchParams_cuda, stretchParams_cpu, 0))
           {
               throw std::runtime_error("cudaHostGetDevicePointer failed");
           }
           if (cudaSuccess != cudaHostAlloc(&image_cpu, IMAGE_SIZE * sizeof(uint8_t), cudaHostAllocMapped))
           {
               throw std::runtime_error("cudaHostAlloc failed");
           }
           if (cudaSuccess != cudaHostGetDevicePointer(&image_cuda, image_cpu, 0))
           {
               throw std::runtime_error("cudaHostGetDevicePointer failed");
           }
   

3. A CUDA stream is a sequence of operations that execute in issue-order on the GPU. The cudaStreamCreateWithFlags() CUDA API creates an asynchronous stream. cuPVA CUDA Interoperability does not support the CUDA default stream, so the CUDA stream handle must be created with the cudaStreamNonBlocking flag. After the stream is created, CUDA interoperability automatically makes it available to cuPVA. Applications do not call the deprecated cuda::CreateStream() API or create and manage a separate cupva::Stream object. Direct submission to a CUDA stream requires PVA driver version 2007 or later and is not supported in Native mode.

           if (cudaSuccess != cudaStreamCreateWithFlags(&cudaStream, cudaStreamNonBlocking))
           {
               throw std::runtime_error("cudaStreamCreateWithFlags failed");
           }
   

4. By default, PVA commands submitted directly to a CUDA stream use cupva::VPU_ANY and can be scheduled on any VPU. The cuda::SetAffinity() API sets the VPU affinity for all PVA submissions on that CUDA stream. This tutorial selects cupva::VPU0 before submitting any PVA commands.

           cuda::SetAffinity(cudaStream, VPU0);
   

5. The PVA program that computes the image dynamic range is created. The steps involving the CmdProgram initialization are very similar to the previous contrast stretching based tutorials. CUDA allocated parameters and image pointers are used as the input and output for the program, respectively.

           if (ReadImageBuffer(inputImageName.c_str(), assetsDirectory, image_cpu, IMAGE_SIZE) != 0)
           {
               cudaFreeHost(image_cpu);
               cudaFreeHost(stretchParams_cpu);
               err = -1;
               return err;
           }
   
           ContrastStretchParams algParams = {
               .inputLowPixelValue           = 0,
               .outputLowPixelValue          = 0,
               .inputHighPixelValue          = 0,
               .outputHighPixelValue         = 255,
               .saturationHistogramCountLow  = IMAGE_SIZE * SATURATED_PIXEL_PERCENTAGE_LOW_INTENSITY / 100,
               .saturationHistogramCountHigh = IMAGE_SIZE * SATURATED_PIXEL_PERCENTAGE_HIGH_INTENSITY / 100};
   
           memcpy(stretchParams_cpu, &algParams, sizeof(ContrastStretchParams));
   
           Executable execComputeDynamicRange =
               Executable::Create(PVA_EXECUTABLE_DATA(cuda_interoperability_compute_dynamic_range_dev),
                                  PVA_EXECUTABLE_SIZE(cuda_interoperability_compute_dynamic_range_dev));
   
           CmdProgram progComputeDynamicRange =
               CreateComputeDynamicRangeProg(execComputeDynamicRange, image_cuda, IMAGE_WIDTH, IMAGE_HEIGHT, TILE_WIDTH,
                                             TILE_HEIGHT, stretchParams_cuda, stretchParams_cpu);
   

6. The cuda::Submit() API submits progComputeDynamicRange directly to cudaStream. No explicit stream import is required. PVA and CUDA commands enqueued on the same CUDA stream are ordered together according to submission order. PVA commands follow the cupva::OrderType value passed to cuda::Submit(), while CUDA commands execute according to the configuration of cudaStream. This example does not submit a cupva::CmdRequestFences command or call cupva::Fence::wait() after submission. Synchronization is achieved with CUDA APIs, as shown in the following steps.

           cuda::Submit(cudaStream, {&progComputeDynamicRange});
   

7. The StretchContrast() call is submitted to the same CUDA stream. Therefore, it is executed after the progComputeDynamicRange PVA program completes.

           StretchContrast(image_cuda, stretchParams_cuda, IMAGE_SIZE, cudaStream);
   

   The stretch_kernel CUDA function performs image contrast stretching on the GPU. stretch_kernel is submitted to cudaStream and runs asynchronously.

   __global__ void stretch_kernel(uint8_t *image, ContrastStretchParams *stretchParams_cuda, int32_t imageSize)
   {
       int32_t inputDynamicRange  = stretchParams_cuda->inputHighPixelValue - stretchParams_cuda->inputLowPixelValue + 1;
       int32_t outputDynamicRange = 256;
   
       int index  = blockIdx.x * blockDim.x + threadIdx.x;
       int stride = blockDim.x * gridDim.x;
       int32_t outputPixelValue;
   
       for (int i = index; i < imageSize; i += stride)
       {
           outputPixelValue = outputDynamicRange * (image[i] - stretchParams_cuda->inputLowPixelValue);
           outputPixelValue /= inputDynamicRange;
           outputPixelValue = max(0, min(outputPixelValue, 255));
           image[i]         = (uint8_t)outputPixelValue;
       }
   }
   
   void StretchContrast(uint8_t *image, ContrastStretchParams *stretchParams_cuda, int32_t imageSize,
                        cudaStream_t cudaStream)
   {
       int blockSize  = BLOCK_SIZE;
       int blockCount = BLOCK_COUNT;
       stretch_kernel<<<blockCount, blockSize, 0, cudaStream>>>(image, stretchParams_cuda, imageSize);
   }
   

8. The cudaStreamSynchronize() call blocks the host until all PVA and CUDA work enqueued on cudaStream completes.

           if (cudaSuccess != cudaStreamSynchronize(cudaStream))
           {
               throw std::runtime_error("cudaStreamSynchronize failed");
           }
   

9. The enhanced image is written to the output file and allocated resources are freed in the last step.

           if (WriteImageBuffer(outputImageName.c_str(), ".", image_cpu, IMAGE_SIZE) != 0)
           {
               err = -1;
           }
       }
       catch (cupva::Exception const &e)
       {
           std::cout << "Caught a cuPVA exception with message: " << e.what() << std::endl;
           err = 1;
       }
       catch (const std::runtime_error &e)
       {
           std::cout << "Caught a CUDA exception with message: " << e.what() << std::endl;
       }
   
       cudaFreeHost(image_cpu);
       cudaFreeHost(stretchParams_cpu);
       cudaStreamDestroy(cudaStream);
       return err;
   }
   

C

1. The main function starts by initializing the image buffer and parameter pointers to NULL.

   int main(int argc, char **argv)
   {
       int32_t err = 0;
   
       /* Submitting PVA commands directly to a CUDA stream requires PVA driver version 2007 or later. */
       uint32_t driverVersion;
       if (CupvaGetDriverVersion(&driverVersion) != CUPVA_ERROR_NONE)
       {
           return 1;
       }
       if (driverVersion < 2007U)
       {
           printf("Unsupported platform, skipping execution...\n");
           return 0;
       }
   
       if (GetAssetsDirectory(argc, argv, assetsDirectory, MAX_IMAGE_PATH_LENGTH) != 0)
       {
           return 1;
       }
   
       uint8_t *image_cpu                        = NULL;
       uint8_t *image_cuda                       = NULL;
       ContrastStretchParams *stretchParams_cpu  = NULL;
       ContrastStretchParams *stretchParams_cuda = NULL;
       cudaStream_t cudaStream                   = NULL;
   
   

2. The second stage, i.e., contrast stretching stage, of the algorithm is executed on the GPU. Contrast stretching stage uses the algorithm parameters computed on PVA in the first stage.

   Memory for algorithm parameters structure and image buffer are allocated using the cudaHostAlloc() CUDA API call. cudaHostAlloc() allocates a page-locked memory on the host and maps the allocation into device space when the cudaHostAllocMapped option is selected.

   On supported platforms, cuPVA can access CUDA device pointers directly, so they can be passed to cuPVA APIs without explicit import. The PVA program that executes in the first stage loads the image and stores the computed image dynamic range at those memory locations.

       CHECK_CUDA_ERROR_GOTO(
           cudaHostAlloc((void **)&stretchParams_cpu, sizeof(ContrastStretchParams), cudaHostAllocMapped), err,
           MemAllocFailed);
   
       CHECK_CUDA_ERROR_GOTO(cudaHostGetDevicePointer((void **)&stretchParams_cuda, (void *)stretchParams_cpu, 0), err,
                             MemAllocFailed);
   
       CHECK_CUDA_ERROR_GOTO(cudaHostAlloc((void **)&image_cpu, IMAGE_SIZE * sizeof(uint8_t), cudaHostAllocMapped), err,
                             MemAllocFailed);
   
       CHECK_CUDA_ERROR_GOTO(cudaHostGetDevicePointer((void **)&image_cuda, (void *)image_cpu, 0), err, MemAllocFailed);
   
   

3. A CUDA stream is a sequence of operations that execute in issue-order on the GPU. The cudaStreamCreateWithFlags() CUDA API creates an asynchronous stream. cuPVA CUDA Interoperability does not support the CUDA default stream, so the CUDA stream handle must be created with the cudaStreamNonBlocking flag. After the stream is created, CUDA interoperability automatically makes it available to cuPVA. Applications do not call the deprecated CupvaCudaCreateStream() API or create and manage a separate cuPVA Stream object. Direct submission to a CUDA stream requires PVA driver version 2007 or later and is not supported in Native mode.

       CHECK_CUDA_ERROR_GOTO(cudaStreamCreateWithFlags(&cudaStream, cudaStreamNonBlocking), err, MemAllocFailed);
   

4. By default, PVA commands submitted directly to a CUDA stream use CUPVA_VPU_ANY and can be scheduled on any VPU. The CupvaCudaStreamSetAffinity() API sets the VPU affinity for all PVA submissions on that CUDA stream. This tutorial selects CUPVA_VPU0 before submitting any PVA commands.

       CHECK_ERROR_GOTO(CupvaCudaStreamSetAffinity(cudaStream, CUPVA_VPU0), err, MemAllocFailed);
   

5. The PVA program that computes the image dynamic range is created. The steps involving the CmdProgram initialization are very similar to the previous contrast stretching based tutorials. CUDA allocated parameters and image pointers are used as the input and output for the program, respectively.

       if (ReadImageBuffer(INPUT_IMAGE_NAME, assetsDirectory, image_cpu, IMAGE_SIZE) != 0)
       {
           err = -1;
           goto MemAllocFailed;
       }
   
       ContrastStretchParams algParams = {
           .inputLowPixelValue           = 0,
           .outputLowPixelValue          = 0,
           .inputHighPixelValue          = 0,
           .outputHighPixelValue         = 255,
           .saturationHistogramCountLow  = IMAGE_SIZE * SATURATED_PIXEL_PERCENTAGE_LOW_INTENSITY / 100,
           .saturationHistogramCountHigh = IMAGE_SIZE * SATURATED_PIXEL_PERCENTAGE_HIGH_INTENSITY / 100};
   
       memcpy(stretchParams_cpu, &algParams, sizeof(ContrastStretchParams));
   
       cupvaExecutable_t execComputeDynamicRange;
       CHECK_ERROR_GOTO(CupvaExecutableCreate(&execComputeDynamicRange,
                                              PVA_EXECUTABLE_DATA(cuda_interoperability_compute_dynamic_range_dev),
                                              PVA_EXECUTABLE_SIZE(cuda_interoperability_compute_dynamic_range_dev)),
                        err, ExecutableCreateFailed);
   
       int32_t createdCmdProgramCount = 0;
       cupvaCmd_t progComputeDynamicRange;
   
       CHECK_ERROR_GOTO(CreateComputeDynamicRangeProg(&progComputeDynamicRange, &execComputeDynamicRange, image_cuda,
                                                      IMAGE_WIDTH, IMAGE_HEIGHT, TILE_WIDTH, TILE_HEIGHT,
                                                      stretchParams_cuda, stretchParams_cpu, &createdCmdProgramCount),
                        err, CmdProgramCreateFailed);
   

6. The CupvaCudaStreamSubmit() API submits progComputeDynamicRange directly to cudaStream. No explicit stream import is required. PVA and CUDA commands enqueued on the same CUDA stream are ordered together according to submission order. PVA commands follow the cupvaOrderType_t value passed to CupvaCudaStreamSubmit(), while CUDA commands execute according to the configuration of cudaStream. This example does not submit a fence-request command or call CupvaFenceWait() after submission. Synchronization is achieved with CUDA APIs, as shown in the following steps.

       cupvaCmd_t const *cmd[1] = {&progComputeDynamicRange};
       CHECK_ERROR_GOTO(CupvaCudaStreamSubmit(cudaStream, cmd, NULL, 1, CUPVA_IN_ORDER, -1, -1), err,
                        DeallocateAllResources);
   

7. The StretchContrast() call is submitted to the same CUDA stream. Therefore, it is executed after the progComputeDynamicRange PVA program completes.

       StretchContrast(image_cuda, stretchParams_cuda, IMAGE_SIZE, cudaStream);
   

   The stretch_kernel CUDA function performs image contrast stretching on the GPU. stretch_kernel is submitted to cudaStream and runs asynchronously.

   __global__ void stretch_kernel(uint8_t *image, ContrastStretchParams *stretchParams_cuda, int32_t imageSize)
   {
       int32_t inputDynamicRange  = stretchParams_cuda->inputHighPixelValue - stretchParams_cuda->inputLowPixelValue + 1;
       int32_t outputDynamicRange = 256;
   
       int index  = blockIdx.x * blockDim.x + threadIdx.x;
       int stride = blockDim.x * gridDim.x;
       int32_t outputPixelValue;
   
       for (int i = index; i < imageSize; i += stride)
       {
           outputPixelValue = outputDynamicRange * (image[i] - stretchParams_cuda->inputLowPixelValue);
           outputPixelValue /= inputDynamicRange;
           outputPixelValue = max(0, min(outputPixelValue, 255));
           image[i]         = (uint8_t)outputPixelValue;
       }
   }
   
   void StretchContrast(uint8_t *image, ContrastStretchParams *stretchParams_cuda, int32_t imageSize,
                        cudaStream_t cudaStream)
   {
       int blockSize  = BLOCK_SIZE;
       int blockCount = BLOCK_COUNT;
       stretch_kernel<<<blockCount, blockSize, 0, cudaStream>>>(image, stretchParams_cuda, imageSize);
   }
   

8. The cudaStreamSynchronize() call blocks the host until all PVA and CUDA work enqueued on cudaStream completes.

       CHECK_CUDA_ERROR_GOTO(cudaStreamSynchronize(cudaStream), err, DeallocateAllResources);
   

9. The enhanced image is written to the output file and allocated resources are freed in the last step.

       if (WriteImageBuffer(OUTPUT_IMAGE_NAME, ".", image_cpu, IMAGE_SIZE) != 0)
       {
           err = -1;
           goto DeallocateAllResources;
       }
   
   DeallocateAllResources:
   CmdProgramCreateFailed:
       if (createdCmdProgramCount > 0)
       {
           CupvaCmdDestroy(&progComputeDynamicRange);
       }
       CupvaExecutableDestroy(execComputeDynamicRange);
   ExecutableCreateFailed:
   MemAllocFailed:
       cudaFreeHost(image_cpu);
       cudaFreeHost(stretchParams_cpu);
       cudaStreamDestroy(cudaStream);
       return err;
   }
   

Output

Make sure to install CUDA toolkit and CUDA cross-compile packages for your target before building this tutorial. CUDAToolkit_ROOT should be set manually in the CMakeLists.txt file for the QNX targets.

The path to the Tutorial assets directory containing the input image file “low-contrast-kodim08-768x512-grayscale.data” should be provided as an argument. The enhanced image output file “unsharp-masked-kodim08-768x512-grayscale.data” is written to the current working directory.

C++

$ ./cuda_interoperability_cpp -a <Tutorial Assets Directory Path>
Read 393216 bytes from <Tutorial Assets Directory Path>/low-contrast-kodim08-768x512-grayscale.data
Wrote 393216 bytes to ./contrast-stretched-kodim08-768x512-grayscale.data

C

$ ./cuda_interoperability_c -a <Tutorial Assets Directory Path>
Read 393216 bytes from <Tutorial Assets Directory Path>/low-contrast-kodim08-768x512-grayscale.data
Wrote 393216 bytes to ./contrast-stretched-kodim08-768x512-grayscale.data