diff options
| -rw-r--r-- | include/clang/Basic/DiagnosticParseKinds.td | 1 | ||||
| -rw-r--r-- | lib/Parse/ParseExpr.cpp | 65 | ||||
| -rw-r--r-- | lib/Parse/Parser.cpp | 2 | ||||
| -rw-r--r-- | test/PCH/cuda-kernel-call.cu | 25 | ||||
| -rw-r--r-- | test/Parser/cuda-kernel-call.cu | 9 | ||||
| -rw-r--r-- | test/SemaCUDA/cuda.h | 12 | ||||
| -rw-r--r-- | test/SemaCUDA/kernel-call.cu | 15 |
7 files changed, 119 insertions, 10 deletions
diff --git a/include/clang/Basic/DiagnosticParseKinds.td b/include/clang/Basic/DiagnosticParseKinds.td index 4142fcca65..0b14ed7024 100644 --- a/include/clang/Basic/DiagnosticParseKinds.td +++ b/include/clang/Basic/DiagnosticParseKinds.td @@ -95,6 +95,7 @@ def err_expected_lsquare : Error<"expected '['">; def err_expected_rsquare : Error<"expected ']'">; def err_expected_rbrace : Error<"expected '}'">; def err_expected_greater : Error<"expected '>'">; +def err_expected_ggg : Error<"expected '>>>'">; def err_expected_semi_declaration : Error< "expected ';' at end of declaration">; def err_expected_semi_decl_list : Error< diff --git a/lib/Parse/ParseExpr.cpp b/lib/Parse/ParseExpr.cpp index 5928871987..55d2ba222f 100644 --- a/lib/Parse/ParseExpr.cpp +++ b/lib/Parse/ParseExpr.cpp @@ -1093,24 +1093,68 @@ Parser::ParsePostfixExpressionSuffix(ExprResult LHS) { break; } - case tok::l_paren: { // p-e: p-e '(' argument-expression-list[opt] ')' + case tok::l_paren: // p-e: p-e '(' argument-expression-list[opt] ')' + case tok::lesslessless: { // p-e: p-e '<<<' argument-expression-list '>>>' + // '(' argument-expression-list[opt] ')' + tok::TokenKind OpKind = Tok.getKind(); InMessageExpressionRAIIObject InMessage(*this, false); + Expr *ExecConfig = 0; + + if (OpKind == tok::lesslessless) { + ExprVector ExecConfigExprs(Actions); + CommaLocsTy ExecConfigCommaLocs; + SourceLocation LLLLoc, GGGLoc; + + LLLLoc = ConsumeToken(); + + if (ParseExpressionList(ExecConfigExprs, ExecConfigCommaLocs)) { + LHS = ExprError(); + } + + if (LHS.isInvalid()) { + SkipUntil(tok::greatergreatergreater); + } else if (Tok.isNot(tok::greatergreatergreater)) { + MatchRHSPunctuation(tok::greatergreatergreater, LLLLoc); + LHS = ExprError(); + } else { + GGGLoc = ConsumeToken(); + } + + if (!LHS.isInvalid()) { + if (ExpectAndConsume(tok::l_paren, diag::err_expected_lparen, "")) + LHS = ExprError(); + else + Loc = PrevTokLocation; + } + + if (!LHS.isInvalid()) { + ExprResult ECResult = Actions.ActOnCUDAExecConfigExpr(getCurScope(), + LLLLoc, move_arg(ExecConfigExprs), GGGLoc); + if (ECResult.isInvalid()) + LHS = ExprError(); + else + ExecConfig = ECResult.get(); + } + } else { + Loc = ConsumeParen(); + } + ExprVector ArgExprs(Actions); CommaLocsTy CommaLocs; - - Loc = ConsumeParen(); if (Tok.is(tok::code_completion)) { Actions.CodeCompleteCall(getCurScope(), LHS.get(), 0, 0); ConsumeCodeCompletionToken(); } - - if (Tok.isNot(tok::r_paren)) { - if (ParseExpressionList(ArgExprs, CommaLocs, &Sema::CodeCompleteCall, - LHS.get())) { - SkipUntil(tok::r_paren); - LHS = ExprError(); + + if (OpKind == tok::l_paren || !LHS.isInvalid()) { + if (Tok.isNot(tok::r_paren)) { + if (ParseExpressionList(ArgExprs, CommaLocs, &Sema::CodeCompleteCall, + LHS.get())) { + SkipUntil(tok::r_paren); + LHS = ExprError(); + } } } @@ -1125,7 +1169,8 @@ Parser::ParsePostfixExpressionSuffix(ExprResult LHS) { ArgExprs.size()-1 == CommaLocs.size())&& "Unexpected number of commas!"); LHS = Actions.ActOnCallExpr(getCurScope(), LHS.take(), Loc, - move_arg(ArgExprs), Tok.getLocation()); + move_arg(ArgExprs), Tok.getLocation(), + ExecConfig); ConsumeParen(); } diff --git a/lib/Parse/Parser.cpp b/lib/Parse/Parser.cpp index 820f703068..bb0966111e 100644 --- a/lib/Parse/Parser.cpp +++ b/lib/Parse/Parser.cpp @@ -126,6 +126,8 @@ SourceLocation Parser::MatchRHSPunctuation(tok::TokenKind RHSTok, case tok::r_brace : LHSName = "{"; DID = diag::err_expected_rbrace; break; case tok::r_square: LHSName = "["; DID = diag::err_expected_rsquare; break; case tok::greater: LHSName = "<"; DID = diag::err_expected_greater; break; + case tok::greatergreatergreater: + LHSName = "<<<"; DID = diag::err_expected_ggg; break; } Diag(Tok, DID); Diag(LHSLoc, diag::note_matching) << LHSName; diff --git a/test/PCH/cuda-kernel-call.cu b/test/PCH/cuda-kernel-call.cu new file mode 100644 index 0000000000..ef12c59207 --- /dev/null +++ b/test/PCH/cuda-kernel-call.cu @@ -0,0 +1,25 @@ +// RUN: %clang_cc1 -emit-pch -o %t %s +// RUN: %clang_cc1 -include-pch %t -fsyntax-only %s + +#ifndef HEADER +#define HEADER +// Header. + +#include "../SemaCUDA/cuda.h" + +void kcall(void (*kp)()) { + kp<<<1, 1>>>(); +} + +__global__ void kern() { +} + +#else +// Using the header. + +void test() { + kcall(kern); + kern<<<1, 1>>>(); +} + +#endif diff --git a/test/Parser/cuda-kernel-call.cu b/test/Parser/cuda-kernel-call.cu new file mode 100644 index 0000000000..f95ae9e619 --- /dev/null +++ b/test/Parser/cuda-kernel-call.cu @@ -0,0 +1,9 @@ +// RUN: %clang_cc1 -fsyntax-only -verify %s + +void foo(void) { + foo<<<1; // expected-error {{expected '>>>'}} expected-note {{to match this '<<<'}} + + foo<<<1,1>>>; // expected-error {{expected '('}} + + foo<<<>>>(); // expected-error {{expected expression}} +} diff --git a/test/SemaCUDA/cuda.h b/test/SemaCUDA/cuda.h index c503747820..e3aeb99ed2 100644 --- a/test/SemaCUDA/cuda.h +++ b/test/SemaCUDA/cuda.h @@ -1,7 +1,19 @@ /* Minimal declarations for CUDA support. Testing purposes only. */ +#include <stddef.h> + #define __constant__ __attribute__((constant)) #define __device__ __attribute__((device)) #define __global__ __attribute__((global)) #define __host__ __attribute__((host)) #define __shared__ __attribute__((shared)) + +struct dim3 { + unsigned x, y, z; + 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); diff --git a/test/SemaCUDA/kernel-call.cu b/test/SemaCUDA/kernel-call.cu new file mode 100644 index 0000000000..6d51695522 --- /dev/null +++ b/test/SemaCUDA/kernel-call.cu @@ -0,0 +1,15 @@ +// RUN: %clang_cc1 -fsyntax-only -verify %s + +#include "cuda.h" + +__global__ void g1(int x) {} + +template <typename T> void t1(T arg) { + g1<<<arg, arg>>>(1); +} + +int main(void) { + g1<<<1, 1>>>(42); + + t1(1); +} |