support cudaKernel grammar

cxx-squid/src/main/java/org/sonar/cxx/parser/CxxGrammarImpl.java

@@ -91,6 +91,7 @@ public enum CxxGrammarImpl implements GrammarRuleKey {
   par_expression,
   expression,
   constantExpression,
+  cudaKernel,
 
   // Statements
   statement,
@@ -469,10 +470,10 @@ private static void expressions(LexerfulGrammarBuilder b) {
 
     b.rule(postfixExpression).is(
       b.firstOf(
-            b.sequence(simpleTypeSpecifier, "(", b.optional(expressionList), ")"),
+            b.sequence(simpleTypeSpecifier, b.optional(cudaKernel), "(", b.optional(expressionList), ")"),
             b.sequence(simpleTypeSpecifier, bracedInitList),
             b.sequence(simpleTypeSpecifier, "::", "typeid"),
-            b.sequence(typenameSpecifier, "(", b.optional(expressionList), ")"),
+            b.sequence(typenameSpecifier, b.optional(cudaKernel), "(", b.optional(expressionList), ")"),
             b.sequence(typenameSpecifier, bracedInitList),
             b.sequence(typenameSpecifier, "::", "typeid"),
 
@@ -511,6 +512,8 @@ private static void expressions(LexerfulGrammarBuilder b) {
         )
         ).skipIfOneChild();
 
+    b.rule(cudaKernel).is(b.sequence("<<", "<", b.optional(expressionList), ">>", ">"));
+
         b.rule(typeIdEnclosed).is(b.firstOf(
           b.sequence("<", typeId, ">"),
           b.sequence("<", innerTypeId, ">>")));

cxx-squid/src/test/java/org/sonar/cxx/parser/CxxParserTest.java

@@ -35,7 +35,7 @@
 
 public class CxxParserTest extends ParserBaseTest {
   String errSources = "/parser/bad/error_recovery_declaration.cc";
-  String[] goodFiles = {"own", "examples", "cli"};
+  String[] goodFiles = {"own", "examples", "cli", "cuda"};
   String[] cCompatibilityFiles = {"c-compat"};
   String rootDir = "src/test/resources/parser";
   File erroneousSources = null;

cxx-squid/src/test/java/org/sonar/cxx/parser/ExpressionTest.java

@@ -149,6 +149,7 @@ public void postfixExpression() {
     g.rule(CxxGrammarImpl.typeId).mock();
     g.rule(CxxGrammarImpl.deleteExpression).mock();
     g.rule(CxxGrammarImpl.pseudoDestructorName).mock();
+    g.rule(CxxGrammarImpl.cudaKernel).mock();
 
     assertThat(p).matches("primaryExpression");
     assertThat(p).matches("primaryExpression [ expression ]");
@@ -175,6 +176,7 @@ public void postfixExpression() {
     assertThat(p).matches("simpleTypeSpecifier :: typeid");
     assertThat(p).matches("typenameSpecifier :: typeid");
 
+    assertThat(p).matches("simpleTypeSpecifier cudaKernel ( expressionList )");
   }
 
   @Test
@@ -189,6 +191,8 @@ public void postfixExpression_reallife() {
 
     assertThat(p).matches("G::typeid");
     assertThat(p).matches("int::typeid");
+
+    assertThat(p).matches("kernel<<<gridDim,blockDim,0>>>(d_data, height, width)");
   }
 
   @Test

cxx-squid/src/test/resources/parser/cuda/cuda-defines.h

@@ -0,0 +1,11 @@
+// define project specific values to reduce parsing errors
+#define __global__
+#define __host__
+#define __device__
+#define __noinline__ 
+#define __forceinline__
+#define __constant__ 
+#define __shared__
+#dfeine __managed__
+#define __restrict__ 
+#define __align__(a)

cxx-squid/src/test/resources/parser/cuda/kernel_function.cpp

@@ -0,0 +1,41 @@
+#include "cuda-defines.h"
+
+texture<float, 2, cudaReadModeElementType> tex;
+
+void foo()
+{
+  cudaArray* cu_array;
+
+  // Allocate array
+  cudaChannelFormatDesc description = cudaCreateChannelDesc<float>();
+  cudaMallocArray(&cu_array, &description, width, height);
+
+  // Copy image data to array
+  cudaMemcpyToArray(cu_array, image, width*height*sizeof(float), cudaMemcpyHostToDevice);
+
+  // Set texture parameters (default)
+  tex.addressMode[0] = cudaAddressModeClamp;
+  tex.addressMode[1] = cudaAddressModeClamp;
+  tex.filterMode = cudaFilterModePoint;
+  tex.normalized = false; // do not normalize coordinates
+
+  // Bind the array to the texture
+  cudaBindTextureToArray(tex, cu_array);
+
+  // Run kernel
+  dim3 blockDim(16, 16, 1);
+  dim3 gridDim((width + blockDim.x - 1)/ blockDim.x, (height + blockDim.y - 1) / blockDim.y, 1);
+  kernel<<< gridDim, blockDim, 0 >>>(d_data, height, width);
+  // Unbind the array from the texture
+  cudaUnbindTexture(tex);
+} //end foo()
+
+__global__ void kernel(float* odata, int height, int width)
+{
+   unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
+   unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
+   if (x < width && y < height) {
+      float c = tex2D(tex, x, y);
+      odata[y*width+x] = c;
+   }
+}