Cvb/CppCudaMemory

// ----------------------------------------------------------------------------
/// \brief  Demonstrates how to pass user-allocated memory as the target
///         buffer(s) for image acquisition from a camera. In this particular
///         example, the allocated memory is either CUDA managed memory or 
///         host mapped memory, enabling automatic page migration or "zero-copy"
///         access to the acquired images respectively and thus reducing manual
///         copy overhead.
// ----------------------------------------------------------------------------
 
#include <iostream>
#include <memory>
#include <string>
#include <vector>
#include <cstdlib>
#include <iomanip>
#include <sstream>
 
#include <cvb/device_factory.hpp>
#include <cvb/global.hpp>
#include <cvb/driver/driver.hpp>
#include <cvb/image.hpp>
#include <cvb/composite.hpp>
#include <cvb/utilities/utilities.hpp>
#include <cvb/driver/image_stream.hpp>
 
#include <cuda_runtime.h>
 
/* 
 * Controls whether managed memory ("automatic page migration") or
 * "zero-copy" mapped host memory should be used.
 * 
 * In the \b true case cudaMallocManaged will be used to allocate the memory.
 * The pointer is written to by the acquisition engine and the same pointer
 * can be used for CUDA device kernels.
 * In the \b false case, cudaAllocHost with cudaHostAllocWriteCombined | cudaHostAllocMapped
 * is used for the acquisition engine buffers.
 * This is effectively "zero-copy" memory, as the GPU will be able to work on the host memory
 * by getting the corresponding device pointer from cudaHostGetDevicePointer.
 * Note: cudaHostAllocWriteCombined leads to disabled host caching of the data, which in turn
 * leads to slow reads on CPU, so the acquired buffers should mainly be consumed on the GPU.
 * The output image is manually managed memory in the latter case.
 */
static const constexpr bool USE_MANAGED_MEMORY = true;
 
static const constexpr auto TIMEOUT                = std::chrono::milliseconds(3000);
static const constexpr int NUM_ELEMENTS_TO_ACQUIRE = 4;
static const constexpr int NUM_BUFFERS             = 4;
static const constexpr int SOBEL_RADIUS            = 1;
__constant__ const int sobelKernel[9]              = {-1, 0, 1, -2, 0, 2, -1, 0, 1};
 
#define CUDA_CHECK(ans)                    \
{                                        \
  CudaAssert((ans), __FILE__, __LINE__); \
}
inline void CudaAssert(cudaError_t code, const char *file, int line, bool abort = true)
{
  if (code != cudaSuccess)
  {
    std::cerr << "CUDA Error: " << cudaGetErrorString(code) << " " << file << ":" << line << std::endl;
    if (abort)
      exit(code);
  }
}
 
void SetupCudaDevice()
{
  int gpuCount = 0;
  CUDA_CHECK(cudaGetDeviceCount(&gpuCount));
  if(gpuCount == 0)
    throw std::runtime_error{"No CUDA device present."};
 
  cudaDeviceProp deviceProps;
  CUDA_CHECK(cudaGetDeviceProperties(&deviceProps, 0));
  std::cout << "Using CUDA device: " << deviceProps.name << std::endl;
 
  if (USE_MANAGED_MEMORY)
  {
    if(!deviceProps.managedMemory)
      throw std::runtime_error{"CUDA device doesn't support managed memory"};
  }
  else
  {
    if (!deviceProps.canMapHostMemory)
      throw std::runtime_error{"CUDA device doesn't support mapping host memory"};
    CUDA_CHECK(cudaSetDeviceFlags(cudaDeviceMapHost));
  }
}
 
__global__ void gpuSobel(const std::uint8_t *inputImage, int width, int height, std::uint8_t *targetImage)
{
  extern __shared__ std::uint8_t inputShMem[];
 
  // calculate the global index of the filtered output image
  const int block_start_x = blockDim.x * blockIdx.x;
  const int block_start_y = blockDim.y * blockIdx.y;
  const int o_x           = blockDim.x * blockIdx.x + threadIdx.x;
  const int o_y           = blockDim.y * blockIdx.y + threadIdx.y;
  const int output_index  = o_y * (width - 2 * SOBEL_RADIUS) + o_x;
  const int s_x           = threadIdx.x + SOBEL_RADIUS;
  const int s_y           = threadIdx.y + SOBEL_RADIUS;
  const int shmem_width   = blockDim.x + 2 * SOBEL_RADIUS + 1;
  const int shmem_height  = blockDim.y + 2 * SOBEL_RADIUS;
  const int shmem_index   = s_y * (shmem_width) + s_x;
  const int numberOfPix   = (width - 2 * SOBEL_RADIUS) * (height - 2 * SOBEL_RADIUS);
 
  // calculate global index of the input image
  const int i_x         = o_x + SOBEL_RADIUS;  // o_x = output_index % aoi_width
  const int i_y         = o_y + SOBEL_RADIUS;  // o_y = output_index / aoi_width
  const int input_index = i_y * width + i_x;
 
  inputShMem[shmem_index] = (input_index < width * height) ? inputImage[input_index] : 0;
  if (threadIdx.x < 2)
  {  // l & r
    const int idx = i_y * width + block_start_x + threadIdx.x * (shmem_width - 2);
    inputShMem[s_y * shmem_width + threadIdx.x * (shmem_width - 2)] = (idx < width * height) ? inputImage[idx] : 0;
  }
 
  if (threadIdx.y < 2)
  {  // t & b
    const int idx = (block_start_y + threadIdx.y * (shmem_height - 1)) * width + i_x;
    inputShMem[threadIdx.y * (shmem_height - 1) * shmem_width + s_x] = (idx < width * height) ? inputImage[idx] : 0;
  }
 
  if (threadIdx.y == 0 && threadIdx.x == 0)
  {  // corners
    const int idxTopLeft  = block_start_y * width + block_start_x;
    const int idxTopRight = block_start_y * width + block_start_x + shmem_width - 2;
    const int idxBotLeft  = (block_start_y + shmem_height - 1) * width + block_start_x;
    const int idxBotRight = (block_start_y + shmem_height - 1) * width + block_start_x + shmem_width - 2;
 
    inputShMem[0]                                = (idxTopLeft < width * height) ? inputImage[idxTopLeft] : 0;
    inputShMem[shmem_width - 2]                  = (idxTopRight < width * height) ? inputImage[idxTopRight] : 0;
    inputShMem[(shmem_height - 1) * shmem_width] = (idxBotLeft < width * height) ? inputImage[idxBotLeft] : 0;
    inputShMem[(shmem_height - 1) * shmem_width + shmem_width - 2] =
        (idxBotRight < width * height) ? inputImage[idxBotRight] : 0;
  }
 
  __syncthreads();  // make sure every thread sees inputShMem.
 
  // test, if the index is still in range
  if (output_index < numberOfPix)
  {
    // variables for solbel filtering
    std::uint16_t sumX = 0, sumY = 0;
 
    // filter in x-direction
    sumX = inputShMem[shmem_index - shmem_width - 1] * sobelKernel[0]
            + inputShMem[shmem_index - shmem_width + 1] * sobelKernel[2] + inputShMem[shmem_index - 1] * sobelKernel[3]
            + inputShMem[shmem_index + 1] * sobelKernel[5] + inputShMem[shmem_index + shmem_width - 1] * sobelKernel[6]
            + inputShMem[shmem_index + shmem_width + 1] * sobelKernel[8];
 
    // filter in y-direction
    sumY = inputShMem[shmem_index - shmem_width - 1] * sobelKernel[0]
            + inputShMem[shmem_index - shmem_width] * sobelKernel[3]
            + inputShMem[shmem_index - shmem_width + 1] * sobelKernel[6]
            + inputShMem[shmem_index + shmem_width - 1] * sobelKernel[2]
            + inputShMem[shmem_index + shmem_width] * sobelKernel[5]
            + inputShMem[shmem_index + shmem_width + 1] * sobelKernel[8];
 
    targetImage[output_index] = min((sumX + sumY) / 2, std::numeric_limits<std::uint8_t>::max());
  }
}
 
namespace Tutorial
{
 
  /* FlowSetPool does not store the FlowSets for further
   * use, therefore we create a subclass to release the
   * buffer later
   */
  class UserFlowSetPool final : public Cvb::Driver::FlowSetPool
  {
 
    using UserFlowSetPoolPtr = std::shared_ptr<UserFlowSetPool>;
 
 
  public:
    UserFlowSetPool(const std::vector<Cvb::FlowInfo> &flowInfo) noexcept
        : Cvb::FlowSetPool(flowInfo, Cvb::FlowSetPool::ProtectedTag{})
    {
    }
 
