From f708fbea7cf72118046d48cf4fba597ebff41d43 Mon Sep 17 00:00:00 2001 From: Bert Date: Fri, 1 Jan 2016 10:08:20 +0100 Subject: [PATCH] support cudaKernel grammar - added optional rule to postfixExpression - "b.rule(cudaKernel).is(b.sequence("<<", "<", b.optional(expressionList), ">>", ">"));" --- .../org/sonar/cxx/parser/CxxGrammarImpl.java | 7 +++- .../org/sonar/cxx/parser/CxxParserTest.java | 2 +- .../org/sonar/cxx/parser/ExpressionTest.java | 4 ++ .../test/resources/parser/cuda/cuda-defines.h | 11 +++++ .../resources/parser/cuda/kernel_function.cpp | 41 +++++++++++++++++++ 5 files changed, 62 insertions(+), 3 deletions(-) create mode 100644 cxx-squid/src/test/resources/parser/cuda/cuda-defines.h create mode 100644 cxx-squid/src/test/resources/parser/cuda/kernel_function.cpp diff --git a/cxx-squid/src/main/java/org/sonar/cxx/parser/CxxGrammarImpl.java b/cxx-squid/src/main/java/org/sonar/cxx/parser/CxxGrammarImpl.java index 0a1e71039f..4626c68c4e 100644 --- a/cxx-squid/src/main/java/org/sonar/cxx/parser/CxxGrammarImpl.java +++ b/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, ">>"))); diff --git a/cxx-squid/src/test/java/org/sonar/cxx/parser/CxxParserTest.java b/cxx-squid/src/test/java/org/sonar/cxx/parser/CxxParserTest.java index d003a79aac..c789903874 100644 --- a/cxx-squid/src/test/java/org/sonar/cxx/parser/CxxParserTest.java +++ b/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; diff --git a/cxx-squid/src/test/java/org/sonar/cxx/parser/ExpressionTest.java b/cxx-squid/src/test/java/org/sonar/cxx/parser/ExpressionTest.java index 0ef68b019d..3ff1c8ef2b 100644 --- a/cxx-squid/src/test/java/org/sonar/cxx/parser/ExpressionTest.java +++ b/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<<>>(d_data, height, width)"); } @Test diff --git a/cxx-squid/src/test/resources/parser/cuda/cuda-defines.h b/cxx-squid/src/test/resources/parser/cuda/cuda-defines.h new file mode 100644 index 0000000000..14c70c3c3e --- /dev/null +++ b/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) diff --git a/cxx-squid/src/test/resources/parser/cuda/kernel_function.cpp b/cxx-squid/src/test/resources/parser/cuda/kernel_function.cpp new file mode 100644 index 0000000000..8f59dac255 --- /dev/null +++ b/cxx-squid/src/test/resources/parser/cuda/kernel_function.cpp @@ -0,0 +1,41 @@ +#include "cuda-defines.h" + +texture tex; + +void foo() +{ + cudaArray* cu_array; + + // Allocate array + cudaChannelFormatDesc description = cudaCreateChannelDesc(); + 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; + } +}