gandhi21299 created this revision. gandhi21299 added a reviewer: yaxunl. gandhi21299 requested review of this revision. Herald added a project: clang. Herald added a subscriber: cfe-commits.
Allow (implicit) address space casting between LLVM-equivalent target address spaces. Repository: rG LLVM Github Monorepo 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,15 @@ +// 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 __global__ __attribute__((global)) +#define __shared__ __attribute__((shared)) +#define __device__ __attribute__((device)) +typedef __attribute__((address_space(0))) float *GP; +typedef __attribute__((address_space(3))) float *LP; + +// CHECK-LABEL: test_ds_atomic_add_f32 +__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 @@ -6547,7 +6547,9 @@ // 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) + if (ArgAS != LangAS::Default && + getASTContext().getTargetAddressSpace(ArgAS) != + getASTContext().getTargetAddressSpace(ParamAS)) continue; // First, ensure that the Arg is an RValue.
Index: clang/test/CodeGenCUDA/builtins-unsafe-atomics-gfx90a.cu =================================================================== --- /dev/null +++ clang/test/CodeGenCUDA/builtins-unsafe-atomics-gfx90a.cu @@ -0,0 +1,15 @@ +// 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 __global__ __attribute__((global)) +#define __shared__ __attribute__((shared)) +#define __device__ __attribute__((device)) +typedef __attribute__((address_space(0))) float *GP; +typedef __attribute__((address_space(3))) float *LP; + +// CHECK-LABEL: test_ds_atomic_add_f32 +__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 @@ -6547,7 +6547,9 @@ // 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) + if (ArgAS != LangAS::Default && + getASTContext().getTargetAddressSpace(ArgAS) != + getASTContext().getTargetAddressSpace(ParamAS)) continue; // First, ensure that the Arg is an RValue.
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits