trueno-gpu 0.4.29

Pure Rust PTX generation for NVIDIA CUDA - no LLVM, no nvcc
Documentation 
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

use super::*;
use crate::kernels::Kernel;

#[test]
fn test_rope_kernel_name() {
    let kernel = RopeKernel::new(32, 64, 10000.0);
    assert_eq!(kernel.name(), "rope");
}

#[test]
fn test_rope_ptx_generation() {
    let kernel = RopeKernel::new(32, 64, 10000.0);
    let ptx = kernel.emit_ptx();

    assert!(ptx.contains(".entry rope"));
    assert!(ptx.contains("sin.approx.f32"));
    assert!(ptx.contains("cos.approx.f32"));
}

#[test]
fn test_rope_indirect_kernel_name() {
    let kernel = RopeIndirectKernel::new(32, 64, 10000.0);
    assert_eq!(kernel.name(), "rope_indirect");
}

#[test]
fn test_rope_indirect_ptx_generation() {
    let kernel = RopeIndirectKernel::new(32, 64, 10000.0);
    let ptx = kernel.emit_ptx();

    assert!(ptx.contains(".entry rope_indirect"));
    assert!(ptx.contains(".param .u64 pos_ptr"));
    assert!(ptx.contains("ld.global.u32"));
}

#[test]
fn test_rope_neox_kernel_name() {
    let kernel = RopeNeoxKernel::new(32, 64, 1_000_000.0);
    assert_eq!(kernel.name(), "rope_neox");
}

#[test]
fn test_rope_neox_indirect_kernel_name() {
    let kernel = RopeNeoxIndirectKernel::new(32, 64, 1_000_000.0);
    assert_eq!(kernel.name(), "rope_neox_indirect");
}

#[test]
fn test_batched_rope_kernel_name() {
    let kernel = BatchedRopeKernel::new(32, 64, 4, 10000.0);
    assert_eq!(kernel.name(), "batched_rope");
}

#[test]
fn test_batched_rope_ptx_generation() {
    let kernel = BatchedRopeKernel::new(32, 64, 4, 10000.0);
    let ptx = kernel.emit_ptx();

    assert!(ptx.contains(".entry batched_rope"));
    assert!(ptx.contains(".param .u64 positions_ptr"));
}

#[test]
fn test_precise_rope_kernel_name() {
    let kernel = PreciseRopeKernel::new(32, 64, 1_000_000.0);
    assert_eq!(kernel.name(), "rope_precise");
}

#[test]
fn test_precise_rope_indirect_kernel_name() {
    let kernel = PreciseRopeIndirectKernel::new(32, 64, 1_000_000.0);
    assert_eq!(kernel.name(), "rope_precise_indirect");
}