Add CUDA support for MSVC (#426)

Tested building https://github.com/trolleyman/cuda-macros with this (the `cuda-macros-test` crate) and it builds & links correctly. Haven't had a chance to test that this runs yet, but will in the morning.

I wasn't sure that this way is the most elegant, but this seemed like the way that did the least amount of changes. I am also not sure that this is the correct way of doing this, especially regarding cross-compiling, but it gets it up and running at least.

To test you can do a `cargo build` in the root of the repo linked above. The build stuff is a bit hacky, but essentially it generates the CUDA function below & calls it.

```c
extern "C" __global__ void hello(int32_t* x, int32_t y) {
    printf("Hello from block %d, thread %d (y=%d)\n", blockIdx.x, threadIdx.x, y);
    *x = 2;
}
```

```rust
extern "C" unsafe fn hello(x: *mut i32, y: i32);
```

Author: trolleyman

src/lib.rs

@@ -209,7 +209,7 @@ impl ToolFamily {
     fn add_debug_flags(&self, cmd: &mut Tool) {
         match *self {
             ToolFamily::Msvc { .. } => {
-                cmd.push_cc_arg("/Z7".into());
+                cmd.push_cc_arg("-Z7".into());
             }
             ToolFamily::Gnu | ToolFamily::Clang => {
                 cmd.push_cc_arg("-g".into());
@@ -218,26 +218,10 @@ impl ToolFamily {
         }
     }
 
-    /// What the flag to include directories into header search path looks like
-    fn include_flag(&self) -> &'static str {
-        match *self {
-            ToolFamily::Msvc { .. } => "/I",
-            ToolFamily::Gnu | ToolFamily::Clang => "-I",
-        }
-    }
-
-    /// What the flag to request macro-expanded source output looks like
-    fn expand_flag(&self) -> &'static str {
-        match *self {
-            ToolFamily::Msvc { .. } => "/E",
-            ToolFamily::Gnu | ToolFamily::Clang => "-E",
-        }
-    }
-
     /// What the flags to enable all warnings
     fn warnings_flags(&self) -> &'static str {
         match *self {
-            ToolFamily::Msvc { .. } => "/W4",
+            ToolFamily::Msvc { .. } => "-W4",
             ToolFamily::Gnu | ToolFamily::Clang => "-Wall",
         }
     }
@@ -253,29 +237,11 @@ impl ToolFamily {
     /// What the flag to turn warning into errors
     fn warnings_to_errors_flag(&self) -> &'static str {
         match *self {
-            ToolFamily::Msvc { .. } => "/WX",
+            ToolFamily::Msvc { .. } => "-WX",
             ToolFamily::Gnu | ToolFamily::Clang => "-Werror",
         }
     }
 
-    /// NVCC-specific. Device code debug info flag. This is separate from the
-    /// debug info flag passed to the C++ compiler.
-    fn nvcc_debug_flag(&self) -> &'static str {
-        match *self {
-            ToolFamily::Msvc { .. } => unimplemented!(),
-            ToolFamily::Gnu | ToolFamily::Clang => "-G",
-        }
-    }
-
-    /// NVCC-specific. Redirect the following flag to the underlying C++
-    /// compiler.
-    fn nvcc_redirect_flag(&self) -> &'static str {
-        match *self {
-            ToolFamily::Msvc { .. } => unimplemented!(),
-            ToolFamily::Gnu | ToolFamily::Clang => "-Xcompiler",
-        }
-    }
-
     fn verbose_stderr(&self) -> bool {
         *self == ToolFamily::Clang
     }
