[CIR][CUDA] Fix address space values for NVPTX (#1445)

Based on
https://github.com/llvm/clangir/blob/7f66a204c4ba1f674cfe0e16e2c9c6b65ca70bc8/clang/lib/Basic/Targets/NVPTX.h#L27,
the current address space values are incorrect. This PR fixes these
values.

Author: anominos

clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/NVPTX.cpp

@@ -52,9 +52,9 @@ class NVPTXTargetLoweringInfo : public TargetLoweringInfo {
     case Kind::offload_global:
       return 1;
     case Kind::offload_constant:
-      return 2;
-    case Kind::offload_generic:
       return 4;
+    case Kind::offload_generic:
+      return 0;
     default:
       cir_cconv_unreachable("Unknown CIR address space for this target");
     }

clang/test/CIR/CodeGen/CUDA/addrspace-lowering.cu

@@ -0,0 +1,19 @@
+#include "../Inputs/cuda.h"
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fclangir \
+// RUN:            -fcuda-is-device -emit-llvm -target-sdk-version=12.3 \
+// RUN:            %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM-DEVICE --input-file=%t.ll %s
+
+
+__shared__ int a;
+
+// LLVM-DEVICE: @a = addrspace(3) {{.*}}
+
+__device__ int b;
+
+// LLVM-DEVICE: @b = addrspace(1) {{.*}}
+
+// __constant__ int c;
+
+// XFAIL-LLVM-DEVICE: @c = addrspace(4) {{.*}}