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

[Benyei, Guy]

Hi Anton,
It's good to resume this effort. I think Clang is very close to fully support OpenCL, so I hope we'll be able to provide full support soon.

I agree, this metadata format you propose is flexible, and effective. I'll use your function in my updated patch; I'll send it soon in another mail.

BTW, 

+  if(S.getLangOptions().OpenCL && ASIdx == LangAS::opencl_constant) {
+    Type = S.Context.getConstType(Type);
+  }

wasn't meant to be part of this patch (even without the braces), but it prevents assignments to constant address space variables. It can be interesting too, but probably as part of another patch.

Thanks
    Guy


-----Original Message-----
From: Anton Lokhmotov [mailto:Anton.Lokhmotov at arm.com] 
Sent: Tuesday, January 03, 2012 19:54
To: cfe-dev at cs.uiuc.edu
Cc: Benyei, Guy; pekka.jaaskelainen at tut.fi
Subject: RE: cfe-dev Digest, Vol 55, Issue 4

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?



---------------------------------------------------------------------
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.