Skip to content

Commit 9fa0dfb

Browse files
committed
powers of alpha device fix
1 parent 02295e5 commit 9fa0dfb

File tree

2 files changed

+6
-4
lines changed

2 files changed

+6
-4
lines changed

core/src/stark/quotient.rs

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -114,6 +114,8 @@ where
114114
)
115115
.unwrap();
116116
quotient_flat.set_max_width();
117+
let powers_of_folding_challenge_device =
118+
powers_of_folding_challenge.to_device_async(stream).unwrap();
117119
quotient_gpu::compute_values(
118120
operations_device.as_ptr(),
119121
operations.len(),
@@ -128,7 +130,7 @@ where
128130
main_evaluations.view(),
129131
permutation_evaluations.view(),
130132
permutation_challenges.as_ptr(),
131-
powers_of_folding_challenge.as_ptr(),
133+
powers_of_folding_challenge_device.as_ptr(),
132134
public_values.as_ptr(),
133135
trace_domain_generator,
134136
quotient_domain_generator,

cuda/quotient/quotient.cuh

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -77,8 +77,8 @@ __global__ void computeValues(
7777
folder.isLastRow = isLastRow;
7878
folder.isTransition = isTransition;
7979
folder.powersOfAlpha = powersOfAlpha;
80-
folder.accumulator = Challenge::zero();
8180
folder.constraintIndex = 0;
81+
folder.accumulator = Challenge::zero();
8282
folder.quotientIdx = quotientIdx;
8383
folder.quotientSize = quotientSize;
8484
folder.nextStep = nextStep;
@@ -569,12 +569,12 @@ __global__ void computeValues(
569569

570570
case 59:
571571
DEBUG("FAssertZero: %d\n", instr.a);
572-
folder.accumulator += folder.powersOfAlpha[folder.constraintIndex] * expr_f[instr.a];
572+
folder.accumulator += (folder.powersOfAlpha[folder.constraintIndex] * expr_f[instr.a]);
573573
folder.constraintIndex++;
574574
break;
575575
case 60:
576576
DEBUG("EAssertZero: %d\n", instr.a);
577-
folder.accumulator += folder.powersOfAlpha[folder.constraintIndex] * expr_ef[instr.a];
577+
folder.accumulator += (folder.powersOfAlpha[folder.constraintIndex] * expr_ef[instr.a]);
578578
folder.constraintIndex++;
579579
break;
580580
}

0 commit comments

Comments
 (0)