Skip to content

Adding a new offload_args intrinsic, which only maps arguments #150683

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
ZuseZ4 wants to merge 6 commits into
base: main
Choose a base branch
from
Open

Adding a new offload_args intrinsic, which only maps arguments #150683

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
7794484
move initialization of omp/ol runtimes into global_ctor/dtor
ZuseZ4 Jan 9, 2026
788a31e
Add an early return if handling multiple offload calls
ZuseZ4 Jan 9, 2026
7f83187
adjust testcase for new logic
ZuseZ4 Jan 9, 2026
22f1c8d
implement offload_args
ZuseZ4 Jan 11, 2026
58d1a17
add test for new offload_args intrinsic
ZuseZ4 Jan 14, 2026
1a22802
Update rustc to handle mangled functions in offload_args, adjust test
ZuseZ4 Jan 14, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion compiler/rustc_codegen_llvm/src/base.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -99,7 +99,7 @@ pub(crate) fn compile_codegen_unit(
.unstable_opts
.offload
.iter()
.any(|o| matches!(o, Offload::Host(_) | Offload::Test));
.any(|o| matches!(o, Offload::Host(_) | Offload::Test | Offload::Args));
if has_host_offload && !cx.sess().target.is_like_gpu {
cx.offload_globals.replace(Some(OffloadGlobals::declare(&cx)));
}
Expand Down
13 changes: 0 additions & 13 deletions compiler/rustc_codegen_llvm/src/builder.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -173,19 +173,6 @@ impl<'a, 'll, CX: Borrow<SCx<'ll>>> GenericBuilder<'a, 'll, CX> {
load
}
}

fn memset(&mut self, ptr: &'ll Value, fill_byte: &'ll Value, size: &'ll Value, align: Align) {
unsafe {
llvm::LLVMRustBuildMemSet(
self.llbuilder,
ptr,
align.bytes() as c_uint,
fill_byte,
size,
false,
);
}
}
}

/// Empty string, to be used where LLVM expects an instruction name, indicating
Expand Down
245 changes: 160 additions & 85 deletions compiler/rustc_codegen_llvm/src/builder/gpu_offload.rs
View file
Open in desktop

Large diffs are not rendered by default.

