clow 0.1.4

Lower-Level wrapper around Cudarc pointer types
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81

//! Demonstrates ClowPtr operations: offset, cast, null checks, and from_raw_parts.

use clow::prelude::*;
use cudarc::driver::{CudaContext, DevicePtr, LaunchConfig, PushKernelArg};
use cudarc::nvrtc::compile_ptx;
use std::sync::Arc;

fn main() -> Result<(), Box<dyn std::error::Error>> {
    let n = 1024;

    let ctx = Arc::new(CudaContext::new(0)?);
    let stream = ctx.default_stream();

    // Compile a kernel that takes a pre-offset pointer
    let ptx = compile_ptx(
        r#"
extern "C" __global__ void fill_offset(float* data, float value, int count) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < count) {
        data[idx] = value;
    }
}
"#,
    )?;
    let module = ctx.load_module(ptx)?;

    // Allocate a buffer of f32
    let buffer = stream.clow_alloc_zeros::<f32>(n)?;
    let base_ptr = buffer.as_device_ptr();

    // ClowPtr::null and is_null
    let null_ptr = ClowPtr::<f32>::null();
    assert!(null_ptr.is_null());
    assert!(!base_ptr.is_null());

    // Use offset to get a pointer to the middle of the buffer
    let half = n / 2;
    let middle_ptr = unsafe { base_ptr.offset(half as isize) };

    // Launch kernel writing to the second half of the buffer
    let kernel = module.load_function("fill_offset")?;
    let block_size = 256u32;
    let grid_size = (half as u32 + block_size - 1) / block_size;

    unsafe {
        let mut builder = stream.launch_builder(&kernel);
        builder.arg(&middle_ptr);
        builder.arg(&42.0f32);
        builder.arg(&half);
        builder.launch(LaunchConfig {
            block_dim: (block_size, 1, 1),
            grid_dim: (grid_size, 1, 1),
            shared_mem_bytes: 0,
        })?;
    }

    stream.synchronize()?;

    let host = stream.clow_clone_dtoh(&buffer)?;

    // First half should be zeros, second half should be 42.0
    for i in 0..half {
        assert_eq!(host[i], 0.0);
    }
    for i in half..n {
        assert_eq!(host[i], 42.0);
    }

    println!("Pointer offset demo: first half zeros, second half 42.0");

    // Demonstrate cast: treat f32 buffer as u32 (same size, different interpretation)
    let as_u32 = base_ptr.cast::<u32>();
    println!("Cast f32 pointer to u32: {:?} -> {:?}", base_ptr, as_u32);

    // Demonstrate from_raw_parts with Cudarc's raw pointer
    let (raw_ptr, _) = buffer.stream().alloc_zeros::<f32>(1)?.device_ptr(&stream);
    let clow_from_raw = ClowPtr::<f32>::from_raw_parts(raw_ptr);
    println!("Created ClowPtr from raw CUdeviceptr: {:?}", clow_from_raw);

    Ok(())
}