Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Minor gpu fixes #130

Merged
merged 1 commit into from
Dec 3, 2024
Merged

Minor gpu fixes #130

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 0 additions & 1 deletion cli/src/commands/pil_helpers.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -202,7 +202,6 @@ impl PilHelpersCmd {
.symbols
.iter()
.filter(|symbol| {
println!("{:?}", symbol);
symbol.air_group_id.is_some()
&& symbol.air_group_id.unwrap() == airgroup_id as u32
&& ((symbol.air_id.is_some() && symbol.air_id.unwrap() == air_id as u32)
Expand Down
7 changes: 4 additions & 3 deletions pil2-stark/src/starkpil/expressions_ctx.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -14,10 +14,11 @@ struct Params {
uint64_t polsMapId;
uint64_t rowOffsetIndex;
bool inverse = false;
bool batch = true;
opType op;
uint64_t value;

Params(ParserParams& params, bool inverse_ = false) : parserParams(params), inverse(inverse_) {
Params(ParserParams& params, bool inverse_ = false, bool batch_ = true) : parserParams(params), inverse(inverse_), batch(batch_) {
dim = params.destDim;
op = opType::tmp;
}
Expand All @@ -41,8 +42,8 @@ struct Dest {

Dest(Goldilocks::Element *dest_, uint64_t offset_ = false) : dest(dest_), offset(offset_) {}

void addParams(ParserParams& parserParams_, bool inverse_ = false) {
params.push_back(Params(parserParams_, inverse_));
void addParams(ParserParams& parserParams_, bool inverse_ = false, bool batch_ = true) {
params.push_back(Params(parserParams_, inverse_, batch_));
uint64_t dimExp = parserParams_.destDim;
dim = std::max(dim, dimExp);
}
Expand Down
24 changes: 18 additions & 6 deletions pil2-stark/src/starkpil/expressions_pack.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -170,10 +170,16 @@ class ExpressionsPack : public ExpressionsCtx {
}
}

inline void copyPolynomial(Goldilocks::Element* destVals, bool inverse, uint64_t dim, Goldilocks::Element* tmp) {
inline void copyPolynomial(Goldilocks::Element* destVals, bool inverse, bool batch, uint64_t dim, Goldilocks::Element* tmp) {
if(dim == 1) {
if(inverse) {
Goldilocks::batchInverse(&destVals[0], &tmp[0], nrowsPack);
if(batch) {
Goldilocks::batchInverse(&destVals[0], &tmp[0], nrowsPack);
} else {
for(uint64_t i = 0; i < nrowsPack; ++i) {
Goldilocks::inv(destVals[i], tmp[i]);
}
}
} else {
Goldilocks::copy_pack(nrowsPack, &destVals[0], &tmp[0]);
}
Expand All @@ -183,7 +189,13 @@ class ExpressionsPack : public ExpressionsCtx {
Goldilocks::copy_pack(nrowsPack, &buff[0], uint64_t(FIELD_EXTENSION), &tmp[0]);
Goldilocks::copy_pack(nrowsPack, &buff[1], uint64_t(FIELD_EXTENSION), &tmp[nrowsPack]);
Goldilocks::copy_pack(nrowsPack, &buff[2], uint64_t(FIELD_EXTENSION), &tmp[2*nrowsPack]);
Goldilocks3::batchInverse((Goldilocks3::Element *)buff, (Goldilocks3::Element *)buff, nrowsPack);
if(batch) {
Goldilocks3::batchInverse((Goldilocks3::Element *)buff, (Goldilocks3::Element *)buff, nrowsPack);
} else {
for(uint64_t i = 0; i < nrowsPack; ++i) {
Goldilocks3::inv((Goldilocks3::Element &)buff[i*FIELD_EXTENSION], (Goldilocks3::Element &)buff[i*FIELD_EXTENSION]);
}
}
Goldilocks::copy_pack(nrowsPack, &destVals[0], &buff[0], uint64_t(FIELD_EXTENSION));
Goldilocks::copy_pack(nrowsPack, &destVals[nrowsPack], &buff[1], uint64_t(FIELD_EXTENSION));
Goldilocks::copy_pack(nrowsPack, &destVals[2*nrowsPack], &buff[2], uint64_t(FIELD_EXTENSION));
Expand Down Expand Up @@ -334,7 +346,7 @@ class ExpressionsPack : public ExpressionsCtx {
uint64_t openingPointIndex = dests[j].params[k].rowOffsetIndex;
uint64_t buffPos = ns*openingPointIndex + dests[j].params[k].stage;
uint64_t stagePos = dests[j].params[k].stagePos;
copyPolynomial(&destVals[j][k*FIELD_EXTENSION*nrowsPack], dests[j].params[k].inverse, dests[j].params[k].dim, &bufferT_[(nColsStagesAcc[buffPos] + stagePos)*nrowsPack]);
copyPolynomial(&destVals[j][k*FIELD_EXTENSION*nrowsPack], dests[j].params[k].inverse, dests[j].params[k].batch, dests[j].params[k].dim, &bufferT_[(nColsStagesAcc[buffPos] + stagePos)*nrowsPack]);
continue;
} else if(dests[j].params[k].op == opType::number) {
uint64_t val = dests[j].params[k].inverse ? Goldilocks::inv(Goldilocks::fromU64(dests[j].params[k].value)).fe : dests[j].params[k].value;
Expand Down Expand Up @@ -788,9 +800,9 @@ class ExpressionsPack : public ExpressionsCtx {
assert(i_args == dests[j].params[k].parserParams.nArgs);

if(dests[j].params[k].parserParams.destDim == 1) {
copyPolynomial(&destVals[j][k*FIELD_EXTENSION*nrowsPack], dests[j].params[k].inverse, dests[j].params[k].parserParams.destDim, &tmp1[dests[j].params[k].parserParams.destId*nrowsPack]);
copyPolynomial(&destVals[j][k*FIELD_EXTENSION*nrowsPack], dests[j].params[k].inverse, dests[j].params[k].batch, dests[j].params[k].parserParams.destDim, &tmp1[dests[j].params[k].parserParams.destId*nrowsPack]);
} else {
copyPolynomial(&destVals[j][k*FIELD_EXTENSION*nrowsPack], dests[j].params[k].inverse, dests[j].params[k].parserParams.destDim, &tmp3[dests[j].params[k].parserParams.destId*FIELD_EXTENSION*nrowsPack]);
copyPolynomial(&destVals[j][k*FIELD_EXTENSION*nrowsPack], dests[j].params[k].inverse, dests[j].params[k].batch, dests[j].params[k].parserParams.destDim, &tmp3[dests[j].params[k].parserParams.destId*FIELD_EXTENSION*nrowsPack]);
}
}

Expand Down
40 changes: 12 additions & 28 deletions pil2-stark/src/starkpil/gen_recursive_proof.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -26,15 +26,14 @@ void *genRecursiveProof(SetupCtx& setupCtx, json& globalInfo, uint64_t airgroupI

Goldilocks::Element* evals = new Goldilocks::Element[setupCtx.starkInfo.evMap.size() * FIELD_EXTENSION];
Goldilocks::Element* challenges = new Goldilocks::Element[setupCtx.starkInfo.challengesMap.size() * FIELD_EXTENSION];
Goldilocks::Element* airgroupValues = new Goldilocks::Element[setupCtx.starkInfo.airgroupValuesMap.size() * FIELD_EXTENSION];

vector<bool> airgroupValuesCalculated(setupCtx.starkInfo.airgroupValuesMap.size(), false);
StepsParams params = {
trace: pAddress,
pols : pols,
publicInputs : publicInputs,
challenges : challenges,
airgroupValues : airgroupValues,
airgroupValues : nullptr,
evals : evals,
xDivXSub : nullptr,
pConstPolsAddress: pConstPols,
Expand Down Expand Up @@ -82,47 +81,33 @@ void *genRecursiveProof(SetupCtx& setupCtx, json& globalInfo, uint64_t airgroupI
}

uint64_t N = 1 << setupCtx.starkInfo.starkStruct.nBits;
Goldilocks::Element *num = new Goldilocks::Element[N*FIELD_EXTENSION];
Goldilocks::Element *den = new Goldilocks::Element[N*FIELD_EXTENSION];
Goldilocks::Element *res = new Goldilocks::Element[N*FIELD_EXTENSION];
Goldilocks::Element *gprod = new Goldilocks::Element[N*FIELD_EXTENSION];

Hint gprod_hint = setupCtx.expressionsBin.hints[0];
auto denField = std::find_if(gprod_hint.fields.begin(), gprod_hint.fields.end(), [](const HintField& hintField) {
return hintField.name == "denominator";
});
auto numField = std::find_if(gprod_hint.fields.begin(), gprod_hint.fields.end(), [](const HintField& hintField) {
return hintField.name == "numerator";
});
auto gprodField = std::find_if(gprod_hint.fields.begin(), gprod_hint.fields.end(), [](const HintField& hintField) {
return hintField.name == "reference";
});
uint64_t gprodFieldId = setupCtx.expressionsBin.hints[0].fields[0].values[0].id;
uint64_t numFieldId = setupCtx.expressionsBin.hints[0].fields[1].values[0].id;
uint64_t denFieldId = setupCtx.expressionsBin.hints[0].fields[2].values[0].id;


Dest numStruct(num);
numStruct.addParams(setupCtx.expressionsBin.expressionsInfo[numField->values[0].id]);
Dest denStruct(den);
denStruct.addParams(setupCtx.expressionsBin.expressionsInfo[denField->values[0].id], true);
std::vector<Dest> dests = {numStruct, denStruct};
Dest destStruct(res);
destStruct.addParams(setupCtx.expressionsBin.expressionsInfo[numFieldId]);
destStruct.addParams(setupCtx.expressionsBin.expressionsInfo[denFieldId], true);
std::vector<Dest> dests = {destStruct};

expressionsCtx.calculateExpressions(params, setupCtx.expressionsBin.expressionsBinArgsExpressions, dests, uint64_t(1 << setupCtx.starkInfo.starkStruct.nBits));


Goldilocks3::copy((Goldilocks3::Element *)&gprod[0], &Goldilocks3::one());
for(uint64_t i = 1; i < N; ++i) {
Goldilocks::Element res[3];
Goldilocks3::mul((Goldilocks3::Element *)res, (Goldilocks3::Element *)&num[(i - 1) * FIELD_EXTENSION], (Goldilocks3::Element *)&den[(i - 1) * FIELD_EXTENSION]);
Goldilocks3::mul((Goldilocks3::Element *)&gprod[i * FIELD_EXTENSION], (Goldilocks3::Element *)&gprod[(i - 1) * FIELD_EXTENSION], (Goldilocks3::Element *)res);
Goldilocks3::mul((Goldilocks3::Element *)&gprod[i * FIELD_EXTENSION], (Goldilocks3::Element *)&gprod[(i - 1) * FIELD_EXTENSION], (Goldilocks3::Element *)&res[(i - 1) * FIELD_EXTENSION]);
}

Polinomial gprodTransposedPol;
setupCtx.starkInfo.getPolynomial(gprodTransposedPol, params.pols, "cm", setupCtx.starkInfo.cmPolsMap[gprodField->values[0].id], false);
setupCtx.starkInfo.getPolynomial(gprodTransposedPol, params.pols, "cm", setupCtx.starkInfo.cmPolsMap[gprodFieldId], false);
#pragma omp parallel for
for(uint64_t j = 0; j < N; ++j) {
std::memcpy(gprodTransposedPol[j], &gprod[j*FIELD_EXTENSION], FIELD_EXTENSION * sizeof(Goldilocks::Element));
}

delete num;
delete den;
delete res;
delete gprod;

TimerStart(CALCULATE_IM_POLS);
Expand Down Expand Up @@ -249,7 +234,6 @@ void *genRecursiveProof(SetupCtx& setupCtx, json& globalInfo, uint64_t airgroupI

delete challenges;
delete evals;
delete airgroupValues;

nlohmann::json jProof = proof.proof.proof2json();
nlohmann::json zkin = proof2zkinStark(jProof, setupCtx.starkInfo);
Expand Down
Loading