How to pass multiple byte arrays to CUDA kernel for parallel SHA-512 hashing in Rust?

[stevefan1999-personal]

#![allow(improper_ctypes_definitions, clippy::missing_safety_doc)]

use cuda_std::prelude::*;
use hmac_sha512::Hash;

#[kernel]
pub unsafe fn sha512(input: &[&[u8]], output: *mut [[u8; 64]]) {
    match thread::index_1d() as usize {
        idx if idx < input.len() => {
            (*output)[idx] = Hash::hash(input[idx]);
        }
        _ => {}
    }
}

How do I get it to do sha512 on multiple input of byte array?

use cust::prelude::*;
use std::error::Error;

static PTX: &str = include_str!(concat!(env!("OUT_DIR"), "/path.ptx"));

fn main() -> Result<(), Box<dyn Error>> {
    let _ctx = cust::quick_init()?;
    let module = Module::from_ptx(PTX, &[])?;
    let stream = Stream::new(StreamFlags::NON_BLOCKING, None)?;
    let hash_input = vec![*b"hello world", *b"hello world",];
    let hash_input_buf = hash_input.as_dbuf()?;

    let mut out = vec![[0u8; 64]; hash_input.len()];
    let out_buf = out.as_dbuf()?;

    let sha512 = module.get_function("sha512")?;
    let (_, block_size) = sha512.suggested_launch_configuration(0, 0.into())?;

    let grid_size = (hash_input.len() as u32).div_ceil(block_size);

    println!(
        "using {} blocks and {} threads per block",
        grid_size, block_size
    );

    unsafe {
        launch!(
            sha512<<<grid_size, block_size, 0, stream>>>(
                hash_input_buf.as_device_ptr(),
                hash_input_buf.len(),
                out_buf.as_device_ptr(),
                out_buf.len(),
            )
        )?;
    }

    stream.synchronize()?;
    out_buf.copy_to(&mut out)?;

    for hash in out {
        println!("{}", hex::encode(hash));
    }

    Ok(())
}

But then I got:

   Compiling serde_json v1.0.140
   Compiling rustc_codegen_nvvm v0.3.0 (https://github.com/Rust-GPU/Rust-CUDA#afb147ed)
   Compiling cust v0.3.2 (https://github.com/Rust-GPU/Rust-CUDA#afb147ed)
   Compiling cuda_builder v0.3.0 (https://github.com/Rust-GPU/Rust-CUDA#afb147ed)
   Compiling gpu-playground v0.1.0 (F:\rust\gpu-playground)
    Finished `dev` profile [unoptimized + debuginfo] target(s) in 24.89s
     Running `target\debug\gpu-playground.exe`
using 1 blocks and 768 threads per block
Error: IllegalAddress
error: process didn't exit successfully: `target\debug\gpu-playground.exe` (exit code: 1)

[LegNeato]

I think I recently fix this...we were emitting a trap() for unsupported operations on i128. We now have a sha256/512 example in the repo.