[cfe-dev] MaxAtomicInlineWidth of SPIR64TargetInfo

[Craig, Ben via cfe-dev]

When adding C11 atomics support to the Hexagon backend, I ended up 
adding almost exactly the same line.  I suspect that the reason it 
wasn't already there is that the people adding C11 and C++11 didn't want 
to try and modify backends they were unfamiliar with.

On 1/14/2016 10:02 PM, Luo, Xionghu via cfe-dev wrote:
>
> 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
>
>
>
> _______________________________________________
> cfe-dev mailing list
> cfe-dev at lists.llvm.org
> http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-dev

-- 
Employee of Qualcomm Innovation Center, Inc.
Qualcomm Innovation Center, Inc. is a member of Code Aurora Forum, a Linux Foundation Collaborative Project

-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/cfe-dev/attachments/20160115/baabf9b6/attachment.html>