@@ -454,12 +420,19 @@ impl Build {
 
         let mut cmd = compiler.to_command();
         let is_arm = target.contains("aarch64") || target.contains("arm");
-        command_add_output_file(&mut cmd, &obj, target.contains("msvc"), false, is_arm);
+        command_add_output_file(
+            &mut cmd,
+            &obj,
+            self.cuda,
+            target.contains("msvc"),
+            false,
+            is_arm,
+        );
 
         // We need to explicitly tell msvc not to link and create an exe
         // in the root directory of the crate
-        if target.contains("msvc") {
-            cmd.arg("/c");
+        if target.contains("msvc") && !self.cuda {
+            cmd.arg("-c");
         }
 
         cmd.arg(&src);
@@ -500,7 +473,6 @@ impl Build {
     ///     .shared_flag(true)
     ///     .compile("libfoo.so");
     /// ```
-
     pub fn shared_flag(&mut self, shared_flag: bool) -> &mut Build {
         self.shared_flag = Some(shared_flag);
         self
@@ -595,7 +567,7 @@ impl Build {
     /// Set warnings flags.
     ///
     /// Adds some flags:
-    /// - "/Wall" for MSVC.
+    /// - "-Wall" for MSVC.
     /// - "-Wall", "-Wextra" for GNU and Clang.
     ///
     /// Enabled by default.
@@ -1008,10 +980,10 @@ impl Build {
             )
         };
         let is_arm = target.contains("aarch64") || target.contains("arm");
-        command_add_output_file(&mut cmd, &obj.dst, msvc, is_asm, is_arm);
+        command_add_output_file(&mut cmd, &obj.dst, self.cuda, msvc, is_asm, is_arm);
         // armasm and armasm64 don't requrie -c option
         if !msvc || !is_asm || !is_arm {
-            cmd.arg(if msvc { "/c" } else { "-c" });
+            cmd.arg("-c");
         }
         cmd.arg(&obj.src);
 
@@ -1026,7 +998,7 @@ impl Build {
         for &(ref a, ref b) in self.env.iter() {
             cmd.env(a, b);
         }
-        cmd.arg(compiler.family.expand_flag());
+        cmd.arg("-E");
 
         assert!(
             self.files.len() <= 1,
@@ -1116,7 +1088,7 @@ impl Build {
         }
 
         for directory in self.include_directories.iter() {
-            cmd.args.push(cmd.family.include_flag().into());
+            cmd.args.push("-I".into());
             cmd.args.push(directory.into());
         }
 
@@ -1153,15 +1125,10 @@ impl Build {
         }
 
         for &(ref key, ref value) in self.definitions.iter() {
-            let lead = if let ToolFamily::Msvc { .. } = cmd.family {
-                "/"
-            } else {
-                "-"
-            };
             if let Some(ref value) = *value {
-                cmd.args.push(format!("{}D{}={}", lead, key, value).into());
+                cmd.args.push(format!("-D{}={}", key, value).into());
             } else {
-                cmd.args.push(format!("{}D{}", lead, key).into());
+                cmd.args.push(format!("-D{}", key).into());
             }
         }
 
@@ -1183,32 +1150,29 @@ impl Build {
         // If the flag is not conditioned on target variable, it belongs here :)
         match cmd.family {
             ToolFamily::Msvc { .. } => {
-                assert!(!self.cuda,
-                    "CUDA C++ compilation not supported for MSVC, yet... but you are welcome to implement it :)");
-
-                cmd.args.push("/nologo".into());
+                cmd.push_cc_arg("-nologo".into());
 
                 let crt_flag = match self.static_crt {
-                    Some(true) => "/MT",
-                    Some(false) => "/MD",
+                    Some(true) => "-MT",
+                    Some(false) => "-MD",
                     None => {
                         let features = self
                             .getenv("CARGO_CFG_TARGET_FEATURE")
                             .unwrap_or(String::new());
                         if features.contains("crt-static") {
-                            "/MT"
+                            "-MT"
                         } else {
-                            "/MD"
+                            "-MD"
                         }
                     }
                 };
-                cmd.args.push(crt_flag.into());
+                cmd.push_cc_arg(crt_flag.into());
 
                 match &opt_level[..] {
                     // Msvc uses /O1 to enable all optimizations that minimize code size.
-                    "z" | "s" | "1" => cmd.push_opt_unless_duplicate("/O1".into()),
+                    "z" | "s" | "1" => cmd.push_opt_unless_duplicate("-O1".into()),
                     // -O3 is a valid value for gcc and clang compilers, but not msvc. Cap to /O2.
-                    "2" | "3" => cmd.push_opt_unless_duplicate("/O2".into()),
+                    "2" | "3" => cmd.push_opt_unless_duplicate("-O2".into()),
                     _ => {}
                 }
             }
@@ -1226,7 +1190,10 @@ impl Build {
                     cmd.push_cc_arg("-fdata-sections".into());
                 }
                 // Disable generation of PIC on RISC-V for now: rust-lld doesn't support this yet
-                if self.pic.unwrap_or(!target.contains("windows-gnu") && !target.contains("riscv")) {
+                if self
+                    .pic
+                    .unwrap_or(!target.contains("windows-gnu") && !target.contains("riscv"))
+                {
                     cmd.push_cc_arg("-fPIC".into());
                     // PLT only applies if code is compiled with PIC support,
                     // and only for ELF targets.
@@ -1239,8 +1206,8 @@ impl Build {
 
         if self.get_debug() {
             if self.cuda {
-                let nvcc_debug_flag = cmd.family.nvcc_debug_flag().into();
-                cmd.args.push(nvcc_debug_flag);
+                // NVCC debug flag
+                cmd.args.push("-G".into());
             }
             let family = cmd.family;
             family.add_debug_flags(cmd);
@@ -1257,13 +1224,13 @@ impl Build {
                         cmd.args.push("-m64".into());
                     } else if target.contains("86") {
                         cmd.args.push("-m32".into());
-                        cmd.args.push("/arch:IA32".into());
+                        cmd.push_cc_arg("-arch:IA32".into());
                     } else {
-                        cmd.args.push(format!("--target={}", target).into());
+                        cmd.push_cc_arg(format!("--target={}", target).into());
                     }
                 } else {
                     if target.contains("i586") {
-                        cmd.args.push("/ARCH:IA32".into());
+                        cmd.push_cc_arg("-arch:IA32".into());
                     }
                 }
 
@@ -1278,7 +1245,7 @@ impl Build {
                 // Windows SDK it is required.
                 if target.contains("arm") || target.contains("thumb") {
                     cmd.args
-                        .push("/D_ARM_WINAPI_PARTITION_DESKTOP_SDK_AVAILABLE=1".into());
+                        .push("-D_ARM_WINAPI_PARTITION_DESKTOP_SDK_AVAILABLE=1".into());
                 }
             }
             ToolFamily::Gnu => {
@@ -1502,18 +1469,18 @@ impl Build {
         };
         let mut cmd = windows_registry::find(&target, tool).unwrap_or_else(|| self.cmd(tool));
         for directory in self.include_directories.iter() {
-            cmd.arg("/I").arg(directory);
+            cmd.arg("-I").arg(directory);
         }
         for &(ref key, ref value) in self.definitions.iter() {
             if let Some(ref value) = *value {
-                cmd.arg(&format!("/D{}={}", key, value));
+                cmd.arg(&format!("-D{}={}", key, value));
             } else {
-                cmd.arg(&format!("/D{}", key));
+                cmd.arg(&format!("-D{}", key));
             }
         }
 
         if target.contains("i686") || target.contains("i586") {
-            cmd.arg("/safeseh");
+            cmd.arg("-safeseh");
         }
         for flag in self.flags.iter() {
             cmd.arg(flag);
@@ -1531,9 +1498,9 @@ impl Build {
         let target = self.get_target()?;
         if target.contains("msvc") {
             let (mut cmd, program) = self.get_ar()?;
-            let mut out = OsString::from("/OUT:");
+            let mut out = OsString::from("-out:");
             out.push(dst);
-            cmd.arg(out).arg("/nologo");
+            cmd.arg(out).arg("-nologo");
 
             // Similar to https://github.com/rust-lang/rust/pull/47507
             // and https://github.com/rust-lang/rust/pull/48548
@@ -1632,19 +1599,21 @@ impl Build {
             }
         };
 
-        let min_version = std::env::var("IPHONEOS_DEPLOYMENT_TARGET")
-                .unwrap_or_else(|_| "7.0".into());
+        let min_version =
+            std::env::var("IPHONEOS_DEPLOYMENT_TARGET").unwrap_or_else(|_| "7.0".into());
 
         let sdk = match arch {
             ArchSpec::Device(arch) => {
                 cmd.args.push("-arch".into());
                 cmd.args.push(arch.into());
-                cmd.args.push(format!("-miphoneos-version-min={}", min_version).into());
+                cmd.args
+                    .push(format!("-miphoneos-version-min={}", min_version).into());
                 "iphoneos"
             }
             ArchSpec::Simulator(arch) => {
                 cmd.args.push(arch.into());
-                cmd.args.push(format!("-mios-simulator-version-min={}", min_version).into());
+                cmd.args
+                    .push(format!("-mios-simulator-version-min={}", min_version).into());
                 "iphonesimulator"
             }
         };
@@ -1776,13 +1745,14 @@ impl Build {
                     }
                 } else if target.contains("cloudabi") {
                     format!("{}-{}", target, traditional)
-                } else if target == "wasm32-wasi" ||
-                          target == "wasm32-unknown-wasi" ||
-                          target == "wasm32-unknown-unknown" {
+                } else if target == "wasm32-wasi"
+                    || target == "wasm32-unknown-wasi"
+                    || target == "wasm32-unknown-unknown"
+                {
                     "clang".to_string()
                 } else if target.contains("vxworks") {
-		    "wr-c++".to_string()
-		} else if self.get_host()? != target {
+                    "wr=c++".to_string()
+                } else if self.get_host()? != target {
                     // CROSS_COMPILE is of the form: "arm-linux-gnueabi-"
                     let cc_env = self.getenv("CROSS_COMPILE");
                     let cross_compile = cc_env.as_ref().map(|s| s.trim_right_matches('-'));
@@ -1880,6 +1850,7 @@ impl Build {
             nvcc_tool
                 .args
                 .push(format!("-ccbin={}", tool.path.display()).into());
+            nvcc_tool.family = tool.family;
             nvcc_tool
         } else {
             tool
@@ -2185,7 +2156,7 @@ impl Tool {
     /// with a "-Xcompiler" flag to get passed to the underlying C++ compiler.
     fn push_cc_arg(&mut self, flag: OsString) {
         if self.cuda {
-            self.args.push(self.family.nvcc_redirect_flag().into());
+            self.args.push("-Xcompiler".into());
         }
         self.args.push(flag);
     }
@@ -2441,13 +2412,16 @@ fn fail(s: &str) -> ! {
     std::process::exit(1);
 }
 
-fn command_add_output_file(cmd: &mut Command, dst: &Path, msvc: bool, is_asm: bool, is_arm: bool) {
-    if msvc && is_asm && is_arm {
-        cmd.arg("-o").arg(&dst);
-    } else if msvc && is_asm {
-        cmd.arg("/Fo").arg(dst);
-    } else if msvc {
-        let mut s = OsString::from("/Fo");
+fn command_add_output_file(
+    cmd: &mut Command,
+    dst: &Path,
+    cuda: bool,
+    msvc: bool,
+    is_asm: bool,
+    is_arm: bool,
+) {
+    if msvc && !cuda && !(is_asm && is_arm) {
+        let mut s = OsString::from("-Fo");
         s.push(&dst);
         cmd.arg(s);
     } else {