Matching CUDA code

Hello everybody,

as someone probably already knows, I'm working on some CUDA source-to
-source transformations using clang 3.7

I'm using the mailing list, even if I received a lot of good answers
from the IRC chat, both for having opinions by a different set of
people and leave a trace on the web, since it seems there's a lot of
interest in this topic (and a lot of previous discussions).

My final goal will be the developing of a tool using the LibTooling
library that will be able to translate CUDA code in (probably) C++11,
with the basic idea being the translation of the kernels in
std::threads.

First of all, I'm trying to use all the clang support for CUDA that is
already available.

As Artem Belevich answered here
http://lists.cs.uiuc.edu/pipermail/cfe-dev/2015-June/043466.html and
looking at the source it seems that all the magic relies in defining
some header like the one in test/SemaCUDA/Inputs/ called cuda.h and
including it in the CUDA code files that you want to parse. So, it
seems that there is no support at all for CUDA syntax in the "real"
clang parser for now.

Of course one would like to use ASTs when doing s2s transformations,
and luckily you can find tree nodes referring to specific CUDA syntax
in the ast-dumps generated by clang.

Enough with the boring introduction. I want to use this thread like a
sort of list of what you can do with the AST matchers about CUDA
(hoping I will not be the only one asking questions).

First question!

Let's match CUDA qualifiers, like for example __global__, in order to
find kernel declarations.

If I try to use a Matcher defined within an ASTConsumer, it will not
work. For example something like:

class MyASTConsumer : public ASTConsumer {
public:
  MyASTConsumer(Rewriter &R) : THandler(R) {
   
Matcher.addMatcher(functionDecl(hasAttr(clang::attr::CUDAGlobal)).bind
("cex"),
        &THandler);
  }

  void HandleTranslationUnit(ASTContext &Context) override {
    Matcher.matchAST(Context);

  }

private:
  MatchFinder Matcher;
  ThreadHandler THandler;
};

Otherwise, using a solution similar the one I found here
http://lists.cs.uiuc.edu/pipermail/cfe-dev/2011-June/015574.html it
will work.

void HandleTranslationUnit(ASTContext &Context) override {
  Matcher.matchAST(Context);

  for(DeclContext::decl_iterator D = Context.getTranslationUnitDecl()
->decls_begin(), DEnd = Context.getTranslationUnitDecl()->decls_end();
D!=DEnd; ++D){
   FunctionDecl* FD;
   if((FD=dyn_cast<FunctionDecl> (*D)) != 0){
    if(FD->hasAttr<CUDAGlobalAttr>())
      std::cout << "found " << FD->getNameAsString() << " fun
decl\n";
   }
  }

}

What do you think about this behaviour? Am I putting the wrong value
in hasAttr() in the first case?

Thanks,
Luca

Hello everybody,

[...]

First question!

Let's match CUDA qualifiers, like for example __global__, in order to
find kernel declarations.

If I try to use a Matcher defined within an ASTConsumer, it will not
work. For example something like:

class MyASTConsumer : public ASTConsumer {
public:
  MyASTConsumer(Rewriter &R) : THandler(R) {
   
Matcher.addMatcher(functionDecl(hasAttr(clang::attr::CUDAGlobal)).bi
nd
("cex"),
        &THandler);
  }

  void HandleTranslationUnit(ASTContext &Context) override {
    Matcher.matchAST(Context);

  }

private:
  MatchFinder Matcher;
  ThreadHandler THandler;
};

Actually this solution works fine, the problem seems to be how do you
define attributes in the input file. The matcher is not capable to
expand macros, so it works only when we have something like

__attribute__((global)) void kernel(...){...}

but not defining

#define __global__ __attribute__((global))

Probably we have to interact in some way with clang's preprocessor.