Skip to content

Instantly share code, notes, and snippets.

Last active January 17, 2022 19:45
  • Star 0 You must be signed in to star a gist
  • Fork 0 You must be signed in to fork a gist
Star You must be signed in to star a gist
Save neoblizz/0a7dcebac76ab6c703d502b70e18a2e2 to your computer and use it in GitHub Desktop.
Envisioning `__ignore__` support in NVCC with a simple example.

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 {
  __host__ void resize(std::size_t n) {

  __host__ void set_ptr() {
    m_raw_ptr =;
  __device__ void get_ptr() {
    // add checks for null_ptr
    return m_raw_ptr;

  __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;
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment