Ensure that we still parse preprocessed CUDA files as CUDA when we use -save-temps option.

[tra]

Summary

Fixes PR22926.

[echristo]

Seems pretty good. Pending the driver aspect rewrite we might want a -### test that verifies the particular phases and commands that are coming out?

-eric

[tra]

Complete test for cuda-related driver pipeline construction is part of the
driver changes I'll send in a bit. Stay tuned.

[echristo]

I think that test will make sense when we get the pipeline up and going so LGTM for this part.

-eric

[tra]

Simplified/cleaned up included test case.

[echristo]

As we mentioned we'll need to split the testcase out to avoid leaving temporary files etc as well. Use -### and split accordingly to test the various features.

-eric

[tra]

Updated test case to avoid creation of temporary files.

Updating D8383: Ensure that we still parse preprocessed CUDA files as CUDA when we use

-save-temps option.

[tra]

Removed unnecessary preprocessing and run compiler checks on the file combined with -x cuda-cpp-output.

Updating D8383: Ensure that we still parse preprocessed CUDA files as CUDA when we use

-save-temps option.

[echristo]

Much better, some inline comments.

Thanks!

-eric

test/Driver/cuda-simple.cu

5

Probably best to avoid most of this comment and just say what you're testing here.

21

Can you move these up by the test that they're from?

[tra]

Cleaned up the test case.

Updating D8383: Ensure that we still parse preprocessed CUDA files as CUDA when we use

-save-temps option.

[tra]

Done.

[echristo]

LGTM. Thanks for all of the revisions!

-eric

[tra]

Committed in http://reviews.llvm.org/rL232737

include/clang/Driver/Types.def

 //  A - The type's temporary suffix should be appended when generating
 //      outputs of this type.
 
 
 // C family source language (with and without preprocessing).
 TYPE("cpp-output",               PP_C,         INVALID,         "i",     "u")
 TYPE("c",                        C,            PP_C,            "c",     "u")
 TYPE("cl",                       CL,           PP_C,            "cl",    "u")
-TYPE("cuda",                     CUDA,         PP_CXX,          "cpp",   "u")
+TYPE("cuda-cpp-output",          PP_CUDA,      INVALID,         "cui",   "u")
+TYPE("cuda",                     CUDA,         PP_CUDA,         "cu",    "u")
 TYPE("objective-c-cpp-output",   PP_ObjC,      INVALID,         "mi",    "u")
 TYPE("objc-cpp-output",          PP_ObjC_Alias, INVALID,        "mi",    "u")
 TYPE("objective-c",              ObjC,         PP_ObjC,         "m",     "u")
 TYPE("c++-cpp-output",           PP_CXX,       INVALID,         "ii",    "u")
 TYPE("c++",                      CXX,          PP_CXX,          "cpp",   "u")
 TYPE("objective-c++-cpp-output", PP_ObjCXX,    INVALID,         "mii",   "u")
 TYPE("objc++-cpp-output",        PP_ObjCXX_Alias, INVALID,      "mii",   "u")
 TYPE("objective-c++",            ObjCXX,       PP_ObjCXX,       "mm",    "u")

include/clang/Frontend/FrontendOptions.h

