Thrust inside user written kernels

by

As it was originally written, Thrust is purely a host side abstraction. It cannot be used inside kernels. You can pass the device memory encapsulated inside a thrust::device_vector to your own kernel like this:

thrust::device_vector< Foo > fooVector;
// Do something thrust-y with fooVector

Foo* fooArray = thrust::raw_pointer_cast( fooVector.data() );

// Pass raw array and its size to kernel
someKernelCall<<< x, y >>>( fooArray, fooVector.size() );

and you can also use device memory not allocated by thrust within thrust algorithms by instantiating a thrust::device_ptr with the bare cuda device memory pointer.

Edited four and half years later to add that as per @JackOLantern’s answer, thrust 1.8 adds a sequential execution policy which means you can run single threaded versions of thrust’s alogrithms on the device. Note that it still isn’t possible to directly pass a thrust device vector to a kernel and device vectors can’t be directly used in device code.

Note that it is also possible to use the thrust::device execution policy in some cases to have parallel thrust execution launched by a kernel as a child grid. This requires separate compilation/device linkage and hardware which supports dynamic parallelism. I am not certain whether this is actually supported in all thrust algorithms or not, but certainly works with some.

More Related Contents:

  1. Polymorphism and derived classes in CUDA / CUDA Thrust
  2. What is the canonical way to check for errors using the CUDA runtime API?
  3. Unspecified launch failure on Memcpy
  4. Different CUDA versions shown by nvcc and NVIDIA-smi
  5. GPU Emulator for CUDA programming without the hardware [closed]
  6. How to measure the inner kernel time in NVIDIA CUDA?
  7. How do CUDA blocks/warps/threads map onto CUDA cores?
  8. Understanding CUDA grid dimensions, block dimensions and threads organization (simple explanation) [closed]
  9. How are 2D / 3D CUDA blocks divided into warps?
  10. What can I do against ‘CUDA driver version is insufficient for CUDA runtime version’?
  11. How to get the CUDA version?
  12. Can anyone provide sample code demonstrating the use of 16 bit floating point in cuda?
  13. CUDA and nvcc: using the preprocessor to choose between float or double
  14. What kind of variables consume registers in CUDA?
  15. CUDA: How many concurrent threads in total?
  16. CUDA apps time out & fail after several seconds – how to work around this?
  17. How do I select which GPU to run a job on?
  18. multi-GPU basic usage
  19. How to use 2D Arrays in CUDA?
  20. How can I make tensorflow run on a GPU with capability 2.x?
  21. How to implement device side CUDA virtual functions?
  22. How to create a CUDA context?
  23. CUBLAS: Incorrect inversion for matrix with zero pivot
  24. Using maximum shared memory in Cuda
  25. Why does my CUDA kernel crash (unspecified launch failure) with a different dataset size?
  26. CUDA determining threads per block, blocks per grid
  27. How does CUDA assign device IDs to GPUs?
  28. How to disable a specific nvcc compiler warnings
  29. CUDA function pointers
  30. Using an array of device function pointers

Leave a Comment