diff --git a/crates/rustc_codegen_nvvm/src/intrinsic.rs b/crates/rustc_codegen_nvvm/src/intrinsic.rs
index f68dcc87..cc0e949c 100644
--- a/crates/rustc_codegen_nvvm/src/intrinsic.rs
+++ b/crates/rustc_codegen_nvvm/src/intrinsic.rs
@@ -605,7 +605,7 @@ impl<'ll, 'tcx> IntrinsicCallBuilderMethods<'tcx> for Builder<'_, 'll, 'tcx> {
                 }
                 let (size, signed) = ty.int_size_and_signed(self.tcx);
                 let width = size.bits();
-                if name == sym::saturating_add || name == sym::saturating_sub {
+                let llret = if name == sym::saturating_add || name == sym::saturating_sub {
                     saturating_intrinsic_impl(
                         self,
                         width as u32,
@@ -673,7 +673,18 @@ impl<'ll, 'tcx> IntrinsicCallBuilderMethods<'tcx> for Builder<'_, 'll, 'tcx> {
                         }
                         _ => unreachable!(),
                     }
-                }
+                };
+                // `ctpop`, `ctlz`, `cttz` and their `_nonzero` variants are declared in `core` as
+                // returning `u32` for *every* operand width, while the corresponding LLVM
+                // intrinsics return a value as wide as their operand. Without this cast the
+                // oversized value is stored into the 4-byte result slot, where it is either
+                // dropped as an out-of-bounds store (`u64`/`u128`, yielding a constant `0`) or
+                // leaves the high bytes uninitialized (`u8`/`u16`).
+                //
+                // Casting to an identical type is a no-op, so the intrinsics in this arm that do
+                // return the operand width (`bswap`, `bitreverse`, `rotate_*`, `saturating_*`)
+                // are unaffected, and any intrinsic added here later is covered automatically.
+                self.intcast(llret, result.layout.llvm_type(self), false)
             }
             sym::raw_eq => {
                 use rustc_codegen_ssa::common::IntPredicate;
diff --git a/guide/src/features.md b/guide/src/features.md
index 8a3fe259..be9ba630 100644
--- a/guide/src/features.md
+++ b/guide/src/features.md
@@ -28,7 +28,7 @@ around to adding it yet.
 | Match | ✔️ |
 | Proc Macros | ✔️ |
 | Try (`?`) | ✔️ |
-| 128 bit integers | 🟨 | Basic ops should work (and are emulated), advanced intrinsics like `ctpop`, `rotate`, etc are unsupported. |
+| 128 bit integers | 🟨 | Basic ops should work (and are emulated). The bit-manipulation intrinsics (`count_ones`, `leading_zeros`, `trailing_zeros`, `swap_bytes`, `reverse_bits`, `rotate_left`/`rotate_right`) are also emulated, over the two 64-bit halves. |
 | Unions | ✔️ |
 | Iterators | ✔️ |
 | Dynamic Dispatch | ✔️ |
diff --git a/tests/compiletests/ui/dis/count_ones_u128.rs b/tests/compiletests/ui/dis/count_ones_u128.rs
new file mode 100644
index 00000000..732f5d8e
--- /dev/null
+++ b/tests/compiletests/ui/dis/count_ones_u128.rs
@@ -0,0 +1,18 @@
+// build-pass
+// compile-flags: -Cllvm-args=--disassemble-entry=count_ones_u128 --error-format=human -Cdebuginfo=0
+
+// Regression test for the emulated 128-bit path: `handle_128_bit_intrinsic` returns an `i128`,
+// which must be truncated to the `u32` that `core::intrinsics::ctpop` is declared to return.
+// Without the truncation the 16-byte store into the 4-byte result slot is discarded and both
+// `popc`s are eliminated as dead, so `u128::count_ones` silently returns `0` on device.
+//
+// The kernel below must contain two `popc.b64`s and the `add` that combines them.
+
+use cuda_std::kernel;
+
+#[kernel]
+pub unsafe fn count_ones_u128(buffer: *const u128, out: *mut u32) {
+    unsafe {
+        *out = (*buffer).count_ones();
+    }
+}
diff --git a/tests/compiletests/ui/dis/count_ones_u128.stderr b/tests/compiletests/ui/dis/count_ones_u128.stderr
new file mode 100644
index 00000000..bc43da90
--- /dev/null
+++ b/tests/compiletests/ui/dis/count_ones_u128.stderr
@@ -0,0 +1,25 @@
+.visible .entry count_ones_u128(
+    .param .u64 count_ones_u128_param_0,
+    .param .u64 count_ones_u128_param_1
+)
+{
+    .reg .b32     %r<3>;
+    .reg .b64     %rd<12>;
+
+
+    ld.param.u64     %rd1, [count_ones_u128_param_0];
+    ld.param.u64     %rd2, [count_ones_u128_param_1];
+    cvta.to.global.u64     %rd3, %rd2;
+    cvta.to.global.u64     %rd4, %rd1;
+    ld.global.v2.u64     {%rd5, %rd6}, [%rd4];
+    popc.b64     %r1, %rd5;
+    cvt.u64.u32     %rd9, %r1;
+    popc.b64     %r2, %rd6;
+    cvt.u64.u32     %rd10, %r2;
+    add.s64     %rd11, %rd9, %rd10;
+    st.global.u32     [%rd3], %rd11;
+    ret;
+
+}
+
+
diff --git a/tests/compiletests/ui/dis/count_ones_u64.rs b/tests/compiletests/ui/dis/count_ones_u64.rs
new file mode 100644
index 00000000..e8122599
--- /dev/null
+++ b/tests/compiletests/ui/dis/count_ones_u64.rs
@@ -0,0 +1,18 @@
+// build-pass
+// compile-flags: -Cllvm-args=--disassemble-entry=count_ones_u64 --error-format=human -Cdebuginfo=0
+
+// Regression test: `core::intrinsics::ctpop` returns `u32` for every operand width, so the
+// `i64` result of `llvm.ctpop.i64` has to be truncated before it is stored into the `u32`
+// result slot. Without that truncation the oversized store is dropped and the `popc` is
+// eliminated as dead, so `count_ones` silently returns `0` on device.
+//
+// The kernel below must contain a `popc.b64`.
+
+use cuda_std::kernel;
+
+#[kernel]
+pub unsafe fn count_ones_u64(buffer: *const u64, out: *mut u32) {
+    unsafe {
+        *out = (*buffer).count_ones();
+    }
+}
diff --git a/tests/compiletests/ui/dis/count_ones_u64.stderr b/tests/compiletests/ui/dis/count_ones_u64.stderr
new file mode 100644
index 00000000..d5366c8c
--- /dev/null
+++ b/tests/compiletests/ui/dis/count_ones_u64.stderr
@@ -0,0 +1,22 @@
+.visible .entry count_ones_u64(
+    .param .u64 count_ones_u64_param_0,
+    .param .u64 count_ones_u64_param_1
+)
+{
+    .reg .b32     %r<2>;
+    .reg .b64     %rd<7>;
+
+
+    ld.param.u64     %rd1, [count_ones_u64_param_0];
+    ld.param.u64     %rd2, [count_ones_u64_param_1];
+    cvta.to.global.u64     %rd3, %rd2;
+    cvta.to.global.u64     %rd4, %rd1;
+    ld.global.u64     %rd5, [%rd4];
+    popc.b64     %r1, %rd5;
+    cvt.u64.u32     %rd6, %r1;
+    st.global.u32     [%rd3], %rd6;
+    ret;
+
+}
+
+
diff --git a/tests/compiletests/ui/dis/leading_zeros_u64.rs b/tests/compiletests/ui/dis/leading_zeros_u64.rs
new file mode 100644
index 00000000..ce7279ad
--- /dev/null
+++ b/tests/compiletests/ui/dis/leading_zeros_u64.rs
@@ -0,0 +1,18 @@
+// build-pass
+// compile-flags: -Cllvm-args=--disassemble-entry=leading_zeros_u64 --error-format=human -Cdebuginfo=0
+
+// Regression test: like `ctpop`, `core::intrinsics::ctlz` returns `u32` for every operand
+// width, so the `i64` result of `llvm.ctlz.i64` has to be truncated before it is stored into
+// the `u32` result slot. Without that truncation the oversized store is dropped and the
+// `clz` is eliminated as dead, so `leading_zeros` silently returns `0` on device.
+//
+// The kernel below must contain a `clz.b64`.
+
+use cuda_std::kernel;
+
+#[kernel]
+pub unsafe fn leading_zeros_u64(buffer: *const u64, out: *mut u32) {
+    unsafe {
+        *out = (*buffer).leading_zeros();
+    }
+}
diff --git a/tests/compiletests/ui/dis/leading_zeros_u64.stderr b/tests/compiletests/ui/dis/leading_zeros_u64.stderr
new file mode 100644
index 00000000..e65c6a2a
--- /dev/null
+++ b/tests/compiletests/ui/dis/leading_zeros_u64.stderr
@@ -0,0 +1,22 @@
+.visible .entry leading_zeros_u64(
+    .param .u64 leading_zeros_u64_param_0,
+    .param .u64 leading_zeros_u64_param_1
+)
+{
+    .reg .b32     %r<2>;
+    .reg .b64     %rd<7>;
+
+
+    ld.param.u64     %rd1, [leading_zeros_u64_param_0];
+    ld.param.u64     %rd2, [leading_zeros_u64_param_1];
+    cvta.to.global.u64     %rd3, %rd2;
+    cvta.to.global.u64     %rd4, %rd1;
+    ld.global.u64     %rd5, [%rd4];
+    clz.b64     %r1, %rd5;
+    cvt.u64.u32     %rd6, %r1;
+    st.global.u32     [%rd3], %rd6;
+    ret;
+
+}
+
+