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);

03.01.2012 19:54, Anton Lokhmotov kirjoitti:

What do you think?

It looks like an improvement to the current proposal (including my
original named metadata for "wg dimension info"). Do you have
a full patch at hand that implements this version?

BR,

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

Hi Guy,

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.

Just a heads-up that we are working on merging your patch (where we really
like your clever handling of the vec_type_hint attribute!) with ours (which
contains extended error checking), and should submit it shortly, so let's
not duplicate the efforts.

Best, Anton.

Hi Anton,
I've sent out another update before I've seen this message, but I'll wait for your patch before any further action.

Thanks
    Guy

Hi Guy,

We've done the merge (including your most recent test using regular
expressions). Please review if anything is missing.

I don't quite understand this change (in lib/Sema/SemaType.cpp):

+ } else if (S.getLangOptions().OpenCL) {
+ S.Diag(DeclLoc, diag::err_missing_type_specifier_opencl)
+ << DS.getSourceRange();
+ declarator.setInvalidType(true);

Could you please explain how it relates to the kernel attributes?

Cheers,
Anton.

kernel_attributes.patch (20.1 KB)

Guy & Anton,

I'll just review this merged patch. Here is my feedback:

- Remove the type specifier diagnostic. I don't think its needed (can you just use the existing one?) and its unrelated to the metadata.

- I'd like to suggest changing how the metadata is organized. So it would be like this:
NamedMetaData Node - opencl.kernels
  MDNode Pair of Function*, MDNode which is a list of all the metadata
    MDNode list that is a key/value pair. So for each attribute you have key (ie. "vec_type_hint"), and then the value(s).

Its just a slight variation on how you have things (where you have a list with the first arg being the key). This makes it pretty obvious what the layout is in my opinion.

Does that seem reasonable?

- The DummyType is a creative solution. I think we'll need Doug to sign off on that change though.

- Add some comments to the function EmitOpenCLKernelMetadata. Maybe one sentence to say what the metadata is and reference to spec.

- Nit pick: Add periods to your comments in AttributeList.h,

Thanks,
Tanya

Guy & Anton,

I'll just review this merged patch. Here is my feedback:

- Remove the type specifier diagnostic. I don't think its needed (can you just use the existing one?) and its unrelated to the metadata.

- I'd like to suggest changing how the metadata is organized. So it would be like this:
NamedMetaData Node - opencl.kernels
  MDNode Pair of Function*, MDNode which is a list of all the metadata
    MDNode list that is a key/value pair. So for each attribute you have key (ie. "vec_type_hint"), and then the value(s).

Its just a slight variation on how you have things (where you have a list with the first arg being the key). This makes it pretty obvious what the layout is in my opinion.

Does that seem reasonable?

Actually, let me clarify this a bit.

So I think the layout you have is this:

NamedMDNode(Fn*, X number of MDNodes..), where each MDNode node is a ("key" (MDString), X number of value*).

If so, then that is what I'm suggesting as well except I was nesting the lists (at each level) in an MDNode. Not sure if thats really beneficial now that I think about it :slight_smile:

So, I retract my opinion here.

Thanks,
Tanya

Hi Anton,
The error I've added in SemaType.cpp is related to the case, when the user defines an unknown identifier as vec_type_hint attribute. In this case in OpenCL, ConvertDeclSpecToType will do something like this:

try.cl:1:39: error: type name requires a specifier or qualifier
__kernel __attribute__((vec_type_hint(mmm))) void foo( void ){}
                                      ^
try.cl:1:39: warning: type specifier missing, defaults to 'int'
__kernel __attribute__((vec_type_hint(mmm))) void foo( void ){}
                                      ^~~
try.cl:1:44: error: expected identifier or '('
__kernel __attribute__((vec_type_hint(mmm))) void foo( void ){}
                                           ^
1 warning and 2 errors generated.

I think the warning above is totally wrong in the context of OpenCL, and should be an error, as written in the C99 specs. Anyhow, this case leads to a failure, but I thought it would make the messages a bit more understandable.

I have only one remark for this patch:

@@ -1356,6 +1356,10 @@ public:

   void GenerateCode(GlobalDecl GD, llvm::Function *Fn,
                     const CGFunctionInfo &FnInfo);

Hi Tanya and Guy,

- Add some comments to the function EmitOpenCLKernelMetadata. Maybe one
sentence to say what the metadata is and reference to spec.

Done.

- Remove the type specifier diagnostic. I don't think its needed (can
you just use the existing one?) and its unrelated to the metadata.

While I appreciate Guy's concern that:

__kernel __attribute__((vec_type_hint(mmm))) void foo( void ){}

should generate a better error message, with his code I still see three
other confusing messages:

kernel-attributes-invalid.cl:7:39: error: type name requires a specifier or
qualifier
__kernel __attribute__((vec_type_hint(mm))) void foo( void ){}
                                      ^
kernel-attributes-invalid.cl:7:39: error: OpenCL requires a type specifier
for all declarations
__kernel __attribute__((vec_type_hint(mm))) void foo( void ){}
                                      ^~~
