clang++ cuda: Lambda capture fails to initialize memory when variable only used in #ifdef __CUDA_ARCH__

[rgreenblatt]

template <typename F> __global__ void call(F f) { f(); }

int main(int argc, char *argv[]) {
  int a = 0;

  call<<<1, 1>>>([=] __device__() {
  // Uncomment line below for code to work
  // (void)a;

#ifdef __CUDA_ARCH__
    printf("%d\n", a);
#endif
  });

  cudaDeviceSynchronize();

  return 0;
}

Build with clang++ -std=c++17 --cuda-gpu-arch=sm_75 -L/usr/local/cuda/lib64 -lcudart test.cu changing options as needed. The issue also occurs on c++11/14. The value printed is garbage instead of 0. Using the variable outside the ifdef or explicity capturing makes the issue go away. The issue occurs regardless of optimization settings as far as I can tell.

My testing is on trunk, but I think this still occurs with clang 9/10

[zygoloid]

On the host side, the lambda does not capture a, so the lambda has a different representation between host and device. This is a bug in your code, just as if you had

struct Lambda {
#ifdef __CUDA_ARCH__
  int a;
#endif
};

and shared the type Lambda between host and device.

Workaround: either don't use #ifdef __CUDA_ARCH__ (at all, it can lead to lots of problems), or in this specific case, don't use implicit lambda captures ([=]) and use an explicit [a] capture instead.

[rgreenblatt]

Thanks for the help. Didn't really think this through.

[llvmbot]

@llvm/issue-subscribers-c-11

Author: Ryan Greenblatt (rgreenblatt)

```cuda template <typename F> __global__ void call(F f) { f(); }

int main(int argc, char *argv[]) {
int a = 0;

call<<<1, 1>>>([=] device() {
// Uncomment line below for code to work
// (void)a;

#ifdef CUDA_ARCH
printf("%d\n", a);
#endif
});

cudaDeviceSynchronize();

return 0;
}

Build with `clang++ -std=c++17 --cuda-gpu-arch=sm_75 -L/usr/local/cuda/lib64 -lcudart test.cu` changing options as needed. The issue also occurs on c++11/14. The value printed is garbage instead of 0. Using the variable outside the ifdef or explicity capturing makes the issue go away. The issue occurs regardless of optimization settings as far as I can tell.

My testing is on trunk, but I think this still occurs with clang 9/10
</details>