[cfe-dev] cfe-dev Digest, Vol 55, Issue 4
Benyei, Guy
guy.benyei at intel.com
Thu Jan 19 01:21:11 PST 2012
Hi Bill,
I went over your proposal, and it looks very clear and flexible; however I couldn't find an easy way to use it to represent all the
OpenCL kernel and kernel argument specific metadata.
I think it can store module level information, but for most of the interesting OpenCL metadata (which is function/kernel level) it might
only act as entry point, replacing the "opencl.kernels" named metadata node. Trying to store all the kernel metadata in separate nodes
would make the metadata lookup very inefficient and difficult. I think Anton's proposal below is more suitable for this purpose.
Anton/Bill,
If you have ideas how to use Bill's proposal for OpenCL metadata, I think there should be an OpenCL specific proposal with examples
to show it's feasible.
Thanks
Guy
-----Original Message-----
From: cfe-dev-bounces at cs.uiuc.edu [mailto:cfe-dev-bounces at cs.uiuc.edu] On Behalf Of Bill Wendling
Sent: Thursday, January 19, 2012 00:34
To: Anton.Lokhmotov at arm.com
Cc: cfe-dev at cs.uiuc.edu
Subject: Re: [cfe-dev] cfe-dev Digest, Vol 55, Issue 4
Hi Anton,
I just sent out a proposal for "module flags metadata". It's similar to what you're proposing here, but more general. If you think it's a benefit to what you would like to do or needs improvements, please send feedback. :-)
-bw
On Jan 3, 2012, at 9:54 AM, Anton Lokhmotov wrote:
> Hi Guy,
>
> Great to see new OpenCL patches coming in! I hope we'll work together on
> providing full support for OpenCL C in Clang/LLVM.
>
> We'd like to propose a different way of representing kernel function
> qualifiers using metadata.
>
> Example code:
>
> __kernel __attribute__(( work_group_size_hint(1,15,92) )) void hint_1_15_92(
> ) { }
>
> __kernel __attribute__(( reqd_work_group_size(12,72,256) )) void
> reqd_12_72_256( ) { }
>
> __kernel void no_attr( ) { }
>
> __kernel __attribute__(( reqd_work_group_size(12,72,256), vec_type_hint
> (float8))) void mix_attr( ) { }
>
> Example metadata:
>
> !opencl.kernels = !{!0, !2, !4, !5}
>
> !0 = metadata !{void ()* @hint_1_15_92, metadata !1}
> !1 = metadata !{metadata !"work_group_size_hint", i32 1, i32 15, i32 92}
> !2 = metadata !{void ()* @reqd_12_72_256, metadata !3}
> !3 = metadata !{metadata !"reqd_work_group_size", i32 12, i32 72, i32 256}
> !4 = metadata !{void ()* @no_attr}
> !5 = metadata !{void ()* @mix_attr, metadata !6, metadata !3}
> !6 = metadata !{metadata !"vec_type_hint", <8 x float> undef}
>
> The only named metadata node "opencl.kernels" references metadata objects
> for kernel functions. The first object provides the kernel signature (in
> fact, LLVM::Function*), and the following refer to optional kernel
> attributes. Note that the number and order of the optional attributes are
> unspecified, which allows future extensions. Note also that the metadata
> node !3 is shared between the 'mix_attr' and 'reqd_12_72_256' kernels.
>
> The most unusual feature here is representing the 'vec_type_hint' attribute
> not as string metadata (e.g. "float8"), but as an undefined value of the
> corresponding LLVM type. This avoids the need to parse the metadata string
> in the backend to get the type of this attribute.
>
> The following function implements emitting this representation:
>
> void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD,
> llvm::Function *Fn)
> {
> if (FD->hasAttr<OpenCLKernelAttr>()) {
> llvm::SmallVector <llvm::Value*, 5> MDArgs;
> llvm::LLVMContext &Context = getLLVMContext();
> MDArgs.push_back(Fn);
>
> if (FD->hasAttr<VecTypeHintAttr>()) {
> llvm::SmallVector <llvm::Value*, 5> attrMDArgs;
> attrMDArgs.
> push_back(llvm::MDString::get(Context, "vec_type_hint"));
>
> llvm::Type *type_hint =
> CGM.getTypes().
> ConvertType(FD->getAttr<VecTypeHintAttr>()->getTypeHint());
> attrMDArgs.push_back(llvm::UndefValue::get(type_hint));
> MDArgs.push_back(llvm::MDNode::get(Context, attrMDArgs));
> }
>
> if (FD->hasAttr<WorkGroupSizeHintAttr>()) {
> llvm::SmallVector <llvm::Value*, 5> attrMDArgs;
> attrMDArgs.
> push_back(llvm::MDString::get(Context, "work_group_size_hint"));
> llvm::Type *iTy = llvm::IntegerType::get(Context, 32);
> WorkGroupSizeHintAttr *attr = FD->getAttr<WorkGroupSizeHintAttr>();
> attrMDArgs.push_back(llvm::ConstantInt::get(iTy,
> llvm::APInt(32, (uint64_t)attr->getXDim())));
> attrMDArgs.push_back(llvm::ConstantInt::get(iTy,
> llvm::APInt(32, (uint64_t)attr->getYDim())));
> attrMDArgs.push_back(llvm::ConstantInt::get(iTy,
> llvm::APInt(32, (uint64_t)attr->getZDim())));
> MDArgs.push_back(llvm::MDNode::get(Context, attrMDArgs));
> }
>
> if (FD->hasAttr<ReqdWorkGroupSizeAttr>()) {
> llvm::SmallVector <llvm::Value*, 5> attrMDArgs;
> attrMDArgs.
> push_back(llvm::MDString::get(Context, "reqd_work_group_size"));
> llvm::Type *iTy = llvm::IntegerType::get(Context, 32);
> ReqdWorkGroupSizeAttr *attr = FD->getAttr<ReqdWorkGroupSizeAttr>();
> attrMDArgs.push_back(llvm::ConstantInt::get(iTy,
> llvm::APInt(32, (uint64_t)attr->getXDim())));
> attrMDArgs.push_back(llvm::ConstantInt::get(iTy,
> llvm::APInt(32, (uint64_t)attr->getYDim())));
> attrMDArgs.push_back(llvm::ConstantInt::get(iTy,
> llvm::APInt(32, (uint64_t)attr->getZDim())));
> MDArgs.push_back(llvm::MDNode::get(Context, attrMDArgs));
> }
>
> llvm::MDNode *kernelMDNode = llvm::MDNode::get(Context, MDArgs);
> llvm::NamedMDNode *OpenCLMetadata =
> CGM.getModule().getOrInsertNamedMetadata("opencl.kernels");
>
> OpenCLMetadata->addOperand(kernelMDNode);
> }
> }
>
> What do you think?
>
> Best wishes,
> Anton.
>
>
> P.S. Please also find some niggles below.
>
> +/// DummyTypeExpr - Not a real expression, but a simple container for a
> type
> +/// Should be used to pass type arguments to attributes, like OpenCL's
> +/// vec_type_hint
>
> Please terminate sentences with dots.
>
> + /// vec_type_hint(...)
> + ExprResult ActOnDummyTypeExpr(ParsedType ParsedTy);
> +
> Either expand or remove the comment?
>
> - } else {
> + } if (S.getLangOptions().OpenCL) {
>
> Do you mean "} else if (S.getLangOptions().OpenCL) {" ???
>
> + if(S.getLangOptions().OpenCL && ASIdx == LangAS::opencl_constant) {
> + Type = S.Context.getConstType(Type);
> + }
> Remove curly braces?
>
>
>
>
> _______________________________________________
> cfe-dev mailing list
> cfe-dev at cs.uiuc.edu
> http://lists.cs.uiuc.edu/mailman/listinfo/cfe-dev
_______________________________________________
cfe-dev mailing list
cfe-dev at cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-dev
---------------------------------------------------------------------
Intel Israel (74) Limited
This e-mail and any attachments may contain confidential material for
the sole use of the intended recipient(s). Any review or distribution
by others is strictly prohibited. If you are not the intended
recipient, please contact the sender and delete all copies.
More information about the cfe-dev
mailing list