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.
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. 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 << ") ";
+ }
+ }
+
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()) << " ";
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__ ";
+ }
+
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 += ")))";
+ }
}
if (Qualifiers::GC gc = getObjCGCAttr()) {
if (!S.empty()) S += ' ';
signature.asc
Description: Digital signature
_______________________________________________ cfe-commits mailing list [email protected] http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
