From c78b572957ba0506e86196b5eb629e40d90ae6b0 Mon Sep 17 00:00:00 2001
From: nik <nlebedenko@hotmail.com>
Date: Fri, 12 Jun 2026 22:04:04 +0000
Subject: [PATCH 1/2] mir-importer: support pointer distance intrinsics

Signed-off-by: nik <nlebedenko@hotmail.com>
---
 .../terminator/intrinsics/memory.rs           | 327 +++++++++++++-
 .../src/translator/terminator/mod.rs          |  31 ++
 crates/mir-importer/src/translator/types.rs   |   7 +
 .../examples/ptr_offset_from/Cargo.lock       | 424 ++++++++++++++++++
 .../examples/ptr_offset_from/Cargo.toml       |  21 +
 .../examples/ptr_offset_from/src/main.rs      |  85 ++++
 6 files changed, 893 insertions(+), 2 deletions(-)
 create mode 100644 crates/rustc-codegen-cuda/examples/ptr_offset_from/Cargo.lock
 create mode 100644 crates/rustc-codegen-cuda/examples/ptr_offset_from/Cargo.toml
 create mode 100644 crates/rustc-codegen-cuda/examples/ptr_offset_from/src/main.rs

diff --git a/crates/mir-importer/src/translator/terminator/intrinsics/memory.rs b/crates/mir-importer/src/translator/terminator/intrinsics/memory.rs
index 2a2fdaa3ea..fd8c3e5029 100644
--- a/crates/mir-importer/src/translator/terminator/intrinsics/memory.rs
+++ b/crates/mir-importer/src/translator/terminator/intrinsics/memory.rs
@@ -9,15 +9,16 @@
 
 use super::super::helpers::{emit_goto, emit_store_result_and_goto};
 use crate::error::{TranslationErr, TranslationResult};
-use crate::translator::rvalue;
 use crate::translator::values::ValueMap;
+use crate::translator::{rvalue, types};
 use dialect_mir::attributes::MirCastKindAttr;
-use dialect_mir::ops::MirCastOp;
+use dialect_mir::ops::{MirCastOp, MirConstantOp, MirDivOp, MirSubOp};
 use dialect_nvvm::ops::{
     CvtF32x2Bf16x2Op, StmatrixM8n8X2Op, StmatrixM8n8X2TransOp, StmatrixM8n8X4Op,
     StmatrixM8n8X4TransOp,
 };
 use pliron::basic_block::BasicBlock;
+use pliron::builtin::attributes::IntegerAttr;
 use pliron::builtin::types::{IntegerType, Signedness};
 use pliron::context::{Context, Ptr};
 use pliron::input_err;
@@ -25,7 +26,10 @@ use pliron::location::{Located, Location};
 use pliron::op::Op;
 use pliron::operation::Operation;
 use pliron::r#type::Typed;
+use pliron::utils::apint::APInt;
+use pliron::value::Value;
 use rustc_public::mir;
+use std::num::NonZeroUsize;
 /// Emits `stmatrix.m8n8.x4`: Warp-cooperative matrix store (4 tiles).
 ///
 /// Stores 4 matrix tiles (32 columns) to shared memory using the warp-cooperative
@@ -381,6 +385,325 @@ pub fn emit_cvt_f32x2_bf16x2(
     )
 }
 
