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
diff --git a/lib/AST/DeclPrinter.cpp b/lib/AST/DeclPrinter.cpp
index 08a1ab5..ebc2937 100644
--- a/lib/AST/DeclPrinter.cpp
+++ b/lib/AST/DeclPrinter.cpp
@@ -596,7 +596,7 @@ void DeclPrinter::VisitLabelDecl(LabelDecl *D) {
void DeclPrinter::VisitVarDecl(VarDecl *D) {
- if (!Policy.SuppressSpecifiers && D->getStorageClass() != SC_None)
+ if (!Policy.SuppressSpecifiers && D->getStorageClassAsWritten() != SC_None)
Out << VarDecl::getStorageClassSpecifierString(D->getStorageClass()) << " ";
if (!Policy.SuppressSpecifiers && D->isThreadSpecified())
diff --git a/lib/AST/TypePrinter.cpp b/lib/AST/TypePrinter.cpp
index fb7b918..14a1c29 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 += ")))";
+ }
}
if (Qualifiers::GC gc = getObjCGCAttr()) {
if (!S.empty()) S += ' ';
diff --git a/test/SemaOpenCL/qualifiers_print.cl b/test/SemaOpenCL/qualifiers_print.cl
index e69de29..1e13bd8 100644
--- a/test/SemaOpenCL/qualifiers_print.cl
+++ b/test/SemaOpenCL/qualifiers_print.cl
@@ -0,0 +1,11 @@
+// RUN: %clang_cc1 %s -ast-print > %t.out
+// RUN: grep __constant %t.out
+// RUN: grep __global %t.out
+// RUN: grep __local %t.out
+
+__constant float f;
+
+__kernel void g1(__global int *in) {
+ __local int i;
+}
+
signature.asc
Description: Digital signature
_______________________________________________ cfe-commits mailing list [email protected] http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
