[CUDA, NVPTX] LLVM crash with "Cannot cast between two non-generic address spaces"

Valid CUDA code results in LLVM crash due to an attempt to generate an impossible addrspacecast in a known-false conditional branch.

Reproducer: https://godbolt.org/z/Yjjsdvj1r

Source:

#include <stdint.h>

// No crash if the function is not inlined, and thus does not know 
// at compile time which pointer it will be handling
//__device__ uintptr_t f(void *p) __noinline__;

__device__ uintptr_t f(void *p) {
   if (__isGlobal(p))
    return __cvta_generic_to_global(p);
   if (__isShared(p))
     return __cvta_generic_to_shared(p);
   return (uintptr_t)p;
}

__shared__ int shared_data;
__device__ int global_data;
__constant__ int const_data = 3;

__global__ void square(uintptr_t* out, int n) {
  out[0] = f(&shared_data);
  out[1] = f(&global_data);
}

IR:

define dso_local void @square(unsigned long*, int)(ptr nocapture noundef writeonly %out, i32 noundef %n) local_unnamed_addr #1 {
entry:
  %0 = tail call i1 @llvm.nvvm.isspacep.global(ptr addrspacecast (ptr addrspace(3) @shared_data to ptr))
  %1 = tail call i1 @llvm.nvvm.isspacep.shared(ptr addrspacecast (ptr addrspace(3) @shared_data to ptr))
  %. = select i1 %1, i64 ptrtoint (ptr addrspace(3) @shared_data to i64), i64 ptrtoint (ptr addrspacecast (ptr addrspace(3) @shared_data to ptr) to i64)
  %retval.0.i = select i1 %0, i64 ptrtoint (ptr addrspace(1) addrspacecast (ptr addrspace(3) @shared_data to ptr addrspace(1)) to i64), i64 %.
  store i64 %retval.0.i, ptr %out, align 8, !tbaa !8
  ret void
}

The culprit is addrspacecast (ptr addrspace(3) @shared_data to ptr addrspace(1)) here:

  %retval.0.i = select i1 %0, i64 ptrtoint (ptr addrspace(1) addrspacecast (ptr addrspace(3) @shared_data to ptr addrspace(1)) to i64),

To add insult to injury, the impossible cast is still going to be executed, and will likely result in a runtime error trying to convert the pointer in the wrong address space.

We need to make sure that __cvta_generic_to_global() is never executed if __isGlobal() is false. Same for the conversions from shared and constant AS.

llvm / llvm-project

[CUDA, NVPTX] LLVM crash with "Cannot cast between two non-generic address spaces" #112760