+#[derive(Clone, Copy)]
+enum PtrOffsetFromResult {
+    Signed,
+    Unsigned,
+}
+
+impl PtrOffsetFromResult {
+    fn result_type(
+        self,
+        ctx: &mut Context,
+    ) -> pliron::r#type::TypePtr<pliron::builtin::types::IntegerType> {
+        match self {
+            Self::Signed => types::get_isize_type(ctx),
+            Self::Unsigned => types::get_usize_type(ctx),
+        }
+    }
+
+    fn intrinsic_name(self) -> &'static str {
+        match self {
+            Self::Signed => "ptr_offset_from",
+            Self::Unsigned => "ptr_offset_from_unsigned",
+        }
+    }
+
+    fn missing_target_message(self) -> &'static str {
+        match self {
+            Self::Signed => "ptr_offset_from call without target block",
+            Self::Unsigned => "ptr_offset_from_unsigned call without target block",
+        }
+    }
+}
+
+/// Emits `core::intrinsics::ptr_offset_from::<T>(this, other) -> isize`.
+///
+/// Computes `(this.addr() - other.addr()) / size_of::<T>()`.
+#[allow(clippy::too_many_arguments)]
+pub fn emit_ptr_offset_from(
+    ctx: &mut Context,
+    body: &mir::Body,
+    args: &[mir::Operand],
+    destination: &mir::Place,
+    target: &Option<usize>,
+    block_ptr: Ptr<BasicBlock>,
+    prev_op: Option<Ptr<Operation>>,
+    value_map: &mut ValueMap,
+    block_map: &[Ptr<BasicBlock>],
+    loc: Location,
+) -> TranslationResult<Ptr<Operation>> {
+    emit_ptr_offset_from_with_result(
+        ctx,
+        body,
+        args,
+        destination,
+        target,
+        block_ptr,
+        prev_op,
+        value_map,
+        block_map,
+        loc,
+        PtrOffsetFromResult::Signed,
+    )
+}
+
+/// Emits `core::intrinsics::ptr_offset_from_unsigned::<T>(this, other) -> usize`.
+///
+/// Computes `(this.addr() - other.addr()) / size_of::<T>()`. The intrinsic
+/// contract guarantees `this >= other` and an exact multiple.
+#[allow(clippy::too_many_arguments)]
+pub fn emit_ptr_offset_from_unsigned(
+    ctx: &mut Context,
+    body: &mir::Body,
+    args: &[mir::Operand],
+    destination: &mir::Place,
+    target: &Option<usize>,
+    block_ptr: Ptr<BasicBlock>,
+    prev_op: Option<Ptr<Operation>>,
+    value_map: &mut ValueMap,
+    block_map: &[Ptr<BasicBlock>],
+    loc: Location,
+) -> TranslationResult<Ptr<Operation>> {
+    emit_ptr_offset_from_with_result(
+        ctx,
+        body,
+        args,
+        destination,
+        target,
+        block_ptr,
+        prev_op,
+        value_map,
+        block_map,
+        loc,
+        PtrOffsetFromResult::Unsigned,
+    )
+}
+
+#[allow(clippy::too_many_arguments)]
+fn emit_ptr_offset_from_with_result(
+    ctx: &mut Context,
+    body: &mir::Body,
+    args: &[mir::Operand],
+    destination: &mir::Place,
+    target: &Option<usize>,
+    block_ptr: Ptr<BasicBlock>,
+    prev_op: Option<Ptr<Operation>>,
+    value_map: &mut ValueMap,
+    block_map: &[Ptr<BasicBlock>],
+    loc: Location,
+    result_kind: PtrOffsetFromResult,
+) -> TranslationResult<Ptr<Operation>> {
+    if args.len() != 2 {
+        return input_err!(
+            loc.clone(),
+            TranslationErr::unsupported(format!(
+                "{} expects 2 arguments (this, other), got {}",
+                result_kind.intrinsic_name(),
+                args.len()
+            ))
+        );
+    }
+
+    let elem_size = pointee_size_bytes(body, &args[0], result_kind.intrinsic_name(), loc.clone())?;
+    let result_ty = result_kind.result_type(ctx);
+    let result_type = result_ty.to_ptr();
+
+    let (this_ptr, op_after_this) = rvalue::translate_operand(
+        ctx,
+        body,
+        &args[0],
+        value_map,
+        block_ptr,
+        prev_op,
+        loc.clone(),
+    )?;
+    let this_addr = emit_pointer_expose_address(
+        ctx,
+        this_ptr,
+        result_type,
+        op_after_this,
+        block_ptr,
+        loc.clone(),
+    );
+
+    let (other_ptr, op_after_other) = rvalue::translate_operand(
+        ctx,
+        body,
+        &args[1],
+        value_map,
+        block_ptr,
+        Some(this_addr),
+        loc.clone(),
+    )?;
+    let other_addr = emit_pointer_expose_address(
+        ctx,
+        other_ptr,
+        result_type,
+        op_after_other,
+        block_ptr,
+        loc.clone(),
+    );
+
+    let this_addr_val = this_addr.deref(ctx).get_result(0);
+    let other_addr_val = other_addr.deref(ctx).get_result(0);
+    let sub_op = Operation::new(
+        ctx,
+        MirSubOp::get_concrete_op_info(),
+        vec![result_type],
+        vec![this_addr_val, other_addr_val],
+        vec![],
+        0,
+    );
+    sub_op.deref_mut(ctx).set_loc(loc.clone());
+    sub_op.insert_after(ctx, other_addr);
+    let byte_diff = sub_op.deref(ctx).get_result(0);
+
+    let size_const = emit_integer_constant(ctx, result_ty, elem_size, sub_op, loc.clone())?;
+    let size_val = size_const.deref(ctx).get_result(0);
+
+    let div_op = Operation::new(
+        ctx,
+        MirDivOp::get_concrete_op_info(),
+        vec![result_type],
+        vec![byte_diff, size_val],
+        vec![],
+        0,
+    );
+    div_op.deref_mut(ctx).set_loc(loc.clone());
+    div_op.insert_after(ctx, size_const);
+    let result = div_op.deref(ctx).get_result(0);
+
+    emit_store_result_and_goto(
+        ctx,
+        destination,
+        result,
+        target,
+        block_ptr,
+        div_op,
+        value_map,
+        block_map,
+        loc,
+        result_kind.missing_target_message(),
+    )
+}
+
+fn pointee_size_bytes(
+    body: &mir::Body,
+    operand: &mir::Operand,
+    intrinsic_name: &str,
+    loc: Location,
+) -> TranslationResult<u64> {
+    use rustc_public::ty::{RigidTy, TyKind};
+
+    let operand_ty = match operand {
+        mir::Operand::Copy(place) | mir::Operand::Move(place) => place.ty(body.locals()).ok(),
+        mir::Operand::Constant(constant) => Some(constant.const_.ty()),
+        mir::Operand::RuntimeChecks(_) => None,
+    };
+    let Some(operand_ty) = operand_ty else {
+        return input_err!(
+            loc.clone(),
+            TranslationErr::unsupported(format!(
+                "{intrinsic_name}: cannot determine pointer operand type"
+            ))
+        );
+    };
+
+    let pointee = match operand_ty.kind() {
+        TyKind::RigidTy(RigidTy::RawPtr(pointee, _))
+        | TyKind::RigidTy(RigidTy::Ref(_, pointee, _)) => pointee,
+        other => {
+            return input_err!(
+                loc.clone(),
+                TranslationErr::unsupported(format!(
+                    "{intrinsic_name}: expected raw pointer or reference operand, got {other:?}"
+                ))
+            );
+        }
+    };
+
+    let layout = match pointee.layout() {
+        Ok(layout) => layout,
+        Err(err) => {
+            return input_err!(
+                loc.clone(),
+                TranslationErr::unsupported(format!(
+                    "{intrinsic_name}: failed to query pointee layout: {err:?}"
+                ))
+            );
+        }
+    };
+    let size = layout.shape().size.bytes() as u64;
+    if size == 0 {
+        return input_err!(
+            loc,
+            TranslationErr::unsupported(format!(
+                "{intrinsic_name}: zero-sized pointee type has no element distance"
+            ))
+        );
+    }
+    Ok(size)
+}
+
+fn emit_pointer_expose_address(
+    ctx: &mut Context,
+    ptr: Value,
+    result_type: Ptr<pliron::r#type::TypeObj>,
+    insert_after: Option<Ptr<Operation>>,
+    block_ptr: Ptr<BasicBlock>,
+    loc: Location,
+) -> Ptr<Operation> {
+    let cast_op = Operation::new(
+        ctx,
+        MirCastOp::get_concrete_op_info(),
+        vec![result_type],
+        vec![ptr],
+        vec![],
+        0,
+    );
+    cast_op.deref_mut(ctx).set_loc(loc);
+    MirCastOp::new(cast_op).set_attr_cast_kind(ctx, MirCastKindAttr::PointerExposeAddress);
+    if let Some(prev) = insert_after {
+        cast_op.insert_after(ctx, prev);
+    } else {
+        cast_op.insert_at_front(block_ptr, ctx);
+    }
+    cast_op
+}
+
+fn emit_integer_constant(
+    ctx: &mut Context,
+    ty: pliron::r#type::TypePtr<IntegerType>,
+    value: u64,
+    insert_after: Ptr<Operation>,
+    loc: Location,
+) -> TranslationResult<Ptr<Operation>> {
+    let value_i64 = i64::try_from(value).map_err(|_| {
+        pliron::input_error!(
+            loc.clone(),
+            TranslationErr::unsupported(format!(
+                "ptr_offset_from pointee size {value} does not fit in isize"
+            ))
+        )
+    })?;
+    let bits = ty.deref(ctx).width() as usize;
+    let apint = APInt::from_i64(value_i64, NonZeroUsize::new(bits).unwrap());
+    let size_attr = IntegerAttr::new(ty, apint);
+    let const_op = Operation::new(
+        ctx,
+        MirConstantOp::get_concrete_op_info(),
+        vec![ty.to_ptr()],
+        vec![],
+        vec![],
+        0,
+    );
+    const_op.deref_mut(ctx).set_loc(loc);
+    MirConstantOp::new(const_op).set_attr_value(ctx, size_attr);
+    const_op.insert_after(ctx, insert_after);
+    Ok(const_op)
+}
+
 /// Emits `SharedArray::index()`: Compute pointer to element in shared memory.
 ///
 /// Computes `base_ptr + index` to get a pointer to the indexed element.
