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

[Peter Collingbourne]

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.

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

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

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

Thanks,
-- 
Peter