Index: include/clang/Driver/Types.def
===================================================================
--- include/clang/Driver/Types.def
+++ include/clang/Driver/Types.def
@@ -42,7 +42,8 @@
 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")
Index: include/clang/Frontend/FrontendOptions.h
===================================================================
--- include/clang/Frontend/FrontendOptions.h
+++ include/clang/Frontend/FrontendOptions.h
@@ -71,6 +71,7 @@
   IK_PreprocessedObjCXX,
   IK_OpenCL,
   IK_CUDA,
+  IK_PreprocessedCuda,
   IK_AST,
   IK_LLVM_IR
 };
Index: lib/Driver/Types.cpp
===================================================================
--- lib/Driver/Types.cpp
+++ lib/Driver/Types.cpp
@@ -85,7 +85,7 @@
   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:
@@ -122,7 +122,7 @@
   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;
   }
 }
@@ -153,6 +153,7 @@
            .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)
Index: lib/Frontend/CompilerInvocation.cpp
===================================================================
--- lib/Frontend/CompilerInvocation.cpp
+++ lib/Frontend/CompilerInvocation.cpp
@@ -990,6 +990,7 @@
       .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)
@@ -1193,6 +1194,7 @@
       LangStd = LangStandard::lang_opencl;
       break;
     case IK_CUDA:
+    case IK_PreprocessedCuda:
       LangStd = LangStandard::lang_cuda;
       break;
     case IK_Asm:
@@ -1245,7 +1247,8 @@
     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;
@@ -1360,6 +1363,7 @@
             << 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";
Index: lib/Frontend/FrontendActions.cpp
===================================================================
--- lib/Frontend/FrontendActions.cpp
+++ lib/Frontend/FrontendActions.cpp
@@ -690,6 +690,7 @@
   case IK_None:
   case IK_Asm:
   case IK_PreprocessedC:
+  case IK_PreprocessedCuda:
   case IK_PreprocessedCXX:
   case IK_PreprocessedObjC:
   case IK_PreprocessedObjCXX:
Index: lib/Frontend/FrontendOptions.cpp
===================================================================
--- lib/Frontend/FrontendOptions.cpp
+++ lib/Frontend/FrontendOptions.cpp
@@ -18,6 +18,7 @@
     .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)
Index: test/Driver/cuda-simple.cu
===================================================================
--- /dev/null
+++ test/Driver/cuda-simple.cu
@@ -0,0 +1,23 @@
+// 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
+//
+// 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
+// CHECK: "-cc1"
+// CHECK: "-E"
+// CHECK: "-x" "cuda"
+// CHECK-NEXT: "-cc1"
+// CHECK: "-x" "cuda-cpp-output"
+//
+// Verify that compiler accepts CUDA syntax 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>>>();
+}
+
Index: test/Driver/lit.local.cfg
===================================================================
--- test/Driver/lit.local.cfg
+++ test/Driver/lit.local.cfg
@@ -1,4 +1,5 @@
-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',