NVPTX backend metabug

[japaric]

The NVPTX backend has been available since: nightly-2017-01-XX

This is a collections of bugs and TODOs related to it.

Documentation

• How to: Run Rust code on your NVIDIA GPU

Bugs

• LLVM assertion when compiling core to PTX. LLVM assertion when using i128 with the NVPTX #38824

• LLVM error when emitting PTX code with debuginfo. LLVM error when emitting PTX code with debuginfo #38785

• NVPTX: No "undefined reference" error is raised when it should be. NVPTX: No "undefined reference" error is raised when it should be #38786

• NVPTX: non-inlined functions can't be used cross crate. NVPTX: non-inlined functions can't be used cross crate #38787

Missing features

• The equivalent to CUDA's __shared__ modifier. Probably needs an RFC to land
  in the compiler as we don't have anything similar to it (AFAIK).

Stabilization

• Stabilize the nvptx targets. IOW, add them to the compiler. Candidates for merging.

• All the non-trivial kernels make use of intrinsics like blockIdx.x. These will have to be stabilized. Right now these intrinsics are implemented as "plaform-intrinsics" but that feature is unstable.

• Stabilize the "ptx-kernel" ABI. Tracking issue for the "ptx-kernel" ABI #38788

cc @rkruppe

[gnzlbg]

It would be better if instead of having to create a new module / crate, adding the #![feature(abi_ptx)] feature, and then declaring the kernels as extern "ptx-kernel" fn foo() {}, we could just handle this using #[target_feature], so that one can add kernels within a non-kernel Rust module. It would be great if this attribute could be applied to closures as well.

That way we could just write our kernels inline with normal Rust code.

#[target_feature(enabled = "ptx_device_fn")] unsafe fn a_device_fn(...) { ... }
#[target_feature(enabled = "ptx_kernel")] unsafe fn a_kernel(...) { ... a_device_fn(...) ... }

