[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