From 32dd41a728ce10113032e20955ba08f8de449857 Mon Sep 17 00:00:00 2001
From: Adrian Kummerlaender
Date: Sun, 12 Sep 2021 14:01:55 +0200
Subject: Start using C++ cuda-api-wrapper instead of raw CUDA

---
 CMakeLists.txt                |   2 +-
 default.nix                   |  24 ++-
 lbm.org                       | 371 +++++++++++++++++++++---------------------
 tangle/LLBM/kernel/executor.h |  37 -----
 tangle/LLBM/lattice.h         |  73 ++++-----
 tangle/LLBM/materials.h       |   2 +-
 tangle/LLBM/memory.h          |  29 ++--
 tangle/LLBM/propagate.h       |  32 ++--
 tangle/benchmark-ldc.cu       |  12 +-
 tangle/channel-with-sphere.cu |   8 +-
 tangle/ldc-2d.cu              |   8 +-
 tangle/ldc-3d.cu              |   8 +-
 tangle/magnus.cu              |  10 +-
 tangle/nozzle.cu              |   8 +-
 tangle/taylor-couette.cu      |   8 +-
 15 files changed, 320 insertions(+), 312 deletions(-)

diff --git a/CMakeLists.txt b/CMakeLists.txt
index ddd0b47..4c3bcbc 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -14,7 +14,7 @@ set(CMAKE_CXX_STANDARD 17)
 set(CMAKE_CXX_STANDARD_REQUIRED ON)
 
 set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --expt-relaxed-constexpr --default-stream per-thread --extended-lambda")
-set(CMAKE_CUDA_ARCHITECTURES 62)
+set(CMAKE_CUDA_ARCHITECTURES 86)
 
 include_directories(
 	${CMAKE_CURRENT_SOURCE_DIR}/tangle
diff --git a/default.nix b/default.nix
index fccf558..1fd61a1 100644
--- a/default.nix
+++ b/default.nix
@@ -7,6 +7,23 @@
 }, ... }:
 
 let