unsafe fn bar() {
   cuda::driver::launch(x, y, z, w, a_kernel);
   cuda::driver::launch(x, y, z, w, #[target_feature(enabled = "ptx_kernel")] |...| a_kernel(...) );
}

This way users can use the typical target_feature facilities to write portable code.

[gnzlbg]

@japaric @rkruppe @alexcrichton I'd like to work on this.

[hanna-kruppe]

You seem to assume a "single source" model? We don't have that currently. You have to compile your kernels as one crate for a nvptx-* target, then embed the resulting PTX asm as string/resource into a crate that's compiled for the host target. The latter crate passes it to the driver for compilation. There's no way to mix host and device code in one compilation unit like nvcc allows.

Supporting that would require novel frontend integration (novel for Rust; clang has something like this already). For example rustc would have to decide for each translation item whether it should be compiled for the host, for the device, or both -- and then combining the resulting PTX and host object files.

Additionally, even if/once we have "single source", target_feature seems inappropriate, since target_feature is for modifying the available instruction set within one target, while this here requires compiling the code for a different target altogether (and in many cases, for two targets as mentioned before). A custom attribute might be more appropriate but since this is really about "entry point for device code or not" rather than "can run on device or not" an ABI seems fine too.

[gnzlbg]

Additionally, even if/once we have "single source", target_feature seems inappropriate, since target_feature is for modifying the available instruction set within one target,

Indeed, makes sense.

Supporting that would require novel frontend integration (novel for Rust; clang has something like this already). For example rustc would have to decide for each translation item whether it should be compiled for the host, for the device, or both -- and then combining the resulting PTX and host object files.

A custom attribute might be more appropriate but since this is really about "entry point for device code or not" rather than "can run on device or not" an ABI seems fine too.

I'd like to work on enabling this via a mixture of the ABI solution to choose, e.g., the ptx or ptx-kernel ABIs, and target_feature to choose, e.g., sm30 vs sm70.

If we could have these multiple ABIs into a single source file, we could have #[target_device] and #[target_device_kernel] procedural macros that just generates copies of a function for different ABIs:

#[inline]
fn baz() { }

#[target_device(host, nvptx(sm = "40"), spirv(version = "1.0"))]
fn bar(...) { ... }

#[target_device_kernel(host, nvptx(sm = "40"), spirv(version = "1.0"))]
fn foo(...) { 
  #[device] bar(...); // device attribute indicates that this fn is a device fn
  baz(); // this function will be used as is
}

that expand to:

fn baz_host(...) { bar_host(...); baz(); }
#[target_feature(enabled = "sm40")]
extern "ptx" fn baz_nvptx(...) { bar_nvptx(...); baz();  }
extern "spriv" fn baz_spirv(...) { bar_spirv(...); baz();  }

fn foo_host(...) { bar_host(...); baz(); }
#[target_feature(enabled = "sm40")]
extern "ptx-kernel" fn foo_nvptx(...) { bar_nvptx(...); baz();  }
extern "spirv-kernell" fn foo_spirv(...) { bar_device(...); baz();  }

And then just launch the kernels using another procedural macro, e.g., kernel_launch!(foo, args...).

One cool feature of clang and nvcc is to allow whoever builds the library to easily choose the devices to target. Procedural macros could allow these via feature flags: --features target_device_nvptx_sm_35, nvptx_sm_70, ... (it's not nice, but should be at least doable).

This approach leaves the door open to doing something nicer in the language in the future, while allowing libraries to experiment with better APIs. I wonder whether these two building blocks (extern ABIs in a single source file, and #[target_feature]) are a good way to do this, and whether this is something that a nicer approach is going to need in one form or another, or whether this approach is completely wrong and there is a better alternative.

[hanna-kruppe]

Frankly, I don't see how any of the tools we have in the language now (target_feature, proc macros, ABIs) can help at all with single source support. Right now, one crate is compiled for one target, period. Subsets of the crate can tweak some parts of the target (e.g., use non-standard ABIs or enable/disable instruction set extensions) but that's a far cry from slicing out a subset of the crate, compiling it for a completely different target, and then stitching the results back together -- and that's precisely what is necessary for single-source offloading (not just CUDA, but also everything else along these lines that I've seen).

In fact the assumption that one crate == one target goes as far as rustc being literally unable to store more than one target per compilation session (Session). Not to mention how you need to generate a whole separate LLVM module for the PTX code.

Even if one attempts to minimize the amount of compiler changes needed during prototyping for faster iteration (generally a good idea) by e.g. splitting the crate into two crates with an external tool and invoking rustc twice, there is ample room for compiler hacking. Even the bare minimum of single source support requires name resolution information, and being able to use generic library code will require type system integration as well.

So in my opinion, this is a rather big feature with at least as much need for experimentation and compiler hacking and design work as SIMD intrinsics. I say this not to discourage you but because your posts so far ignore the technical challenges that are, in my opinion, the biggest obstacle to single source support.

I'm also rather puzzled by the priorities here. Before experimenting with the best way to allow users to compile their single-source applications not just for multiple CUDA devices but also for entirely different targets like SPIR-V, basic features like an equivalent to __shared__ seem like a simpler and more important first step. Again, not trying to discourage, but even the greatest most ergonomic portable offloading solution seems pointless if the kernels can't even use group-shared memory.

[gnzlbg]

Frankly, I don't see how any of the tools we have in the language now (target_feature, proc macros, ABIs) can help at all with single source support.

Oh no, I think I expressed myself wrong. I meant that once we get single source support using extern ABIs, the combination of tools that we have already available in the language can allow for pretty nice APIs.

basic features like an equivalent to shared seem like a simpler and more important first step.

I think that __shared__ can be implemented as a core::intrinsic and is thus not a big deal: fn shared() -> *mut u8 would do.

[hanna-kruppe]

I meant that once we get single source support using extern ABIs,

What does this mean? I am not aware of any plans for any kind of single source support. And what does "using extern ABIs" mean? It seems to presuppose some strategy for single source support but it's not clear to me which one (and it doesn't sound like any of the strategies that I am aware of). Finally, assuming I'm correct that single source support is not on the horizon, I'm puzzled why we're hashing out details of how it could be exposed better to the user if the basic technical prerequsites aren't even on the horizon.

I think that shared can be implemented as a core::intrinsic and is thus not a big deal: fn shared() -> *mut u8 would do.

IIUC such an intrinsic would be basically like an alloca intrinsic, which Rust has rejected in favor of better support for unsized (DST) values. So while this is a possible strategy (although I can think of some technical challenges as well) that a prototype implementation might choose, it is far from clear to me that it's the approach we'd want to adopt.

[gnzlbg]

IIUC such an intrinsic would be basically like an alloca intrinsic

How so? The kernel does not allocate anything: __shared__ just initializes a pointer with a value.

Finally, assuming I'm correct that single source support is not on the horizon

I'd like to work on enabling single source support and I'd like to enable it in such a way that it is useful.

And what does "using extern ABIs" mean?

In a single source model:

extern "ptx-kernel" unsafe fn foo(...) { ... is compiled to a ptx kernel ... }
fn bar(...) { ... is compiled for the host ... }

[hanna-kruppe]

How so? The kernel does not allocate anything: shared just initializes a pointer with a value.

It's a storage specifies in C parlance. You declare variables to live in shared memory as opposed to thread-private memory or global memory or constant memory. In pointer types it's just an optional hint that the pointee lives in shared memory, that aspect isn't even needed. What is absolutely necessary is to be able to do declare locals like __shared__ float As[BLOCK_SIZE][BLOCK_SIZE]; (from http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#shared-memory) which is like let As: [[f32; BLOCK_SIZE]; BLOCK_SIZE]; except there's just one array per thread group, accessed by all threads in the group.

[gnzlbg]

It's a storage specifies in C parlance.

Sure, but what's the point of making it a storage specifier in Rust? You can't have two variables on shared memory, that is, the following is not valid CUDA C:

__global__ void foo(float* foo) {
  __shared__ a float[];
  __shared__ b float[];  // ERROR: you can only have one pointer to shared memory per kernel
  foo[0] = a[0] + b[t0]; 
}

What is absolutely necessary is to be able to do declare locals like shared float As[BLOCK_SIZE][BLOCK_SIZE]; (from http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#shared-memory)

In particular, here you are not allocating a float [BLOCK_SIZE][BLOCK_SIZE]; region of shared memory, but assigning the pointer to shared memory to a float [BLOCK_SIZE][BLOCK_SIZE] array (C arrays are just pointers). This might lead to UB if no memory actually has been allocated, or if some other kernel reinterpreted it to have another type, or if less memory than the one required here was allocated, etc.

So IMO, independently of what the spec says, we should focus on the actual semantics of __shared__. The only thing __shared__ does is initializing some pointer to point to the shared memory region with the restriction that one cannot initialize two pointers to it. Nothing more, nothing less. It does not allocate anything, it does not guarantee that any memory exists, or was allocated, or that the data in the memory has some "type", ... no nothin'.

This

extern __shared__ a float[];

is just:

float* a = __get_ptr_to_shared_memory();

[hanna-kruppe]

I want to point out that we're pretty badly derailing this metabug. You should probably open a thread on internals.rlo if you want to take this discussion much further.

But first I would invite you to double check your facts. Many things you've said go against everything I've ever heard and seen about CUDA (and other offloading solutions, for that matter).

Sure, but what's the point of making it a storage specifier in Rust?

I'm not saying it should be a storage specifier in Rust. We don't even have such a concept at the moment. I'm just saying, it's fundamentally a variation on variable declaration (like DST locals), not something about pointers.

You can't have two variables on shared memory, that is, the following is not valid CUDA C:

???? Check the code I linked earlier, it has at least two such variables. Or pretty a random non-trivial CUDA program using shared memory. I don't know why the code you give is rejected, but I've never even seen this syntax so I'm not even sure what it means.

For all you know, As[0] will fail because no shared memory actually has been allocated.

Are you saying Nvidia's examples, as well as all the other programs using shared memory, are effectively borked? That can't be right. In fact, I'm pretty sure the size of the is recorded (provided it has a size -- again I don't know wth __shared__ a float[]; is) in the binary and the scheduler makes sure to only place as many thread groups on one core as the available shared memory permits. I know this because it means the amount of shared memory you use impacts occupancy and thus performance.

That reminds me, another way in which such intrinsic for __shared__ would be weird is that it should probably require the size to be a compile time constant (unlike alloca).