[PATCH] D69818: [HIP] Fix pointer type kernel arg for amdgpu

Artem Belevich via Phabricator via cfe-commits cfe-commits at lists.llvm.org
Mon Nov 4 13:51:24 PST 2019


tra added inline comments.


================
Comment at: clang/lib/CodeGen/CGCall.cpp:1172
     if (isa<llvm::PointerType>(Ty))
-      return CGF.Builder.CreateBitCast(Val, Ty, "coerce.val");
+      return CGF.Builder.CreatePointerCast(Val, Ty, "coerce.val");
 
----------------
What's the purpose of this change?


================
Comment at: clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu:15-16
+// CHECK: define amdgpu_kernel void  @_Z7kernel3PU3AS2iPU3AS1i(i32 addrspace(2)* %x, i32 addrspace(1)* %y)
+__global__ void kernel3(__attribute__((address_space(2))) int *x,
+                        __attribute__((address_space(1))) int *y) {
+  y[0] = x[0];
----------------
Interesting. Clang used to crash on explicit address space attribute in the past. I'm glad it works now.


CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D69818/new/

https://reviews.llvm.org/D69818





More information about the cfe-commits mailing list