kernel-attributes-invalid.cl:7:39: error: expected ')'
__kernel __attribute__((vec_type_hint(mm))) void foo( void ){}
                                      ^
                                      )
kernel-attributes-invalid.cl:7:44: error: expected identifier or '('
__kernel __attribute__((vec_type_hint(mm))) void foo( void ){}

I think we should try to come up with a better solution...

- The DummyType is a creative solution. I think we'll need Doug to sign
off on that change though.

Please see below an excerpt from an alternative solution handling the
vec_type_hint parameter as an identifier (yuck!). Hopefully, it's enough to
convince Doug ;).

Cheers,
Anton.

+static void handleOpenCLVecTypeHint(Sema &S, Decl *D, const AttributeList
&Attr)
+{
+ if (!Attr.getParameterName()) {
+ S.Diag(Attr.getLoc(), diag::err_attribute_wrong_number_arguments) << 1;
+ return;
+ }

kernel_optional_attribute_qualifiers.patch (18.7 KB)

Hi Tanya and Guy,

- Add some comments to the function EmitOpenCLKernelMetadata. Maybe one
sentence to say what the metadata is and reference to spec.

Done.

- Remove the type specifier diagnostic. I don't think its needed (can
you just use the existing one?) and its unrelated to the metadata.

While I appreciate Guy's concern that:

__kernel __attribute__((vec_type_hint(mmm))) void foo( void ){}

should generate a better error message, with his code I still see three
other confusing messages:

kernel-attributes-invalid.cl:7:39: error: type name requires a specifier or
qualifier
__kernel __attribute__((vec_type_hint(mm))) void foo( void ){}
                                     ^
kernel-attributes-invalid.cl:7:39: error: OpenCL requires a type specifier
for all declarations
__kernel __attribute__((vec_type_hint(mm))) void foo( void ){}
                                     ^~~
kernel-attributes-invalid.cl:7:39: error: expected ')'
__kernel __attribute__((vec_type_hint(mm))) void foo( void ){}
                                     ^
                                     )
kernel-attributes-invalid.cl:7:44: error: expected identifier or '('
__kernel __attribute__((vec_type_hint(mm))) void foo( void ){}

I think we should try to come up with a better solution...

- The DummyType is a creative solution. I think we'll need Doug to sign
off on that change though.

Please see below an excerpt from an alternative solution handling the
vec_type_hint parameter as an identifier (yuck!). Hopefully, it's enough to
convince Doug ;).

I'm not particularly thrilled with either solution, because both involve special-purpose hacks.

Let's first characterize how best to approach this problem. First of all, the syntax of the vec_type_hint attribute implies that what we really want to do is parse the attribute argument as a type name (via Parse::ParseTypeName). That's exactly what the attached patch does, which is great. I was hoping that the patch would also eliminate the BuiltinType stuff from Parser::ParseGNUAttributeArgs.

The next step is to get that parsed type information from the parser to semantic analysis. The best way to do so would be for AttributeList to properly support type arguments, because then we could directly pass the type through. A refactor of AttributeList et al to support type arguments would benefit vec_type_hint as well as several other arguments that accept types (e.g., iboutletcollection), and would be the best way to approach this problem. It would maintain type-source information properly, and wouldn't require the allocation of an expression that will effectively be leaked.

DummyTypeExpr will encode the type as an expression to get it through the AttributeList machinery without any refactoring. It will work, and it is expedient, but it's not the right long-term solution for Clang. If we have to go in this direction, we shouldn't introduce a new Expr kind for it. Rather, we should just use OpaqueValueExpr, which can fulfill the same role.

  - Doug

DummyTypeExpr will encode the type as an expression to get it through
the AttributeList machinery without any refactoring. It will work, and
it is expedient, but it's not the right long-term solution for Clang.
If we have to go in this direction, we shouldn't introduce a new Expr
kind for it. Rather, we should just use OpaqueValueExpr, which can
fulfill the same role.

Agree.

Guy, is it something you could try to refactor?

Cheers,
Anton.

I can check this direction, but it will take time to finish this work.

Anyhow, I think the short term solution with the OpaqueValueExpr instead DummyTypeExpr could be committed to solve this issue for now.

Thanks
    Guy

I can check this direction, but it will take time to finish this work.

Understood.

Anyhow, I think the short term solution with the OpaqueValueExpr instead DummyTypeExpr could be committed to solve this issue for now.

In the interest of rectifying more of the various vendors' OpenCL implementations, I can live with this in the short term.

I feel obligated to make a code owner comment: Clang is great to work with specifically because we refactor until our new features can be dropped in beautifully. If we compromise on that approach too often, or for too long, our code base will be seriously damaged.

  - Doug

Hi Doug,
I'm looking into the attributes stuff, and it seems like there was some effort to make it tablegenned all the way.
Currently even with the Attr.td file in place, each attribute must be added also in AttributeList.h, in AttributeList.cpp, and also there is specific handling for every attribute in SemaDeclAttr.cpp.

Is there any reason for this?
I would imagine a simpler way to add attributes by only inserting them in the Attr.td file, and getting the rest for free.

Thanks
     Guy

That was indeed the design; it simply hasn't happened yet. The person who was working on this went on to work on other things.

John.