[PATCH] D111734: [HIP] Relax conditions for address space cast in builtin args
Anshil Gandhi via Phabricator via cfe-commits
cfe-commits at lists.llvm.org
Wed Oct 13 09:55:44 PDT 2021
gandhi21299 updated this revision to Diff 379444.
gandhi21299 marked an inline comment as done.
gandhi21299 added a comment.
adding codegen test
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D111734/new/
https://reviews.llvm.org/D111734
Files:
clang/lib/Sema/SemaExpr.cpp
clang/test/CodeGenCUDA/builtins-unsafe-atomics-gfx90a.cu
Index: clang/test/CodeGenCUDA/builtins-unsafe-atomics-gfx90a.cu
===================================================================
--- /dev/null
+++ clang/test/CodeGenCUDA/builtins-unsafe-atomics-gfx90a.cu
@@ -0,0 +1,20 @@
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx90a -x hip \
+// RUN: -aux-triple x86_64-unknown-linux-gnu -fcuda-is-device -emit-llvm %s \
+// RUN: -o - | FileCheck %s
+
+#define __device__ __attribute__((device))
+typedef __attribute__((address_space(3))) float *LP;
+
+// CHECK-LABEL: test_ds_atomic_add_f32
+// CHECK: %[[ADDR_ADDR:.*]] = alloca float*, align 8, addrspace(5)
+// CHECK: %[[ADDR_ADDR_ASCAST_PTR:.*]] = addrspacecast float* addrspace(5)* %[[ADDR_ADDR]] to float**
+// CHECK: store float* %addr, float** %[[ADDR_ADDR_ASCAST_PTR]], align 8
+// CHECK: %[[ADDR_ADDR_ASCAST:.*]] = load float*, float** %[[ADDR_ADDR_ASCAST_PTR]], align 8
+// CHECK: %[[AS_CAST:.*]] = addrspacecast float* %[[ADDR_ADDR_ASCAST]] to float addrspace(3)*
+// CHECK: %3 = call contract float @llvm.amdgcn.ds.fadd.f32(float addrspace(3)* %[[AS_CAST]]
+// CHECK: %4 = load float*, float** %rtn.ascast, align 8
+// CHECK: store float %3, float* %4, align 4
+__device__ void test_ds_atomic_add_f32(float *addr, float val) {
+ float *rtn;
+ *rtn = __builtin_amdgcn_ds_faddf((LP)addr, val, 0, 0, 0);
+}
Index: clang/lib/Sema/SemaExpr.cpp
===================================================================
--- clang/lib/Sema/SemaExpr.cpp
+++ clang/lib/Sema/SemaExpr.cpp
@@ -6545,9 +6545,14 @@
auto ArgPtTy = ArgTy->getPointeeType();
auto ArgAS = ArgPtTy.getAddressSpace();
- // Only allow implicit casting from a non-default address space pointee
- // type to a default address space pointee type
- if (ArgAS != LangAS::Default || ParamAS == LangAS::Default)
+ // Only allow implicit casting when either the argument has a pointee in
+ // a non-default address space and the target address spaces of the argument
+ // pointee is different from the target address space of the pointee of the
+ // parameter, or the parameter is in the default address space.
+ if ((ArgAS != LangAS::Default &&
+ getASTContext().getTargetAddressSpace(ArgAS) !=
+ getASTContext().getTargetAddressSpace(ParamAS)) ||
+ ParamAS == LangAS::Default)
continue;
// First, ensure that the Arg is an RValue.
-------------- next part --------------
A non-text attachment was scrubbed...
Name: D111734.379444.patch
Type: text/x-patch
Size: 2432 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20211013/770e4641/attachment.bin>
More information about the cfe-commits
mailing list