diff --git a/crates/mir-importer/src/translator/terminator/mod.rs b/crates/mir-importer/src/translator/terminator/mod.rs
index 2949f868f9..31b454db85 100644
--- a/crates/mir-importer/src/translator/terminator/mod.rs
+++ b/crates/mir-importer/src/translator/terminator/mod.rs
@@ -1785,6 +1785,37 @@ fn try_dispatch_intrinsic(
             )?))
         }
 
+        "core::intrinsics::ptr_offset_from" | "std::intrinsics::ptr_offset_from" => {
+            Ok(Some(intrinsics::memory::emit_ptr_offset_from(
+                ctx,
+                body,
+                args,
+                destination,
+                target,
+                block_ptr,
+                prev_op,
+                value_map,
+                block_map,
+                loc,
+            )?))
+        }
+
+        "core::intrinsics::ptr_offset_from_unsigned"
+        | "std::intrinsics::ptr_offset_from_unsigned" => {
+            Ok(Some(intrinsics::memory::emit_ptr_offset_from_unsigned(
+                ctx,
+                body,
+                args,
+                destination,
+                target,
+                block_ptr,
+                prev_op,
+                value_map,
+                block_map,
+                loc,
+            )?))
+        }
+
         // =================================================================
         // Thread/Block Position Intrinsics
         // Support both re-exported (cuda_device::) and full paths (cuda_device::thread::)
