ICE when using nalgebra Matrix3 in rust-cuda kernel: "failed to parse bitcode for LTO module" #222

Activity

1. LegNeato

   added

   bugSomething isn't workingSomething isn't working

   on May 31, 2025
2. LegNeato commented on May 31, 2025

   

   LegNeato

   on May 31, 2025

   Contributor

   More actions

   Thanks for the report!
3. Neutron3529 commented on Aug 22, 2025

   

   Neutron3529

   on Aug 22, 2025

   Last edited by Neutron3529

   More actions

   I wonder whether you're still learning writting CUDA code with Rust.

   I happened to have some spare time to try compiling the code with my small crate cuda_min

   The GPU side code is shown as follows:

   #![cfg(target_arch = "nvptx64")]
   #![no_std]
   #![no_main]
   #![feature(abi_ptx, stdarch_nvptx, asm_experimental_arch)]
   
   use core::arch::nvptx::*;
   use nalgebra::Matrix3;
   
   #[allow(unused_imports)] // for its panic-handler
   use cuda_min;
   
   #[unsafe(no_mangle)]
   pub unsafe extern "ptx-kernel" fn mat_mul(
       a: *const Matrix3<u8>,
       b: *const Matrix3<u8>,
       c: *mut Matrix3<u8>,
   ) {
       unsafe {
           let x = _thread_idx_x();
           let y = _thread_idx_y();
           let a = &*a;
           let b = &*b;
           let c = &mut *c;
   
           c[(x as usize, y as usize)] = (a.row(x as usize) * b.column(y as usize)).sum();
       }
   }

   which compiles and generate following PTX code

    
           PTX Code:
           //
           // Generated by LLVM NVPTX Back-End
           //
       .version 8.7
       .target sm_120
       .address_size 64
       
               // .globl       mat_mul                 // -- Begin function mat_mul
       .extern .func  (.param .b32 func_retval0) vprintf
       (
               .param .b64 vprintf_param_0,
               .param .b64 vprintf_param_1
       )
       ;
       .func _ZN4core9panicking9panic_fmt17ha964c0166b5de554E
       (
               .param .b64 _ZN4core9panicking9panic_fmt17ha964c0166b5de554E_param_0,
               .param .b64 _ZN4core9panicking9panic_fmt17ha964c0166b5de554E_param_1
       )
       .noreturn;
       .func  (.param .b32 func_retval0) memcmp
       (
               .param .b64 memcmp_param_0,
               .param .b64 memcmp_param_1,
               .param .b64 memcmp_param_2
       )
       ;
       .global .align 1 .b8 anon_$_6995e9b5afb94c728c982fc65adc2dcb_$_0[29] = {77, 97, 116, 114, 105, 120, 32, 115, 108, 105, 99, 105, 110, 103, 32, 111, 117, 116, 32, 111, 102, 32, 98, 111, 117, 110, 100, 115, 46};
       .global .align 8 .u64 anon_$_6995e9b5afb94c728c982fc65adc2dcb_$_1[2] = {generic(anon_$_6995e9b5afb94c728c982fc65adc2dcb_$_0), 29};
       .global .align 1 .b8 anon_$_6995e9b5afb94c728c982fc65adc2dcb_$_2[111] = {47, 104, 111, 109, 101, 47, 110, 101, 117, 116, 114, 111, 110, 47, 46, 99, 97, 114, 103, 111, 47, 114, 101, 103, 105, 115, 116, 114, 121, 47, 115, 114, 99, 47, 109, 105, 114, 114, 111, 114, 115, 46, 117, 115, 116, 99, 46, 101, 100, 117, 46, 99, 110, 45, 51, 56, 100, 48, 101, 53, 101, 98, 53, 100, 97, 50, 97, 98, 97, 101, 47, 110, 97, 108, 103, 101, 98, 114, 97, 45, 48, 46, 51, 52, 46, 48, 47, 115, 114, 99, 47, 98, 97, 115, 101, 47, 109, 97, 116, 114, 105, 120, 95, 118, 105, 101, 119, 46, 114, 115};
       .global .align 8 .u64 anon_$_6995e9b5afb94c728c982fc65adc2dcb_$_3[3] = {generic(anon_$_6995e9b5afb94c728c982fc65adc2dcb_$_2), 110, 38654705954};
       .global .align 8 .u64 anon_$_6995e9b5afb94c728c982fc65adc2dcb_$_4[3] = {generic(anon_$_6995e9b5afb94c728c982fc65adc2dcb_$_2), 110, 38654705958};
       .global .align 1 .b8 anon_$_a44653bcb419bc498496063559ae8ca0_$_0[51] = {98, 108, 111, 99, 107, 32, 40, 37, 100, 44, 32, 37, 100, 44, 32, 37, 100, 41, 32, 116, 104, 114, 101, 97, 100, 32, 40, 37, 100, 44, 32, 37, 100, 44, 32, 37, 100, 41, 32, 112, 97, 110, 105, 99, 115, 58, 32, 37, 115, 10};
       .global .align 1 .b8 anon_$_a44653bcb419bc498496063559ae8ca0_$_1[55] = {98, 108, 111, 99, 107, 32, 40, 37, 100, 44, 32, 37, 100, 44, 32, 37, 100, 41, 32, 116, 104, 114, 101, 97, 100, 32, 40, 37, 100, 44, 32, 37, 100, 44, 32, 37, 100, 41, 58, 32, 112, 97, 110, 105, 99, 32, 111, 99, 99, 111, 114, 115, 46, 10};
       .global .align 1 .b8 anon_$_a44653bcb419bc498496063559ae8ca0_$_2[75] = {98, 108, 111, 99, 107, 32, 40, 37, 100, 44, 32, 37, 100, 44, 32, 37, 100, 41, 32, 116, 104, 114, 101, 97, 100, 32, 40, 37, 100, 44, 32, 37, 100, 44, 32, 37, 100, 41, 58, 32, 32, 97, 116, 32, 102, 105, 108, 101, 58, 32, 37, 115, 44, 32, 108, 105, 110, 101, 58, 32, 37, 100, 32, 99, 111, 108, 117, 109, 110, 58, 32, 37, 100, 10};
                                               // @mat_mul
       .visible .entry mat_mul(
               .param .u64 .ptr .align 1 mat_mul_param_0,
               .param .u64 .ptr .align 1 mat_mul_param_1,
               .param .u64 .ptr .align 1 mat_mul_param_2
       )
       {
               .local .align 8 .b8     __local_depot0[48];
               .reg .b64       %SP;
               .reg .b64       %SPL;
               .reg .pred      %p<3>;
               .reg .b16       %rs<10>;
               .reg .b32       %r<4>;
               .reg .b64       %rd<33>;
       
       // %bb.0:
               mov.u64         %SPL, __local_depot0;
               cvta.local.u64  %SP, %SPL;
               ld.param.u64    %rd6, [mat_mul_param_0];
               ld.param.u64    %rd7, [mat_mul_param_2];
               cvta.to.global.u64      %rd1, %rd7;
               ld.param.u64    %rd8, [mat_mul_param_1];
               cvta.to.global.u64      %rd2, %rd8;
               cvta.to.global.u64      %rd3, %rd6;
               add.u64         %rd9, %SP, 0;
               add.u64         %rd4, %SPL, 0;
               add.u64         %rd5, %SPL, 0;
               mov.u32         %r1, %tid.x;
               setp.lt.u32     %p1, %r1, 3;
               @%p1 bra        $L__BB0_2;
               bra.uni         $L__BB0_1;
       $L__BB0_2:
               mov.u32         %r2, %tid.y;
               setp.lt.u32     %p2, %r2, 3;
               @%p2 bra        $L__BB0_4;
               bra.uni         $L__BB0_3;
       $L__BB0_4:                              // %.peel.begin
               cvt.u64.u32     %rd27, %r1;
               add.s64         %rd28, %rd3, %rd27;
               mul.lo.s32      %r3, %r2, 3;
               cvt.u64.u32     %rd29, %r3;
               add.s64         %rd30, %rd2, %rd29;
               ld.global.u8    %rs1, [%rd30];
               ld.global.u8    %rs2, [%rd28];
               ld.global.u8    %rs3, [%rd30+1];
               ld.global.u8    %rs4, [%rd28+3];
               ld.global.u8    %rs5, [%rd30+2];
               ld.global.u8    %rs6, [%rd28+6];
               mul.lo.s16      %rs7, %rs2, %rs1;
               mad.lo.s16      %rs8, %rs4, %rs3, %rs7;
               mad.lo.s16      %rs9, %rs6, %rs5, %rs8;
               add.s64         %rd31, %rd1, %rd29;
               add.s64         %rd32, %rd31, %rd27;
               st.global.u8    [%rd32], %rs9;
               ret;
       $L__BB0_1:
               mov.u64         %rd11, anon_$_6995e9b5afb94c728c982fc65adc2dcb_$_1;
               cvta.global.u64         %rd12, %rd11;
               st.local.u64    [%rd5], %rd12;
               mov.b64         %rd13, 1;
               st.local.u64    [%rd5+8], %rd13;
               mov.b64         %rd14, 0;
               st.local.u64    [%rd5+32], %rd14;
               mov.b64         %rd15, 8;
               st.local.u64    [%rd5+16], %rd15;
               st.local.u64    [%rd5+24], %rd14;
               mov.u64         %rd16, anon_$_6995e9b5afb94c728c982fc65adc2dcb_$_3;
               cvta.global.u64         %rd17, %rd16;
               { // callseq 0, 0
               .param .b64 param0;
               st.param.b64    [param0], %rd9;
               .param .b64 param1;
               st.param.b64    [param1], %rd17;
               call.uni 
               _ZN4core9panicking9panic_fmt17ha964c0166b5de554E, 
               (
               param0, 
               param1
               );
               } // callseq 0
       $L__BB0_3:
               mov.u64         %rd19, anon_$_6995e9b5afb94c728c982fc65adc2dcb_$_1;
               cvta.global.u64         %rd20, %rd19;
               st.local.u64    [%rd4], %rd20;
               mov.b64         %rd21, 1;
               st.local.u64    [%rd4+8], %rd21;
               mov.b64         %rd22, 0;
               st.local.u64    [%rd4+32], %rd22;
               mov.b64         %rd23, 8;
               st.local.u64    [%rd4+16], %rd23;
               st.local.u64    [%rd4+24], %rd22;
               mov.u64         %rd24, anon_$_6995e9b5afb94c728c982fc65adc2dcb_$_4;
               cvta.global.u64         %rd25, %rd24;
               { // callseq 1, 0
               .param .b64 param0;
               st.param.b64    [param0], %rd9;
               .param .b64 param1;
               st.param.b64    [param1], %rd25;
               call.uni 
               _ZN4core9panicking9panic_fmt17ha964c0166b5de554E, 
               (
               param0, 
               param1
               );
               } // callseq 1
                                               // -- End function
       }
       .func _RNvCsars8JMXwbW5_7___rustc17rust_begin_unwind(
               .param .b64 _RNvCsars8JMXwbW5_7___rustc17rust_begin_unwind_param_0
       )
       .noreturn                               // -- Begin function _RNvCsars8JMXwbW5_7___rustc17rust_begin_unwind
                                               // @_RNvCsars8JMXwbW5_7___rustc17rust_begin_unwind
       {
               .local .align 8 .b8     __local_depot1[1064];
               .reg .b64       %SP;
               .reg .b64       %SPL;
               .reg .pred      %p<17>;
               .reg .b16       %rs<9>;
               .reg .b32       %r<42>;
               .reg .b64       %rd<84>;
       
       // %bb.0:
               mov.u64         %SPL, __local_depot1;
               cvta.local.u64  %SP, %SPL;
               ld.param.u64    %rd28, [_RNvCsars8JMXwbW5_7___rustc17rust_begin_unwind_param_0];
               add.u64         %rd30, %SP, 0;
               add.u64         %rd1, %SPL, 0;
               add.u64         %rd31, %SP, 40;
               add.u64         %rd2, %SPL, 40;
               add.u64         %rd3, %SPL, 40;
               add.u64         %rd4, %SPL, 0;
               add.u64         %rd5, %SPL, 40;
               add.u64         %rd6, %SPL, 40;
               add.u64         %rd7, %SPL, 40;
               ld.u64  %rd37, [%rd28];
               ld.u64  %rd29, [%rd37+8];
               ld.u64  %rd9, [%rd37+24];
               setp.eq.s64     %p1, %rd29, 1;
               mov.u32         %r36, %tid.z;
               mov.u32         %r37, %tid.y;
               mov.u32         %r38, %tid.x;
               mov.u32         %r39, %ctaid.z;
               mov.u32         %r40, %ctaid.y;
               mov.u32         %r41, %ctaid.x;
               @%p1 bra        $L__BB1_17;
       // %bb.1:
               setp.ne.s64     %p2, %rd29, 0;
               @%p2 bra        $L__BB1_3;
       // %bb.2:
               setp.eq.s64     %p4, %rd9, 0;
               mov.b64         %rd11, 1;
               mov.b64         %rd76, 0;
               @%p4 bra        $L__BB1_4;
               bra.uni         $L__BB1_3;
       $L__BB1_17:
               setp.ne.s64     %p3, %rd9, 0;
               @%p3 bra        $L__BB1_3;
               bra.uni         $L__BB1_18;
       $L__BB1_3:
               st.local.u32    [%rd7], %r41;
               st.local.u32    [%rd7+4], %r40;
               st.local.u32    [%rd7+8], %r39;
               st.local.u32    [%rd7+12], %r38;
               st.local.u32    [%rd7+16], %r37;
               st.local.u32    [%rd7+20], %r36;
               mov.u64         %rd59, anon_$_a44653bcb419bc498496063559ae8ca0_$_1;
               cvta.global.u64         %rd60, %rd59;
               { // callseq 5, 0
               .param .b64 param0;
               st.param.b64    [param0], %rd60;
               .param .b64 param1;
               st.param.b64    [param1], %rd31;
               .param .b32 retval0;
               call.uni (retval0), 
               vprintf, 
               (
               param0, 
               param1
               );
               ld.param.b32    %r24, [retval0];
               } // callseq 5
               bra.uni         $L__BB1_20;
       $L__BB1_18:
               ld.u64  %rd8, [%rd37];
               ld.u64  %rd11, [%rd8];
               ld.u64  %rd76, [%rd8+8];
       $L__BB1_4:
               mov.b32         %r9, 0;
               st.local.u32    [%rd3], %r9;
               setp.eq.s64     %p5, %rd76, 0;
               @%p5 bra        $L__BB1_6;
       // %bb.5:
               add.s64         %rd40, %rd11, %rd76;
               add.s64         %rd41, %rd40, -1;
               { // callseq 2, 0
               .param .b64 param0;
               st.param.b64    [param0], %rd31;
               .param .b64 param1;
               st.param.b64    [param1], %rd41;
               .param .b64 param2;
               st.param.b64    [param2], 1;
               .param .b32 retval0;
               call.uni (retval0), 
               memcmp, 
               (
               param0, 
               param1, 
               param2
               );
               ld.param.b32    %r10, [retval0];
               } // callseq 2
               setp.eq.s32     %p6, %r10, 0;
               @%p6 bra        $L__BB1_19;
       $L__BB1_6:                              // %loadstoreloop2.preheader
               mov.b64         %rd78, 0;
       $L__BB1_7:                              // %loadstoreloop2
                                               // =>This Inner Loop Header: Depth=1
               add.s64         %rd47, %rd5, %rd78;
               mov.b16         %rs1, 0;
               st.local.u8     [%rd47], %rs1;
               add.s64         %rd78, %rd78, 1;
               setp.lt.u64     %p7, %rd78, 1024;
               @%p7 bra        $L__BB1_7;
       // %bb.8:                               // %split1
               setp.lt.u64     %p8, %rd76, 1024;
               @%p8 bra        $L__BB1_12;
               bra.uni         $L__BB1_9;
       $L__BB1_12:
               @%p5 bra        $L__BB1_15;
       // %bb.13:                              // %loop-memcpy-expansion5.preheader
               mov.b64         %rd80, 0;
       $L__BB1_14:                             // %loop-memcpy-expansion5
                                               // =>This Inner Loop Header: Depth=1
               add.s64         %rd52, %rd11, %rd80;
               ld.u8   %rs3, [%rd52];
               add.s64         %rd53, %rd5, %rd80;
               st.local.u8     [%rd53], %rs3;
               add.s64         %rd80, %rd80, 1;
               setp.lt.u64     %p11, %rd80, %rd76;
               @%p11 bra       $L__BB1_14;
       $L__BB1_15:                             // %post-loop-memcpy-expansion4
               add.s64         %rd54, %rd5, %rd76;
               mov.b16         %rs4, 0;
               st.local.u8     [%rd54], %rs4;
               bra.uni         $L__BB1_16;
       $L__BB1_9:
               mov.b64         %rd79, 0;
       $L__BB1_10:                             // %load-store-loop8
                                               // =>This Inner Loop Header: Depth=1
               add.s64         %rd49, %rd11, %rd79;
               ld.u8   %rs2, [%rd49];
               add.s64         %rd50, %rd5, %rd79;
               st.local.u8     [%rd50], %rs2;
               add.s64         %rd79, %rd79, 1;
               setp.lt.u64     %p9, %rd79, 1020;
               @%p9 bra        $L__BB1_10;
       // %bb.11:                              // %memcpy-split7
               mov.b32         %r14, 0;
               st.local.u8     [%rd5+1023], %r14;
               mov.b32         %r15, 46;
               st.local.u8     [%rd5+1022], %r15;
               st.local.u8     [%rd5+1021], %r15;
               st.local.u8     [%rd5+1020], %r15;
       $L__BB1_16:
               st.local.v2.u32         [%rd4], {%r41, %r40};
               st.local.v2.u32         [%rd4+8], {%r39, %r38};
               st.local.v2.u32         [%rd4+16], {%r37, %r36};
               st.local.u64    [%rd4+24], %rd31;
               mov.u64         %rd56, anon_$_a44653bcb419bc498496063559ae8ca0_$_0;
               cvta.global.u64         %rd57, %rd56;
               { // callseq 4, 0
               .param .b64 param0;
               st.param.b64    [param0], %rd57;
               .param .b64 param1;
               st.param.b64    [param1], %rd30;
               .param .b32 retval0;
               call.uni (retval0), 
               vprintf, 
               (
               param0, 
               param1
               );
               ld.param.b32    %r16, [retval0];
               } // callseq 4
               bra.uni         $L__BB1_20;
       $L__BB1_19:
               st.local.v2.u32         [%rd6], {%r41, %r40};
               st.local.v2.u32         [%rd6+8], {%r39, %r38};
               st.local.v2.u32         [%rd6+16], {%r37, %r36};
               st.local.u64    [%rd6+24], %rd11;
               mov.u64         %rd43, anon_$_a44653bcb419bc498496063559ae8ca0_$_0;
               cvta.global.u64         %rd44, %rd43;
               { // callseq 3, 0
               .param .b64 param0;
               st.param.b64    [param0], %rd44;
               .param .b64 param1;
               st.param.b64    [param1], %rd31;
               .param .b32 retval0;
               call.uni (retval0), 
               vprintf, 
               (
               param0, 
               param1
               );
               ld.param.b32    %r12, [retval0];
               } // callseq 3
       $L__BB1_20:                             // %loadstoreloop.preheader
               ld.u64  %rd63, [%rd28+8];
               ld.u64  %rd20, [%rd63];
               ld.u64  %rd21, [%rd63+8];
               ld.v2.u32       {%r7, %r8}, [%rd63+16];
               mov.b64         %rd81, 0;
       $L__BB1_21:                             // %loadstoreloop
                                               // =>This Inner Loop Header: Depth=1
               add.s64         %rd64, %rd2, %rd81;
               mov.b16         %rs5, 0;
               st.local.u8     [%rd64], %rs5;
               add.s64         %rd81, %rd81, 1;
               setp.lt.u64     %p12, %rd81, 1024;
               @%p12 bra       $L__BB1_21;
       // %bb.22:                              // %split
               setp.lt.u64     %p13, %rd21, 1024;
               @%p13 bra       $L__BB1_26;
               bra.uni         $L__BB1_23;
       $L__BB1_26:
               setp.eq.s64     %p15, %rd21, 0;
               @%p15 bra       $L__BB1_29;
       // %bb.27:                              // %loop-memcpy-expansion.preheader
               mov.b64         %rd83, 0;
       $L__BB1_28:                             // %loop-memcpy-expansion
                                               // =>This Inner Loop Header: Depth=1
               add.s64         %rd69, %rd20, %rd83;
               ld.u8   %rs7, [%rd69];
               add.s64         %rd70, %rd2, %rd83;
               st.local.u8     [%rd70], %rs7;
               add.s64         %rd83, %rd83, 1;
               setp.lt.u64     %p16, %rd83, %rd21;
               @%p16 bra       $L__BB1_28;
       $L__BB1_29:                             // %post-loop-memcpy-expansion
               add.s64         %rd71, %rd2, %rd21;
               mov.b16         %rs8, 0;
               st.local.u8     [%rd71], %rs8;
               bra.uni         $L__BB1_30;
       $L__BB1_23:
               mov.b64         %rd82, 0;
       $L__BB1_24:                             // %load-store-loop
                                               // =>This Inner Loop Header: Depth=1
               add.s64         %rd66, %rd20, %rd82;
               ld.u8   %rs6, [%rd66];
               add.s64         %rd67, %rd2, %rd82;
               st.local.u8     [%rd67], %rs6;
               add.s64         %rd82, %rd82, 1;
               setp.lt.u64     %p14, %rd82, 1020;
               @%p14 bra       $L__BB1_24;
       // %bb.25:                              // %memcpy-split
               mov.b32         %r26, 0;
               st.local.u8     [%rd2+1023], %r26;
               mov.b32         %r27, 46;
               st.local.u8     [%rd2+1022], %r27;
               st.local.u8     [%rd2+1021], %r27;
               st.local.u8     [%rd2+1020], %r27;
       $L__BB1_30:
               st.local.v2.u32         [%rd1], {%r41, %r40};
               st.local.v2.u32         [%rd1+8], {%r39, %r38};
               st.local.v2.u32         [%rd1+16], {%r37, %r36};
               st.local.u64    [%rd1+24], %rd31;
               st.local.v2.u32         [%rd1+32], {%r7, %r8};
               mov.u64         %rd73, anon_$_a44653bcb419bc498496063559ae8ca0_$_2;
               cvta.global.u64         %rd74, %rd73;
               { // callseq 6, 0
               .param .b64 param0;
               st.param.b64    [param0], %rd74;
               .param .b64 param1;
               st.param.b64    [param1], %rd30;
               .param .b32 retval0;
               call.uni (retval0), 
               vprintf, 
               (
               param0, 
               param1
               );
               ld.param.b32    %r34, [retval0];
               } // callseq 6
               // begin inline asm
               trap;
               // end inline asm
                                               // -- End function
       }
       .func _ZN4core9panicking9panic_fmt17ha964c0166b5de554E(
               .param .b64 _ZN4core9panicking9panic_fmt17ha964c0166b5de554E_param_0,
               .param .b64 _ZN4core9panicking9panic_fmt17ha964c0166b5de554E_param_1
       )
       .noreturn                               // -- Begin function _ZN4core9panicking9panic_fmt17ha964c0166b5de554E
                                               // @_ZN4core9panicking9panic_fmt17ha964c0166b5de554E
       {
               .local .align 8 .b8     __local_depot2[24];
               .reg .b64       %SP;
               .reg .b64       %SPL;
               .reg .b16       %rs<3>;
               .reg .b64       %rd<5>;
       
       // %bb.0:
               mov.u64         %SPL, __local_depot2;
               cvta.local.u64  %SP, %SPL;
               ld.param.u64    %rd1, [_ZN4core9panicking9panic_fmt17ha964c0166b5de554E_param_0];
               ld.param.u64    %rd2, [_ZN4core9panicking9panic_fmt17ha964c0166b5de554E_param_1];
               add.u64         %rd3, %SP, 0;
               add.u64         %rd4, %SPL, 0;
               st.local.u64    [%rd4], %rd1;
               st.local.u64    [%rd4+8], %rd2;
               mov.b16         %rs1, 0;
               mov.b16         %rs2, 1;
               st.local.v2.u8  [%rd4+16], {%rs2, %rs1};
               { // callseq 7, 0
               .param .b64 param0;
               st.param.b64    [param0], %rd3;
               call.uni 
               _RNvCsars8JMXwbW5_7___rustc17rust_begin_unwind, 
               (
               param0
               );
               } // callseq 7
                                               // -- End function
       }
       .func  (.param .b32 func_retval0) memcmp(
               .param .b64 memcmp_param_0,
               .param .b64 memcmp_param_1,
               .param .b64 memcmp_param_2
       )                                       // -- Begin function memcmp
                                               // @memcmp
       {
               .reg .pred      %p<4>;
               .reg .b16       %rs<3>;
               .reg .b32       %r<8>;
               .reg .b64       %rd<13>;
       
       // %bb.0:
               ld.param.u64    %rd10, [memcmp_param_2];
               ld.param.u64    %rd12, [memcmp_param_1];
               ld.param.u64    %rd11, [memcmp_param_0];
               setp.eq.s64     %p1, %rd10, 0;
               mov.b32         %r7, 0;
               @%p1 bra        $L__BB3_4;
       $L__BB3_2:                              // %.preheader
                                               // =>This Inner Loop Header: Depth=1
               ld.u8   %rs1, [%rd11];
               ld.u8   %rs2, [%rd12];
               setp.eq.s16     %p2, %rs1, %rs2;
               @%p2 bra        $L__BB3_1;
               bra.uni         $L__BB3_3;
       $L__BB3_1:                              //   in Loop: Header=BB3_2 Depth=1
               add.s64         %rd12, %rd12, 1;
               add.s64         %rd11, %rd11, 1;
               add.s64         %rd10, %rd10, -1;
               setp.eq.s64     %p3, %rd10, 0;
               @%p3 bra        $L__BB3_4;
               bra.uni         $L__BB3_2;
       $L__BB3_3:
               cvt.u32.u16     %r1, %rs1;
               cvt.u32.u16     %r2, %rs2;
               sub.s32         %r7, %r1, %r2;
       $L__BB3_4:                              // %.loopexit
               st.param.b32    [func_retval0], %r7;
               ret;
                                               // -- End function
       }
   
   

   In the Rust side, the code is a little bit complex since my crate is designed to mainly concern about 1d problems (e.g., verify whether a number has a divisor in some range.)

   use nalgebra::{Const, Matrix3};
   use std::time::Instant;
   
   use cuda_min::{Device, Param};
   fn main() {
       const A: &'static str = include_str!(concat!(env!("OUT_DIR"), "/gpu_ptx_code.ptx"));
       println!("PTX Code:");
       println!("{A}");
       let device = Device::init();
   
       let module = device.compile(A).unwrap();
       let func = module.get_function("mat_mul").unwrap();
       println!(
           "max thread per block is {}",
           func.get_max_thread_per_block().unwrap()
       );
       const LEN: usize = 1;
       let now = Instant::now();
       let mut ret = vec![Matrix3::<u8>::zeros(); LEN]; // ret could also be input, and its length decided the max avaliable tasks.
       let a = vec![
           Matrix3::<u8>::from_row_slice_generic(
               Const::<3>,
               Const::<3>,
               &[1u8, 2, 3, 4, 5, 6, 7, 8, 9]
           );
           LEN
       ];
       let b = vec![
           Matrix3::<u8>::from_row_slice_generic(
               Const::<3>,
               Const::<3>,
               &[9u8, 8, 7, 6, 5, 4, 3, 2, 1]
           );
           LEN
       ];
       let mut param = Param::new(&mut ret).block_size(1).push(&a).push(&b);
       param = unsafe { param.set_block_grid_size((3, 3, 1), (1, 1, 1)) };
   
       let res = func.call(param).unwrap();
       res.sync().unwrap();
       println!("{:?}", now.elapsed());
       println!("{ret:?}");
       println!("{:?}", now.elapsed());
   }

   Such program generate a result

   [[[30, 84, 138], [24, 69, 114], [18, 54, 90]]]
   

   If you're interested in 1d jobs, maybe using my crate is another choice