Skip to content

Commit 5c18976

Browse files
robik75robik75
committed
fix(gpu_prover): fix compilation issues when compiling with the MSVC toolchain
1 parent e4711f4 commit 5c18976

File tree

3 files changed

+50
-24
lines changed

3 files changed

+50
-24
lines changed

‎gpu_prover/native/witness/trace.cuh‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -16,8 +16,8 @@ struct TimestampData {
1616
static constexpr unsigned TIMESTAMP_COLUMNS_NUM_BITS_MASK = (1u << TIMESTAMP_COLUMNS_NUM_BITS) - 1;
1717
static constexpr unsigned NUM_EMPTY_BITS_FOR_RAM_TIMESTAMP = 2;
1818
static constexpr u32 TOTAL_TIMESTAMP_BITS = TIMESTAMP_COLUMNS_NUM_BITS * NUM_TIMESTAMP_COLUMNS_FOR_RAM;
19-
static constexpr TimestampScalar TIMESTAMP_STEP = 1ul << NUM_EMPTY_BITS_FOR_RAM_TIMESTAMP;
20-
static constexpr TimestampScalar MAX_INITIAL_TIMESTAMP = (1ul << TOTAL_TIMESTAMP_BITS) - TIMESTAMP_STEP * 2;
19+
static constexpr TimestampScalar TIMESTAMP_STEP = 1ull << NUM_EMPTY_BITS_FOR_RAM_TIMESTAMP;
20+
static constexpr TimestampScalar MAX_INITIAL_TIMESTAMP = (1ull << TOTAL_TIMESTAMP_BITS) - TIMESTAMP_STEP * 2;
2121

2222
u16 limbs[NUM_TIMESTAMP_DATA_LIMBS];
2323

‎gpu_prover/native/witness/trace_delegation.cuh‎

Lines changed: 37 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -61,10 +61,34 @@ struct KeccakSpecial5AbiDescription {
6161

6262
template <typename DESCRIPTION> struct DelegationWitness {
6363
const TimestampScalar write_timestamp;
64-
const RegisterOrIndirectReadWriteData reg_accesses[DESCRIPTION::REG_ACCESSES];
65-
const RegisterOrIndirectReadData indirect_reads[DESCRIPTION::INDIRECT_READS];
66-
const RegisterOrIndirectReadWriteData indirect_writes[DESCRIPTION::INDIRECT_WRITES];
67-
const u16 variables_offsets[DESCRIPTION::VARIABLE_OFFSETS];
64+
// instead of
65+
// const RegisterOrIndirectReadWriteData reg_accesses[DESCRIPTION::REG_ACCESSES];
66+
// const RegisterOrIndirectReadData indirect_reads[DESCRIPTION::INDIRECT_READS];
67+
// const RegisterOrIndirectReadWriteData indirect_writes[DESCRIPTION::INDIRECT_WRITES];
68+
// const u16 variables_offsets[DESCRIPTION::VARIABLE_OFFSETS];
69+
// we implement this as a single byte array to avoid compilation errors when some of the arrays would be zero-size
70+
const u8 contents[DESCRIPTION::REG_ACCESSES * sizeof(RegisterOrIndirectReadWriteData) + DESCRIPTION::INDIRECT_READS * sizeof(RegisterOrIndirectReadData) +
71+
DESCRIPTION::INDIRECT_WRITES * sizeof(RegisterOrIndirectReadWriteData) + DESCRIPTION::VARIABLE_OFFSETS * sizeof(u16)];
72+
73+
DEVICE_FORCEINLINE const RegisterOrIndirectReadWriteData *reg_accesses() const { return reinterpret_cast<const RegisterOrIndirectReadWriteData *>(contents); }
74+
75+
DEVICE_FORCEINLINE const RegisterOrIndirectReadData *indirect_reads() const {
76+
constexpr size_t offset = DESCRIPTION::REG_ACCESSES * sizeof(RegisterOrIndirectReadWriteData);
77+
return reinterpret_cast<const RegisterOrIndirectReadData *>(contents + offset);
78+
}
79+
80+
DEVICE_FORCEINLINE const RegisterOrIndirectReadWriteData *indirect_writes() const {
81+
constexpr size_t offset =
82+
DESCRIPTION::REG_ACCESSES * sizeof(RegisterOrIndirectReadWriteData) + DESCRIPTION::INDIRECT_READS * sizeof(RegisterOrIndirectReadData);
83+
return reinterpret_cast<const RegisterOrIndirectReadWriteData *>(contents + offset);
84+
}
85+
86+
DEVICE_FORCEINLINE const u16 *variables_offsets() const {
87+
constexpr size_t offset = DESCRIPTION::REG_ACCESSES * sizeof(RegisterOrIndirectReadWriteData) +
88+
DESCRIPTION::INDIRECT_READS * sizeof(RegisterOrIndirectReadData) +
89+
DESCRIPTION::INDIRECT_WRITES * sizeof(RegisterOrIndirectReadWriteData);
90+
return reinterpret_cast<const u16 *>(contents + offset);
91+
}
6892
};
6993

7094
template <typename DESCRIPTION> struct DelegationTrace {
@@ -79,20 +103,20 @@ template <typename DESCRIPTION> struct DelegationTrace {
79103
const auto cycle_data = tracing_data + trace_row;
80104
switch (placeholder.tag) {
81105
case DelegationRegisterReadValue: {
82-
return cycle_data->reg_accesses[reg_offset].read_value;
106+
return cycle_data->reg_accesses()[reg_offset].read_value;
83107
}
84108
case DelegationRegisterWriteValue: {
85-
return cycle_data->reg_accesses[reg_offset].write_value;
109+
return cycle_data->reg_accesses()[reg_offset].write_value;
86110
}
87111
case DelegationIndirectReadValue: {
88-
return DESCRIPTION::use_read_indirects(register_index) ? cycle_data->indirect_reads[word_index].read_value
89-
: cycle_data->indirect_writes[word_index].read_value;
112+
return DESCRIPTION::use_read_indirects(register_index) ? cycle_data->indirect_reads()[word_index].read_value
113+
: cycle_data->indirect_writes()[word_index].read_value;
90114
}
91115
case DelegationIndirectWriteValue: {
92116
if (DESCRIPTION::use_read_indirects(register_index)) {
93117
__trap();
94118
}
95-
return cycle_data->indirect_writes[word_index].write_value;
119+
return cycle_data->indirect_writes()[word_index].write_value;
96120
}
97121
default:
98122
__trap();
@@ -110,7 +134,7 @@ template <typename DESCRIPTION> struct DelegationTrace {
110134
case DelegationIndirectAccessVariableOffset: {
111135
const u32 variable_index = placeholder.payload[0];
112136
const auto cycle_data = tracing_data + trace_row;
113-
return cycle_data->variables_offsets[variable_index];
137+
return cycle_data->variables_offsets()[variable_index];
114138
}
115139
default:
116140
__trap();
@@ -138,11 +162,11 @@ template <typename DESCRIPTION> struct DelegationTrace {
138162
case DelegationWriteTimestamp:
139163
return TimestampData::from_scalar(cycle_data->write_timestamp);
140164
case DelegationRegisterReadTimestamp: {
141-
return cycle_data->reg_accesses[reg_offset].timestamp;
165+
return cycle_data->reg_accesses()[reg_offset].timestamp;
142166
}
143167
case DelegationIndirectReadTimestamp: {
144-
return DESCRIPTION::use_read_indirects(register_index) ? cycle_data->indirect_reads[word_index].timestamp
145-
: cycle_data->indirect_writes[word_index].timestamp;
168+
return DESCRIPTION::use_read_indirects(register_index) ? cycle_data->indirect_reads()[word_index].timestamp
169+
: cycle_data->indirect_writes()[word_index].timestamp;
146170
}
147171
default:
148172
__trap();

‎gpu_prover/native/witness/witness_generation.cuh‎

Lines changed: 11 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -390,26 +390,28 @@ template <class R> struct WitnessProxy {
390390
if (VAR(S).inner) { \
391391
T \
392392
}
393-
#define SET_MEMORY_PLACE(IDX, V) p.template set_memory_place(IDX, VAR(V));
394-
#define SET_WITNESS_PLACE(IDX, V) p.template set_witness_place(IDX, VAR(V));
393+
#define SET_MEMORY_PLACE(IDX, V) p.set_memory_place(IDX, VAR(V));
394+
#define SET_WITNESS_PLACE(IDX, V) p.set_witness_place(IDX, VAR(V));
395395
#define SET_SCRATCH_PLACE(IDX, V) p.set_scratch_place(IDX, VAR(V));
396396

397397
#define FN_BEGIN(N) template <class R> DEVICE_FORCEINLINE void fn_##N(const WitnessProxy<R> p) {
398398
#define FN_END }
399399

400400
#define FN_CALL(N) fn_##N(p);
401401

402+
// NOLINTBEGIN
402403
// clang-format off
403-
#define INCLUDE_PREFIX ../../../../circuit_defs/ // whitespace! NOLINT
404-
#define UNROLLED_INCLUDE_PREFIX ../../../../circuit_defs/unrolled_circuits/ // whitespace! NOLINT
405-
#define INCLUDE_SUFFIX /generated/witness_generation_fn.cuh
404+
#define INCLUDE_PREFIX ../../../../circuit_defs
405+
#define UNROLLED_INCLUDE_PREFIX ../../../../circuit_defs/unrolled_circuits
406+
#define INCLUDE_SUFFIX generated/witness_generation_fn.cuh
407+
#define PATH_CAT(a, b, c) a/b/c
406408
// clang-format on
409+
// NOLINTEND
407410
#define STRINGIFY(X) STRINGIFY2(X)
408411
#define STRINGIFY2(X) #X
409-
#define IDENT(x) x
410-
#define CAT_3(x, y, z) IDENT(x) IDENT(y) IDENT(z)
411-
#define CIRCUIT_INCLUDE(NAME) STRINGIFY(CAT_3(INCLUDE_PREFIX, NAME, INCLUDE_SUFFIX))
412-
#define UNROLLED_CIRCUIT_INCLUDE(NAME) STRINGIFY(CAT_3(UNROLLED_INCLUDE_PREFIX, NAME, INCLUDE_SUFFIX))
412+
#define CIRCUIT_INCLUDE(NAME) STRINGIFY(PATH_CAT(INCLUDE_PREFIX, NAME, INCLUDE_SUFFIX))
413+
#define UNROLLED_CIRCUIT_INCLUDE(NAME) STRINGIFY(PATH_CAT(UNROLLED_INCLUDE_PREFIX, NAME, INCLUDE_SUFFIX))
414+
413415
#define KERNEL_NAME(NAME) ab_generate_witness_values_##NAME##_kernel
414416
#define KERNEL(NAME, ORACLE) \
415417
EXTERN __global__ void KERNEL_NAME(NAME)(const __grid_constant__ ORACLE oracle, const wrapped_f *const __restrict__ generic_lookup_tables, \

0 commit comments

Comments
 (0)