principles from introductory CUDA examples to an existing project that has a meaningful amount of non-trivial code. • Provide some guidance to people about to embark on using CUDA to speed up existing software float* y) { for (int i = 0; i < n; i++) y[i] = x[i] + y[i]; } TEST_CASE("cppcon-0", "[CUDA]") { int N = 1 << 20; float* x = new float[N]; float* y = new float[N]; for (int add_cpu(N, x, y); delete[] x; delete[] y; } An Even Easier Introduction to CUDA 4 |TEST_CASE("cppcon-1", "[CUDA]") { int N = 1 << 20; float* x; float* y; cudaMallocManaged(&x, N*sizeof(float));

the Gap: Writing Portable Programs for CPU and GPU using CUDA Thomas Mejstrik Sebastian Woblistin 2/66Content 1 Motivation Audience etc.. Cuda crash course Quiz time 2 Patterns Oldschool host device everywhere Conditional function body constexpr everything Disable Cuda warnings host device template 3 The dark path Function dispatch triple 4 Cuda proposal Conditional host device Forbid bad cross function dark path Cuda proposal Thank you Motivation 1 Motivation Audience etc.. Cuda crash course Quiz time 2 Patterns 3 The dark path 4 Cuda proposal5/66 Motivation Patterns The dark path Cuda proposal

Polling D C 1 #include 2 #include cuda.hpp> 3 4 taro::Taro taro{NUM_THREADS}; 5 auto cuda = taro.cuda_scheduler(NUM_STREAMS); 6 29Taro’s Programming Model – Example com/dian-lun-lin/taro A B Callback Wait Polling D C 7 auto task_a = taro.emplace([&]() { 8 cuda.wait([&](cudaStream_t stream) { 9 kernel_a1<<<32, 256, 0, stream>>>(); 10 }); // synchronize 11 7 auto task_a = taro.emplace([&]() { 8 cuda.wait([&](cudaStream_t stream) { 9 kernel_a1<<<32, 256, 0, stream>>>(); 10 }); // synchronize 11 }); CUDA stream for offloading GPU kernels 32Taro’s

very fast intra-node transfers GPU GPU Fast Intra- Node Fabric DataGPU Communication Libraries CUDA-Aware MPI NVSHMEM ROC_SHMEM - Communication libraries offering increasing support for GPU-to-GPU will utilize both GPUDirect RDMA and NVLink GASNet-EX Memory KindsGPU Communication Libraries CUDA-Aware MPI NVSHMEM ROC_SHMEM - Communication libraries offering increasing support for GPU-to-GPU = BCL::broadcast(ptr, 0); ptr[BCL::rank()] = BCL::rank(); BCL::cuda::ptr ptr = nullptr; if (BCL::rank() == 0) { ptr = BCL::cuda::alloc(BCL::nprocs()); } ptr = BCL::broadcast(ptr, 0); ptr[BCL::rank()]

CMake Presets support 3. ARM and ARM64 support (Raspberry Pi, Surface Pro X, Apple Silicon) 4. CUDA IntelliSense and GPU debuggingVisual Studio Code What’s new? 1. GitHub Codespaces (coding from your CMake Presets support 3. ARM and ARM64 support (Raspberry Pi, Surface Pro X, Apple Silicon) 4. CUDA IntelliSense and GPU debugging 5. Disassembly View while debugging Preview!Visual Studio Code What’s CMake Presets support 3. ARM and ARM64 support (Raspberry Pi, Surface Pro X, Apple Silicon) 4. CUDA IntelliSense and GPU debugging 5. Disassembly View while debugging Preview!Visual Studio Code What’s

TensorFlow. These are built using optimized, hardware-specific libraries (such as Intel’s MKL or NVIDIA’s CUDA) which speed up performance without code changes. Read more about how conda supports data scientists corresponds to the package. The currently supported list of virtual packages includes: • __cuda: Maximum version of CUDA supported by the display driver. • __osx: OSX version if applicable. • __glibc: Version post8+8f640d35a conda-build version : 3.17.8 python version : 3.7.2.final.0 virtual packages : __cuda=10.0 base environment : /Users/demo/dev/conda/devenv (writable) channel URLs : https://repo.anaconda

TensorFlow. These are built using optimized, hardware-specific libraries (such as Intel’s MKL or NVIDIA’s CUDA) which speed up performance without code changes. Read more about how conda supports data scientists corresponds to the package. The currently supported list of virtual packages includes: • __cuda: Maximum version of CUDA supported by the display driver. • __osx: OSX version if applicable. • __glibc: Version post8+8f640d35a conda-build version : 3.17.8 python version : 3.7.2.final.0 virtual packages : __cuda=10.0 base environment : /Users/demo/dev/conda/devenv (writable) channel URLs : https://repo.anaconda

TensorFlow. These are built using optimized, hardware-specific libraries (such as Intel’s MKL or NVIDIA’s CUDA) which speed up performance without code changes. Read more about how conda supports data scientists corresponds to the package. The currently supported list of virtual packages includes: • __cuda: Maximum version of CUDA supported by the display driver. • __osx: OSX version if applicable. • __glibc: Version post8+8f640d35a conda-build version : 3.17.8 python version : 3.7.2.final.0 virtual packages : __cuda=10.0 base environment : /Users/demo/dev/conda/devenv (writable) channel URLs : https://repo.anaconda

corresponds to the package. The currently supported list of virtual packages includes: • __cuda: Maximum version of CUDA supported by the display driver. • __osx: OSX version if applicable. • __glibc: Version post8+8f640d35a conda-build version : 3.17.8 python version : 3.7.2.final.0 virtual packages : __cuda=10.0 base environment : /Users/demo/dev/conda/devenv (writable) channel URLs : https://repo.anaconda using an environment variable. Supported variables include: • CONDA_OVERRIDE_CUDA - CUDA version number or set to "" for no CUDA detected. • CONDA_OVERRIDE_OSX - OSX version number or set to "" for no OSX

corresponds to the package. The currently supported list of virtual packages includes: • __cuda: Maximum version of CUDA supported by the display driver. • __osx: OSX version if applicable. • __glibc: Version post8+8f640d35a conda-build version : 3.17.8 python version : 3.7.2.final.0 virtual packages : __cuda=10.0 base environment : /Users/demo/dev/conda/devenv (writable) channel URLs : https://repo.anaconda using an environment variable. Supported variables include: • CONDA_OVERRIDE_CUDA - CUDA version number or set to "" for no CUDA detected. • CONDA_OVERRIDE_OSX - OSX version number or set to "" for no OSX