PTX support

- `--emit=asm --target=nvptx64-nvidia-cuda` can be used to turn a crate
  into a PTX module (a `.s` file).

- intrinsics like `__syncthreads` and `blockIdx.x` are exposed as
  `"platform-intrinsics"`.

- "cabi" has been implemented for the nvptx and nvptx64 architectures.
  i.e. `extern "C"` works.

- a new ABI, `"ptx-kernel"`. That can be used to generate "global"
  functions. Example: `extern "ptx-kernel" fn kernel() { .. }`. All
  other functions are "device" functions.

2 files changed, 18 insertions, 1 deletions

diff --git a/src/libsyntax/abi.rs b/src/libsyntax/abi.rs
index a39cac8db99..0cc62fb43a5 100644
--- a/src/libsyntax/abi.rs
+++ b/src/libsyntax/abi.rs
@@ -41,6 +41,7 @@ pub enum Abi {
     Aapcs,
     Win64,
     SysV64,
+    PtxKernel,
 
     // Multiplatform / generic ABIs
     Rust,
@@ -82,6 +83,7 @@ const AbiDatas: &'static [AbiData] = &[
     AbiData {abi: Abi::Aapcs, name: "aapcs", generic: false },
     AbiData {abi: Abi::Win64, name: "win64", generic: false },
     AbiData {abi: Abi::SysV64, name: "sysv64", generic: false },
+    AbiData {abi: Abi::PtxKernel, name: "ptx-kernel", generic: false },
 
     // Cross-platform ABIs
     AbiData {abi: Abi::Rust, name: "Rust", generic: true },
diff --git a/src/libsyntax/feature_gate.rs b/src/libsyntax/feature_gate.rs
index e04cc11f15e..625af803458 100644
--- a/src/libsyntax/feature_gate.rs
+++ b/src/libsyntax/feature_gate.rs
@@ -318,6 +318,9 @@ declare_features! (
 
     // Allow safe suggestions for potential type conversions.
     (active, safe_suggestion, "1.0.0", Some(37384)),
+
+    // `extern "ptx-*" fn()`
+    (active, abi_ptx, "1.15.0", None),
 );
 
 declare_features! (
@@ -986,7 +989,19 @@ impl<'a> PostExpansionVisitor<'a> {
                 gate_feature_post!(&self, abi_sysv64, span,
                                    "sysv64 ABI is experimental and subject to change");
             },
-            _ => {}
+            Abi::PtxKernel => {
+                gate_feature_post!(&self, abi_ptx, span,
+                                   "PTX ABIs are experimental and subject to change");
+            }
+            // Stable
+            Abi::Cdecl |
+            Abi::Stdcall |
+            Abi::Fastcall |
+            Abi::Aapcs |
+            Abi::Win64 |
+            Abi::Rust |
+            Abi::C |
+            Abi::System => {}
         }
     }
 }