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
_______________________________________________
Beignet mailing list
[email protected]
http://lists.freedesktop.org/mailman/listinfo/beignet
