sp1-gpu-sys 6.2.2

FFI bindings and CUDA build system for SP1-GPU
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

#pragma once

extern "C" void* branching_program_kernel();
extern "C" void* transition_kernel();
extern "C" void* interpolateAndObserve_kernel();

// A range of values where the start is inclusive and the end is exclusive.
struct Range {
    int start;
    int end;

    __device__ bool in_range(int val) { return val >= start && val < end; }
};

// All the values of the BitState struct:
// https://github.com/succinctlabs/slop/blob/783136d30dc2b5e9ce558385b333dad93a89fd29/crates/jagged/src/poly.rs#L73
enum BitState {
    ROW_0__INDEX_0__CURR_PS_0__NEXT_PS_0,
    ROW_0__INDEX_0__CURR_PS_0__NEXT_PS_1,
    ROW_0__INDEX_0__CURR_PS_1__NEXT_PS_0,
    ROW_0__INDEX_0__CURR_PS_1__NEXT_PS_1,
    ROW_0__INDEX_1__CURR_PS_0__NEXT_PS_0,
    ROW_0__INDEX_1__CURR_PS_0__NEXT_PS_1,
    ROW_0__INDEX_1__CURR_PS_1__NEXT_PS_0,
    ROW_0__INDEX_1__CURR_PS_1__NEXT_PS_1,
    ROW_1__INDEX_0__CURR_PS_0__NEXT_PS_0,
    ROW_1__INDEX_0__CURR_PS_0__NEXT_PS_1,
    ROW_1__INDEX_0__CURR_PS_1__NEXT_PS_0,
    ROW_1__INDEX_0__CURR_PS_1__NEXT_PS_1,
    ROW_1__INDEX_1__CURR_PS_0__NEXT_PS_0,
    ROW_1__INDEX_1__CURR_PS_0__NEXT_PS_1,
    ROW_1__INDEX_1__CURR_PS_1__NEXT_PS_0,
    ROW_1__INDEX_1__CURR_PS_1__NEXT_PS_1,
    BIT_STATE_COUNT,
};

// All the values of the MemoryState struct:
// https://github.com/succinctlabs/slop/blob/783136d30dc2b5e9ce558385b333dad93a89fd29/crates/jagged/src/poly.rs#L39
// and the StateOrFail enum:
// https://github.com/succinctlabs/slop/blob/783136d30dc2b5e9ce558385b333dad93a89fd29/crates/jagged/src/poly.rs#L64
enum MemoryState {
    COMP_SO_FAR_0__CARRY_0,
    COMP_SO_FAR_0__CARRY_1,
    COMP_SO_FAR_1__CARRY_0,
    COMP_SO_FAR_1__CARRY_1,
    FAIL,
    MEMORY_STATE_COUNT,
};

// The success memory state:
// https://github.com/succinctlabs/slop/blob/783136d30dc2b5e9ce558385b333dad93a89fd29/crates/jagged/src/poly.rs#L53
__device__ constexpr int SUCCESS_STATE = COMP_SO_FAR_1__CARRY_0;

__device__ constexpr int INITIAL_MEMORY_STATE = COMP_SO_FAR_0__CARRY_0;

__constant__ constexpr const MemoryState TRANSITIONS[BIT_STATE_COUNT][MEMORY_STATE_COUNT] = {
    {COMP_SO_FAR_0__CARRY_0, FAIL, COMP_SO_FAR_1__CARRY_0, FAIL, FAIL},
    {COMP_SO_FAR_1__CARRY_0, FAIL, COMP_SO_FAR_1__CARRY_0, FAIL, FAIL},
    {FAIL, COMP_SO_FAR_0__CARRY_1, FAIL, COMP_SO_FAR_1__CARRY_1, FAIL},
    {FAIL, COMP_SO_FAR_1__CARRY_1, FAIL, COMP_SO_FAR_1__CARRY_1, FAIL},
    {FAIL, COMP_SO_FAR_0__CARRY_0, FAIL, COMP_SO_FAR_0__CARRY_0, FAIL},
    {FAIL, COMP_SO_FAR_0__CARRY_0, FAIL, COMP_SO_FAR_1__CARRY_0, FAIL},
    {COMP_SO_FAR_0__CARRY_0, FAIL, COMP_SO_FAR_0__CARRY_0, FAIL, FAIL},
    {COMP_SO_FAR_0__CARRY_0, FAIL, COMP_SO_FAR_1__CARRY_0, FAIL, FAIL},
    {FAIL, COMP_SO_FAR_0__CARRY_1, FAIL, COMP_SO_FAR_1__CARRY_1, FAIL},
    {FAIL, COMP_SO_FAR_1__CARRY_1, FAIL, COMP_SO_FAR_1__CARRY_1, FAIL},
    {COMP_SO_FAR_0__CARRY_1, FAIL, COMP_SO_FAR_1__CARRY_1, FAIL, FAIL},
    {COMP_SO_FAR_1__CARRY_1, FAIL, COMP_SO_FAR_1__CARRY_1, FAIL, FAIL},
    {COMP_SO_FAR_0__CARRY_0, FAIL, COMP_SO_FAR_0__CARRY_0, FAIL, FAIL},
    {COMP_SO_FAR_0__CARRY_0, FAIL, COMP_SO_FAR_1__CARRY_0, FAIL, FAIL},
    {FAIL, COMP_SO_FAR_0__CARRY_1, FAIL, COMP_SO_FAR_0__CARRY_1, FAIL},
    {FAIL, COMP_SO_FAR_0__CARRY_1, FAIL, COMP_SO_FAR_1__CARRY_1, FAIL},
};