audit every kernel argument in `include/nvexec` for accessibility from device

[ericniebler]

also do the inverse: make sure references or pointers to device memory are never accessed by the host.

[gevtushenko]

Sketch of a helper function:

#include <cuda_runtime.h>
#include <optional>
#include <iostream>

std::optional<bool> is_device_accessible(void* ptr) {
  struct cudaPointerAttributes attributes;

  if (cudaError_t err = cudaPointerGetAttributes(&attributes, ptr); err != cudaSuccess)  {
    return {};
  }

  return attributes.type == cudaMemoryTypeDevice || attributes.type == cudaMemoryTypeManaged ||
         attributes.type == cudaMemoryTypeHost;
}

int main() {
  char *managed_ptr{};
  char *pinned_ptr{};
  char *device_ptr{};
  char *host_ptr_1{};
  char host_ptr_2[1024];

  cudaMallocManaged(&managed_ptr, 1024);
  cudaMallocHost(&pinned_ptr, 1024);
  cudaMalloc(&device_ptr, 1024);
  host_ptr_1 = new char[1024];

  for (int offset = 0; offset < 1024; offset++)
  {
    if (*is_device_accessible(managed_ptr + offset) != true) {
      std::cerr << "managed error" << std::endl;
    }
    if (*is_device_accessible(pinned_ptr + offset) != true) {
      std::cerr << "pinned error" << std::endl;
    }
    if (*is_device_accessible(device_ptr + offset) != true) {
      std::cerr << "device error" << std::endl;
    }
    if (*is_device_accessible(host_ptr_1 + offset) == true) {
      std::cerr << "host 1 error" << std::endl;
    }
    if (*is_device_accessible(host_ptr_2 + offset) == true) {
      std::cerr << "host 2 error" << std::endl;
    }
  }
}