enum InputKind {
   IK_ObjC,
   IK_ObjCXX,
   IK_PreprocessedC,
   IK_PreprocessedCXX,
   IK_PreprocessedObjC,
   IK_PreprocessedObjCXX,
   IK_OpenCL,
   IK_CUDA,
+  IK_PreprocessedCuda,
   IK_AST,
   IK_LLVM_IR
 };
 
   
 /// \brief An input file for the front end.
 class FrontendInputFile {
   /// \brief The file name, or "-" to read from standard input.

lib/Driver/Types.cpp

 bool types::isAcceptedByClang(ID Id) {
   switch (Id) {
   default:
     return false;
 
   case TY_Asm:
   case TY_C: case TY_PP_C:
   case TY_CL:
-  case TY_CUDA:
+  case TY_CUDA: case TY_PP_CUDA:
   case TY_ObjC: case TY_PP_ObjC: case TY_PP_ObjC_Alias:
   case TY_CXX: case TY_PP_CXX:
   case TY_ObjCXX: case TY_PP_ObjCXX: case TY_PP_ObjCXX_Alias:
   case TY_CHeader: case TY_PP_CHeader:
   case TY_CLHeader:
   case TY_ObjCHeader: case TY_PP_ObjCHeader:
   case TY_CXXHeader: case TY_PP_CXXHeader:
   case TY_ObjCXXHeader: case TY_PP_ObjCXXHeader:
bool types::isCXX(ID Id) {
   switch (Id) {
   default:
     return false;
 
   case TY_CXX: case TY_PP_CXX:
   case TY_ObjCXX: case TY_PP_ObjCXX: case TY_PP_ObjCXX_Alias:
   case TY_CXXHeader: case TY_PP_CXXHeader:
   case TY_ObjCXXHeader: case TY_PP_ObjCXXHeader:
-  case TY_CUDA:
+  case TY_CUDA: case TY_PP_CUDA:
     return true;
   }
 }
 
 types::ID types::lookupTypeForExtension(const char *Ext) {
   return llvm::StringSwitch<types::ID>(Ext)
            .Case("c", TY_C)
            .Case("i", TY_PP_C)
return llvm::StringSwitch<types::ID>(Ext)
            .Case("mi", TY_PP_ObjC)
            .Case("mm", TY_ObjCXX)
            .Case("bc", TY_LLVM_BC)
            .Case("cc", TY_CXX)
            .Case("CC", TY_CXX)
            .Case("cl", TY_CL)
            .Case("cp", TY_CXX)
            .Case("cu", TY_CUDA)
+           .Case("cui", TY_PP_CUDA)
            .Case("hh", TY_CXXHeader)
            .Case("ll", TY_LLVM_IR)
            .Case("hpp", TY_CXXHeader)
            .Case("ads", TY_Ada)
            .Case("adb", TY_Ada)
            .Case("ast", TY_AST)
            .Case("c++", TY_CXX)
            .Case("C++", TY_CXX)

lib/Frontend/CompilerInvocation.cpp

DashX = llvm::StringSwitch<InputKind>(A->getValue())
       .Case("cl", IK_OpenCL)
       .Case("cuda", IK_CUDA)
       .Case("c++", IK_CXX)
       .Case("objective-c", IK_ObjC)
       .Case("objective-c++", IK_ObjCXX)
       .Case("cpp-output", IK_PreprocessedC)
       .Case("assembler-with-cpp", IK_Asm)
       .Case("c++-cpp-output", IK_PreprocessedCXX)
+      .Case("cuda-cpp-output", IK_PreprocessedCuda)
       .Case("objective-c-cpp-output", IK_PreprocessedObjC)
       .Case("objc-cpp-output", IK_PreprocessedObjC)
       .Case("objective-c++-cpp-output", IK_PreprocessedObjCXX)
       .Case("objc++-cpp-output", IK_PreprocessedObjCXX)
       .Case("c-header", IK_C)
       .Case("cl-header", IK_OpenCL)
       .Case("objective-c-header", IK_ObjC)
       .Case("c++-header", IK_CXX)
if (LangStd == LangStandard::lang_unspecified) {
     case IK_None:
     case IK_AST:
     case IK_LLVM_IR:
       llvm_unreachable("Invalid input kind!");
     case IK_OpenCL:
       LangStd = LangStandard::lang_opencl;
       break;
     case IK_CUDA:
+    case IK_PreprocessedCuda:
       LangStd = LangStandard::lang_cuda;
       break;
     case IK_Asm:
     case IK_C:
     case IK_PreprocessedC:
     case IK_ObjC:
     case IK_PreprocessedObjC:
       LangStd = LangStandard::lang_gnu11;
void CompilerInvocation::setLangDefaults(LangOptions &Opts, InputKind IK,
   if (Opts.OpenCL) {
     Opts.AltiVec = 0;
     Opts.CXXOperatorNames = 1;
     Opts.LaxVectorConversions = 0;
     Opts.DefaultFPContract = 1;
     Opts.NativeHalfType = 1;
   }
 
-  Opts.CUDA = LangStd == LangStandard::lang_cuda || IK == IK_CUDA;
+  Opts.CUDA = IK == IK_CUDA || IK == IK_PreprocessedCuda ||
+              LangStd == LangStandard::lang_cuda;
 
   // OpenCL and C++ both have bool, true, false keywords.
   Opts.Bool = Opts.OpenCL || Opts.CPlusPlus;
 
   // OpenCL has half keyword
   Opts.Half = Opts.OpenCL;
 
   // C++ has wchar_t keyword.
else {
             << A->getAsString(Args) << "C++/ObjC++";
         break;
       case IK_OpenCL:
         if (!Std.isC99())
           Diags.Report(diag::err_drv_argument_not_allowed_with)
             << A->getAsString(Args) << "OpenCL";
         break;
       case IK_CUDA:
+      case IK_PreprocessedCuda:
         if (!Std.isCPlusPlus())
           Diags.Report(diag::err_drv_argument_not_allowed_with)
             << A->getAsString(Args) << "CUDA";
         break;
       default:
         break;
       }
     }

lib/Frontend/FrontendActions.cpp

void PrintPreambleAction::ExecuteAction() {
   case IK_ObjCXX:
   case IK_OpenCL:
   case IK_CUDA:
     break;
       
   case IK_None:
   case IK_Asm:
   case IK_PreprocessedC:
+  case IK_PreprocessedCuda:
   case IK_PreprocessedCXX:
   case IK_PreprocessedObjC:
   case IK_PreprocessedObjCXX:
   case IK_AST:
   case IK_LLVM_IR:
     // We can't do anything with these.
     return;
   }

lib/Frontend/FrontendOptions.cpp

 
 InputKind FrontendOptions::getInputKindForExtension(StringRef Extension) {
   return llvm::StringSwitch<InputKind>(Extension)
     .Cases("ast", "pcm", IK_AST)
     .Case("c", IK_C)
     .Cases("S", "s", IK_Asm)
     .Case("i", IK_PreprocessedC)
     .Case("ii", IK_PreprocessedCXX)
+    .Case("cui", IK_PreprocessedCuda)
     .Case("m", IK_ObjC)
     .Case("mi", IK_PreprocessedObjC)
     .Cases("mm", "M", IK_ObjCXX)
     .Case("mii", IK_PreprocessedObjCXX)
     .Cases("C", "cc", "cp", IK_CXX)
     .Cases("cpp", "CPP", "c++", "cxx", "hpp", IK_CXX)
     .Case("cl", IK_OpenCL)
     .Case("cu", IK_CUDA)
     .Cases("ll", "bc", IK_LLVM_IR)
     .Default(IK_C);
 }

test/Driver/cuda-simple.cu

+// Verify that we can parse a simple CUDA file with or without -save-temps
+// http://llvm.org/PR22936
+// RUN: %clang  -Werror -fsyntax-only -c %s
+//
+// We can't use -fsyntax-only with -save-temps because it would create
+// a temporary file after preprocessor stage. Instead we'll verify that we
+// pass -x cuda-cpp-output to compiler after preprocessing a CUDA file
+// RUN: %clang  -Werror -### -save-temps -c %s 2>&1 | FileCheck %s
+//
+// .. and then we'll verify that compiler accepts CUDA syntax in files
+// compiled with "-x cuda-cpp-output".
+// RUN: %clang -Werror -fsyntax-only -x cuda-cpp-output -c %s
+
+int cudaConfigureCall(int, int);
+__attribute__((global)) void kernel() {}
+
+void func() {
+     kernel<<<1,1>>>();
+}
+
+// CHECK: "-cc1"
+// CHECK: "-E"
+// CHECK: "-x" "cuda"
+// CHECK-NEXT: "-cc1"
+// CHECK: "-x" "cuda-cpp-output"

test/Driver/lit.local.cfg

-config.suffixes = ['.c', '.cpp', '.h', '.m', '.mm', '.S', '.s', '.f90', '.f95']
+config.suffixes = ['.c', '.cpp', '.h', '.m', '.mm', '.S', '.s', '.f90', '.f95',
+                   '.cu']
 config.substitutions = list(config.substitutions)
 config.substitutions.insert(0,
     ('%clang_cc1',
      """*** Do not use 'clang -cc1' in Driver tests. ***""") )