[cfe-dev] cfe-dev Digest, Vol 55, Issue 4

[Bill Wendling]

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