diff --git a/crates/mir-importer/src/translator/types.rs b/crates/mir-importer/src/translator/types.rs
index 6698db6881..4f74084273 100644
--- a/crates/mir-importer/src/translator/types.rs
+++ b/crates/mir-importer/src/translator/types.rs
@@ -84,6 +84,13 @@ fn closure_upvar_tys(substs: &rustc_public::ty::GenericArgs) -> Option<Vec<rustc
     Some(upvar_tys)
 }
 
+/// Returns the `isize` type (i64 on 64-bit targets).
+pub fn get_isize_type(
+    ctx: &mut Context,
+) -> pliron::r#type::TypePtr<pliron::builtin::types::IntegerType> {
+    pliron::builtin::types::IntegerType::get(ctx, 64, pliron::builtin::types::Signedness::Signed)
+}
+
 /// Returns the 32-bit floating point type.
 pub fn get_f32_type(
     ctx: &mut Context,
diff --git a/crates/rustc-codegen-cuda/examples/ptr_offset_from/Cargo.lock b/crates/rustc-codegen-cuda/examples/ptr_offset_from/Cargo.lock
new file mode 100644
index 0000000000..cff073d381
--- /dev/null
+++ b/crates/rustc-codegen-cuda/examples/ptr_offset_from/Cargo.lock
@@ -0,0 +1,424 @@
+# This file is automatically @generated by Cargo.
+# It is not intended for manual editing.
+version = 4
+
+[[package]]
+name = "aho-corasick"
+version = "1.1.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ddd31a130427c27518df266943a5308ed92d4b226cc639f5a8f1002816174301"
+dependencies = [
+ "memchr",
+]
+
+[[package]]
+name = "anyhow"
+version = "1.0.102"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "7f202df86484c868dbad7eaa557ef785d5c66295e41b460ef922eca0723b842c"
+
+[[package]]
+name = "bindgen"
+version = "0.71.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "5f58bf3d7db68cfbac37cfc485a8d711e87e064c3d0fe0435b92f7a407f9d6b3"
+dependencies = [
+ "bitflags",
+ "cexpr",
+ "clang-sys",
+ "itertools",
+ "log",
+ "prettyplease",
+ "proc-macro2",
+ "quote",
+ "regex",
+ "rustc-hash",
+ "shlex",
+ "syn",
+]
+
+[[package]]
+name = "bitflags"
+version = "2.13.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b4388bee8683e3d04af747c73422af53102d2bd24d9eadb6cbc100baef4b43f8"
+
+[[package]]
+name = "cexpr"
+version = "0.6.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "6fac387a98bb7c37292057cffc56d62ecb629900026402633ae9160df93a8766"
+dependencies = [
+ "nom",
+]
+
+[[package]]
+name = "cfg-if"
+version = "1.0.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "9330f8b2ff13f34540b44e946ef35111825727b38d33286ef986142615121801"
+
+[[package]]
+name = "clang-sys"
+version = "1.8.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0b023947811758c97c59bf9d1c188fd619ad4718dcaa767947df1cadb14f39f4"
+dependencies = [
+ "glob",
+ "libc",
+ "libloading",
+]
+
+[[package]]
+name = "crc32fast"
+version = "1.5.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "9481c1c90cbf2ac953f07c8d4a58aa3945c425b7185c9154d67a65e4230da511"
+dependencies = [
+ "cfg-if",
+]
+
+[[package]]
+name = "crunchy"
+version = "0.2.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "460fbee9c2c2f33933d720630a6a0bac33ba7053db5344fac858d4b8952d77d5"
+
+[[package]]
+name = "cuda-bindings"
+version = "0.2.1"
+dependencies = [
+ "bindgen",
+]
+
+[[package]]
+name = "cuda-core"
+version = "0.2.1"
+dependencies = [
+ "anyhow",
+ "cuda-bindings",
+ "half",
+ "oxide-artifacts",
+]
+
+[[package]]
+name = "cuda-device"
+version = "0.2.1"
+dependencies = [
+ "cuda-macros",
+]
+
+[[package]]
+name = "cuda-host"
+version = "0.2.1"
+dependencies = [
+ "cuda-core",
+ "cuda-macros",
+ "half",
+ "libnvvm-sys",
+ "nvjitlink-sys",
+ "thiserror",
+]
+
+[[package]]
+name = "cuda-macros"
+version = "0.2.1"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "reserved-oxide-symbols",
+ "syn",
+]
+
+[[package]]
+name = "either"
+version = "1.16.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "91622ff5e7162018101f2fea40d6ebf4a78bbe5a49736a2020649edf9693679e"
+
+[[package]]
+name = "equivalent"
+version = "1.0.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "877a4ace8713b0bcf2a4e7eec82529c029f1d0619886d18145fea96c3ffe5c0f"
+
+[[package]]
+name = "foldhash"
+version = "0.1.5"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d9c4f5dac5e15c24eb999c26181a6ca40b39fe946cbe4c263c7209467bc83af2"
+
+[[package]]
+name = "glob"
+version = "0.3.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0cc23270f6e1808e30a928bdc84dea0b9b4136a8bc82338574f23baf47bbd280"
+
+[[package]]
+name = "half"
+version = "2.7.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "6ea2d84b969582b4b1864a92dc5d27cd2b77b622a8d79306834f1be5ba20d84b"
+dependencies = [
+ "cfg-if",
+ "crunchy",
+ "zerocopy",
+]
+
+[[package]]
+name = "hashbrown"
+version = "0.15.5"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "9229cfe53dfd69f0609a49f65461bd93001ea1ef889cd5529dd176593f5338a1"
+dependencies = [
+ "foldhash",
+]
+
+[[package]]
+name = "hashbrown"
+version = "0.17.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ed5909b6e89a2db4456e54cd5f673791d7eca6732202bbf2a9cc504fe2f9b84a"
+
+[[package]]
+name = "indexmap"
+version = "2.14.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d466e9454f08e4a911e14806c24e16fba1b4c121d1ea474396f396069cf949d9"
+dependencies = [
+ "equivalent",
+ "hashbrown 0.17.1",
+]
+
+[[package]]
+name = "itertools"
+version = "0.13.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "413ee7dfc52ee1a4949ceeb7dbc8a33f2d6c088194d9f922fb8318faf1f01186"
+dependencies = [
+ "either",
+]
+
+[[package]]
+name = "libc"
+version = "0.2.186"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "68ab91017fe16c622486840e4c83c9a37afeff978bd239b5293d61ece587de66"
+
+[[package]]
+name = "libloading"
+version = "0.8.9"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d7c4b02199fee7c5d21a5ae7d8cfa79a6ef5bb2fc834d6e9058e89c825efdc55"
+dependencies = [
+ "cfg-if",
+ "windows-link",
+]
+
+[[package]]
+name = "libnvvm-sys"
+version = "0.2.1"
+dependencies = [
+ "libloading",
+ "thiserror",
+]
+
+[[package]]
+name = "log"
+version = "0.4.32"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "953f07c43838f8e6f9758cab68bf5bed85465e7587ebe0b823f1bcd81978ad3a"
+
+[[package]]
+name = "memchr"
+version = "2.8.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "88904434abc2901f197fe8cc55f0445e7ded921dba5911dad2e2b39b48e663c4"
+
+[[package]]
+name = "minimal-lexical"
+version = "0.2.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "68354c5c6bd36d73ff3feceb05efa59b6acb7626617f4962be322a825e61f79a"
+
+[[package]]
+name = "nom"
+version = "7.1.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d273983c5a657a70a3e8f2a01329822f3b8c8172b73826411a55751e404a0a4a"
+dependencies = [
+ "memchr",
+ "minimal-lexical",
+]
+
+[[package]]
+name = "nvjitlink-sys"
+version = "0.2.1"
+dependencies = [
+ "libloading",
+ "thiserror",
+]
+
+[[package]]
+name = "object"
+version = "0.36.7"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "62948e14d923ea95ea2c7c86c71013138b66525b86bdc08d2dcc262bdb497b87"
+dependencies = [
+ "crc32fast",
+ "hashbrown 0.15.5",
+ "indexmap",
+ "memchr",
+]
+
+[[package]]
+name = "oxide-artifacts"
+version = "0.2.1"
+dependencies = [
+ "object",
+]
+
+[[package]]
+name = "prettyplease"
+version = "0.2.37"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "479ca8adacdd7ce8f1fb39ce9ecccbfe93a3f1344b3d0d97f20bc0196208f62b"
+dependencies = [
+ "proc-macro2",
+ "syn",
+]
+
+[[package]]
+name = "proc-macro2"
+version = "1.0.106"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "8fd00f0bb2e90d81d1044c2b32617f68fcb9fa3bb7640c23e9c748e53fb30934"
+dependencies = [
+ "unicode-ident",
+]
+
+[[package]]
+name = "ptr_offset_from"
+version = "0.1.0"
+dependencies = [
+ "cuda-core",
+ "cuda-device",
+ "cuda-host",
+]
+
+[[package]]
+name = "quote"
+version = "1.0.45"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "41f2619966050689382d2b44f664f4bc593e129785a36d6ee376ddf37259b924"
+dependencies = [
+ "proc-macro2",
+]
+
+[[package]]
+name = "regex"
+version = "1.12.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "f1292b7759ae1cb9ec195452d1390a074f0cd8541ab7a5a8c31cd6db45d4a6ba"
+dependencies = [
+ "aho-corasick",
+ "memchr",
+ "regex-automata",
+ "regex-syntax",
+]
+
+[[package]]
+name = "regex-automata"
+version = "0.4.14"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "6e1dd4122fc1595e8162618945476892eefca7b88c52820e74af6262213cae8f"
+dependencies = [
+ "aho-corasick",
+ "memchr",
+ "regex-syntax",
+]
+
+[[package]]
+name = "regex-syntax"
+version = "0.8.11"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d6f6ff9a378485b298a5286656da665ba74413d36db0979633275d2e708145d4"
+
+[[package]]
+name = "reserved-oxide-symbols"
+version = "0.2.1"
+
+[[package]]
+name = "rustc-hash"
+version = "2.1.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "94300abf3f1ae2e2b8ffb7b58043de3d399c73fa6f4b73826402a5c457614dbe"
+
+[[package]]
+name = "shlex"
+version = "1.3.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0fda2ff0d084019ba4d7c6f371c95d8fd75ce3524c3cb8fb653a3023f6323e64"
+
+[[package]]
+name = "syn"
+version = "2.0.117"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e665b8803e7b1d2a727f4023456bbbbe74da67099c585258af0ad9c5013b9b99"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "unicode-ident",
+]
+
+[[package]]
+name = "thiserror"
+version = "2.0.18"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "4288b5bcbc7920c07a1149a35cf9590a2aa808e0bc1eafaade0b80947865fbc4"
+dependencies = [
+ "thiserror-impl",
+]
+
+[[package]]
+name = "thiserror-impl"
+version = "2.0.18"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ebc4ee7f67670e9b64d05fa4253e753e016c6c95ff35b89b7941d6b856dec1d5"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "syn",
+]
+
+[[package]]
+name = "unicode-ident"
+version = "1.0.24"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e6e4313cd5fcd3dad5cafa179702e2b244f760991f45397d14d4ebf38247da75"
+
+[[package]]
+name = "windows-link"
+version = "0.2.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "f0805222e57f7521d6a62e36fa9163bc891acd422f971defe97d64e70d0a4fe5"
+
+[[package]]
+name = "zerocopy"
+version = "0.8.52"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ce1022995ff5ff5d841ad7d994facc23098cd40152f2c1d11cd607c6f530653f"
+dependencies = [
+ "zerocopy-derive",
+]
+
+[[package]]
+name = "zerocopy-derive"
+version = "0.8.52"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "1ae7f38b72ec2a254e2b87ef277cf2cd4fb97cbebf944faa6f33354da0867930"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "syn",
+]
diff --git a/crates/rustc-codegen-cuda/examples/ptr_offset_from/Cargo.toml b/crates/rustc-codegen-cuda/examples/ptr_offset_from/Cargo.toml
new file mode 100644
index 0000000000..4d12083f1b
--- /dev/null
+++ b/crates/rustc-codegen-cuda/examples/ptr_offset_from/Cargo.toml
@@ -0,0 +1,21 @@
+[package]
+name = "ptr_offset_from"
+version = "0.1.0"
+edition = "2024"
+license = "Apache-2.0"
+
+# Standalone crate, not part of the parent workspace.
+[workspace]
+
+# Exercises raw pointer distance intrinsics:
+# - `ptr_offset_from_unsigned::<T>` via `offset_from_unsigned`
+# - `ptr_offset_from::<T>` via `offset_from`
+# - byte-oriented wrappers via `byte_offset_from*`
+#
+# Build and run with:
+#   cargo oxide run ptr_offset_from
+
+[dependencies]
+cuda-device = { path = "../../../cuda-device" }
+cuda-host = { path = "../../../cuda-host" }
+cuda-core = { path = "../../../cuda-core" }
diff --git a/crates/rustc-codegen-cuda/examples/ptr_offset_from/src/main.rs b/crates/rustc-codegen-cuda/examples/ptr_offset_from/src/main.rs
new file mode 100644
index 0000000000..944dca5de2
--- /dev/null
+++ b/crates/rustc-codegen-cuda/examples/ptr_offset_from/src/main.rs
@@ -0,0 +1,85 @@
+// Copyright (c) 2024-2026 NVIDIA CORPORATION. All rights reserved.
+// SPDX-License-Identifier: Apache-2.0
+
+//! Regression coverage for raw pointer distance intrinsics.
+//!
+//! The stable raw pointer methods lower through rustc intrinsics in MIR:
+//! `offset_from_unsigned` uses `ptr_offset_from_unsigned`, while
+//! `offset_from` and the byte-oriented wrappers use `ptr_offset_from`.
+//! cuda-oxide must compute the byte address difference and divide by the
+//! pointee size, preserving signed results for negative distances.
+
+use cuda_core::{CudaContext, DeviceBuffer, LaunchConfig};
+use cuda_device::{DisjointSlice, cuda_module, kernel, thread};
+
+#[cuda_module]
+mod kernels {
+    use super::*;
+
+    #[kernel]
+    pub fn pointer_distances(values: &[u32], lo: usize, hi: usize, mut out: DisjointSlice<i64>) {
+        if thread::index_1d().get() == 0 {
+            let base = values.as_ptr();
+            unsafe {
+                let lo_ptr = base.add(lo);
+                let hi_ptr = base.add(hi);
+
+                *out.get_unchecked_mut(0) = hi_ptr.offset_from_unsigned(lo_ptr) as i64;
+                *out.get_unchecked_mut(1) = lo_ptr.offset_from(hi_ptr) as i64;
+                *out.get_unchecked_mut(2) = hi_ptr.byte_offset_from(lo_ptr) as i64;
+                *out.get_unchecked_mut(3) = hi_ptr.byte_offset_from_unsigned(lo_ptr) as i64;
+            }
+        }
+    }
+}
+
+const SLOTS: usize = 4;
+
+fn main() {
+    let ctx = CudaContext::new(0).expect("ctx");
+    let stream = ctx.default_stream();
+    let module = kernels::load(&ctx).expect("load");
+
+    let values: Vec<u32> = (0..16).map(|v| v * 11).collect();
+    let lo = 3usize;
+    let hi = 9usize;
+
+    let values_dev = DeviceBuffer::from_host(&stream, &values).unwrap();
+    let mut out_dev = DeviceBuffer::<i64>::zeroed(&stream, SLOTS).unwrap();
+
+    module
+        .pointer_distances(
+            &stream,
+            LaunchConfig::for_num_elems(1),
+            &values_dev,
+            lo,
+            hi,
+            &mut out_dev,
+        )
+        .expect("launch pointer_distances");
+
+    let got = out_dev.to_host_vec(&stream).unwrap();
+    let element_distance = (hi - lo) as i64;
+    let byte_distance = ((hi - lo) * core::mem::size_of::<u32>()) as i64;
+    let expected = [
+        element_distance,
+        -element_distance,
+        byte_distance,
+        byte_distance,
+    ];
+
+    let mut errors = 0;
+    for (i, (&actual, &want)) in got.iter().zip(expected.iter()).enumerate() {
+        if actual != want {
+            eprintln!("slot {i}: expected {want}, got {actual}");
+            errors += 1;
+        }
+    }
+
+    if errors == 0 {
+        println!("SUCCESS: pointer distances are correct");
+    } else {
+        println!("FAILURE: {errors} pointer distance mismatches");
+        std::process::exit(1);
+    }
+}

