Hi,
The patch adds ASTMatchers for matching CUDA declarations.
Kindly asking for review.
Jacques
--
Jacques Pienaar | Platforms | [email protected] | 765-430-6883
Index: include/clang/ASTMatchers/ASTMatchers.h
===================================================================
--- include/clang/ASTMatchers/ASTMatchers.h (revision 214460)
+++ include/clang/ASTMatchers/ASTMatchers.h (working copy)
@@ -3643,6 +3643,48 @@
return InnerMatcher.matches(*Node.getLHS(), Finder, Builder);
}
+/// \brief Matches CUDA kernel call expression.
+///
+/// Example matches kernel<<<i,j>>>()
+/// \code
+/// kernel<<<i,j>>>();
+/// \endcode
+const internal::VariadicDynCastAllOfMatcher<Stmt, CUDAKernelCallExpr>
+ CUDAKernelCallExpr;
+
+/// \brief Matches declaration that has CUDA device attribute.
+///
+/// Given
+/// \code
+/// __attribute__((device)) void f() { ... }
+/// \endcode
+/// matches the function declaration of f.
+AST_MATCHER(Decl, hasCudaDeviceAttr) {
+ return Node.hasAttr<clang::CUDADeviceAttr>();
+}
+
+/// \brief Matches declaration that has CUDA host attribute.
+///
+/// Given
+/// \code
+/// __attribute__((host)) void f() { ... }
+/// \endcode
+/// matches the function declaration of f.
+AST_MATCHER(Decl, hasCudaHostAttr) {
+ return Node.hasAttr<clang::CUDAHostAttr>();
+}
+
+/// \brief Matches declaration that has CUDA global attribute.
+///
+/// Given
+/// \code
+/// __attribute__((global)) void f() { ... }
+/// \endcode
+/// matches the function declaration of f.
+AST_MATCHER(Decl, hasCudaGlobalAttr) {
+ return Node.hasAttr<clang::CUDAGlobalAttr>();
+}
+
} // end namespace ast_matchers
} // end namespace clang
Index: unittests/ASTMatchers/ASTMatchersTest.cpp
===================================================================
--- unittests/ASTMatchers/ASTMatchersTest.cpp (revision 214460)
+++ unittests/ASTMatchers/ASTMatchersTest.cpp (working copy)
@@ -649,6 +649,18 @@
EXPECT_TRUE(matches("void f() { int i; }", CannotMemoize));
}
+TEST(DeclarationMatcher, MatchCudaDecl) {
+ EXPECT_TRUE(matchesWithCuda("__global__ void f() { }"
+ "void g() { f<<<1, 2>>>(); }",
+ CUDAKernelCallExpr()));
+ EXPECT_TRUE(matchesWithCuda("__attribute__((device)) void f() {}",
+ hasCudaDeviceAttr()));
+ EXPECT_TRUE(matchesWithCuda("__attribute__((host)) void f() {}",
+ hasCudaHostAttr()));
+ EXPECT_TRUE(matchesWithCuda("__attribute__((global)) void f() {}",
+ hasCudaGlobalAttr()));
+}
+
// Implements a run method that returns whether BoundNodes contains a
// Decl bound to Id that can be dynamically cast to T.
// Optionally checks that the check succeeded a specific number of times.
Index: unittests/ASTMatchers/ASTMatchersTest.h
===================================================================
--- unittests/ASTMatchers/ASTMatchersTest.h (revision 214460)
+++ unittests/ASTMatchers/ASTMatchersTest.h (working copy)
@@ -103,7 +103,73 @@
return matchesConditionally(Code, AMatcher, false, "-std=c++11");
}
+// Function based on matchesConditionally with "-x cuda" argument added and
+// small CUDA header prepended to the code string.
template <typename T>
+testing::AssertionResult matchesConditionallyWithCuda(
+ const std::string &Code, const T &AMatcher, bool ExpectMatch,
+ llvm::StringRef CompileArg) {
+ const std::string CudaHeader =
+ "typedef unsigned int size_t;\n"
+ "#define __constant__ __attribute__((constant))\n"
+ "#define __device__ __attribute__((device))\n"
+ "#define __global__ __attribute__((global))\n"
+ "#define __host__ __attribute__((host))\n"
+ "#define __shared__ __attribute__((shared))\n"
+ "struct dim3 {"
+ " unsigned x, y, z;"
+ " __host__ __device__ dim3(unsigned x, unsigned y = 1, unsigned z = 1)"
+ " : x(x), y(y), z(z) {}"
+ "};"
+ "typedef struct cudaStream *cudaStream_t;"
+ "int cudaConfigureCall(dim3 gridSize, dim3 blockSize,"
+ " size_t sharedSize = 0,"
+ " cudaStream_t stream = 0);";
+
+ bool Found = false, DynamicFound = false;
+ MatchFinder Finder;
+ VerifyMatch VerifyFound(nullptr, &Found);
+ Finder.addMatcher(AMatcher, &VerifyFound);
+ VerifyMatch VerifyDynamicFound(nullptr, &DynamicFound);
+ if (!Finder.addDynamicMatcher(AMatcher, &VerifyDynamicFound))
+ return testing::AssertionFailure() << "Could not add dynamic matcher";
+ std::unique_ptr<FrontendActionFactory> Factory(
+ newFrontendActionFactory(&Finder));
+ // Some tests use typeof, which is a gnu extension.
+ std::vector<std::string> Args({"-xcuda", CompileArg});
+ if (!runToolOnCodeWithArgs(Factory->create(),
+ CudaHeader + Code, Args)) {
+ return testing::AssertionFailure() << "Parsing error in \"" << Code << "\"";
+ }
+ if (Found != DynamicFound) {
+ return testing::AssertionFailure() << "Dynamic match result ("
+ << DynamicFound
+ << ") does not match static result ("
+ << Found << ")";
+ }
+ if (!Found && ExpectMatch) {
+ return testing::AssertionFailure()
+ << "Could not find match in \"" << Code << "\"";
+ } else if (Found && !ExpectMatch) {
+ return testing::AssertionFailure()
+ << "Found unexpected match in \"" << Code << "\"";
+ }
+ return testing::AssertionSuccess();
+}
+
+template <typename T>
+testing::AssertionResult matchesWithCuda(const std::string &Code,
+ const T &AMatcher) {
+ return matchesConditionallyWithCuda(Code, AMatcher, true, "-std=c++11");
+}
+
+template <typename T>
+testing::AssertionResult notMatchesWithCuda(const std::string &Code,
+ const T &AMatcher) {
+ return matchesConditionallyWithCuda(Code, AMatcher, false, "-std=c++11");
+}
+
+template <typename T>
testing::AssertionResult
matchAndVerifyResultConditionally(const std::string &Code, const T &AMatcher,
BoundNodesCallback *FindResultVerifier,
_______________________________________________
cfe-commits mailing list
[email protected]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
