[HIP] Relax conditions for address space cast in builtin args

[gandhi21299]

Summary

Allow (implicit) address space casting between LLVM-equivalent
target address spaces.

[gandhi21299]

removed irrelevant lines in the test

[yaxunl]

we need a codegen test for the case

clang/lib/Sema/SemaExpr.cpp

6550

we still need to keep this condition so it should be:

if ((ArgAS != LangAS::Default &&
        getASTContext().getTargetAddressSpace(ArgAS) !=
        getASTContext().getTargetAddressSpace(ParamAS))
   || ParamAS == LangAS::Default)

[gandhi21299]

adding codegen test

[tra]

I think for a Sema change there should be a Sema test, too.

[gandhi21299]

Passed internal CI, still working on a Sema test.

[gandhi21299]

adding sema test

[gandhi21299]

installed clang-format, refreshing patch

[gandhi21299]

removed unused diagnostic sema note

[tra]

clang/lib/Sema/SemaExpr.cpp

6548–6553

I'd rephrase it in terms of add ASC if target AS is different.

Prevent assumes that we've already decided to add an ASC, while it is not the case. All we know is that we've got a pointer argument and we are figuring out whether ASC is needed.

clang/test/SemaCUDA/builtins-unsafe-atomics-gfx90a.cu

4

You do not want to produce any output. I think -fsyntax-only -verify along with // expected-no-diagnostics should do the trick.

9

Not needed w/o FileCheck.

[gandhi21299]

corrected sema test, as requested

[yaxunl]

LGTM. Thanks.

[gandhi21299]

Thanks for the review!

[tra]

LGTM in general, modulo a readability suggestion.

clang/lib/Sema/SemaExpr.cpp

6548–6553

Hmm. I guess that does not quite match what we do here.
Oringal code would add ASC for generic->specific only.
The test cases you've addedd are for specific->specific AS, AFAICT, and that makes me wonder if I understand what's going on here.

So the test case currently does fail, because we end up w/o implicit ASC and end up trying to pass wrong types of pointers. https://godbolt.org/z/MaWMbGb5z

Now we actually want to allow implicit ASC for the pointer types that are in the *same* target AS, so my comment suggestion was actually wrong.

This confusion also suggests that the condition is hard to understand. Perhaps we can rewrite it.

E.g.

bool NeedImplicitASC = 
      ParamAS != LangAS::Default &&   // Pointers params in generic AS don't need special handling.
      ( ArgAs == LangAS::Default  ||         // We do allow implicit conversion from generic AS 
        // or from specific AS which has target AS matching that of Param.
        getASTContext().getTargetAddressSpace(ArgAS) == getASTContext().getTargetAddressSpace(ParamAS));
if (!NeedImplicitASC)
   continue;

[gandhi21299]

@tra I see, we sure can rewrite that segment for readability.

[tra]

In D111734#3067624, @gandhi21299 wrote:

@tra I see, we sure can rewrite that segment for readability.

It could've been done with a follow-up clean-up commit. revert/reland works, too. Thank you for following up on the comment, either way.

[gandhi21299]

cleaned up code for readability

clang/lib/Sema/SemaExpr.cpp

if (getLangOpts().HIP && getLangOpts().CUDAIsDevice && FD &&
             !ArgExprs[Idx]->getType()->isPointerType())
           continue;
 
         auto ParamAS = Param->getType()->getPointeeType().getAddressSpace();
         auto ArgTy = ArgExprs[Idx]->getType();
         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)
+        // Add address space cast if target address spaces are different
+        if ((ArgAS != LangAS::Default &&
+             getASTContext().getTargetAddressSpace(ArgAS) !=
+                 getASTContext().getTargetAddressSpace(ParamAS)) ||
+            ParamAS == LangAS::Default)
           continue;
 
         // First, ensure that the Arg is an RValue.
         if (ArgExprs[Idx]->isGLValue()) {
           ArgExprs[Idx] = ImplicitCastExpr::Create(
               Context, ArgExprs[Idx]->getType(), CK_NoOp, ArgExprs[Idx],
               nullptr, VK_PRValue, FPOptionsOverride());
         }
 

clang/test/CodeGenCUDA/builtins-unsafe-atomics-gfx90a.cu

+// 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);
+}

clang/test/SemaCUDA/builtins-unsafe-atomics-gfx90a.cu

+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx90a -x hip \
+// RUN:  -aux-triple x86_64-unknown-linux-gnu -fcuda-is-device %s \
+// RUN:  -fsyntax-only -verify
+// expected-no-diagnostics
+
+#define __device__ __attribute__((device))
+typedef __attribute__((address_space(3))) float *LP;
+
+__device__ void test_ds_atomic_add_f32(float *addr, float val) {
+  float *rtn;
+  *rtn = __builtin_amdgcn_ds_faddf((LP)addr, val, 0, 0, 0);
+}