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

[Anton Lokhmotov]

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?