baracuda-driver 0.0.1-alpha.40

Safe Rust wrappers for the CUDA Driver API (devices, contexts, streams, events, memory, kernels, graphs).
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

//
// Hand-written PTX for tests/kernels/vector_add.cu.
// Equivalent CUDA C:
//   extern "C" __global__ void vector_add(const float* a, const float* b,
//                                         float* c, unsigned int n) {
//     unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
//     if (i < n) c[i] = a[i] + b[i];
//   }
//
// Compatible with CUDA 11.4+ drivers and compute capability 5.0+.
//

.version 7.4
.target sm_50
.address_size 64

.visible .entry vector_add(
    .param .u64 vector_add_param_0,
    .param .u64 vector_add_param_1,
    .param .u64 vector_add_param_2,
    .param .u32 vector_add_param_3
)
{
    .reg .pred %p<2>;
    .reg .b32 %r<6>;
    .reg .f32 %f<4>;
    .reg .b64 %rd<11>;

    ld.param.u64  %rd1, [vector_add_param_0];
    ld.param.u64  %rd2, [vector_add_param_1];
    ld.param.u64  %rd3, [vector_add_param_2];
    ld.param.u32  %r1,  [vector_add_param_3];
    mov.u32       %r2, %ntid.x;
    mov.u32       %r3, %ctaid.x;
    mov.u32       %r4, %tid.x;
    mad.lo.s32    %r5, %r3, %r2, %r4;
    setp.ge.u32   %p1, %r5, %r1;
    @%p1 bra      $L__BB0_2;

    cvta.to.global.u64 %rd4, %rd1;
    mul.wide.u32       %rd5, %r5, 4;
    add.s64            %rd6, %rd4, %rd5;
    ld.global.f32      %f1, [%rd6];
    cvta.to.global.u64 %rd7, %rd2;
    add.s64            %rd8, %rd7, %rd5;
    ld.global.f32      %f2, [%rd8];
    add.f32            %f3, %f1, %f2;
    cvta.to.global.u64 %rd9, %rd3;
    add.s64            %rd10, %rd9, %rd5;
    st.global.f32      [%rd10], %f3;

$L__BB0_2:
    ret;
}