Polymorphism and derived classes in CUDA / CUDA Thrust

by

I am not going to attempt to answer everything in this question, it is just too large. Having said that here are some observations about the code you posted which might help:

  • The GPU side new operator allocates memory from a private runtime heap. As of CUDA 6, that memory cannot be accessed by the host side CUDA APIs. You can access the memory from within kernels and device functions, but that memory cannot be accessed by the host. So using new inside a thrust device functor is a broken design that can never work. That is why your “vector of pointers” model fails.
  • Thrust is fundamentally intended to allow data parallel versions of typical STL algorithms to be applied to POD types. Building a codebase using complex polymorphic objects and trying to cram those through Thrust containers and algorithms might be made to work, but it isn’t what Thrust was designed for, and I wouldn’t recommend it. Don’t be surprised if you break thrust in unexpected ways if you do.
  • CUDA supports a lot of C++ features, but the compilation and object models are much simpler than even the C++98 standard upon which they are based. CUDA lacks several key features (RTTI for example) which make complex polymorphic object designs workable in C++. My suggestion is use C++ features sparingly. Just because you can do something in CUDA doesn’t mean you should. The GPU is a simple architecture and simple data structures and code are almost always more performant than functionally similar complex objects.

Having skim read the code you posted, my overall recommendation is to go back to the drawing board. If you want to look at some very elegant CUDA/C++ designs, spend some time reading the code bases of CUB and CUSP. They are both very different, but there is a lot to learn from both (and CUSP is built on top of Thrust, which makes it even more relevant to your usage case, I suspect).

More Related Contents:

  1. Thrust inside user written kernels
  2. CUDA compute capability requirements
  3. Cuda atomics change flag
  4. How is CUDA memory managed?
  5. Modifying registry to increase GPU timeout, windows 7
  6. How can I implement a custom atomic function involving several variables?
  7. Passing Host Function as a function pointer in __global__ OR __device__ function in CUDA
  8. How to verify CuDNN installation?
  9. What is the purpose of using multiple “arch” flags in Nvidia’s NVCC compiler?
  10. How are 2D / 3D CUDA blocks divided into warps?
  11. What is a bank conflict? (Doing Cuda/OpenCL programming)
  12. nvidia-smi Volatile GPU-Utilization explanation?
  13. How do I use Nvidia Multi-process Service (MPS) to run multiple non-MPI CUDA applications?
  14. Any particular function to initialize GPU other than the first cudaMalloc call?
  15. CUDA and nvcc: using the preprocessor to choose between float or double
  16. Copy an object to device?
  17. CUDA: How many concurrent threads in total?
  18. How do I select which GPU to run a job on?
  19. multi-GPU basic usage
  20. device function pointers
  21. How can I make tensorflow run on a GPU with capability 2.x?
  22. cudaMemset() – does it set bytes or integers?
  23. How to implement device side CUDA virtual functions?
  24. How to disable a specific nvcc compiler warnings
  25. Inconsistency of IDs between ‘nvidia-smi -L’ and cuDeviceGetName()
  26. CUDA allocation alignment is 256 bytes – seriously?
  27. printf() in my CUDA kernel doesn’t result produce any output
  28. Multiply Rectangular Matrices in CUDA
  29. Which Compute Capability is supported by which CUDA versions?
  30. “invalid configuration argument ” error for the call of CUDA kernel?

Leave a Comment