Browse Source

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);
```
wip-new-parallel
Callum Tolley 5 years ago
committed by Alex Crichton
parent
commit
0187762feb
  1. 162
      src/lib.rs
  2. 28
      src/windows_registry.rs
  3. 20
      tests/test.rs

162
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 {

28
src/windows_registry.rs

@ -180,8 +180,8 @@ mod impl_ {
use std::ffi::OsString;
use std::fs::File;
use std::io::Read;
use std::mem;
use std::iter;
use std::mem;
use std::path::{Path, PathBuf};
use Tool;
@ -218,7 +218,7 @@ mod impl_ {
}
}
fn vs16_instances() -> Box<Iterator<Item=PathBuf>> {
fn vs16_instances() -> Box<Iterator<Item = PathBuf>> {
let instances = if let Some(instances) = vs15_instances() {
instances
} else {
@ -236,17 +236,19 @@ mod impl_ {
}
fn find_tool_in_vs16_path(tool: &str, target: &str) -> Option<Tool> {
vs16_instances().filter_map(|path| {
let path = path.join(tool);
if !path.is_file() {
return None;
}
let mut tool = Tool::new(path);
if target.contains("x86_64") {
tool.env.push(("Platform".into(), "X64".into()));
}
Some(tool)
}).next()
vs16_instances()
.filter_map(|path| {
let path = path.join(tool);
if !path.is_file() {
return None;
}
let mut tool = Tool::new(path);
if target.contains("x86_64") {
tool.env.push(("Platform".into(), "X64".into()));
}
Some(tool)
})
.next()
}
fn find_msbuild_vs16(target: &str) -> Option<Tool> {

20
tests/test.rs

@ -290,11 +290,11 @@ fn msvc_smoke() {
test.gcc().file("foo.c").compile("foo");
test.cmd(0)
.must_have("/O2")
.must_have("-O2")
.must_have("foo.c")
.must_not_have("/Z7")
.must_have("/c")
.must_have("/MD");
.must_not_have("-Z7")
.must_have("-c")
.must_have("-MD");
test.cmd(1).must_have(test.td.path().join("foo.o"));
}
@ -303,14 +303,14 @@ fn msvc_opt_level_0() {
let test = Test::msvc();
test.gcc().opt_level(0).file("foo.c").compile("foo");
test.cmd(0).must_not_have("/O2");
test.cmd(0).must_not_have("-O2");
}
#[test]
fn msvc_debug() {
let test = Test::msvc();
test.gcc().debug(true).file("foo.c").compile("foo");
test.cmd(0).must_have("/Z7");
test.cmd(0).must_have("-Z7");
}
#[test]
@ -318,7 +318,7 @@ fn msvc_include() {
let test = Test::msvc();
test.gcc().include("foo/bar").file("foo.c").compile("foo");
test.cmd(0).must_have("/I").must_have("foo/bar");
test.cmd(0).must_have("-I").must_have("foo/bar");
}
#[test]
@ -330,7 +330,7 @@ fn msvc_define() {
.file("foo.c")
.compile("foo");
test.cmd(0).must_have("/DFOO=bar").must_have("/DBAR");
test.cmd(0).must_have("-DFOO=bar").must_have("-DBAR");
}
#[test]
@ -338,7 +338,7 @@ fn msvc_static_crt() {
let test = Test::msvc();
test.gcc().static_crt(true).file("foo.c").compile("foo");
test.cmd(0).must_have("/MT");
test.cmd(0).must_have("-MT");
}
#[test]
@ -346,5 +346,5 @@ fn msvc_no_static_crt() {
let test = Test::msvc();
test.gcc().static_crt(false).file("foo.c").compile("foo");
test.cmd(0).must_have("/MD");
test.cmd(0).must_have("-MD");
}

Loading…
Cancel
Save