+  cuda-api-wrappers = pkgs.stdenv.mkDerivation rec {
+    name = "cuda-api-wrappers";
+    version = "0.4.3";
+
+    src = pkgs.fetchFromGitHub {
+      owner = "eyalroz";
+      repo = "cuda-api-wrappers";
+      rev = "v${version}";
+      sha256 = "plsjzIeNjgOoJJCbrSVQp7TK30Bh/ka1SWAakYgQR2s=";
+    };
+
+    buildInputs = with pkgs; [
+      cmake
+      cudatoolkit_11
+    ];
+  };
+
   cuda-samples-common-headers = pkgs.stdenv.mkDerivation rec {
     name = "cuda-samples-common-headers";
     version = "11.1";
@@ -25,7 +42,7 @@ let
       cp -r $src/Common/* $out/include/cuda-samples/Common
     '';
   };
-    
+
   imgui-sfml = pkgs.stdenv.mkDerivation rec {
     name = "imgui-sfml";
     version = "2.1";
@@ -35,7 +52,7 @@ let
       repo = "imgui-sfml";
       rev = "v${version}";
       sha256 = "1g8gqly156miv12ajapnhmxfcv9i3fqhdmdy45gmdw47kh8ly5zj";
-    }; 
+    };
 
     buildInputs = with pkgs; [
       cmake
@@ -82,8 +99,9 @@ in pkgs.stdenv.mkDerivation rec {
   in with pkgs; [
     local-python
     cudatoolkit_11
+    cuda-api-wrappers
     cuda-samples-common-headers
-    linuxPackages.nvidia_x11 
+    linuxPackages.nvidia_x11
     libGL
     sfml
     imgui-sfml
diff --git a/lbm.org b/lbm.org
index 059a579..168427b 100644
--- a/lbm.org
+++ b/lbm.org
@@ -16,7 +16,7 @@ to a set of interesting examples. All of this runs on GPUs using CUDA near the m
 
 *This document is a [[*Open tasks][work in progress]].*
 #+BEGIN_EXPORT html
-<video style="width:100%" src="https://literatelb.org/talk/fire.webm" onloadeddata="this.play();" playsinline loop muted/>
+<video style="width:100%" src="https://literatelb.org/media/fire.webm" playsinline controls loop muted/>
 #+END_EXPORT
 #+END_ABSTRACT
 #+TOC: headlines 2
@@ -1349,6 +1349,7 @@ to access it.
 #include "kernel/propagate.h"
 
 #include <cuda.h>
+#include <cuda/runtime_api.hpp>
 #+END_SRC
 
 Our code employs the /Periodic Shift (PS)/ cite:kummerlanderImplicitPropagationDirectly2021 propagation
@@ -1404,6 +1405,7 @@ protected:
 
 public:
   CyclicPopulationBuffer(descriptor::Cuboid<DESCRIPTOR> cuboid);
+  ~CyclicPopulationBuffer();
 
   LatticeView<DESCRIPTOR,S> view() {
     return LatticeView<DESCRIPTOR,S>{ _cuboid, _population.device() };
@@ -1420,15 +1422,13 @@ out that this can be done at virtually no cost by using the in-hardware virtual
 logic. Doing so requires the array sizes to be exact multiples of the device page size.
 
 #+BEGIN_SRC cpp :tangle tangle/LLBM/propagate.h
-std::size_t getDevicePageSize(int device_id=-1) {
-  if (device_id == -1) {
-    cudaGetDevice(&device_id);
-  }
+std::size_t getDevicePageSize() {
+  auto device = cuda::device::current::get();
   std::size_t granularity = 0;
   CUmemAllocationProp prop = {};
   prop.type = CU_MEM_ALLOCATION_TYPE_PINNED;
   prop.location.type = CU_MEM_LOCATION_TYPE_DEVICE;
-  prop.location.id = device_id;
+  prop.location.id = device.id();
   cuMemGetAllocationGranularity(&granularity, &prop, CU_MEM_ALLOC_GRANULARITY_MINIMUM);
   return granularity;
 }
@@ -1457,8 +1457,7 @@ physical array in virtual memory.
 To do this we first need to know which device is currently selected.
 
 #+BEGIN_SRC cpp :tangle tangle/LLBM/propagate.h
-  int device_id = -1;
-  cudaGetDevice(&device_id);
+  auto device = cuda::device::current::get();
 #+END_SRC
 
 Using this device ID, a device-pinned address area large enough to fit two full views of the
@@ -1467,7 +1466,7 @@ lattice can be reserved.
 #+BEGIN_SRC cpp :tangle tangle/LLBM/propagate.h
   _prop.type = CU_MEM_ALLOCATION_TYPE_PINNED;
   _prop.location.type = CU_MEM_LOCATION_TYPE_DEVICE;
-  _prop.location.id = device_id;
+  _prop.location.id = device.id();
   cuMemAddressReserve(&_ptr, 2 * _volume * DESCRIPTOR::q, 0, 0, 0);
 #+END_SRC
 
@@ -1488,7 +1487,7 @@ the correct access flags
 
 #+BEGIN_SRC cpp :tangle tangle/LLBM/propagate.h
   _access.location.type = CU_MEM_LOCATION_TYPE_DEVICE;
-  _access.location.id = 0;
+  _access.location.id = device.id();
   _access.flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE;
   cuMemSetAccess(_ptr, 2 * _volume * DESCRIPTOR::q, &_access, 1);
 #+END_SRC
@@ -1507,7 +1506,7 @@ we also need to store those somewhere.
 
 #+BEGIN_SRC cpp :tangle tangle/LLBM/propagate.h
   for (unsigned iPop=0; iPop < DESCRIPTOR::q; ++iPop) {
-    _base[iPop] = device() + iPop * 2 * (_volume / sizeof(S));
+    _base[iPop] = this->device() + iPop * 2 * (_volume / sizeof(S));
     _population[iPop] = _base[iPop] + iPop * ((_volume / sizeof(S)) / DESCRIPTOR::q);
   }
 
@@ -1516,6 +1515,20 @@ we also need to store those somewhere.
 }
 #+END_SRC
 
+Finally, once the population buffer is no longer needed, we should also release the
+mapped memory again.
+
+#+BEGIN_SRC cpp :tangle tangle/LLBM/propagate.h
+template <typename DESCRIPTOR, typename S>
+CyclicPopulationBuffer<DESCRIPTOR,S>::~CyclicPopulationBuffer() {
+  cuMemUnmap(_ptr, 2 * _volume * DESCRIPTOR::q);
+  for (unsigned iPop=0; iPop < DESCRIPTOR::q; ++iPop) {
+    cuMemRelease(_handle[iPop]);
+  }
+  cuMemAddressFree(_ptr, 2 * _volume * DESCRIPTOR::q);
+}
+#+END_SRC
+
 ** Access
 The common interface of most of out GPU kernels is to accept an array of current propulations and
 write the new populations to another array. This way we can control where the populations are read
@@ -1558,7 +1571,9 @@ for (unsigned iPop=0; iPop < DESCRIPTOR::q; ++iPop) {
 #+BEGIN_SRC cpp :tangle tangle/LLBM/propagate.h :noweb no
 template <typename DESCRIPTOR, typename S>
 void CyclicPopulationBuffer<DESCRIPTOR,S>::stream() {
-  propagate<DESCRIPTOR,S><<<1,1>>>(view(), _base.device(), _volume / sizeof(S));
+  cuda::launch(propagate<DESCRIPTOR,S>,
+               cuda::launch_configuration_t(1,1),
+               view(), _base.device(), _volume / sizeof(S));
 }
 #+END_SRC
 
@@ -1873,6 +1888,8 @@ SignedDistanceBoundary(Lattice<DESCRIPTOR,T,S>&, CellMaterials<DESCRIPTOR>&, SDF
 #include <memory>
 #include <vector>
 #include <cstring>
+
+#include <cuda/runtime_api.hpp>
 #+END_SRC
 
 Most memory of our simulation resides solely on the GPU. While we already defined a data structure for
@@ -1884,7 +1901,8 @@ template <typename T>
 class DeviceBuffer {
 protected:
   const std::size_t _size;
-  T* _data;
+  cuda::device_t _device;
+  cuda::memory::device::unique_ptr<T[]> _data;
 #+END_SRC
 
 Note that the value of the =_data= pointer is going to be a address in GPU memory. We only expose access to it
@@ -1896,19 +1914,15 @@ host.
 #+BEGIN_SRC cpp :tangle tangle/LLBM/memory.h
 public:
   DeviceBuffer(std::size_t size):
-    _size(size) {
-    cudaMalloc(&_data, _size*sizeof(T));
-    cudaMemset(_data, 0, _size*sizeof(T));
-  }
+    _size(size),
+    _device(cuda::device::current::get()),
+    _data(cuda::memory::device::make_unique<T[]>(_device, size))
+  { }
   <<device-buffer-load-from-plain-data>>
   <<device-buffer-load-from-std-vector>>
- 
-  ~DeviceBuffer() {
-    cudaFree(_data);
-  }
 
   T* device() {
-    return _data;
+    return _data.get();
   }
 
   std::size_t size() const {
@@ -1924,7 +1938,7 @@ of a pointer and the size.
 #+BEGIN_SRC cpp :eval no
 DeviceBuffer(const T* data, std::size_t size):
   DeviceBuffer(size) {
-  cudaMemcpy(_data, data, size*sizeof(T), cudaMemcpyHostToDevice);
+  cuda::memory::copy(_data.get(), data, size*sizeof(T));
 }
 #+END_SRC
 
@@ -1947,7 +1961,7 @@ template <typename T>
 class SharedVector : public DeviceBuffer<T> {
 private:
   std::unique_ptr<T[]> _host_data;
- 
+
 public:
   SharedVector(std::size_t size):
     DeviceBuffer<T>(size),
@@ -1964,11 +1978,11 @@ public:
   }
 
   void syncHostFromDevice() {
-    cudaMemcpy(_host_data.get(), this->_data, this->_size*sizeof(T), cudaMemcpyDeviceToHost);
+    cuda::memory::copy(_host_data.get(), this->_data.get(), this->_size*sizeof(T));
   }
 
   void syncDeviceFromHost() {
-    cudaMemcpy(this->_data, _host_data.get(), this->_size*sizeof(T), cudaMemcpyHostToDevice);
+    cuda::memory::copy(this->_data.get(), _host_data.get(), this->_size*sizeof(T));
   }
 
 };
@@ -2024,7 +2038,7 @@ public:
 
   DeviceTexture(descriptor::CuboidD<3> c):
     DeviceTexture(c.nX, c.nY, c.nZ) { }
-  
+
   ~DeviceTexture() {
     cudaFreeArray(_array);
   }
@@ -2525,20 +2539,73 @@ struct post_process_by_spatial_cell_mask { };
 }
 #+END_SRC
 
+Generally speaking we want to avoid control flow branching inside of GPU kernels. However, it can
+make sense to group multiple masked operators into a single lattice pass of a single kernel using some
+small number of branches instead of launching a separate branch-free kernel for each operator.
+
+In fact, the most common type of operator call we will use for the collisions in our simulations is the
+fused application of a set of masked operators /curried/ with their arguments.
+
 #+BEGIN_SRC cpp :noweb no :tangle tangle/LLBM/lattice.h
 template <typename... OPERATOR>
 void apply(OPERATOR... ops) {
   const auto block_size = 32;
   const auto block_count = (_cuboid.volume + block_size - 1) / block_size;
-  kernel::call_operators<DESCRIPTOR,T,S,OPERATOR...><<<block_count,block_size>>>(
-    _population.view(), ops...
-  );
+  cuda::launch(kernel::call_operators<DESCRIPTOR,T,S,OPERATOR...>,
+               cuda::launch_configuration_t(block_count, block_size),
+               _population.view(),
+               ops...);
 }
 #+END_SRC
 
-The dispatching method =apply= calls the appropriate overload by instantiating the type defined by =OPERATOR::call_tag= and passing
-it on to =call_operator= alongside any other arguments. Note that the dispatching function and the overloaded functions have to be of
-different names to prevent infinite recursion as the variadic parameter pack could of course also capture tags.
+In order to generate a single kernel from multiple operators curried with their respective arguments
+we implement a small =Operator= helper class.
+
+#+BEGIN_SRC cpp :noweb no :tangle tangle/LLBM/operator.h
+#pragma once
+
+#include <tuple>
+
+template <typename OPERATOR, typename... ARGS>
+struct Operator {
+  bool* const mask;
+  const std::tuple<ARGS...> config;
+
+  Operator(OPERATOR, DeviceBuffer<bool>& m, ARGS... args):
+    mask(m.device()),
+    config(args...) { }
+#+END_SRC
+
+This =Operator= wrapper stores both the mask pointer and any arguments for a given =OPERATOR= type.
+The arguments will be passed to the actual =OPERATOR::apply= by the following proxy method for
+any masked cell index.
+
+#+BEGIN_SRC cpp :noweb no :tangle tangle/LLBM/operator.h
+  template <typename DESCRIPTOR, typename T, typename S>
+  __device__ bool apply(DESCRIPTOR d, S f_curr[DESCRIPTOR::q], S f_next[DESCRIPTOR::q], std::size_t gid) const {
+    if (mask[gid]) {
+      std::apply([](auto... args) { OPERATOR::template apply<T,S>(args...); },
+                 std::tuple_cat(std::make_tuple(d, f_curr, f_next, gid), config));
+      return true;
+    } else {
+      return false;
+    }
+  }
+};
+#+END_SRC
+
+Finally a deduction guide can be used to remove some of the visual cruft of declaring =Operator= instances in
+the application code.
+
+#+BEGIN_SRC cpp :noweb no :tangle tangle/LLBM/operator.h
+template <typename OPERATOR, typename... ARGS>
+Operator(OPERATOR, DeviceBuffer<bool>&, ARGS... args) -> Operator<OPERATOR,std::remove_reference_t<ARGS>...>;
+#+END_SRC
+
+For single operator application the =apply= dispatcher calls the appropriate overload by instantiating the type defined
+by =OPERATOR::call_tag= and passing it on to =call_operator= alongside any other arguments. Note that the dispatching
+function and the overloaded functions have to be of different names to prevent infinite recursion as the variadic parameter
+pack could of course also capture tags.
 
 #+BEGIN_SRC cpp :noweb no :tangle tangle/LLBM/lattice.h
 template <typename OPERATOR, typename... ARGS>
@@ -2547,47 +2614,39 @@ void apply(ARGS&&... args) {
 }
 #+END_SRC
 
-As cell IDs can be extracted from both a list of them and a mask applied to the full set we can use the different types of device buffer
-specializations to distinguish between masked and list-based applications of cell-ID-accepting kernels.
+One common case is to apply a =call_by_cell_id= operator on a list of such IDs.
 
 #+BEGIN_SRC cpp :noweb no :tangle tangle/LLBM/lattice.h
 template <typename OPERATOR, typename... ARGS>
 void call_operator(tag::call_by_cell_id, DeviceBuffer<std::size_t>& cells, ARGS... args) {
   const auto block_size = 32;
   const auto block_count = (cells.size() + block_size - 1) / block_size;
-  kernel::call_operator<OPERATOR,DESCRIPTOR,T,S,ARGS...><<<block_count,block_size>>>(
-    _population.view(), cells.device(), cells.size(), std::forward<ARGS>(args)...
-  );
-}
-#+END_SRC
-
-#+BEGIN_SRC cpp :noweb no :tangle tangle/LLBM/lattice.h
-template <typename OPERATOR, typename... ARGS>
-void call_operator(tag::call_by_cell_id, DeviceBuffer<bool>& mask, ARGS... args) {
-  const auto block_size = 32;
-  const auto block_count = (_cuboid.volume + block_size - 1) / block_size;
-  kernel::call_operator<OPERATOR,DESCRIPTOR,T,S,ARGS...><<<block_count,block_size>>>(
-    _population.view(), mask.device(), std::forward<ARGS>(args)...
-  );
+  cuda::launch(kernel::call_operator<OPERATOR,DESCRIPTOR,T,S,ARGS...>,
+               cuda::launch_configuration_t(block_count, block_size),
+               _population.view(),
+               cells.device(), cells.size(),
+               std::forward<ARGS>(args)...);
 }
 #+END_SRC
 
-The list-based application doesn't perform any detection of device buffer types but can be thought of as simply iterating
-over =count= elements and passing the current index to the individual device operator. This is useful when calling e.g. the
-interpolated bounce back kernel as there is no 1:1 mapping between threads and cells to be found there.
+The more generic version of this caller is list-based application. These calls can be thought of as simply iterating
+over =count= elements and passing the current index to the individual device operator. This is useful when calling
+e.g. the interpolated bounce back kernel as there is no 1:1 mapping between threads and cells to be found there.
 
 #+BEGIN_SRC cpp :noweb no :tangle tangle/LLBM/lattice.h
 template <typename OPERATOR, typename... ARGS>
 void call_operator(tag::call_by_list_index, std::size_t count, ARGS... args) {
   const auto block_size = 32;
   const auto block_count = (count + block_size - 1) / block_size;
-  kernel::call_operator_using_list<OPERATOR,DESCRIPTOR,T,S,ARGS...><<<block_count,block_size>>>(
-    _population.view(), count, std::forward<ARGS>(args)...
-  );
+  cuda::launch(kernel::call_operator_using_list<OPERATOR,DESCRIPTOR,T,S,ARGS...>,
+               cuda::launch_configuration_t(block_count, block_size),
+               _population.view(),
+               count,
+               std::forward<ARGS>(args)...);
 }
 #+END_SRC
 
-We provide corresponding =inspect= callers for read-only functor kernels.
+Next, we provide corresponding =inspect= callers for read-only functor kernels.
 
 #+BEGIN_SRC cpp :noweb no :tangle tangle/LLBM/lattice.h
 template <typename FUNCTOR, typename... ARGS>
@@ -2596,25 +2655,19 @@ void inspect(ARGS&&... args) {
 }
 #+END_SRC
 
-#+BEGIN_SRC cpp :noweb no :tangle tangle/LLBM/lattice.h
-template <typename FUNCTOR, typename... ARGS>
-void call_functor(tag::call_by_cell_id, DeviceBuffer<std::size_t>& cells, ARGS... args) {
-  const auto block_size = 32;
-  const auto block_count = (cells.size() + block_size - 1) / block_size;
-  kernel::call_functor<FUNCTOR,DESCRIPTOR,T,S,ARGS...><<<block_count,block_size>>>(
-    _population.view(), cells.device(), cells.size(), std::forward<ARGS>(args)...
-  );
-}
-#+END_SRC
+Different from cell-based operators, the common case for functors is to evaluate a =call_by_cell_id=
+functor using a mask (specifically the bulk mask) instead of a list.
 
 #+BEGIN_SRC cpp :noweb no :tangle tangle/LLBM/lattice.h
 template <typename FUNCTOR, typename... ARGS>
 void call_functor(tag::call_by_cell_id, DeviceBuffer<bool>& mask, ARGS... args) {
   const auto block_size = 32;
   const auto block_count = (_cuboid.volume + block_size - 1) / block_size;
-  kernel::call_functor<FUNCTOR,DESCRIPTOR,T,S,ARGS...><<<block_count,block_size>>>(
-    _population.view(), mask.device(), std::forward<ARGS>(args)...
-  );
+  cuda::launch(kernel::call_functor<FUNCTOR,DESCRIPTOR,T,S,ARGS...>,
+               cuda::launch_configuration_t(block_count, block_size),
+               _population.view(),
+               mask.device(),
+               std::forward<ARGS>(args)...);
 }
 #+END_SRC
 
@@ -2629,9 +2682,11 @@ void call_functor(tag::call_by_spatial_cell_mask, DeviceBuffer<bool>& mask, ARGS
   const dim3 grid((_cuboid.nX + block.x - 1) / block.x,
                   (_cuboid.nY + block.y - 1) / block.y,
                   (_cuboid.nZ + block.z - 1) / block.z);
-  kernel::call_spatial_functor<FUNCTOR,DESCRIPTOR,T,S,ARGS...><<<grid,block>>>(
-    _population.view(), mask.device(), std::forward<ARGS>(args)...
-  );
+  cuda::launch(kernel::call_spatial_functor<FUNCTOR,DESCRIPTOR,T,S,ARGS...>,
+               cuda::launch_configuration_t(grid, block),
+               _population.view(),
+               mask.device(),
+               std::forward<ARGS>(args)...);
 }
 #+END_SRC
 
@@ -2651,9 +2706,11 @@ template <typename OPERATOR, typename... ARGS>
 void tagged_helper(tag::post_process_by_list_index, std::size_t count, ARGS... args) {
   const auto block_size = 32;
   const auto block_count = (count + block_size - 1) / block_size;
-  kernel::call_operator_using_list<OPERATOR,DESCRIPTOR,T,S,ARGS...><<<block_count,block_size>>>(
-    DESCRIPTOR(), count, std::forward<ARGS>(args)...
-  );
+  cuda::launch(kernel::call_operator_using_list<OPERATOR,DESCRIPTOR,T,S,ARGS...>,
+               cuda::launch_configuration_t(block_count, block_size),
+               DESCRIPTOR(),
+               count,
+               std::forward<ARGS>(args)...);
 }
 #+END_SRC
 
@@ -2664,62 +2721,16 @@ void tagged_helper(tag::post_process_by_spatial_cell_mask, DeviceBuffer<bool>& m
   const dim3 grid((_cuboid.nX + block.x - 1) / block.x,
                   (_cuboid.nY + block.y - 1) / block.y,
                   (_cuboid.nZ + block.z - 1) / block.z);
-  kernel::call_spatial_operator<OPERATOR,DESCRIPTOR,T,S,ARGS...><<<grid,block>>>(
-    _cuboid, mask.device(), std::forward<ARGS>(args)...
-  );
+  cuda::launch(kernel::call_spatial_operator<OPERATOR,DESCRIPTOR,T,S,ARGS...>,
+               cuda::launch_configuration_t(grid, block),
+               _cuboid,
+               mask.device(),
+               std::forward<ARGS>(args)...);
 }
 
 };
 #+END_SRC
 
-Generally speaking we want to avoid control flow branching inside of GPU kernels. However it can
-make sense to group multiple masked operators into a single lattice pass of a single kernel using some
-small number of branches instead of launching a separate branch-free kernel for each operator.
-
-In order to generate a single kernel from multiple operators curried with their respective arguments
-we implement a small =Operator= helper class.
-
-#+BEGIN_SRC cpp :noweb no :tangle tangle/LLBM/operator.h
-#pragma once
-
-#include <tuple>
-
-template <typename OPERATOR, typename... ARGS>
-struct Operator {
-  bool* const mask;
-  const std::tuple<ARGS...> config;
-
-  Operator(OPERATOR, DeviceBuffer<bool>& m, ARGS... args):
-    mask(m.device()),
-    config(args...) { }
-#+END_SRC
-
-This =Operator= wrapper stores both the mask pointer and any arguments for a given =OPERATOR= type.
-The arguments will be passed to the actual =OPERATOR::apply= by the following proxy method for
-any masked cell index.
-
-#+BEGIN_SRC cpp :noweb no :tangle tangle/LLBM/operator.h
-  template <typename DESCRIPTOR, typename T, typename S>
-  __device__ bool apply(DESCRIPTOR d, S f_curr[DESCRIPTOR::q], S f_next[DESCRIPTOR::q], std::size_t gid) const {
-    if (mask[gid]) {
-      std::apply([](auto... args) { OPERATOR::template apply<T,S>(args...); },
-                 std::tuple_cat(std::make_tuple(d, f_curr, f_next, gid), config));
-      return true;
-    } else {
-      return false;
-    }
-  }
-};
-#+END_SRC
-
-Finally a deduction guide can be used to remove some of the visual cruft of declaring =Operator= instances in
-the application code.
-
-#+BEGIN_SRC cpp :noweb no :tangle tangle/LLBM/operator.h
-template <typename OPERATOR, typename... ARGS>
-Operator(OPERATOR, DeviceBuffer<bool>&, ARGS... args) -> Operator<OPERATOR,std::remove_reference_t<ARGS>...>;
-#+END_SRC
-
 ** Kernel Execution
 #+BEGIN_SRC cpp :tangle tangle/LLBM/kernel/executor.h
 #pragma once
@@ -2758,30 +2769,6 @@ __global__ void call_operator(
 }
 #+END_SRC
 
-Besides applying an operator to each member of a cell list we also want the possibility
-of using a mask instead. Doing this reduces the required memory bandwidth for large
-ratios between masked and unmasked cells.
-
-#+BEGIN_SRC cpp :tangle tangle/LLBM/kernel/executor.h
-template <typename OPERATOR, typename DESCRIPTOR, typename T, typename S, typename... ARGS>
-__global__ void call_operator(
-    LatticeView<DESCRIPTOR,S> lattice
-  , bool* mask
-  , ARGS... args
-) {
-  const std::size_t gid = blockIdx.x * blockDim.x + threadIdx.x;
-  if (!(gid < lattice.cuboid.volume) || !mask[gid]) {
-      return;
-  }
-
-  S f_curr[DESCRIPTOR::q];
-  S f_next[DESCRIPTOR::q];
-  <<read-f-curr>>
-  OPERATOR::template apply<T,S>(DESCRIPTOR(), f_curr, f_next, gid, std::forward<ARGS>(args)...);
-  <<write-f-next>>
-}
-#+END_SRC
-
 Of course the same argument holds for read-only functors where we in most cases only want
 to compute moments for the bulk cells.
 
@@ -2847,25 +2834,6 @@ __global__ void call_operator_using_list(
 }
 #+END_SRC
 
-Even more general is the following caller kernel that doesn't assume anything about the
-lattice data required or written to by the wrapped operator. This is mostly used for easily
-calling post-processing kernels that do not directly operate on lattice data.
-
-#+BEGIN_SRC cpp :tangle tangle/LLBM/kernel/executor.h
-template <typename OPERATOR, typename DESCRIPTOR, typename T, typename S, typename... ARGS>
-__global__ void call_operator_using_list(
-    DESCRIPTOR descriptor
-  , std::size_t count
-  , ARGS... args
-) {
-  const std::size_t index = blockIdx.x * blockDim.x + threadIdx.x;
-  if (!(index < count)) {
-      return;
-  }
-  OPERATOR::template apply<T,S>(descriptor, index, count, std::forward<ARGS>(args)...);
-}
-#+END_SRC
-
 For some post-processing operators and functors it is convenient to apply them in
 a spatial setting where the threads are mapped in a 3D grid by CUDA. e.g. this is
 used when writing data into the 3D textures used by our ray marcher.
@@ -2959,7 +2927,7 @@ public:
   descriptor::Cuboid<DESCRIPTOR> cuboid() const {
     return _cuboid;
   };
- 
+
   <<cell-materials-basic-access>>
   <<cell-materials-fancy-setters>>
   <<cell-materials-get-list>>
@@ -5049,7 +5017,11 @@ using T = float;
 using DESCRIPTOR = descriptor::D2Q9;
 
 int main() {
-cudaSetDevice(0);
+if (cuda::device::count() == 0) {
+  std::cerr << "No CUDA devices on this system" << std::endl;
+  return -1;
+}
+auto current = cuda::device::current::get();
 #+END_SRC
 
 After including the relevant headers and declaring which floating point precision and descriptor
@@ -5127,7 +5099,7 @@ std::size_t iStep = 0;
 while (window.isOpen()) {
   <<ldc-simulation-step>>
   if (iStep % 100 == 0) {
-    cudaDeviceSynchronize();
+    cuda::synchronize(current);
     <<ldc-visualization-step>>
     window.draw([&]() {
       ImGui::Begin("Render");
@@ -5154,7 +5126,11 @@ using T = float;
 using DESCRIPTOR = descriptor::D3Q19;
 
 int main() {
-cudaSetDevice(0);
+if (cuda::device::count() == 0) {
+  std::cerr << "No CUDA devices on this system" << std::endl;
+  return -1;
+}
+auto current = cuda::device::current::get();
 #+END_SRC
 
 After including the relevant headers we construct the D3Q19 lattice.
@@ -5186,7 +5162,7 @@ auto bulk_mask = materials.mask_of_material(1);
 auto wall_mask = materials.mask_of_material(2);
 auto lid_mask  = materials.mask_of_material(3);
 
-cudaDeviceSynchronize();
+cuda::synchronize(current);
 #+END_SRC
 
 In order to model the flow we employ standard BGK collisions in the bulk. The side and bottom
@@ -5232,7 +5208,11 @@ using T = float;
 using DESCRIPTOR = descriptor::D2Q9;
 
 int main() {
-cudaSetDevice(0);
+if (cuda::device::count() == 0) {
+  std::cerr << "No CUDA devices on this system" << std::endl;
+  return -1;
+}
+auto current = cuda::device::current::get();
 #+END_SRC
 
 Adding another non-spinning cylinder allows for direct comparison and also produces a Kármán vortex street.
@@ -5316,7 +5296,7 @@ auto edge_mask = materials.mask_of_material(5);
 #+END_SRC
 
 #+BEGIN_SRC cpp :tangle tangle/magnus.cu
-cudaDeviceSynchronize();
+cuda::synchronize(current);
 #+END_SRC
 
 During each timestep we apply all purely local collision operators in a single fused
@@ -5377,7 +5357,7 @@ std::size_t iStep = 0;
 while (window.isOpen()) {
   <<magnus-simulation-step>>
   if (iStep % 200 == 0) {
-    cudaDeviceSynchronize();
+    cuda::synchronize(current);
     <<magnus-visualization-step>>
     window.draw([&]() {
       ImGui::Begin("Render");
@@ -5412,7 +5392,11 @@ using T = float;
 using DESCRIPTOR = descriptor::D3Q19;
 
 int main() {
-cudaSetDevice(0);
+if (cuda::device::count() == 0) {
+  std::cerr << "No CUDA devices on this system" << std::endl;
+  return -1;
+}
+auto current = cuda::device::current::get();
 #+END_SRC
 
 After including the relevant headers we construct the D3Q19 lattice.
@@ -5476,7 +5460,7 @@ auto inflow_mask  = materials.mask_of_material(4);
 auto outflow_mask = materials.mask_of_material(5);
 auto edge_mask    = materials.mask_of_material(6);
 
-cudaDeviceSynchronize();
+cuda::synchronize(current);
 #+END_SRC
 
 In order to model the flow we employ standard BGK collisions in the bulk. The channel
@@ -5530,7 +5514,11 @@ using T = float;
 using DESCRIPTOR = descriptor::D3Q19;
 
 int main() {
-cudaSetDevice(0);
+if (cuda::device::count() == 0) {
+  std::cerr << "No CUDA devices on this system" << std::endl;
+  return -1;
+}
+auto current = cuda::device::current::get();
 #+END_SRC
 
 After including the relevant headers we construct the D3Q19 lattice.
@@ -5593,7 +5581,7 @@ auto bulk_list = materials.list_of_material(1);
 auto wall_mask = materials.mask_of_material(2);
 auto wall_list = materials.list_of_material(2);
 
-cudaDeviceSynchronize();
+cuda::synchronize(current);
 #+END_SRC
 
 Due to the large quantity of inactive cells in this example we call the collision
@@ -5648,7 +5636,11 @@ using T = float;
 using DESCRIPTOR = descriptor::D3Q19;
 
 int main() {
-cudaSetDevice(0);
+if (cuda::device::count() == 0) {
+  std::cerr << "No CUDA devices on this system" << std::endl;
+  return -1;
+}
+auto current = cuda::device::current::get();
 #+END_SRC
 
 After including the relevant headers we construct the D3Q19 lattice.
@@ -5698,7 +5690,7 @@ auto boundary_mask = materials.mask_of_material(2);
 auto inflow_mask   = materials.mask_of_material(3);
 auto outflow_mask  = materials.mask_of_material(4);
 
-cudaDeviceSynchronize();
+cuda::synchronize(current);
 #+END_SRC
 
 In order to model the turbulence we employ Smagorinsky BGK collisions in the bulk.
@@ -5819,7 +5811,7 @@ auto bulk_mask = materials.mask_of_material(1);
 auto box_mask  = materials.mask_of_material(2);
 auto lid_mask  = materials.mask_of_material(3);
 
-cudaDeviceSynchronize();
+cuda::synchronize(current);
 #+END_SRC
 
 The simulation step consists of BGK collisions in the bulk and bounce back boundaries at the sides.
@@ -5837,7 +5829,7 @@ and the floating point type to be used for lattice data and computations.
 
 #+BEGIN_SRC cpp :tangle tangle/benchmark-ldc.cu
 void simulate(descriptor::Cuboid<DESCRIPTOR> cuboid, std::size_t nStep) {
-  cudaSetDevice(0);
+  auto current = cuda::device::current::get();
 
   <<benchmark-ldc-setup-lattice>>
 
@@ -5845,7 +5837,7 @@ void simulate(descriptor::Cuboid<DESCRIPTOR> cuboid, std::size_t nStep) {
     <<benchmark-ldc-simulation-step>>
   }
 
-  cudaDeviceSynchronize();
+  cuda::synchronize(current);
 
   auto start = timer::now();
 
@@ -5853,7 +5845,7 @@ void simulate(descriptor::Cuboid<DESCRIPTOR> cuboid, std::size_t nStep) {
     <<benchmark-ldc-simulation-step>>
   }
 
-  cudaDeviceSynchronize();
+  cuda::synchronize(current);
 
   auto mlups = timer::mlups(cuboid.volume, nStep, start);
 
@@ -5866,6 +5858,10 @@ some very basic CLI parameter handling in the =main= function.
 
 #+BEGIN_SRC cpp :tangle tangle/benchmark-ldc.cu
 int main(int argc, char* argv[]) {
+  if (cuda::device::count() == 0) {
+    std::cerr << "No CUDA devices on this system" << std::endl;
+    return -1;
+  }
   if (argc != 3) {
     std::cerr << "Invalid parameter count" << std::endl;
     return -1;
@@ -5918,7 +5914,6 @@ done
 :unnumbered: notoc
 :end:
 ** TODO Add more detailed explanations
-** TODO Add CUDA error handling
 ** TODO Expand example documentation
 ** TODO Add literature citations for method sections
 ** TODO Add physical dimensionalization
diff --git a/tangle/LLBM/kernel/executor.h b/tangle/LLBM/kernel/executor.h
index 942918d..cce023b 100644
--- a/tangle/LLBM/kernel/executor.h
+++ b/tangle/LLBM/kernel/executor.h
@@ -30,30 +30,6 @@ __global__ void call_operator(
   }
 }
 
-template <typename OPERATOR, typename DESCRIPTOR, typename T, typename S, typename... ARGS>
-__global__ void call_operator(
-    LatticeView<DESCRIPTOR,S> lattice
-  , bool* mask
-  , ARGS... args
-) {
-  const std::size_t gid = blockIdx.x * blockDim.x + threadIdx.x;
-  if (!(gid < lattice.cuboid.volume) || !mask[gid]) {
-      return;
-  }
-
-  S f_curr[DESCRIPTOR::q];
-  S f_next[DESCRIPTOR::q];
-  S* preshifted_f[DESCRIPTOR::q];
-  for (unsigned iPop=0; iPop < DESCRIPTOR::q; ++iPop) {
-    preshifted_f[iPop] = lattice.pop(iPop, gid);
-    f_curr[iPop] = *preshifted_f[iPop];
-  }
-  OPERATOR::template apply<T,S>(DESCRIPTOR(), f_curr, f_next, gid, std::forward<ARGS>(args)...);
-  for (unsigned iPop=0; iPop < DESCRIPTOR::q; ++iPop) {
-    *preshifted_f[iPop] = f_next[iPop];
-  }
-}
-
 template <typename FUNCTOR, typename DESCRIPTOR, typename T, typename S, typename... ARGS>
 __global__ void call_functor(
     LatticeView<DESCRIPTOR,S> lattice
@@ -110,19 +86,6 @@ __global__ void call_operator_using_list(
   OPERATOR::template apply<T,S>(lattice, index, count, std::forward<ARGS>(args)...);
 }
 
-template <typename OPERATOR, typename DESCRIPTOR, typename T, typename S, typename... ARGS>
-__global__ void call_operator_using_list(
-    DESCRIPTOR descriptor
-  , std::size_t count
-  , ARGS... args
-) {
-  const std::size_t index = blockIdx.x * blockDim.x + threadIdx.x;
-  if (!(index < count)) {
-      return;
-  }
-  OPERATOR::template apply<T,S>(descriptor, index, count, std::forward<ARGS>(args)...);
-}
-
 template <typename FUNCTOR, typename DESCRIPTOR, typename T, typename S, typename... ARGS>
 __global__ void call_spatial_functor(
     LatticeView<DESCRIPTOR,S> lattice
diff --git a/tangle/LLBM/lattice.h b/tangle/LLBM/lattice.h
index 7157c78..d3a1840 100644
--- a/tangle/LLBM/lattice.h
+++ b/tangle/LLBM/lattice.h
@@ -31,9 +31,10 @@ template <typename... OPERATOR>
 void apply(OPERATOR... ops) {
   const auto block_size = 32;
   const auto block_count = (_cuboid.volume + block_size - 1) / block_size;
-  kernel::call_operators<DESCRIPTOR,T,S,OPERATOR...><<<block_count,block_size>>>(
-    _population.view(), ops...
-  );
+  cuda::launch(kernel::call_operators<DESCRIPTOR,T,S,OPERATOR...>,
+               cuda::launch_configuration_t(block_count, block_size),
+               _population.view(),
+               ops...);
 }
 
 template <typename OPERATOR, typename... ARGS>
@@ -45,27 +46,22 @@ template <typename OPERATOR, typename... ARGS>
 void call_operator(tag::call_by_cell_id, DeviceBuffer<std::size_t>& cells, ARGS... args) {
   const auto block_size = 32;
   const auto block_count = (cells.size() + block_size - 1) / block_size;
-  kernel::call_operator<OPERATOR,DESCRIPTOR,T,S,ARGS...><<<block_count,block_size>>>(
-    _population.view(), cells.device(), cells.size(), std::forward<ARGS>(args)...
-  );
-}
-
-template <typename OPERATOR, typename... ARGS>
-void call_operator(tag::call_by_cell_id, DeviceBuffer<bool>& mask, ARGS... args) {
-  const auto block_size = 32;
-  const auto block_count = (_cuboid.volume + block_size - 1) / block_size;
-  kernel::call_operator<OPERATOR,DESCRIPTOR,T,S,ARGS...><<<block_count,block_size>>>(
-    _population.view(), mask.device(), std::forward<ARGS>(args)...
-  );
+  cuda::launch(kernel::call_operator<OPERATOR,DESCRIPTOR,T,S,ARGS...>,
+               cuda::launch_configuration_t(block_count, block_size),
+               _population.view(),
+               cells.device(), cells.size(),
+               std::forward<ARGS>(args)...);
 }
 
 template <typename OPERATOR, typename... ARGS>
 void call_operator(tag::call_by_list_index, std::size_t count, ARGS... args) {
   const auto block_size = 32;
   const auto block_count = (count + block_size - 1) / block_size;
-  kernel::call_operator_using_list<OPERATOR,DESCRIPTOR,T,S,ARGS...><<<block_count,block_size>>>(
-    _population.view(), count, std::forward<ARGS>(args)...
-  );
+  cuda::launch(kernel::call_operator_using_list<OPERATOR,DESCRIPTOR,T,S,ARGS...>,
+               cuda::launch_configuration_t(block_count, block_size),
+               _population.view(),
+               count,
+               std::forward<ARGS>(args)...);
 }
 
 template <typename FUNCTOR, typename... ARGS>
@@ -73,22 +69,15 @@ void inspect(ARGS&&... args) {
   call_functor<FUNCTOR>(typename FUNCTOR::call_tag{}, std::forward<ARGS&&>(args)...);
 }
 
-template <typename FUNCTOR, typename... ARGS>
-void call_functor(tag::call_by_cell_id, DeviceBuffer<std::size_t>& cells, ARGS... args) {
-  const auto block_size = 32;
-  const auto block_count = (cells.size() + block_size - 1) / block_size;
-  kernel::call_functor<FUNCTOR,DESCRIPTOR,T,S,ARGS...><<<block_count,block_size>>>(
-    _population.view(), cells.device(), cells.size(), std::forward<ARGS>(args)...
-  );
-}
-
 template <typename FUNCTOR, typename... ARGS>
 void call_functor(tag::call_by_cell_id, DeviceBuffer<bool>& mask, ARGS... args) {
   const auto block_size = 32;
   const auto block_count = (_cuboid.volume + block_size - 1) / block_size;
-  kernel::call_functor<FUNCTOR,DESCRIPTOR,T,S,ARGS...><<<block_count,block_size>>>(
-    _population.view(), mask.device(), std::forward<ARGS>(args)...
-  );
+  cuda::launch(kernel::call_functor<FUNCTOR,DESCRIPTOR,T,S,ARGS...>,
+               cuda::launch_configuration_t(block_count, block_size),
+               _population.view(),
+               mask.device(),
+               std::forward<ARGS>(args)...);
 }
 
 template <typename FUNCTOR, typename... ARGS>
@@ -97,9 +86,11 @@ void call_functor(tag::call_by_spatial_cell_mask, DeviceBuffer<bool>& mask, ARGS
   const dim3 grid((_cuboid.nX + block.x - 1) / block.x,
                   (_cuboid.nY + block.y - 1) / block.y,
                   (_cuboid.nZ + block.z - 1) / block.z);
-  kernel::call_spatial_functor<FUNCTOR,DESCRIPTOR,T,S,ARGS...><<<grid,block>>>(
-    _population.view(), mask.device(), std::forward<ARGS>(args)...
-  );
+  cuda::launch(kernel::call_spatial_functor<FUNCTOR,DESCRIPTOR,T,S,ARGS...>,
+               cuda::launch_configuration_t(grid, block),
+               _population.view(),
+               mask.device(),
+               std::forward<ARGS>(args)...);
 }
 
 template <typename OPERATOR, typename... ARGS>
@@ -111,9 +102,11 @@ template <typename OPERATOR, typename... ARGS>
 void tagged_helper(tag::post_process_by_list_index, std::size_t count, ARGS... args) {
   const auto block_size = 32;
   const auto block_count = (count + block_size - 1) / block_size;
-  kernel::call_operator_using_list<OPERATOR,DESCRIPTOR,T,S,ARGS...><<<block_count,block_size>>>(
-    DESCRIPTOR(), count, std::forward<ARGS>(args)...
-  );
+  cuda::launch(kernel::call_operator_using_list<OPERATOR,DESCRIPTOR,T,S,ARGS...>,
+               cuda::launch_configuration_t(block_count, block_size),
+               DESCRIPTOR(),
+               count,
+               std::forward<ARGS>(args)...);
 }
 
 template <typename OPERATOR, typename... ARGS>
@@ -122,9 +115,11 @@ void tagged_helper(tag::post_process_by_spatial_cell_mask, DeviceBuffer<bool>& m
   const dim3 grid((_cuboid.nX + block.x - 1) / block.x,
                   (_cuboid.nY + block.y - 1) / block.y,
                   (_cuboid.nZ + block.z - 1) / block.z);
-  kernel::call_spatial_operator<OPERATOR,DESCRIPTOR,T,S,ARGS...><<<grid,block>>>(
-    _cuboid, mask.device(), std::forward<ARGS>(args)...
-  );
+  cuda::launch(kernel::call_spatial_operator<OPERATOR,DESCRIPTOR,T,S,ARGS...>,
+               cuda::launch_configuration_t(grid, block),
+               _cuboid,
+               mask.device(),
+               std::forward<ARGS>(args)...);
 }
 
 };
diff --git a/tangle/LLBM/materials.h b/tangle/LLBM/materials.h
index d782d8d..5734688 100644
--- a/tangle/LLBM/materials.h
+++ b/tangle/LLBM/materials.h
@@ -24,7 +24,7 @@ public:
   descriptor::Cuboid<DESCRIPTOR> cuboid() const {
     return _cuboid;
   };
- 
+
   int get(std::size_t iCell) const {
     return _materials[iCell];
   }
diff --git a/tangle/LLBM/memory.h b/tangle/LLBM/memory.h
index 97fec5c..149034b 100644
--- a/tangle/LLBM/memory.h
+++ b/tangle/LLBM/memory.h
@@ -4,31 +4,30 @@
 #include <vector>
 #include <cstring>
 
+#include <cuda/runtime_api.hpp>
+
 template <typename T>
 class DeviceBuffer {
 protected:
   const std::size_t _size;
-  T* _data;
+  cuda::device_t _device;
+  cuda::memory::device::unique_ptr<T[]> _data;
 
 public:
   DeviceBuffer(std::size_t size):
-    _size(size) {
-    cudaMalloc(&_data, _size*sizeof(T));
-    cudaMemset(_data, 0, _size*sizeof(T));
-  }
+    _size(size),
+    _device(cuda::device::current::get()),
+    _data(cuda::memory::device::make_unique<T[]>(_device, size))
+  { }
   DeviceBuffer(const T* data, std::size_t size):
     DeviceBuffer(size) {
-    cudaMemcpy(_data, data, size*sizeof(T), cudaMemcpyHostToDevice);
+    cuda::memory::copy(_data.get(), data, size*sizeof(T));
   }
   DeviceBuffer(const std::vector<T>& data):
     DeviceBuffer(data.data(), data.size()) { }
- 
-  ~DeviceBuffer() {
-    cudaFree(_data);
-  }
 
   T* device() {
-    return _data;
+    return _data.get();
   }
 
   std::size_t size() const {
@@ -40,7 +39,7 @@ template <typename T>
 class SharedVector : public DeviceBuffer<T> {
 private:
   std::unique_ptr<T[]> _host_data;
- 
+
 public:
   SharedVector(std::size_t size):
     DeviceBuffer<T>(size),
@@ -57,11 +56,11 @@ public:
   }
 
   void syncHostFromDevice() {
-    cudaMemcpy(_host_data.get(), this->_data, this->_size*sizeof(T), cudaMemcpyDeviceToHost);
+    cuda::memory::copy(_host_data.get(), this->_data.get(), this->_size*sizeof(T));
   }
 
   void syncDeviceFromHost() {
-    cudaMemcpy(this->_data, _host_data.get(), this->_size*sizeof(T), cudaMemcpyHostToDevice);
+    cuda::memory::copy(this->_data.get(), _host_data.get(), this->_size*sizeof(T));
   }
 
 };
@@ -103,7 +102,7 @@ public:
 
   DeviceTexture(descriptor::CuboidD<3> c):
     DeviceTexture(c.nX, c.nY, c.nZ) { }
-  
+
   ~DeviceTexture() {
     cudaFreeArray(_array);
   }
diff --git a/tangle/LLBM/propagate.h b/tangle/LLBM/propagate.h
index acb1d6c..91b1131 100644
--- a/tangle/LLBM/propagate.h
+++ b/tangle/LLBM/propagate.h
@@ -5,6 +5,7 @@
 #include "kernel/propagate.h"
 
 #include <cuda.h>
+#include <cuda/runtime_api.hpp>
 
 template <typename DESCRIPTOR, typename S>
 struct LatticeView {
@@ -37,6 +38,7 @@ protected:
 
 public:
   CyclicPopulationBuffer(descriptor::Cuboid<DESCRIPTOR> cuboid);
+  ~CyclicPopulationBuffer();
 
   LatticeView<DESCRIPTOR,S> view() {
     return LatticeView<DESCRIPTOR,S>{ _cuboid, _population.device() };
@@ -46,15 +48,13 @@ public:
 
 };
 
-std::size_t getDevicePageSize(int device_id=-1) {
-  if (device_id == -1) {
-    cudaGetDevice(&device_id);
-  }
+std::size_t getDevicePageSize() {
+  auto device = cuda::device::current::get();
   std::size_t granularity = 0;
   CUmemAllocationProp prop = {};
   prop.type = CU_MEM_ALLOCATION_TYPE_PINNED;
   prop.location.type = CU_MEM_LOCATION_TYPE_DEVICE;
-  prop.location.id = device_id;
+  prop.location.id = device.id();
   cuMemGetAllocationGranularity(&granularity, &prop, CU_MEM_ALLOC_GRANULARITY_MINIMUM);
   return granularity;
 }
@@ -69,12 +69,11 @@ CyclicPopulationBuffer<DESCRIPTOR,S>::CyclicPopulationBuffer(
   _population(DESCRIPTOR::q)
 {
 
-  int device_id = -1;
-  cudaGetDevice(&device_id);
+  auto device = cuda::device::current::get();
 
   _prop.type = CU_MEM_ALLOCATION_TYPE_PINNED;
   _prop.location.type = CU_MEM_LOCATION_TYPE_DEVICE;
-  _prop.location.id = device_id;
+  _prop.location.id = device.id();
   cuMemAddressReserve(&_ptr, 2 * _volume * DESCRIPTOR::q, 0, 0, 0);
 
   for (unsigned iPop=0; iPop < DESCRIPTOR::q; ++iPop) {
@@ -85,7 +84,7 @@ CyclicPopulationBuffer<DESCRIPTOR,S>::CyclicPopulationBuffer(
   }
 
   _access.location.type = CU_MEM_LOCATION_TYPE_DEVICE;
-  _access.location.id = 0;
+  _access.location.id = device.id();
   _access.flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE;
   cuMemSetAccess(_ptr, 2 * _volume * DESCRIPTOR::q, &_access, 1);
 
@@ -95,7 +94,7 @@ CyclicPopulationBuffer<DESCRIPTOR,S>::CyclicPopulationBuffer(
   }
 
   for (unsigned iPop=0; iPop < DESCRIPTOR::q; ++iPop) {
-    _base[iPop] = device() + iPop * 2 * (_volume / sizeof(S));
+    _base[iPop] = this->device() + iPop * 2 * (_volume / sizeof(S));
     _population[iPop] = _base[iPop] + iPop * ((_volume / sizeof(S)) / DESCRIPTOR::q);
   }
 
@@ -103,6 +102,15 @@ CyclicPopulationBuffer<DESCRIPTOR,S>::CyclicPopulationBuffer(
   _population.syncDeviceFromHost();
 }
 
+template <typename DESCRIPTOR, typename S>
+CyclicPopulationBuffer<DESCRIPTOR,S>::~CyclicPopulationBuffer() {
+  cuMemUnmap(_ptr, 2 * _volume * DESCRIPTOR::q);
+  for (unsigned iPop=0; iPop < DESCRIPTOR::q; ++iPop) {
+    cuMemRelease(_handle[iPop]);
+  }
+  cuMemAddressFree(_ptr, 2 * _volume * DESCRIPTOR::q);
+}
+
 template <typename DESCRIPTOR, typename S>
 __device__ __forceinline__
 S* LatticeView<DESCRIPTOR,S>::pop(pop_index_t iPop, std::size_t gid) const {
@@ -111,5 +119,7 @@ S* LatticeView<DESCRIPTOR,S>::pop(pop_index_t iPop, std::size_t gid) const {
 
 template <typename DESCRIPTOR, typename S>
 void CyclicPopulationBuffer<DESCRIPTOR,S>::stream() {
-  propagate<DESCRIPTOR,S><<<1,1>>>(view(), _base.device(), _volume / sizeof(S));
+  cuda::launch(propagate<DESCRIPTOR,S>,
+               cuda::launch_configuration_t(1,1),
+               view(), _base.device(), _volume / sizeof(S));
 }
diff --git a/tangle/benchmark-ldc.cu b/tangle/benchmark-ldc.cu
index 1400547..18b30d1 100644
--- a/tangle/benchmark-ldc.cu
+++ b/tangle/benchmark-ldc.cu
@@ -12,7 +12,7 @@ using T = float;
 using DESCRIPTOR = descriptor::D3Q19;
 
 void simulate(descriptor::Cuboid<DESCRIPTOR> cuboid, std::size_t nStep) {
-  cudaSetDevice(0);
+  auto current = cuda::device::current::get();
 
   Lattice<DESCRIPTOR,T> lattice(cuboid);
   
@@ -30,7 +30,7 @@ void simulate(descriptor::Cuboid<DESCRIPTOR> cuboid, std::size_t nStep) {
   auto box_mask  = materials.mask_of_material(2);
   auto lid_mask  = materials.mask_of_material(3);
   
-  cudaDeviceSynchronize();
+  cuda::synchronize(current);
 
   for (std::size_t iStep=0; iStep < 100; ++iStep) {
     lattice.apply(Operator(BgkCollideO(), bulk_mask, 0.56),
@@ -39,7 +39,7 @@ void simulate(descriptor::Cuboid<DESCRIPTOR> cuboid, std::size_t nStep) {
     lattice.stream();
   }
 
-  cudaDeviceSynchronize();
+  cuda::synchronize(current);
 
   auto start = timer::now();
 
@@ -50,7 +50,7 @@ void simulate(descriptor::Cuboid<DESCRIPTOR> cuboid, std::size_t nStep) {
     lattice.stream();
   }
 
-  cudaDeviceSynchronize();
+  cuda::synchronize(current);
 
   auto mlups = timer::mlups(cuboid.volume, nStep, start);
 
@@ -58,6 +58,10 @@ void simulate(descriptor::Cuboid<DESCRIPTOR> cuboid, std::size_t nStep) {
 }
 
 int main(int argc, char* argv[]) {
+  if (cuda::device::count() == 0) {
+    std::cerr << "No CUDA devices on this system" << std::endl;
+    return -1;
+  }
   if (argc != 3) {
     std::cerr << "Invalid parameter count" << std::endl;
     return -1;
diff --git a/tangle/channel-with-sphere.cu b/tangle/channel-with-sphere.cu
index 29cc7de..280a142 100644
--- a/tangle/channel-with-sphere.cu
+++ b/tangle/channel-with-sphere.cu
@@ -15,7 +15,11 @@ using T = float;
 using DESCRIPTOR = descriptor::D3Q19;
 
 int main() {
-cudaSetDevice(0);
+if (cuda::device::count() == 0) {
+  std::cerr << "No CUDA devices on this system" << std::endl;
+  return -1;
+}
+auto current = cuda::device::current::get();
 
 const descriptor::Cuboid<DESCRIPTOR> cuboid(448, 64, 64);
 Lattice<DESCRIPTOR,T> lattice(cuboid);
@@ -55,7 +59,7 @@ auto inflow_mask  = materials.mask_of_material(4);
 auto outflow_mask = materials.mask_of_material(5);
 auto edge_mask    = materials.mask_of_material(6);
 
-cudaDeviceSynchronize();
+cuda::synchronize(current);
 
 VolumetricExample renderer(cuboid);
 renderer.add<QCriterionS>(lattice, bulk_mask, obstacle);
diff --git a/tangle/ldc-2d.cu b/tangle/ldc-2d.cu
index acba98f..5fa36f9 100644
--- a/tangle/ldc-2d.cu
+++ b/tangle/ldc-2d.cu
@@ -13,7 +13,11 @@ using T = float;
 using DESCRIPTOR = descriptor::D2Q9;
 
 int main() {
-cudaSetDevice(0);
+if (cuda::device::count() == 0) {
+  std::cerr << "No CUDA devices on this system" << std::endl;
+  return -1;
+}
+auto current = cuda::device::current::get();
 
 const descriptor::Cuboid<DESCRIPTOR> cuboid(500, 500);
 Lattice<DESCRIPTOR,T> lattice(cuboid);
@@ -52,7 +56,7 @@ while (window.isOpen()) {
                 Operator(BounceBackMovingWallO(), lid_mask, std::min(iStep*1e-3, 1.0)*u_lid, 0.f));
   lattice.stream();
   if (iStep % 100 == 0) {
-    cudaDeviceSynchronize();
+    cuda::synchronize(current);
     lattice.inspect<CollectMomentsF>(bulk_mask, moments_rho.device(), moments_u.device());
     renderSliceViewToTexture<<<
       dim3(cuboid.nX / 32 + 1, cuboid.nY / 32 + 1),
diff --git a/tangle/ldc-3d.cu b/tangle/ldc-3d.cu
index e9b42f2..c70b12e 100644
--- a/tangle/ldc-3d.cu
+++ b/tangle/ldc-3d.cu
@@ -15,7 +15,11 @@ using T = float;
 using DESCRIPTOR = descriptor::D3Q19;
 
 int main() {
-cudaSetDevice(0);
+if (cuda::device::count() == 0) {
+  std::cerr << "No CUDA devices on this system" << std::endl;
+  return -1;
+}
+auto current = cuda::device::current::get();
 
 const descriptor::Cuboid<DESCRIPTOR> cuboid(100, 100, 100);
 Lattice<DESCRIPTOR,T> lattice(cuboid);
@@ -34,7 +38,7 @@ auto bulk_mask = materials.mask_of_material(1);
 auto wall_mask = materials.mask_of_material(2);
 auto lid_mask  = materials.mask_of_material(3);
 
-cudaDeviceSynchronize();
+cuda::synchronize(current);
 
 auto none = [] __device__ (float3) -> float { return 1; };
 VolumetricExample renderer(cuboid);
diff --git a/tangle/magnus.cu b/tangle/magnus.cu
index 08a4515..a98bcf3 100644
--- a/tangle/magnus.cu
+++ b/tangle/magnus.cu
@@ -13,7 +13,11 @@ using T = float;
 using DESCRIPTOR = descriptor::D2Q9;
 
 int main() {
-cudaSetDevice(0);
+if (cuda::device::count() == 0) {
+  std::cerr << "No CUDA devices on this system" << std::endl;
+  return -1;
+}
+auto current = cuda::device::current::get();
 
 const descriptor::Cuboid<DESCRIPTOR> cuboid(1200, 500);
 Lattice<DESCRIPTOR,T> lattice(cuboid);
@@ -64,7 +68,7 @@ auto inflow_mask  = materials.mask_of_material(3);
 auto outflow_mask = materials.mask_of_material(4);
 auto edge_mask = materials.mask_of_material(5);
 
-cudaDeviceSynchronize();
+cuda::synchronize(current);
 
 RenderWindow window("Magnus");
 cudaSurfaceObject_t colormap;
@@ -83,7 +87,7 @@ while (window.isOpen()) {
   lattice.apply<BouzidiO>(bouzidi.getCount(), bouzidi.getConfig());
   lattice.stream();
   if (iStep % 200 == 0) {
-    cudaDeviceSynchronize();
+    cuda::synchronize(current);
     lattice.inspect<CollectMomentsF>(bulk_mask, moments_rho.device(), moments_u.device());
     renderSliceViewToTexture<<<
       dim3(cuboid.nX / 32 + 1, cuboid.nY / 32 + 1),
diff --git a/tangle/nozzle.cu b/tangle/nozzle.cu
index 03c18f9..5278096 100644
--- a/tangle/nozzle.cu
+++ b/tangle/nozzle.cu
@@ -15,7 +15,11 @@ using T = float;
 using DESCRIPTOR = descriptor::D3Q19;
 
 int main() {
-cudaSetDevice(0);
+if (cuda::device::count() == 0) {
+  std::cerr << "No CUDA devices on this system" << std::endl;
+  return -1;
+}
+auto current = cuda::device::current::get();
 
 const descriptor::Cuboid<DESCRIPTOR> cuboid(500, 80, 80);
 Lattice<DESCRIPTOR,T> lattice(cuboid);
@@ -46,7 +50,7 @@ auto boundary_mask = materials.mask_of_material(2);
 auto inflow_mask   = materials.mask_of_material(3);
 auto outflow_mask  = materials.mask_of_material(4);
 
-cudaDeviceSynchronize();
+cuda::synchronize(current);
 
 VolumetricExample renderer(cuboid);
 renderer.add<CurlNormS>(lattice, bulk_mask, obstacle);
diff --git a/tangle/taylor-couette.cu b/tangle/taylor-couette.cu
index 2e69bfb..6b39afa 100644
--- a/tangle/taylor-couette.cu
+++ b/tangle/taylor-couette.cu
@@ -15,7 +15,11 @@ using T = float;
 using DESCRIPTOR = descriptor::D3Q19;
 
 int main() {
-cudaSetDevice(0);
+if (cuda::device::count() == 0) {
+  std::cerr << "No CUDA devices on this system" << std::endl;
+  return -1;
+}
+auto current = cuda::device::current::get();
 
 const descriptor::Cuboid<DESCRIPTOR> cuboid(500, 96, 96);
 Lattice<DESCRIPTOR,T> lattice(cuboid);
@@ -55,7 +59,7 @@ auto bulk_list = materials.list_of_material(1);
 auto wall_mask = materials.mask_of_material(2);
 auto wall_list = materials.list_of_material(2);
 
-cudaDeviceSynchronize();
+cuda::synchronize(current);
 
 VolumetricExample renderer(cuboid);
 renderer.add<VelocityNormS>(lattice, bulk_mask, inner_cylinder);
-- 
cgit v1.2.3