[cfe-commits] PATCH: support for pretty-printing OpenCL/CUDA keywords/attributes

[Richard Membarth]

On 2011.10.26 19:14, Peter Collingbourne wrote:
> On Tue, Oct 25, 2011 at 04:40:07PM +0200, Richard Membarth wrote:
> > Hello all,
> > 
> > I'm currently working on a source-to-source compiler for CUDA and
> > OpenCL. The CUDA and OpenCL target code is emitted using the
> > pretty-printer in Clang.
> 
> Have you looked at the Rewrite library?  In general this is a
> much better way of doing source to source translation using
> Clang.

I use both, the Rewrite library for the host-code and the pretty
printer for the device-code. So far, this works for me apart from
the missing support for CUDA/OpenCL keywords/qualifiers.

> 
> > Attached is a patch that adds support for CUDA and OpenCL
> > keywords/attributes to the pretty-printer.
> > 
> > Please let me know if this is ok.
> 
> Comments inline.
> 
> > Thanks,
> >   Richard
> > 
> 
> > diff --git a/lib/AST/DeclPrinter.cpp b/lib/AST/DeclPrinter.cpp
> > index 08a1ab5..1e46a7f 100644
> > --- a/lib/AST/DeclPrinter.cpp
> > +++ b/lib/AST/DeclPrinter.cpp
> > @@ -390,6 +390,30 @@ void DeclPrinter::VisitFunctionDecl(FunctionDecl *D) {
> >      if (D->isModulePrivate())    Out << "__module_private__ ";
> >    }
> >  
> > +  if (Policy.LangOpts.OpenCL) {
> > +    if (D->hasAttr<OpenCLKernelAttr>()) Out << "__kernel ";
> > +    if (D->hasAttr<ReqdWorkGroupSizeAttr>()) {
> > +      Out << "__attribute__ ((reqd_work_group_size(";
> > +      ReqdWorkGroupSizeAttr * Attr = D->getAttr<ReqdWorkGroupSizeAttr>();
> > +      Out << Attr->getXDim() << ", ";
> > +      Out << Attr->getYDim() << ", ";
> > +      Out << Attr->getZDim() << "))) ";
> > +    }
> > +  }
> > +
> > +  if (Policy.LangOpts.CUDA) {
> > +    if (D->hasAttr<CUDAHostAttr>()) Out << "__host__ ";
> > +    if (D->hasAttr<CUDADeviceAttr>()) Out << "__device__ ";
> > +    if (D->hasAttr<CUDAGlobalAttr>()) Out << "__global__ ";
> > +    if (D->hasAttr<CUDALaunchBoundsAttr>()) {
> > +      CUDALaunchBoundsAttr *Attr = D->getAttr<CUDALaunchBoundsAttr>();
> > +      Out << "__launch_bounds__ (";
> > +      Out << Attr->getMaxThreads();
> > +      if (Attr->getMinBlocks()) Out << ", " << Attr->getMinBlocks();
> > +      Out << ") ";
> > +    }
> > +  }
> > +
> 
> I'm not so sure that we should add manual support for pretty printing
> every attribute to the decl printer (which is the logical conclusion of
> your changes above), as this increases maintenance burden.  Ideally I'd
> like to see tblgen capable of generating the necessary code for pretry
> printing attributes.

I agree on that. Having tblgen generate the required code for all
attributes is a much cleaner solution. I'm working on this and
sending a separate patch.

> 
> >    PrintingPolicy SubPolicy(Policy);
> >    SubPolicy.SuppressSpecifiers = false;
> >    std::string Proto = D->getNameInfo().getAsString();
> > @@ -596,7 +620,8 @@ void DeclPrinter::VisitLabelDecl(LabelDecl *D) {
> >  
> >  
> >  void DeclPrinter::VisitVarDecl(VarDecl *D) {
> > -  if (!Policy.SuppressSpecifiers && D->getStorageClass() != SC_None)
> > +  if (!Policy.SuppressSpecifiers && D->getStorageClass() != SC_None &&
> > +          D->getStorageClass() != SC_OpenCLWorkGroupLocal)
> >      Out << VarDecl::getStorageClassSpecifierString(D->getStorageClass()) << " ";
> 
> I think that this should use getStorageClassAsWritten, which reflects
> the original storage class used in the source code (which would always
> be SC_None for OpenCL and in particular for __local variables).
> 

You're right, using getStorageClassAsWritten() is sufficient.

> >  
> >    if (!Policy.SuppressSpecifiers && D->isThreadSpecified())
> > @@ -604,6 +629,12 @@ void DeclPrinter::VisitVarDecl(VarDecl *D) {
> >    if (!Policy.SuppressSpecifiers && D->isModulePrivate())
> >      Out << "__module_private__ ";
> >  
> > +  if (Policy.LangOpts.CUDA) {
> > +    if (D->hasAttr<CUDAConstantAttr>()) Out << "__constant__ ";
> > +    if (D->hasAttr<CUDADeviceAttr>()) Out << "__device__ ";
> > +    if (D->hasAttr<CUDASharedAttr>()) Out << "__shared__ ";
> > +  }
> > +
> 
> Same comment as above.
> 
> >    std::string Name = D->getNameAsString();
> >    QualType T = D->getType();
> >    if (ParmVarDecl *Parm = dyn_cast<ParmVarDecl>(D))
> > diff --git a/lib/AST/TypePrinter.cpp b/lib/AST/TypePrinter.cpp
> > index fb7b918..f119f88 100644
> > --- a/lib/AST/TypePrinter.cpp
> > +++ b/lib/AST/TypePrinter.cpp
> > @@ -1157,9 +1157,21 @@ void Qualifiers::getAsStringInternal(std::string &S,
> >    AppendTypeQualList(S, getCVRQualifiers());
> >    if (unsigned addrspace = getAddressSpace()) {
> >      if (!S.empty()) S += ' ';
> > -    S += "__attribute__((address_space(";
> > -    S += llvm::utostr_32(addrspace);
> > -    S += ")))";
> > +    switch (addrspace) {
> > +      case LangAS::opencl_global:
> > +        S += "__global ";
> > +        break;
> > +      case LangAS::opencl_local:
> > +        S += "__local ";
> > +        break;
> > +      case LangAS::opencl_constant:
> > +        S += "__constant ";
> > +        break;
> > +      default:
> > +        S += "__attribute__((address_space(";
> > +        S += llvm::utostr_32(addrspace);
> > +        S += ")))";
> > +    }
> 
> This seems fine, but is the whitespace after the address space names
> necessary?  If you add a test case (testing that diagnostics use the
> OpenCL address spaces), I will commit this.

The whitespace is not required here.

Attached is an updated patch for OpenCL qualifiers and a test case
to check if the OpenCL address space is used.

Thanks,
  Richard

> 
> Thanks,
> -- 
> Peter

-------------- next part --------------
A non-text attachment was scrubbed...
Name: prettyprint_opencl_qualifiers.patch
Type: text/x-diff
Size: 1931 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20111028/9015a03e/attachment.patch>
-------------- next part --------------
A non-text attachment was scrubbed...
Name: signature.asc
Type: application/pgp-signature
Size: 198 bytes
Desc: Digital signature
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20111028/9015a03e/attachment.sig>