4 changes: 4 additions & 0 deletions compiler/rustc_codegen_llvm/src/common.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -124,6 +124,10 @@ impl<'ll, CX: Borrow<SCx<'ll>>> GenericCx<'ll, CX> {
pub(crate) fn const_null(&self, t: &'ll Type) -> &'ll Value {
unsafe { llvm::LLVMConstNull(t) }
}

pub(crate) fn const_struct(&self, elts: &[&'ll Value], packed: bool) -> &'ll Value {
struct_in_context(self.llcx(), elts, packed)
}
}

impl<'ll, 'tcx> ConstCodegenMethods for CodegenCx<'ll, 'tcx> {
Expand Down
54 changes: 48 additions & 6 deletions compiler/rustc_codegen_llvm/src/intrinsic.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,7 +30,9 @@ use tracing::debug;
use crate::abi::FnAbiLlvmExt;
use crate::builder::Builder;
use crate::builder::autodiff::{adjust_activity_to_abi, generate_enzyme_call};
use crate::builder::gpu_offload::{OffloadKernelDims, gen_call_handling, gen_define_handling};
use crate::builder::gpu_offload::{
OffloadKernelDims, gen_call_handling, gen_define_handling, register_offload,
};
use crate::context::CodegenCx;
use crate::declare::declare_raw_fn;
use crate::errors::{
Expand Down Expand Up @@ -215,7 +217,19 @@ impl<'ll, 'tcx> IntrinsicCallBuilderMethods<'tcx> for Builder<'_, 'll, 'tcx> {
let _ = tcx.dcx().emit_almost_fatal(OffloadWithoutFatLTO);
}

codegen_offload(self, tcx, instance, args);
codegen_offload(self, tcx, instance, args, false);
return Ok(());
}
sym::offload_args => {
if tcx.sess.opts.unstable_opts.offload.is_empty() {
let _ = tcx.dcx().emit_almost_fatal(OffloadWithoutEnable);
}

if tcx.sess.lto() != rustc_session::config::Lto::Fat {
let _ = tcx.dcx().emit_almost_fatal(OffloadWithoutFatLTO);
}

codegen_offload(self, tcx, instance, args, true);
return Ok(());
}
sym::is_val_statically_known => {
Expand Down Expand Up @@ -1368,6 +1382,7 @@ fn codegen_offload<'ll, 'tcx>(
tcx: TyCtxt<'tcx>,
instance: ty::Instance<'tcx>,
args: &[OperandRef<'tcx, &'ll Value>],
host: bool,
) {
let cx = bx.cx;
let fn_args = instance.args;
Expand All @@ -1390,8 +1405,18 @@ fn codegen_offload<'ll, 'tcx>(
}
};

let offload_dims = OffloadKernelDims::from_operands(bx, &args[1], &args[2]);
let args = get_args_from_tuple(bx, args[3], fn_target);
let llfn = cx.get_fn(fn_target);
let (offload_dims, args) = if host {
// If we only map arguments to the gpu and otherwise work on host code, there is no need to
// handle block or thread dimensions.
let args = get_args_from_tuple(bx, args[1], fn_target);
(None, args)
} else {
let offload_dims = OffloadKernelDims::from_operands(bx, &args[1], &args[2]);
let args = get_args_from_tuple(bx, args[3], fn_target);
(Some(offload_dims), args)
};

let target_symbol = symbol_name_for_instance_in_crate(tcx, fn_target, LOCAL_CRATE);

let sig = tcx.fn_sig(fn_target.def_id()).skip_binder().skip_binder();
Expand All @@ -1409,8 +1434,25 @@ fn codegen_offload<'ll, 'tcx>(
return;
}
};
let offload_data = gen_define_handling(&cx, &metadata, &types, target_symbol, offload_globals);
gen_call_handling(bx, &offload_data, &args, &types, &metadata, offload_globals, &offload_dims);
register_offload(cx);
let instance = rustc_middle::ty::Instance::mono(tcx, fn_target.def_id());
let fn_abi = cx.fn_abi_of_instance(instance, tcx.mk_type_list(&[]));
let host_fn_ty = fn_abi.llvm_type(cx);

let offload_data =
gen_define_handling(&cx, &metadata, &types, target_symbol.clone(), offload_globals, host);
gen_call_handling(
bx,
&offload_data,
&args,
&types,
&metadata,
offload_globals,
offload_dims.as_ref(),
host,
llfn,
host_fn_ty,
);
}

fn get_args_from_tuple<'ll, 'tcx>(
Expand Down
2 changes: 2 additions & 0 deletions compiler/rustc_hir_analysis/src/check/intrinsic.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -165,6 +165,7 @@ fn intrinsic_operation_unsafety(tcx: TyCtxt<'_>, intrinsic_id: LocalDefId) -> hi
| sym::mul_with_overflow
| sym::needs_drop
| sym::offload
| sym::offload_args
| sym::offset_of
| sym::overflow_checks
| sym::powf16
Expand Down Expand Up @@ -338,6 +339,7 @@ pub(crate) fn check_intrinsic_type(
],
param(2),
),
sym::offload_args => (3, 0, vec![param(0), param(1)], param(2)),
sym::offset => (2, 0, vec![param(0), param(1)], param(0)),
sym::arith_offset => (
1,
Expand Down
5 changes: 4 additions & 1 deletion compiler/rustc_monomorphize/src/collector/autodiff.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -15,7 +15,10 @@ pub(crate) fn collect_autodiff_fn<'tcx>(
intrinsic: IntrinsicDef,
output: &mut MonoItems<'tcx>,
) {
if intrinsic.name != rustc_span::sym::autodiff {
if intrinsic.name != rustc_span::sym::autodiff
&& intrinsic.name != rustc_span::sym::offload
&& intrinsic.name != rustc_span::sym::offload_args
{
return;
};

Expand Down
2 changes: 2 additions & 0 deletions compiler/rustc_session/src/config.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -196,6 +196,8 @@ pub enum Offload {
Device,
/// Second step in the offload pipeline, generates the host code to call kernels.
Host(String),
/// We only map arguments, but still call host (=CPU) code.
Args,
/// Test is similar to Host, but allows testing without a device artifact.
Test,
}
Expand Down
19 changes: 14 additions & 5 deletions compiler/rustc_session/src/options.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -796,7 +796,7 @@ mod desc {
"a comma-separated list of strings, with elements beginning with + or -";
pub(crate) const parse_autodiff: &str = "a comma separated list of settings: `Enable`, `PrintSteps`, `PrintTA`, `PrintTAFn`, `PrintAA`, `PrintPerf`, `PrintModBefore`, `PrintModAfter`, `PrintModFinal`, `PrintPasses`, `NoPostopt`, `LooseTypes`, `Inline`, `NoTT`";
pub(crate) const parse_offload: &str =
"a comma separated list of settings: `Host=<Absolute-Path>`, `Device`, `Test`";
"a comma separated list of settings: `Host=<Absolute-Path>`, `Device`, `Test`, `Args`";
pub(crate) const parse_comma_list: &str = "a comma-separated list of strings";
pub(crate) const parse_opt_comma_list: &str = parse_comma_list;
pub(crate) const parse_number: &str = "a number";
Expand Down Expand Up @@ -1480,6 +1480,13 @@ pub mod parse {
}
Offload::Test
}
"Args" => {
if let Some(_) = arg {
// Args does not accept a value
return false;
}
Offload::Args
}
_ => {
// FIXME(ZuseZ4): print an error saying which value is not recognized
return false;
Expand Down Expand Up @@ -2546,10 +2553,12 @@ options! {
normalize_docs: bool = (false, parse_bool, [TRACKED],
"normalize associated items in rustdoc when generating documentation"),
offload: Vec<crate::config::Offload> = (Vec::new(), parse_offload, [TRACKED],
"a list of offload flags to enable
Mandatory setting:
`=Enable`
Currently the only option available"),
"a list of offload flags to enable:
`=Device`
`=Host(path)`
`=Test`
`=Args`
Multiple options can be combined with commas."),
on_broken_pipe: OnBrokenPipe = (OnBrokenPipe::Default, parse_on_broken_pipe, [TRACKED],
"behavior of std::io::ErrorKind::BrokenPipe (SIGPIPE)"),
osx_rpath_install_name: bool = (false, parse_bool, [TRACKED],
Expand Down
1 change: 1 addition & 0 deletions compiler/rustc_span/src/symbol.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1609,6 +1609,7 @@ symbols! {
of,
off,
offload,
offload_args,
offset,
offset_of,
offset_of_enum,
Expand Down
61 changes: 61 additions & 0 deletions library/core/src/intrinsics/mod.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3384,6 +3384,67 @@ pub const fn copysignf128(x: f128, y: f128) -> f128;
#[rustc_intrinsic]
pub const fn autodiff<F, G, T: crate::marker::Tuple, R>(f: F, df: G, args: T) -> R;

/// This intrinsic maps the given args from the Host(=CPU) to a GPU device. It then calls the given
/// function. Unlike the full `offload` intrinsic, this intrinsic expects a host function, in which
/// we will replace all usages of the given host args with their device version. This enables
/// support for various GPU libraries like `cuBLAS`, `cuDNN`, or `rocBLAS`, which *must* be called
/// from the host, but expect a mixture of host and device arguments.
///
/// Type Parameters:
/// - `F`: The kernel to call. Must be a function item.
/// - `T`: A tuple of arguments passed to `f`.
/// - `R`: The return type of the kernel.
///
/// Arguments:
/// - `f`: The host function to be called.
/// - `args`: A tuple of arguments, will be mapped to the gpu and forwarded to `f`.
///
/// Example usage (pseudocode):
///
/// ```rust,ignore (pseudocode)
/// fn kernel(A: &[f32; 6], x: &[f32; 3], y: &mut [f64; 2]) {
/// core::intrinsics::offload_args(sgemv_wrapper, (A,x,y))
/// }
///
/// #[cfg(target_os = "linux")]
/// extern "C" {
/// pub fn rocblas_sgemv(
/// alpha: *const f32,
/// A: *const f32,
/// x: *const f32,
/// beta: *const f32,
/// y: *mut f32,
/// );
/// }
///
/// #[cfg(not(target_os = "linux"))]
/// fn sgemv_wrapper(A: &[f32; 6], x: &[f32; 3], y: &mut [f64; 2]) {
Copy link
Contributor

@Flakebi Flakebi Jan 21, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

A here is a reference to invalid memory, as the pointer behind A is only valid on the GPU but this code runs on the CPU (also the cfg(not(target_os should probably be cfg(target_os without the not).

It works out fine here because A is never accessed, but I think it violates Rust’s guarantees.
To be safe, offload_args needs to map the type as well. I.e.

offload_args::<_, (&T, &[T])>(|t: *const T, ts: (*const T, usize)| {}, (t, &[t]));
// etc. for other mapped types

Copy link
Member Author

@ZuseZ4 ZuseZ4 Jan 21, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yes you are right, that's what I was mentioning above to Ralf: https://github.com/rust-lang/rust/pull/150683/changes#r2706310846 That means some safe usages can be currently unsound till the follow-up pr lands, but it's a half implemented rustc intrinsic (which end-users aren't supposed to use anyway), so I'd argue it's ok. The alternative is to make this PR bigger, but it clearly already takes a while to review, so I'd rather have that separate.

Generally, I am very willing to put in significant effort to preserve noalias for the GPU kernels, which your type mapping I think does not do. If you watch the slides from my US LLVM dev talk, they show how I intend to keep it even for mut args like &mut [f32].

Copy link
Member Author

@ZuseZ4 ZuseZ4 Jan 21, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Also for the vision, I want these two intrinsics to be the default, which should (thanks to our new llvm opts) almost always be fast enough. However, we'll also offer explicit host2device, device2host memtransfers for users. oli-obk had a nice idea for a thin wrapper type that represents that a type is now on the GPU. In that case we could prevent usages on the CPU, and enforce that people only pass data to the GPU kernel which is already on the GPU.

But I see this third explicit mode (well it's not really a fully orthogonal mode, more like an extension) a bit similar to unsafe, it should hopefully be rarely necessary.

Copy link
Member

@RalfJung RalfJung Jan 21, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

That link doesn't work here so I am not sure what you are referring to. But I definitely didn't realize that the example contains unsound code, or else I would have made a fuzz about it. ;)

some safe usages can be currently unsound

The problem with the example isn't about being safe or not, it's about having UB. Examples with UB without a big fat warning sign are not okay.

Is there a better example that avoids the UB? Like, if you made all these references into raw pointers, would that work?

Also, this affects not just A but also x and y, wouldn't it? They are all GPU pointers now?

Copy link
Member Author

@ZuseZ4 ZuseZ4 Jan 21, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Ah, gh being lovely again. I was referring to my comment here:

@RalfJung I'm somewhat concerned whether people could write host code in the wrapper, which breaks if we map arguments to the GPU.

Yes, A,x,y are GPU pointers. That was why I had asked you about only supporting extern ForeignFn once the follow-up pr lands. The easiest solution for now is indeed to move this to raw pointers, till we have either of the follow-up pr's landed which make it safe.

Copy link
Member

@RalfJung RalfJung Jan 21, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The intrinsic is unsafe, so it's entirely fine if some uses of it cause UB. But the docs for the intrinsic should be clear about what exactly is or isn't UB, or if that is yet to be determined then they should be clear about that.

Also, the example shouldn't be UB. :)

/// // rocblas expects scalars to be passed as host pointers.
/// let alpha = 1.0;
/// let beta = 1.0;
/// unsafe {
/// rocblas_sgemv(
/// // Host ptr
/// &alpha as *const f32,
/// // Replaced by device ptr
/// A.as_ptr(),
/// // Replaced by device ptr
/// x.as_ptr(),
/// // Host ptr
/// &beta as *const f32,
/// // Replaced by device ptr
/// y.as_mut_ptr()
/// );
/// }
/// }
/// ```
///
/// For reference, see the Clang documentation on offloading:
/// <https://clang.llvm.org/docs/OffloadingDesign.html>.
#[rustc_nounwind]
#[rustc_intrinsic]
pub const fn offload_args<F, T: crate::marker::Tuple, R>(f: F, args: T) -> R;
Copy link
Member Author

@ZuseZ4 ZuseZ4 Jan 18, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

TODO: docs

Copy link
Member

@RalfJung RalfJung Jan 19, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yeah that's a blocker for reviews.^^

@rustbot author

Copy link
Member Author

@ZuseZ4 ZuseZ4 Jan 19, 2026 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Details ^^ Here you go.

@RalfJung I'm somewhat concerned whether people could write host code in the wrapper, which breaks if we map arguments to the GPU. I will make a follow-up PR, in which I will add support for forwarding host arguments that should not be mapped. In the case of the rocBLAS test, those would be alpha and beta. Once we have that, I think I could prohibit calling a Rust function, and enforce that f is a ForeignFn to prevent this risk, if you think it's needed?

Copy link
Member

@RalfJung RalfJung Jan 20, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I don't know enough about this to have an opinion on that question.^^

But we can bikeshed the intrinsic as a whole a bit.

  • Why do we need both offload and offload_args?
  • Given that offload also has arguments, the name doesn't seem very descriptive. Maybe offload_with_host_helper or so? Just spitballing here as I have know idea how this all actually gets used. ;)

Copy link
Member Author

@ZuseZ4 ZuseZ4 Jan 20, 2026 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

offload is there to map all arguments to the GPU, and then launch a GPU kernel. It has a couple of extra args, to tell offload (ol) how many threads and blocks we want when launching it. That gpu kernel comes from a Rust function, which we compiled to the GPU.

offload_args also maps all the arguments to the GPU, and then launches a CPU function, which accepts pointers to GPU memory. That's relying on the called CPU function independently launching a GPU kernel. A lot of highly optimized libraries like cuBLAS, rocBLAS, cuDNN, etc. already do that. If you try to launch them with offload you will get a failure, because they are not meant to be launched from the GPU. If you launch them but pass CPU pointers, you will also get a failure, because all arrays are supposed to be in GPU memory. In the past NVIDIA used to partly support calling those functions from the GPU, but they dropped support. I could probably try to unify both offload intrinsics under one name, but with autodiff we also decided to split it into autodiff_forward and autodiff_reverse since they were behaving too differently, which otherwise made it confusing.

Under the hood, both create an offload_begin_mapper which copies data to the GPU, and an offload_end_mapper, which copies them back. The main intrinsic then also has a tgt_target_kernel_launch, which calls the GPU kernel. The offload_args intrinsic just calls the host function, but as you can see in the codegen test, we swap out the pointers before. This gives us the benefit of being able to optimize them both the same, so now both can use all the LLVM profiling/debugging/visualization tools (not yet tested, but should work) and both benefit from the LLVM opts that we started writing for them.

Happy to change the intrinsic name since we expect to have a pretty macro on top of it. My motivation for offload_args was, that this intrinsic only offloads arguments, whereas the full offload intrinsic offloads both the arguments and a function to the GPU.

Does that help?

Copy link
Member

@RalfJung RalfJung Jan 20, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It sounds a bit like offload_args is a lower-level primitive that could be used to implement offload... but it also seems like that's not actually how it works?

Copy link
Member Author

@ZuseZ4 ZuseZ4 Jan 20, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

They share a lot of the implementation, that's why this PR is so small, especially once the first 3 refactor commits from a different PR are merged. If you look at the commits, you'll see that I've added a boolean host flag. If it is set to true, we directly call the host function given to the intrinsic. The compilation pipeline in this case is also shorten, and a few of our optional args are set to None, since we don't need to generate all that much info for offload.

If the host boolean is set to false, we generate a __tgt_target_kernel (I love their naming style) and launch the function given to the intrinsic through that. This requires that we have previously compiled the function to the GPU, and embedded it's bitcode into our IR.

Copy link
Member

@RalfJung RalfJung Jan 20, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Okay, so... I take it that it does not make sense to implement offset on top of offset_args in "userland" (rather than in codegen).

Happy to change the intrinsic name since we expect to have a pretty macro on top of it. My motivation for offload_args was, that this intrinsic only offloads arguments, whereas the full offload intrinsic offloads both the arguments and a function to the GPU.

I see. Happy to take input some more folks on that.

Copy link
Contributor

@Flakebi Flakebi Jan 21, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I’m not that familiar with offload, so I’ll need help on the details, but this is how I understood it:

There are two (and a half?) primitive operations involved:

  1. Converting a CPU pointer to a GPU pointer (and the other half: converting it back)
  2. Launching a kernel on a GPU with some given arguments

The current offload implements both of these. However, for using cuda/rocm libraries, you want to be able to only do 1 without also doing 2, right?

On a low level, 1 and 2 really are orthogonal operations, so it would make sense to split them into two separate things.
Though the details of actually implementing that are more complicated. offload is a rather high-level api that hides things going on underneath, including cpu vs gpu memory/pointers.

The interface of doing the copy to and back around a closure seems nice, it looks like scoped threads and similar patterns.
Does it anything else than copy memory? Then we could also call it offload_copy_mem_args, offload_mem or offload_ptrs (or maybe something with ‘scoped’?).

I think we can the expose the two primitive operations as intrinsics:

/// Transfer memory referenced by args back and forth
pub fn offload_copy_mem_args<F, T: crate::marker::Tuple, R>(f: F, args: T) -> R;

/// Call a kernel.
/// Does **not** copy memory (note: also no restriction on T)
pub fn offload_call_kernel<F>(args: T) -> R;

This composes nicely to implement offload in code:

pub fn offload<F, T: crate::marker::Tuple, R>(f: F, args: T) -> R {
    offload_copy_mem_args(|offloaded_args| {
        offload_call_kernel(f, offloaded_args)
    })
}

And it gets rid of the bool argument to trigger one or the other implementation.

Copy link
Member

@RalfJung RalfJung Jan 21, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Replying here since it seems you posted that in the wrong thread

Also for the vision, I want these two intrinsics to be the default, which should (thanks to our new llvm opts) almost always be fast enough. However, we'll also offer explicit host2device, device2host memtransfers for users. oli-obk had a nice idea for a thin wrapper type that represents that a type is now on the GPU. In that case we could prevent usages on the CPU, and enforce that people only pass data to the GPU kernel which is already on the GPU.

But I see this third explicit mode (well it's not really a fully orthogonal mode, more like an extension) a bit similar to unsafe, it should hopefully be rarely necessary.

This sounds like you are discussing the public, user-visible API. In this PR we are discussing the underlying language primitives (intrinsics) powering that API. Those don't have the same concerns. No user should ever directly call these intrinsics. The intrinsics should be designed for

  • ease of specification
  • ease of implementation in backends


/// Generates the LLVM body of a wrapper function to offload a kernel `f`.
///
/// Type Parameters:
Expand Down
Loading
Loading