gpurs 0.4.0

GPU acceleration/linear algebra crate for scientific computing
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
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100

//! Module for GPU-based mass vector transformation.

mod gpu_transformer;

pub use gpu_transformer::GPUTransformer;

// Default amount of memory slots for transformations in GPUTransformer.
const INIT_MEMORY_CAPACITY: usize = 3;

// List of default kernel names.
const PROGRAM_LIST_FLOAT: [&str; 1] = [
    "bulk_transform",
];
const PROGRAM_LIST_DOUBLE: [&str; 1] = [
    "bulk_transform",
];

// Source code for default kernels.
const PROGRAM_SOURCE_FLOAT: &str = r#"
typedef struct _vec3h {
    float x;
    float y;
    float z;
    float w;
}vec3h;

kernel void bulk_transform (
    global vec3h* output,
    const int N,
    const global float* transform,
    const global vec3h* vectors
) {
    const int globalCol = get_global_id(0);

    float interm_x = transform[0] * vectors[globalCol].x;
    interm_x += transform[1] * vectors[globalCol].y;
    interm_x += transform[2] * vectors[globalCol].z;
    interm_x += transform[3] * vectors[globalCol].w;
    output[globalCol].x = interm_x;

    float interm_y = transform[4] * vectors[globalCol].x;
    interm_y += transform[5] * vectors[globalCol].y;
    interm_y += transform[6] * vectors[globalCol].z;
    interm_y += transform[7] * vectors[globalCol].w;
    output[globalCol].y = interm_y;

    float interm_z = transform[8] * vectors[globalCol].x;
    interm_z += transform[9] * vectors[globalCol].y;
    interm_z += transform[10] * vectors[globalCol].z;
    interm_z += transform[11] * vectors[globalCol].w;
    output[globalCol].z = interm_z;

    float interm_w = transform[12] * vectors[globalCol].x;
    interm_w += transform[13] * vectors[globalCol].y;
    interm_w += transform[14] * vectors[globalCol].z;
    interm_w += transform[15] * vectors[globalCol].w;
    output[globalCol].w = interm_w;
}
"#;
const PROGRAM_SOURCE_DOUBLE: &str = r#"
typedef struct _vec3h {
    double x;
    double y;
    double z;
    double w;
}vec3h;

kernel void bulk_transform (
    global vec3h* output,
    const int N,
    const global double* transform,
    const global vec3h* vectors
) {
    const int globalCol = get_global_id(0);

    double interm_x = transform[0] * vectors[globalCol].x;
    interm_x += transform[1] * vectors[globalCol].y;
    interm_x += transform[2] * vectors[globalCol].z;
    interm_x += transform[3] * vectors[globalCol].w;
    output[globalCol].x = interm_x;

    double interm_y = transform[4] * vectors[globalCol].x;
    interm_y += transform[5] * vectors[globalCol].y;
    interm_y += transform[6] * vectors[globalCol].z;
    interm_y += transform[7] * vectors[globalCol].w;
    output[globalCol].y = interm_y;

    double interm_z = transform[8] * vectors[globalCol].x;
    interm_z += transform[9] * vectors[globalCol].y;
    interm_z += transform[10] * vectors[globalCol].z;
    interm_z += transform[11] * vectors[globalCol].w;
    output[globalCol].z = interm_z;

    double interm_w = transform[12] * vectors[globalCol].x;
    interm_w += transform[13] * vectors[globalCol].y;
    interm_w += transform[14] * vectors[globalCol].z;
    interm_w += transform[15] * vectors[globalCol].w;
    output[globalCol].w = interm_w;
}
"#;