4 ms·

I think CUDA even lets you allocate pinned host memory too now - cuHostMalloc or something like that - so you can skip the double-copy

• my123 4y ago
  CUDA provides a tier significantly above that: unified memory. See: https://on-demand.gputechconf.com/gtc/2017/presentation/s7285-nikolay-sakharnykh-unified-memory-on-pascal-and-volta.pdf https://on-demand.gputechconf.com/gtc/2017/presentation/s728... And: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#um-unified-memory-programming-hd https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.... However, the Windows driver infrastructure's unified memory support is much further behind, with the pre-Pascal feature set. For those you'll have to use Linux. Note that WSL2 is considered as Windows for this, it's a driver infrastructure limitation in Windows.
  • andoma 4y ago
    I've switched to using cudaMallocManaged() exclusively. From what I can tell there's isn't much of a performance difference. A few cudaMemPrefetchAsync() at strategic places will remedy any performance problems. I really love the ability that you can just break with gdb and look around in that memory as well.
  • einpoklum 4y ago
    Unified memory is just _different_, not above or below. It offers on-demand paging. But that comes at a cost (at times) in terms of memory I/O speed.
    • my123 4y ago
      It's a feature tier above, with much more emphasis on ease of use from the programmer's perspective. It also allows for significantly more approachable programming models. For example: https://developer.nvidia.com/blog/accelerating-standard-c-with-gpus-using-stdpar/ https://developer.nvidia.com/blog/accelerating-standard-c-wi...
• boulos 4y ago
  Yeah, sorry if I was unclear: some folks thought that cuHostMalloc et al. and pinned memory were "impure". That you should instead have a unified sense of "allocate" and that it could sometimes be host, sometimes device, sometimes migrate. The unified memory support in CUDA (originally intended for Denver, IIRC) is mostly a response to people finding it too hard to decide (a la mmap, really). So it's not that CUDA doesn't have these. It's that it does, but many people never have to understand anything beyond "there's a thing called malloc, and there's host and device".
• 01100011 4y ago
  Sure, but pinned memory is often a limited resource and requires the GPU to issue PCI transactions. Depending on your needs, it's generally better to copy to/from the GPU explicitly, which can be done asynchronously, hiding the overhead behind other work to a degree.
  • jhj 4y ago
    In CUDA, some transfers involving pageable host memory are completely synchronous from the perspective of the host, even if you use `cudaMemcpyAsync`: https://docs.nvidia.com/cuda/cuda-runtime-api/api-sync-behavior.html#api-sync-behavior__memcpy-async https://docs.nvidia.com/cuda/cuda-runtime-api/api-sync-behav... Pinned memory is typically used to get around the synchronization aspects.