[cfe-commits] r125219 - in /cfe/trunk: include/clang/Basic/DiagnosticParseKinds.td lib/Parse/ParseExpr.cpp lib/Parse/Parser.cpp test/PCH/cuda-kernel-call.cu test/Parser/cuda-kernel-call.cu test/SemaCUDA/cuda.h test/SemaCUDA/kernel-call.cu
Peter Collingbourne
peter at pcc.me.uk
Wed Feb 9 13:12:02 PST 2011
Author: pcc
Date: Wed Feb 9 15:12:02 2011
New Revision: 125219
URL: http://llvm.org/viewvc/llvm-project?rev=125219&view=rev
Log:
Parse: add support for parsing CUDA kernel calls
Added:
cfe/trunk/test/PCH/cuda-kernel-call.cu
cfe/trunk/test/Parser/cuda-kernel-call.cu
cfe/trunk/test/SemaCUDA/kernel-call.cu
Modified:
cfe/trunk/include/clang/Basic/DiagnosticParseKinds.td
cfe/trunk/lib/Parse/ParseExpr.cpp
cfe/trunk/lib/Parse/Parser.cpp
cfe/trunk/test/SemaCUDA/cuda.h
Modified: cfe/trunk/include/clang/Basic/DiagnosticParseKinds.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticParseKinds.td?rev=125219&r1=125218&r2=125219&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticParseKinds.td (original)
+++ cfe/trunk/include/clang/Basic/DiagnosticParseKinds.td Wed Feb 9 15:12:02 2011
@@ -95,6 +95,7 @@
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<
Modified: cfe/trunk/lib/Parse/ParseExpr.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Parse/ParseExpr.cpp?rev=125219&r1=125218&r2=125219&view=diff
==============================================================================
--- cfe/trunk/lib/Parse/ParseExpr.cpp (original)
+++ cfe/trunk/lib/Parse/ParseExpr.cpp Wed Feb 9 15:12:02 2011
@@ -1093,24 +1093,68 @@
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 @@
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();
}
Modified: cfe/trunk/lib/Parse/Parser.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Parse/Parser.cpp?rev=125219&r1=125218&r2=125219&view=diff
==============================================================================
--- cfe/trunk/lib/Parse/Parser.cpp (original)
+++ cfe/trunk/lib/Parse/Parser.cpp Wed Feb 9 15:12:02 2011
@@ -126,6 +126,8 @@
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;
Added: cfe/trunk/test/PCH/cuda-kernel-call.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/PCH/cuda-kernel-call.cu?rev=125219&view=auto
==============================================================================
--- cfe/trunk/test/PCH/cuda-kernel-call.cu (added)
+++ cfe/trunk/test/PCH/cuda-kernel-call.cu Wed Feb 9 15:12:02 2011
@@ -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
Added: cfe/trunk/test/Parser/cuda-kernel-call.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/Parser/cuda-kernel-call.cu?rev=125219&view=auto
==============================================================================
--- cfe/trunk/test/Parser/cuda-kernel-call.cu (added)
+++ cfe/trunk/test/Parser/cuda-kernel-call.cu Wed Feb 9 15:12:02 2011
@@ -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}}
+}
Modified: cfe/trunk/test/SemaCUDA/cuda.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/cuda.h?rev=125219&r1=125218&r2=125219&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/cuda.h (original)
+++ cfe/trunk/test/SemaCUDA/cuda.h Wed Feb 9 15:12:02 2011
@@ -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);
Added: cfe/trunk/test/SemaCUDA/kernel-call.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/kernel-call.cu?rev=125219&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/kernel-call.cu (added)
+++ cfe/trunk/test/SemaCUDA/kernel-call.cu Wed Feb 9 15:12:02 2011
@@ -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);
+}
More information about the cfe-commits
mailing list