[Beignet] MaxAtomicInlineWidth of SPIR64TargetInfo

[Anastasia Stulova]

Yes, I think it makes sense to set the max width in SPIR for atomics to 64 (at least looking at OpenCL and SPIR spec of atomics).

If you send this as a patch to me and cfe-commits, I will review.

Cheers,
Anastasia

From: Song, Ruiling [mailto:ruiling.song at intel.com]
Sent: 27 January 2016 03:22
To: Anastasia Stulova
Cc: beignet at lists.freedesktop.org; Luo, Xionghu; 'cfe-dev at lists.llvm.org'
Subject: RE: MaxAtomicInlineWidth of SPIR64TargetInfo

Hi Anastasia,

Do you have any comment/suggestion on this?

Thanks!
Ruiling

From: cfe-dev [mailto:cfe-dev-bounces at lists.llvm.org] On Behalf Of Luo, Xionghu via cfe-dev
Sent: Friday, January 15, 2016 12:03 PM
To: 'cfe-dev at lists.llvm.org' <cfe-dev at lists.llvm.org<mailto:cfe-dev at lists.llvm.org>>
Cc: beignet at lists.freedesktop.org<mailto:beignet at lists.freedesktop.org>
Subject: [cfe-dev] MaxAtomicInlineWidth of SPIR64TargetInfo

Hello,

We are developing the atomic built-in function for OpenCL 2.0. we use SPIR64 as the target and hope to lower the llvm built-in functions like '__c11_atomic_fetch_and_sub' to instructions like 'atomicrmw', but the MaxAtomicInlineWidth of SPIR64TargetInfo is not set by default, so the EmitAtomicExpr goes to libcall path and got function '__atomic_fetch_sub_4', which is unexpected.

After set the variable MaxAtomicInlineWidth, the EmitAtomicExpr take the non-libcall path and get instruction:  '%1 = atomicrmw volatile sub i32* %arrayinit.begin, i32 1 acquire'

class SPIR64TargetInfo : public SPIRTargetInfo {
public:
   SPIR64TargetInfo(const llvm::Triple &Triple) : SPIRTargetInfo(Triple) {
     PointerWidth = PointerAlign = 64;
     SizeType = TargetInfo::UnsignedLong;
     PtrDiffType = IntPtrType = TargetInfo::SignedLong;
     DataLayoutString = "e-i64:64-v16:16-v24:32-v32:32-v48:64-"
                        "v96:128-v192:256-v256:256-v512:512-v1024:1024";
+    MaxAtomicPromoteWidth = MaxAtomicInlineWidth = 64;
   }
   void getTargetDefines(const LangOptions &Opts,
                         MacroBuilder &Builder) const override {
     DefineStd(Builder, "SPIR64", Opts);
   }

So the question is why some other targets set the MaxAtomicInlineWidth when created but SPIR64TargetInfo did NOT? Shall we add it as above?
Thanks.



The mentioned OpenCL test kernel is as below, command is './clang -cc1 -emit-llvm -triple spir64 -cl-std=CL2.0 atomic_functions.cl -o atomic_functions.spir':

__kernel void atomic_functions()
{
  volatile atomic_uint ptest[2] = {0};
  int test = __c11_atomic_fetch_sub(&ptest[0], 1, 1);
}


Luo Xionghu
Best Regards

-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.freedesktop.org/archives/beignet/attachments/20160127/cc6a10f1/attachment.html>