From 316cb67840a31d16d7c19c19d0c66484e1ea4583 Mon Sep 17 00:00:00 2001
From: nihalpasham <nihalp@nvidia.com>
Date: Tue, 16 Jun 2026 14:47:30 +0530
Subject: [PATCH 2/2] docs: note pointer-distance intrinsics in
 supported-features

---
 cuda-oxide-book/appendix/supported-features.md | 1 +
 1 file changed, 1 insertion(+)

diff --git a/cuda-oxide-book/appendix/supported-features.md b/cuda-oxide-book/appendix/supported-features.md
index 2a9eae24c1..74175128d9 100644
--- a/cuda-oxide-book/appendix/supported-features.md
+++ b/cuda-oxide-book/appendix/supported-features.md
@@ -17,6 +17,7 @@ roadmap, **N/A** = not applicable or no identified need.
 | HMM / Unified Memory Management | **Full** | GPU directly reads/writes host memory without `cudaMemcpy`. Reference captures in closures leverage HMM for host pointer access. Requires Turing+ GPU, Linux 6.1.24+, CUDA 12.2+. |
 | Unified Struct ABI (no `#[repr(C)]`) | **Full** | Device struct layout matches host exactly. The compiler queries rustc's actual layout and reproduces it with explicit padding in LLVM IR. Works with `#[repr(Rust)]` default. |
 | Dynamic Layout Matching | **Full** | Compiler queries rustc's `fields_by_offset_order()` and byte offsets, builds LLVM structs with correct field order and explicit padding bytes. Independent of LLVM's datalayout. |
+| Pointer Distance (`offset_from`) | **Full** | `ptr_offset_from` / `ptr_offset_from_unsigned` intrinsics (and the `offset_from`, `offset_from_unsigned`, `byte_offset_from`, `byte_offset_from_unsigned` methods) lower to an address difference divided by the rustc-reported pointee size, returning `isize` (signed) or `usize` (unsigned). Errors on a zero-sized pointee. |
 
 ## Compiler: Type System