Envisioning `__ignore__` support in NVCC with a simple example.

ignore_on_device.md

How crazy is it to imagine a keyword (NVCC-supported), something like __ignore__, where if you use that in front of an expression (function, variable, object, etc.), it is ignored on the device side (in __device__ and __global__). This solves the issue where complicated containers that support host and device code, and their constructors/destructors that run on host code are all just ignored on device when they are passed as a member of larger class or struct. For example;

__global__ void kernel(foo_t foo) {
  auto idx = threadIdx.x;
  auto ptr = foo.get_ptr();
  ptr[idx] = idx;
}

class foo_t {
public:
  __host__ void resize(std::size_t n) {
    m_vector.resize(n);
  }

  __host__ void set_ptr() {
    m_raw_ptr = m_vector.data().get();
  }
  
  __device__ void get_ptr() {
    // add checks for null_ptr
    return m_raw_ptr;
  }

private:
  __ignore__ thrust::device_vector<int> m_vector;
  int * m_raw_ptr;
}

int main(int argc, char** argv) {
  std::size_t n = 10;
  foo_t foo;
  foo.resize(n);
  foo.set_ptr();
  kernel<<<1,n>>>(foo);
  cudaDeviceSynchronize();
}