    virtual ~UserFlowSetPool()
    {
      for (auto &flowSet : *this)
        for (auto &flow : flowSet)
          if (USE_MANAGED_MEMORY)
          {
            CUDA_CHECK(cudaFree(flow.Buffer));
          }
          else
          {
            CUDA_CHECK(cudaFreeHost(flow.Buffer));
          }
    }
 
    static UserFlowSetPoolPtr Create(const std::vector<Cvb::FlowInfo> &flowInfos)
    {
      return std::make_shared<UserFlowSetPool>(flowInfos);
    }
  };
 
}  // namespace Tutorial
 
template <class Func>
void AcquireData(const Cvb::CompositeStreamPtr &stream, Func &&processFunc)
{
  for (auto i = 0; i < NUM_ELEMENTS_TO_ACQUIRE; i++)
  {
    Cvb::CompositePtr composite;
    Cvb::WaitStatus waitStatus;
    Cvb::NodeMapEnumerator enumerator;
    std::tie(composite, waitStatus, enumerator) = stream->WaitFor(TIMEOUT);
    switch (waitStatus)
    {
      default:
        std::cout << "wait status unknown.\n";
      case Cvb::WaitStatus::Abort:
      case Cvb::WaitStatus::Timeout:
      {
        std::cout << "wait status not ok\n";
        continue;
      }
      case Cvb::WaitStatus::Ok:
      {
        break;
      }
    }
 
    // assume the composites first element is an image
    auto firstElement = composite->ItemAt(0);
 
    if (!Cvb::holds_alternative<Cvb::ImagePtr>(firstElement))
    {
      std::cout << "composite does not contain an image at the first element\n";
      continue;
    }
 
    auto image        = Cvb::get<Cvb::ImagePtr>(firstElement);
    auto linearAccess = image->Plane(0).LinearAccess();
    std::cout << "acquired image: " << i << " at memory location: " << reinterpret_cast<intptr_t>(linearAccess.BasePtr()) << "\n";
 
    // generate file name
    std::stringstream acqFileName;
    acqFileName << "acquired" << std::setw(2) << std::setfill('0') << i << ".png";
    image->Save(acqFileName.str());
    std::stringstream procFileName;
    procFileName << "processed" << std::setw(2) << std::setfill('0') << i << ".png";
 
    // call processing lambda
    processFunc(image, linearAccess, procFileName.str());
  }
}
 
