decuda 0.1.1

CUDA to HIP, SYCL, OpenCL, and Rust GPU migration tool — automatic source-code translator for porting CUDA C++ kernels to AMD ROCm HIP, Intel oneAPI SYCL, Khronos OpenCL, and Rust GPU (cust / rust-gpu)
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

// Generated by decuda.
// Host-side runtime calls are translated to the `cust` crate. Kernels
// are emitted as `TODO(decuda)` blocks: rust-gpu translation requires
// the kernel to be authored as a Rust fn. See `examples/` for a
// scaffolded SPIR-V kernel module you can flesh out.
//
// Add to your Cargo.toml:
//   [dependencies]
//   cust = "0.3"

// Complex fixture: asynchronous pipeline using CUDA streams and events.
//
// Exercises:
//   - cudaStreamCreate / cudaStreamDestroy / cudaStreamSynchronize
//   - cudaEventCreate / cudaEventRecord / cudaEventSynchronize / cudaEventDestroy
//   - cudaMemcpyAsync (async memcpy on a stream)
//   - Multiple kernels launched on different streams
//   - Launch with shared-memory size argument: kernel<<<grid, block, smem, stream>>>
//   - __global__ kernels with __syncthreads
//   - cuda_runtime.h header
// was: #include cuda_runtime.h  ->  cust::cuda_build_setup() /* TODO: import cust crate */

// TODO(decuda): rewrite as rust-gpu kernel fn
 void scale(float a, float* x, int n) {
    int i = block_idx * block_dim + thread_idx;
    if (i < n) {
        x[i] *= a;
    }
}

// TODO(decuda): rewrite as rust-gpu kernel fn
 void add(const float* x, const float* y, float* out, int n) {
    int i = block_idx * block_dim + thread_idx;
    if (i < n) {
        out[i] = x[i] + y[i];
    }
    group.sync();
}

int main(void) {
    const int N = 1 << 16;
    float *dx = nullptr, *dy = nullptr, *dz = nullptr;

    cudaMalloc((void**)&dx, N * sizeof(float));
    cudaMalloc((void**)&dy, N * sizeof(float));
    cudaMalloc((void**)&dz, N * sizeof(float));

    cust::Stream() s1, s2;
    cudaStreamCreate(&s1);
    cudaStreamCreate(&s2);

    cudaEvent_t e1, e2;
    cudaEventCreate(&e1);
    cudaEventCreate(&e2);

    dim3 grid(N / 256);
    dim3 block(256);

    // Async memcpy on stream s1.
    cudaMemcpyAsync(dx, dy, N * sizeof(float), cudaMemcpyDeviceToDevice, s1);

    // Launch with explicit shared-memory size and stream.
    { /* decuda cust launch */ let _kernel = modules.get_function("scale"); unsafe { let _ = launch!( _kernel<<<grid as grid_size, block as block_size, 0 as usize, s1>>>(2.0f, dx, N) ); } };
    cudaEventRecord(e1, s1);

    // Second stream waits on event e1 via host-side synchronize.
    cudaEventSynchronize(e1);
    { /* decuda cust launch */ let _kernel = modules.get_function("add"); unsafe { let _ = launch!( _kernel<<<grid as grid_size, block as block_size, 128 as usize, s2>>>(dx, dy, dz, N) ); } };
    cudaEventRecord(e2, s2);

    cudaStreamSynchronize(s1);
    cudaStreamSynchronize(s2);

    cudaEventDestroy(e1);
    cudaEventDestroy(e2);
    cudaStreamDestroy(s1);
    cudaStreamDestroy(s2);

    cudaFree(dx);
    cudaFree(dy);
    cudaFree(dz);
    return 0;
}