int main()
{
  try
  {
    // check for present CUDA device, print name and check for required device features
    SetupCudaDevice();
 
    // discover transport layers
    auto infoList = Cvb::DeviceFactory::Discover(Cvb::DiscoverFlags::IgnoreVins);
 
    // can't continue the demo if there's no available device
    if (infoList.empty())
      throw std::runtime_error("There is no available device for this demonstration.");
 
    // instantiate the first device in the discovered list
    auto device = Cvb::DeviceFactory::Open<Cvb::GenICamDevice>(infoList[0].AccessToken(), Cvb::AcquisitionStack::GenTL);
 
    if (device->StreamCount() == 0)
      throw std::runtime_error("There is no available stream for this demonstration.");
 
    // get the first stream
    auto stream = device->Stream<Cvb::CompositeStream>(0);
 
    // get the flow set information that is needed for the current stream
    auto flowSetInfo = stream->FlowSetInfo();
 
    // create a subclass of FlowSetPool to store the created buffer
    auto flowSetPoolPtr = Tutorial::UserFlowSetPool::Create(flowSetInfo);
 
    std::generate_n(std::back_inserter(*flowSetPoolPtr), NUM_BUFFERS, [&flowSetInfo]() {
      auto flows = std::vector<void *>(flowSetInfo.size());
      std::transform(flowSetInfo.begin(), flowSetInfo.end(), flows.begin(),
                     [](Cvb::Driver::FlowInfo info) {
                       void *ptr = nullptr;
                       if (USE_MANAGED_MEMORY)
                       {
                         // allocate managed memory, but attach to host
                         CUDA_CHECK(cudaMallocManaged(&ptr, info.Size, cudaMemAttachHost));
                       }
                       else
                       {
                         // allocate host memory, which is readable from the GPU
                         CUDA_CHECK(cudaHostAlloc(&ptr, info.Size, cudaHostAllocWriteCombined | cudaHostAllocMapped));
                       }
                       return ptr;
                     });
      return flows;
    });
 
    // register the user flow set pool
    stream->RegisterExternalFlowSetPool(std::move(flowSetPoolPtr));
 
    std::uint8_t *dTarget = nullptr;
    std::uint8_t *target  = nullptr;
    if (USE_MANAGED_MEMORY)
    {
      // allocate managed memory, but attach to host
      CUDA_CHECK(cudaMallocManaged(&dTarget, flowSetInfo.front().Size, cudaMemAttachHost));
    }
    else
    {
      // allocate manually managed memory on GPU
      CUDA_CHECK(cudaMalloc(&dTarget, flowSetInfo.front().Size));
      // allocate pinned memory on host
      CUDA_CHECK(cudaMallocHost(&target, flowSetInfo.front().Size));
    }
 
    // start the data acquisition for that stream
    stream->EngineStart();
    stream->DeviceStart();
 
    // create a GPU processing stream
    cudaStream_t gpuStream;
    CUDA_CHECK(cudaStreamCreate(&gpuStream));
    if (USE_MANAGED_MEMORY)
    {
      AcquireData(stream, [&](Cvb::ImagePtr &image, Cvb::LinearAccessData &linearAccess, std::string outputName) {
        const dim3 blocks{(static_cast<unsigned>(image->Width() - 2 * SOBEL_RADIUS) + 15) / 16,
                          (static_cast<unsigned>(image->Height() - 2 * SOBEL_RADIUS) + 15) / 16};
        const dim3 threads{16, 16};
        std::uint8_t *dInput = reinterpret_cast<std::uint8_t *>(linearAccess.BasePtr());
        // attach memory to stream -> make writable from GPU
        CUDA_CHECK(cudaStreamAttachMemAsync(gpuStream, dTarget));
        CUDA_CHECK(cudaStreamAttachMemAsync(gpuStream, dInput));
        
        // enqueue CUDA kernel
        gpuSobel<<<blocks, threads, (16 + 2 * SOBEL_RADIUS + 1) * (16 + 2 * SOBEL_RADIUS), gpuStream>>>(
            dInput, image->Width(), image->Height(), dTarget);
        
        // prefetch target buffer
        CUDA_CHECK(cudaStreamAttachMemAsync(gpuStream, dTarget, 0, cudaMemAttachHost));
        
        // detach from stream -> make writable from acquisition engine again
        CUDA_CHECK(cudaStreamAttachMemAsync(gpuStream, dInput, 0, cudaMemAttachHost));
        CUDA_CHECK(cudaStreamSynchronize(gpuStream));
 
        // map managed memory to Cvb image
        auto wrapped = Cvb::WrappedImage::FromGrey8Pixels(dTarget, image->Width() - 2 * SOBEL_RADIUS,
                                                          image->Height() - 2 * SOBEL_RADIUS);
        wrapped->Save(outputName);
      });
    }
    else
      AcquireData(stream, [&](Cvb::ImagePtr &image, Cvb::LinearAccessData &linearAccess, std::string outputName) {
        const dim3 blocks{(static_cast<unsigned>(image->Width()) + 15) / 16,
                          (static_cast<unsigned>(image->Height()) + 15) / 16};
        const dim3 threads{16, 16};
        
        // get GPU pointer for mapped host pointer
        std::uint8_t *dImage;
        CUDA_CHECK(cudaHostGetDevicePointer(&dImage, reinterpret_cast<void *>(linearAccess.BasePtr()), 0));
        
        // enqueue CUDA kernel
        gpuSobel<<<blocks, threads, (16 + 2 * SOBEL_RADIUS + 1) * (16 + 2 * SOBEL_RADIUS), gpuStream>>>(dImage, image->Width(),
                                                                                             image->Height(), dTarget);
        
        // copy output to host memory
        CUDA_CHECK(cudaMemcpyAsync(target, dTarget,
                              (image->Width() - 2 * SOBEL_RADIUS) * (image->Height() - 2 * SOBEL_RADIUS),
                              cudaMemcpyDeviceToHost, gpuStream));
        CUDA_CHECK(cudaStreamSynchronize(gpuStream));
 
        // map host memory to Cvb image
        auto wrapped = Cvb::WrappedImage::FromGrey8Pixels(target, image->Width() - 2 * SOBEL_RADIUS,
                                                          image->Height() - 2 * SOBEL_RADIUS);
        wrapped->Save(outputName);
      });
    CUDA_CHECK(cudaStreamDestroy(gpuStream));
 
    stream->DeviceAbort();
    stream->EngineAbort();
 
    // deregister the user flow set pool to get free buffer (releaseCallback)
    stream->DeregisterFlowSetPool();
  }
  catch (const std::exception &e)
  {
    std::cout << e.what() << std::endl;
  }
 
  return 0;
}