diff --git a/sxt/proof/sumcheck/BUILD b/sxt/proof/sumcheck/BUILD index 99f9cf01..a80a29ed 100644 --- a/sxt/proof/sumcheck/BUILD +++ b/sxt/proof/sumcheck/BUILD @@ -1,200 +1,4 @@ # sxt_cc_component( -# name = "mle_utility", -# impl_deps = [ -# "//sxt/base/container:span_utility", -# "//sxt/base/device:memory_utility", -# "//sxt/base/device:property", -# "//sxt/base/num:divide_up", -# "//sxt/base/num:ceil_log2", -# "//sxt/memory/management:managed_array", -# "//sxt/scalar25/type:element", -# ], -# test_deps = [ -# "//sxt/base/device:stream", -# "//sxt/base/device:synchronization", -# "//sxt/base/test:unit_test", -# "//sxt/memory/management:managed_array", -# "//sxt/memory/resource:managed_device_resource", -# "//sxt/scalar25/type:element", -# "//sxt/scalar25/type:literal", -# ], -# deps = [ -# "//sxt/base/container:span", -# "//sxt/memory/management:managed_array_fwd", -# ], -# ) -# -# sxt_cc_component( -# name = "sum_gpu", -# impl_deps = [ -# ":constant", -# ":device_cache", -# ":mle_utility", -# ":polynomial_mapper", -# ":reduction_gpu", -# "//sxt/algorithm/reduction:kernel_fit", -# "//sxt/algorithm/reduction:thread_reduction", -# "//sxt/base/device:memory_utility", -# "//sxt/base/device:stream", -# "//sxt/base/device:state", -# "//sxt/base/iterator:split", -# "//sxt/base/num:ceil_log2", -# "//sxt/base/num:constexpr_switch", -# "//sxt/execution/async:coroutine", -# "//sxt/execution/async:future", -# "//sxt/execution/device:for_each", -# "//sxt/execution/kernel:kernel_dims", -# "//sxt/memory/resource:device_resource", -# "//sxt/memory/resource:async_device_resource", -# "//sxt/scalar25/operation:add", -# "//sxt/scalar25/operation:mul", -# "//sxt/scalar25/type:element", -# ], -# test_deps = [ -# ":device_cache", -# "//sxt/base/iterator:split", -# "//sxt/base/test:unit_test", -# "//sxt/execution/async:future", -# "//sxt/execution/schedule:scheduler", -# "//sxt/scalar25/operation:overload", -# "//sxt/scalar25/type:element", -# "//sxt/scalar25/type:literal", -# ], -# deps = [ -# "//sxt/base/container:span", -# "//sxt/execution/async:future_fwd", -# ], -# ) -# -# sxt_cc_component( -# name = "sumcheck_transcript", -# with_test = False, -# deps = [ -# "//sxt/base/container:span", -# ], -# ) -# -# sxt_cc_component( -# name = "reference_transcript", -# impl_deps = [ -# "//sxt/scalar25/type:element", -# "//sxt/proof/transcript:transcript_utility", -# ], -# test_deps = [ -# "//sxt/base/test:unit_test", -# "//sxt/scalar25/type:literal", -# ], -# deps = [ -# ":sumcheck_transcript", -# "//sxt/proof/transcript", -# ], -# ) -# -# sxt_cc_component( -# name = "sumcheck_random", -# impl_deps = [ -# "//sxt/base/error:assert", -# "//sxt/base/num:fast_random_number_generator", -# "//sxt/scalar25/random:element", -# "//sxt/scalar25/type:element", -# ], -# with_test = False, -# deps = [ -# ":constant", -# ], -# ) -# -# sxt_cc_component( -# name = "workspace", -# with_test = False, -# ) -# -# sxt_cc_component( -# name = "driver", -# with_test = False, -# deps = [ -# ":workspace", -# "//sxt/base/container:span", -# "//sxt/execution/async:future_fwd", -# ], -# ) -# -# sxt_cc_component( -# name = "driver_test", -# impl_deps = [ -# ":driver", -# ":workspace", -# "//sxt/base/test:unit_test", -# "//sxt/execution/async:future", -# "//sxt/execution/schedule:scheduler", -# "//sxt/scalar25/operation:overload", -# "//sxt/scalar25/type:element", -# "//sxt/scalar25/type:literal", -# ], -# with_test = False, -# ) -# -# sxt_cc_component( -# name = "cpu_driver", -# impl_deps = [ -# ":polynomial_utility", -# "//sxt/base/container:stack_array", -# "//sxt/base/error:panic", -# "//sxt/base/num:ceil_log2", -# "//sxt/execution/async:future", -# "//sxt/memory/management:managed_array", -# "//sxt/scalar25/operation:mul", -# "//sxt/scalar25/operation:sub", -# "//sxt/scalar25/operation:muladd", -# "//sxt/scalar25/type:element", -# "//sxt/scalar25/type:literal", -# ], -# test_deps = [ -# ":driver_test", -# "//sxt/base/test:unit_test", -# ], -# deps = [ -# ":driver", -# ":workspace", -# ], -# ) -# -# sxt_cc_component( -# name = "gpu_driver", -# impl_deps = [ -# ":constant", -# ":polynomial_mapper", -# ":sum_gpu", -# "//sxt/algorithm/iteration:for_each", -# "//sxt/algorithm/reduction:reduction", -# "//sxt/base/container:stack_array", -# "//sxt/base/device:stream", -# "//sxt/base/device:memory_utility", -# "//sxt/base/error:panic", -# "//sxt/base/num:ceil_log2", -# "//sxt/base/num:constexpr_switch", -# "//sxt/execution/async:coroutine", -# "//sxt/execution/async:future", -# "//sxt/execution/device:synchronization", -# "//sxt/memory/management:managed_array", -# "//sxt/memory/resource:async_device_resource", -# "//sxt/memory/resource:device_resource", -# "//sxt/scalar25/operation:mul", -# "//sxt/scalar25/operation:sub", -# "//sxt/scalar25/operation:muladd", -# "//sxt/scalar25/type:element", -# "//sxt/scalar25/type:literal", -# ], -# test_deps = [ -# ":driver_test", -# "//sxt/base/test:unit_test", -# ], -# deps = [ -# ":driver", -# ], -# ) -# -# sxt_cc_component( # name = "polynomial_utility", # impl_deps = [ # "//sxt/scalar25/operation:add", diff --git a/sxt/proof/sumcheck/chunked_gpu_driver.cc b/sxt/proof/sumcheck/chunked_gpu_driver.cc deleted file mode 100644 index 1358c4cc..00000000 --- a/sxt/proof/sumcheck/chunked_gpu_driver.cc +++ /dev/null @@ -1,154 +0,0 @@ -/** Proofs GPU - Space and Time's cryptographic proof algorithms on the CPU and GPU. - * - * Copyright 2025-present Space and Time Labs, Inc. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#include "sxt/proof/sumcheck/chunked_gpu_driver.h" - -#include - -#include "sxt/algorithm/iteration/transform.h" -#include "sxt/base/error/assert.h" -#include "sxt/base/num/ceil_log2.h" -#include "sxt/execution/async/coroutine.h" -#include "sxt/execution/async/future.h" -#include "sxt/memory/management/managed_array.h" -#include "sxt/proof/sumcheck/device_cache.h" -#include "sxt/proof/sumcheck/fold_gpu.h" -#include "sxt/proof/sumcheck/gpu_driver.h" -#include "sxt/proof/sumcheck/mle_utility.h" -#include "sxt/proof/sumcheck/sum_gpu.h" -#include "sxt/scalar25/operation/sub.h" -#include "sxt/scalar25/type/element.h" -#include "sxt/scalar25/type/literal.h" - -namespace sxt::prfsk { -//-------------------------------------------------------------------------------------------------- -// chunked_gpu_workspace -//-------------------------------------------------------------------------------------------------- -namespace { -struct chunked_gpu_workspace final : public workspace { - std::unique_ptr single_gpu_workspace; - - device_cache cache; - memmg::managed_array mles_data; - basct::cspan mles; - unsigned n; - unsigned num_variables; - - chunked_gpu_workspace(basct::cspan> product_table, - basct::cspan product_terms) noexcept - : cache{product_table, product_terms} {} -}; -} // namespace - -//-------------------------------------------------------------------------------------------------- -// try_make_single_gpu_workspace -//-------------------------------------------------------------------------------------------------- -static xena::future<> try_make_single_gpu_workspace(chunked_gpu_workspace& work, - double no_chunk_cutoff) noexcept { - auto gpu_memory_fraction = get_gpu_memory_fraction(work.mles); - if (gpu_memory_fraction > no_chunk_cutoff) { - co_return; - } - - // construct single gpu workspace - auto cache_data = work.cache.clear(); - gpu_driver drv; - work.single_gpu_workspace = - co_await drv.make_workspace(work.mles, std::move(cache_data->product_table), - std::move(cache_data->product_terms), work.n); - - // free data we no longer need - work.mles_data.reset(); - work.mles = {}; -} - -//-------------------------------------------------------------------------------------------------- -// constructor -//-------------------------------------------------------------------------------------------------- -chunked_gpu_driver::chunked_gpu_driver(double no_chunk_cutoff) noexcept - : no_chunk_cutoff_{no_chunk_cutoff} { - SXT_RELEASE_ASSERT(0 <= no_chunk_cutoff_ && no_chunk_cutoff_ <= 1.0); -} - -//-------------------------------------------------------------------------------------------------- -// make_workspace -//-------------------------------------------------------------------------------------------------- -xena::future> -chunked_gpu_driver::make_workspace(basct::cspan mles, - basct::cspan> product_table, - basct::cspan product_terms, - unsigned n) const noexcept { - auto res = std::make_unique(product_table, product_terms); - res->mles = mles; - res->n = n; - res->num_variables = std::max(basn::ceil_log2(n), 1); - auto gpu_memory_fraction = get_gpu_memory_fraction(mles); - if (gpu_memory_fraction <= no_chunk_cutoff_) { - gpu_driver drv; - res->single_gpu_workspace = co_await drv.make_workspace(mles, product_table, product_terms, n); - } - co_return std::unique_ptr(std::move(res)); -} - -//-------------------------------------------------------------------------------------------------- -// sum -//-------------------------------------------------------------------------------------------------- -xena::future<> chunked_gpu_driver::sum(basct::span polynomial, - workspace& ws) const noexcept { - auto& work = static_cast(ws); - if (work.single_gpu_workspace) { - gpu_driver drv; - co_return co_await drv.sum(polynomial, *work.single_gpu_workspace); - } - co_await sum_gpu(polynomial, work.cache, work.mles, work.n); -} - -//-------------------------------------------------------------------------------------------------- -// fold -//-------------------------------------------------------------------------------------------------- -xena::future<> chunked_gpu_driver::fold(workspace& ws, const s25t::element& r) const noexcept { - using s25t::operator""_s25; - auto& work = static_cast(ws); - if (work.single_gpu_workspace) { - gpu_driver drv; - co_return co_await drv.fold(*work.single_gpu_workspace, r); - } - auto n = work.n; - auto mid = 1u << (work.num_variables - 1u); - auto num_mles = work.mles.size() / n; - SXT_RELEASE_ASSERT( - // clang-format off - work.n >= mid && work.mles.size() % n == 0 - // clang-format on - ); - - s25t::element one_m_r = 0x1_s25; - s25o::sub(one_m_r, one_m_r, r); - - // fold - memmg::managed_array mles_p(num_mles * mid); - co_await fold_gpu(mles_p, work.mles, n, r); - - // update - work.n = mid; - --work.num_variables; - work.mles_data = std::move(mles_p); - work.mles = work.mles_data; - - // check if we should fall back to single gpu workspace - co_await try_make_single_gpu_workspace(work, no_chunk_cutoff_); -} -} // namespace sxt::prfsk diff --git a/sxt/proof/sumcheck/chunked_gpu_driver.h b/sxt/proof/sumcheck/chunked_gpu_driver.h deleted file mode 100644 index ed93decb..00000000 --- a/sxt/proof/sumcheck/chunked_gpu_driver.h +++ /dev/null @@ -1,42 +0,0 @@ -/** Proofs GPU - Space and Time's cryptographic proof algorithms on the CPU and GPU. - * - * Copyright 2025-present Space and Time Labs, Inc. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#pragma once - -#include "sxt/proof/sumcheck/driver.h" - -namespace sxt::prfsk { -//-------------------------------------------------------------------------------------------------- -// chunked_gpu_driver -//-------------------------------------------------------------------------------------------------- -class chunked_gpu_driver final : public driver { -public: - explicit chunked_gpu_driver(double no_chunk_cutoff = 0.5) noexcept; - - // driver - xena::future> - make_workspace(basct::cspan mles, - basct::cspan> product_table, - basct::cspan product_terms, unsigned n) const noexcept override; - - xena::future<> sum(basct::span polynomial, workspace& ws) const noexcept override; - - xena::future<> fold(workspace& ws, const s25t::element& r) const noexcept override; - -private: - double no_chunk_cutoff_; -}; -} // namespace sxt::prfsk diff --git a/sxt/proof/sumcheck/chunked_gpu_driver.t.cc b/sxt/proof/sumcheck/chunked_gpu_driver.t.cc deleted file mode 100644 index 8fb4a3aa..00000000 --- a/sxt/proof/sumcheck/chunked_gpu_driver.t.cc +++ /dev/null @@ -1,35 +0,0 @@ -/** Proofs GPU - Space and Time's cryptographic proof algorithms on the CPU and GPU. - * - * Copyright 2025-present Space and Time Labs, Inc. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#include "sxt/proof/sumcheck/chunked_gpu_driver.h" - -#include "sxt/base/test/unit_test.h" -#include "sxt/proof/sumcheck/driver_test.h" - -using namespace sxt; -using namespace sxt::prfsk; - -TEST_CASE("we can perform the primitive operations for sumcheck proofs") { - SECTION("we handle the case when only chunking is used") { - chunked_gpu_driver drv{0.0}; - exercise_driver(drv); - } - - SECTION("we handle the case when the chunked driver falls back to the single gpu driver") { - chunked_gpu_driver drv{1.0}; - exercise_driver(drv); - } -} diff --git a/sxt/proof/sumcheck/cpu_driver.cc b/sxt/proof/sumcheck/cpu_driver.cc deleted file mode 100644 index cc8d6919..00000000 --- a/sxt/proof/sumcheck/cpu_driver.cc +++ /dev/null @@ -1,161 +0,0 @@ -/** Proofs GPU - Space and Time's cryptographic proof algorithms on the CPU and GPU. - * - * Copyright 2024-present Space and Time Labs, Inc. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#include "sxt/proof/sumcheck/cpu_driver.h" - -#include - -#include "sxt/base/container/stack_array.h" -#include "sxt/base/error/panic.h" -#include "sxt/base/num/ceil_log2.h" -#include "sxt/execution/async/future.h" -#include "sxt/memory/management/managed_array.h" -#include "sxt/proof/sumcheck/polynomial_utility.h" -#include "sxt/scalar25/operation/mul.h" -#include "sxt/scalar25/operation/muladd.h" -#include "sxt/scalar25/operation/sub.h" -#include "sxt/scalar25/type/element.h" -#include "sxt/scalar25/type/literal.h" - -namespace sxt::prfsk { -//-------------------------------------------------------------------------------------------------- -// cpu_workspace -//-------------------------------------------------------------------------------------------------- -namespace { -struct cpu_workspace final : public workspace { - memmg::managed_array mles; - basct::cspan> product_table; - basct::cspan product_terms; - unsigned n; - unsigned num_variables; -}; -} // namespace - -//-------------------------------------------------------------------------------------------------- -// make_workspace -//-------------------------------------------------------------------------------------------------- -xena::future> -cpu_driver::make_workspace(basct::cspan mles, - basct::cspan> product_table, - basct::cspan product_terms, unsigned n) const noexcept { - auto res = std::make_unique(); - res->mles = memmg::managed_array{mles.begin(), mles.end()}; - res->product_table = product_table; - res->product_terms = product_terms; - res->n = n; - res->num_variables = std::max(basn::ceil_log2(n), 1); - return xena::make_ready_future>(std::move(res)); -} - -//-------------------------------------------------------------------------------------------------- -// sum -//-------------------------------------------------------------------------------------------------- -xena::future<> cpu_driver::sum(basct::span polynomial, - workspace& ws) const noexcept { - auto& work = static_cast(ws); - auto n = work.n; - auto mid = 1u << (work.num_variables - 1u); - SXT_RELEASE_ASSERT(work.n >= mid); - - auto mles = work.mles.data(); - auto product_table = work.product_table; - auto product_terms = work.product_terms; - - for (auto& val : polynomial) { - val = {}; - } - - // expand paired terms - auto n1 = work.n - mid; - for (unsigned i = 0; i < n1; ++i) { - unsigned term_first = 0; - for (auto [mult, num_terms] : product_table) { - SXT_RELEASE_ASSERT(num_terms < polynomial.size()); - auto terms = product_terms.subspan(term_first, num_terms); - SXT_STACK_ARRAY(p, num_terms + 1u, s25t::element); - expand_products(p, mles + i, n, mid, terms); - for (unsigned term_index = 0; term_index < p.size(); ++term_index) { - s25o::muladd(polynomial[term_index], mult, p[term_index], polynomial[term_index]); - } - term_first += num_terms; - } - } - - // expand terms where the corresponding pair is zero (i.e. n is not a power of 2) - for (unsigned i = n1; i < mid; ++i) { - unsigned term_first = 0; - for (auto [mult, num_terms] : product_table) { - auto terms = product_terms.subspan(term_first, num_terms); - SXT_STACK_ARRAY(p, num_terms + 1u, s25t::element); - partial_expand_products(p, mles + i, n, terms); - for (unsigned term_index = 0; term_index < p.size(); ++term_index) { - s25o::muladd(polynomial[term_index], mult, p[term_index], polynomial[term_index]); - } - term_first += num_terms; - } - } - - return xena::make_ready_future(); -} - -//-------------------------------------------------------------------------------------------------- -// fold -//-------------------------------------------------------------------------------------------------- -xena::future<> cpu_driver::fold(workspace& ws, const s25t::element& r) const noexcept { - using s25t::operator""_s25; - - auto& work = static_cast(ws); - auto n = work.n; - auto mid = 1u << (work.num_variables - 1u); - auto num_mles = work.mles.size() / n; - SXT_RELEASE_ASSERT( - // clang-format off - work.n >= mid && work.mles.size() % n == 0 - // clang-format on - ); - - auto mles = work.mles.data(); - memmg::managed_array mles_p(num_mles * mid); - - s25t::element one_m_r = 0x1_s25; - s25o::sub(one_m_r, one_m_r, r); - auto n1 = work.n - mid; - for (auto mle_index = 0; mle_index < num_mles; ++mle_index) { - auto data = mles + n * mle_index; - auto data_p = mles_p.data() + mid * mle_index; - - // fold paired terms - for (unsigned i = 0; i < n1; ++i) { - auto val = data[i]; - s25o::mul(val, val, one_m_r); - s25o::muladd(val, r, data[mid + i], val); - data_p[i] = val; - } - - // fold terms paired with zero - for (unsigned i = n1; i < mid; ++i) { - auto val = data[i]; - s25o::mul(val, val, one_m_r); - data_p[i] = val; - } - } - - work.n = mid; - --work.num_variables; - work.mles = std::move(mles_p); - return xena::make_ready_future(); -} -} // namespace sxt::prfsk diff --git a/sxt/proof/sumcheck/cpu_driver.h b/sxt/proof/sumcheck/cpu_driver.h deleted file mode 100644 index b017396d..00000000 --- a/sxt/proof/sumcheck/cpu_driver.h +++ /dev/null @@ -1,37 +0,0 @@ -/** Proofs GPU - Space and Time's cryptographic proof algorithms on the CPU and GPU. - * - * Copyright 2024-present Space and Time Labs, Inc. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#pragma once - -#include "sxt/proof/sumcheck/driver.h" - -namespace sxt::prfsk { -//-------------------------------------------------------------------------------------------------- -// cpu_driver -//-------------------------------------------------------------------------------------------------- -class cpu_driver final : public driver { -public: - // driver - xena::future> - make_workspace(basct::cspan mles, - basct::cspan> product_table, - basct::cspan product_terms, unsigned n) const noexcept override; - - xena::future<> sum(basct::span polynomial, workspace& ws) const noexcept override; - - xena::future<> fold(workspace& ws, const s25t::element& r) const noexcept override; -}; -} // namespace sxt::prfsk diff --git a/sxt/proof/sumcheck/cpu_driver.t.cc b/sxt/proof/sumcheck/cpu_driver.t.cc deleted file mode 100644 index 42ae1e7c..00000000 --- a/sxt/proof/sumcheck/cpu_driver.t.cc +++ /dev/null @@ -1,30 +0,0 @@ -/** Proofs GPU - Space and Time's cryptographic proof algorithms on the CPU and GPU. - * - * Copyright 2024-present Space and Time Labs, Inc. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#include "sxt/proof/sumcheck/cpu_driver.h" - -#include - -#include "sxt/base/test/unit_test.h" -#include "sxt/proof/sumcheck/driver_test.h" - -using namespace sxt; -using namespace sxt::prfsk; - -TEST_CASE("we can perform the primitive operations for sumcheck proofs") { - cpu_driver drv; - exercise_driver(drv); -} diff --git a/sxt/proof/sumcheck/driver.cc b/sxt/proof/sumcheck/driver.cc deleted file mode 100644 index 6e46927f..00000000 --- a/sxt/proof/sumcheck/driver.cc +++ /dev/null @@ -1,17 +0,0 @@ -/** Proofs GPU - Space and Time's cryptographic proof algorithms on the CPU and GPU. - * - * Copyright 2024-present Space and Time Labs, Inc. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#include "sxt/proof/sumcheck/driver.h" diff --git a/sxt/proof/sumcheck/driver.h b/sxt/proof/sumcheck/driver.h deleted file mode 100644 index 422bb88a..00000000 --- a/sxt/proof/sumcheck/driver.h +++ /dev/null @@ -1,47 +0,0 @@ -/** Proofs GPU - Space and Time's cryptographic proof algorithms on the CPU and GPU. - * - * Copyright 2024-present Space and Time Labs, Inc. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#pragma once - -#include - -#include "sxt/base/container/span.h" -#include "sxt/execution/async/future_fwd.h" -#include "sxt/proof/sumcheck/workspace.h" - -namespace sxt::s25t { -class element; -} - -namespace sxt::prfsk { -//-------------------------------------------------------------------------------------------------- -// driver -//-------------------------------------------------------------------------------------------------- -class driver { -public: - virtual ~driver() noexcept = default; - - virtual xena::future> - make_workspace(basct::cspan mles, - basct::cspan> product_table, - basct::cspan product_terms, unsigned n) const noexcept = 0; - - virtual xena::future<> sum(basct::span polynomial, - workspace& ws) const noexcept = 0; - - virtual xena::future<> fold(workspace& ws, const s25t::element& r) const noexcept = 0; -}; -} // namespace sxt::prfsk diff --git a/sxt/proof/sumcheck/driver_test.cc b/sxt/proof/sumcheck/driver_test.cc deleted file mode 100644 index 97fa80df..00000000 --- a/sxt/proof/sumcheck/driver_test.cc +++ /dev/null @@ -1,133 +0,0 @@ -/** Proofs GPU - Space and Time's cryptographic proof algorithms on the CPU and GPU. - * - * Copyright 2024-present Space and Time Labs, Inc. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#include "sxt/proof/sumcheck/driver_test.h" - -#include - -#include "sxt/base/test/unit_test.h" -#include "sxt/execution/async/future.h" -#include "sxt/execution/schedule/scheduler.h" -#include "sxt/proof/sumcheck/driver.h" -#include "sxt/proof/sumcheck/workspace.h" -#include "sxt/scalar25/operation/overload.h" -#include "sxt/scalar25/type/element.h" -#include "sxt/scalar25/type/literal.h" - -namespace sxt::prfsk { -using s25t::operator""_s25; - -//-------------------------------------------------------------------------------------------------- -// exercise_driver -//-------------------------------------------------------------------------------------------------- -void exercise_driver(const driver& drv) { - std::vector mles; - std::vector> product_table{ - {0x1_s25, 1}, - }; - std::vector product_terms = {0}; - - std::vector p(2); - - SECTION("we can sum a polynomial with n = 1") { - std::vector mles = {0x123_s25}; - auto ws = drv.make_workspace(mles, product_table, product_terms, 1); - xens::get_scheduler().run(); - auto fut = drv.sum(p, *ws.value()); - xens::get_scheduler().run(); - REQUIRE(fut.ready()); - REQUIRE(p[0] == mles[0]); - REQUIRE(p[1] == -mles[0]); - } - - SECTION("we can sum a polynomial with a non-unity multiplier") { - std::vector mles = {0x123_s25}; - product_table[0].first = 0x2_s25; - auto ws = drv.make_workspace(mles, product_table, product_terms, 1); - xens::get_scheduler().run(); - auto fut = drv.sum(p, *ws.value()); - xens::get_scheduler().run(); - REQUIRE(fut.ready()); - REQUIRE(p[0] == 0x2_s25 * mles[0]); - REQUIRE(p[1] == -0x2_s25 * mles[0]); - } - - SECTION("we can sum a polynomial with n = 2") { - std::vector mles = {0x123_s25, 0x456_s25}; - auto ws = drv.make_workspace(mles, product_table, product_terms, 2); - xens::get_scheduler().run(); - auto fut = drv.sum(p, *ws.value()); - xens::get_scheduler().run(); - REQUIRE(fut.ready()); - REQUIRE(p[0] == mles[0]); - REQUIRE(p[1] == mles[1] - mles[0]); - } - - SECTION("we can sum a polynomial with two MLEs added together") { - std::vector mles = {0x123_s25, 0x456_s25}; - std::vector> product_table{ - {0x1_s25, 1}, - {0x1_s25, 1}, - }; - std::vector product_terms = {0, 1}; - - auto ws = drv.make_workspace(mles, product_table, product_terms, 1); - xens::get_scheduler().run(); - auto fut = drv.sum(p, *ws.value()); - xens::get_scheduler().run(); - REQUIRE(fut.ready()); - REQUIRE(p[0] == mles[0] + mles[1]); - REQUIRE(p[1] == -mles[0] - mles[1]); - } - - SECTION("we can sum a polynomial with two MLEs multiplied together") { - std::vector mles = {0x123_s25, 0x456_s25}; - std::vector> product_table{ - {0x1_s25, 2}, - }; - std::vector product_terms = {0, 1}; - p.resize(3); - - auto ws = drv.make_workspace(mles, product_table, product_terms, 1); - xens::get_scheduler().run(); - auto fut = drv.sum(p, *ws.value()); - xens::get_scheduler().run(); - REQUIRE(fut.ready()); - REQUIRE(p[0] == mles[0] * mles[1]); - REQUIRE(p[1] == -mles[0] * mles[1] - mles[1] * mles[0]); - REQUIRE(p[2] == mles[0] * mles[1]); - } - - SECTION("we can fold mles") { - std::vector mles = {0x123_s25, 0x456_s25, 0x789_s25}; - auto ws = drv.make_workspace(mles, product_table, product_terms, 3); - xens::get_scheduler().run(); - auto r = 0xabc123_s25; - auto fut = drv.fold(*ws.value(), r); - xens::get_scheduler().run(); - REQUIRE(fut.ready()); - fut = drv.sum(p, *ws.value()); - xens::get_scheduler().run(); - REQUIRE(fut.ready()); - - mles[0] = (0x1_s25 - r) * mles[0] + r * mles[2]; - mles[1] = (0x1_s25 - r) * mles[1]; - - REQUIRE(p[0] == mles[0]); - REQUIRE(p[1] == mles[1] - mles[0]); - } -} -} // namespace sxt::prfsk diff --git a/sxt/proof/sumcheck/driver_test.h b/sxt/proof/sumcheck/driver_test.h deleted file mode 100644 index 69c0a53e..00000000 --- a/sxt/proof/sumcheck/driver_test.h +++ /dev/null @@ -1,26 +0,0 @@ -/** Proofs GPU - Space and Time's cryptographic proof algorithms on the CPU and GPU. - * - * Copyright 2024-present Space and Time Labs, Inc. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#pragma once - -namespace sxt::prfsk { -class driver; - -//-------------------------------------------------------------------------------------------------- -// exercise_driver -//-------------------------------------------------------------------------------------------------- -void exercise_driver(const driver& drv); -} // namespace sxt::prfsk diff --git a/sxt/proof/sumcheck/gpu_driver.cc b/sxt/proof/sumcheck/gpu_driver.cc deleted file mode 100644 index de996b48..00000000 --- a/sxt/proof/sumcheck/gpu_driver.cc +++ /dev/null @@ -1,200 +0,0 @@ -/** Proofs GPU - Space and Time's cryptographic proof algorithms on the CPU and GPU. - * - * Copyright 2024-present Space and Time Labs, Inc. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#include "sxt/proof/sumcheck/gpu_driver.h" - -#include -#include - -#include "sxt/algorithm/iteration/for_each.h" -#include "sxt/base/device/memory_utility.h" -#include "sxt/base/device/stream.h" -#include "sxt/base/device/synchronization.h" -#include "sxt/base/error/panic.h" -#include "sxt/base/num/ceil_log2.h" -#include "sxt/base/num/constexpr_switch.h" -#include "sxt/execution/async/coroutine.h" -#include "sxt/execution/async/future.h" -#include "sxt/execution/device/synchronization.h" -#include "sxt/memory/management/managed_array.h" -#include "sxt/memory/resource/async_device_resource.h" -#include "sxt/memory/resource/device_resource.h" -#include "sxt/proof/sumcheck/constant.h" -#include "sxt/proof/sumcheck/sum_gpu.h" -#include "sxt/scalar25/operation/mul.h" -#include "sxt/scalar25/operation/muladd.h" -#include "sxt/scalar25/operation/sub.h" -#include "sxt/scalar25/type/element.h" -#include "sxt/scalar25/type/literal.h" - -namespace sxt::prfsk { -//-------------------------------------------------------------------------------------------------- -// gpu_workspace -//-------------------------------------------------------------------------------------------------- -namespace { -struct gpu_workspace final : public workspace { - memmg::managed_array mles; - memmg::managed_array> product_table; - memmg::managed_array product_terms; - unsigned n; - unsigned num_variables; - - gpu_workspace() noexcept - : mles{memr::get_device_resource()}, product_table{memr::get_device_resource()}, - product_terms{memr::get_device_resource()} {} -}; -} // namespace - -//-------------------------------------------------------------------------------------------------- -// make_workspace -//-------------------------------------------------------------------------------------------------- -xena::future> gpu_driver::make_workspace( - basct::cspan mles, - memmg::managed_array>&& product_table_dev, - memmg::managed_array&& product_terms_dev, unsigned n) const noexcept { - auto ws = std::make_unique(); - - // dimensions - ws->n = n; - ws->num_variables = std::max(basn::ceil_log2(n), 1); - - // mles - ws->mles = memmg::managed_array{ - mles.size(), - memr::get_device_resource(), - }; - basdv::stream mle_stream; - basdv::async_copy_host_to_device(ws->mles, mles, mle_stream); - - // product_table - ws->product_table = std::move(product_table_dev); - - // product_terms - ws->product_terms = std::move(product_terms_dev); - - // await - co_await xendv::await_stream(mle_stream); - co_return ws; -} - -xena::future> -gpu_driver::make_workspace(basct::cspan mles, - basct::cspan> product_table, - basct::cspan product_terms, unsigned n) const noexcept { - auto ws = std::make_unique(); - - // dimensions - ws->n = n; - ws->num_variables = std::max(basn::ceil_log2(n), 1); - - // mles - ws->mles = memmg::managed_array{ - mles.size(), - memr::get_device_resource(), - }; - basdv::stream mle_stream; - basdv::async_copy_host_to_device(ws->mles, mles, mle_stream); - - // product_table - ws->product_table = memmg::managed_array>{ - product_table.size(), - memr::get_device_resource(), - }; - basdv::stream product_table_stream; - basdv::async_copy_host_to_device(ws->product_table, product_table, product_table_stream); - - // product_terms - ws->product_terms = memmg::managed_array{ - product_terms.size(), - memr::get_device_resource(), - }; - basdv::stream product_terms_stream; - basdv::async_copy_host_to_device(ws->product_terms, product_terms, product_terms_stream); - - // await - co_await xendv::await_stream(mle_stream); - co_await xendv::await_stream(product_table_stream); - co_await xendv::await_stream(product_terms_stream); - co_return ws; -} - -//-------------------------------------------------------------------------------------------------- -// sum -//-------------------------------------------------------------------------------------------------- -xena::future<> gpu_driver::sum(basct::span polynomial, - workspace& ws) const noexcept { - auto& work = static_cast(ws); - auto n = work.n; - auto mid = 1u << (work.num_variables - 1u); - SXT_RELEASE_ASSERT( - // clang-format off - work.n >= mid && - polynomial.size() - 1u <= max_degree_v - // clang-format on - ); - co_await sum_gpu(polynomial, work.mles, work.product_table, work.product_terms, n); -} - -//-------------------------------------------------------------------------------------------------- -// fold -//-------------------------------------------------------------------------------------------------- -xena::future<> gpu_driver::fold(workspace& ws, const s25t::element& r) const noexcept { - using s25t::operator""_s25; - auto& work = static_cast(ws); - auto n = work.n; - auto mid = 1u << (work.num_variables - 1u); - auto num_mles = work.mles.size() / n; - SXT_RELEASE_ASSERT( - // clang-format off - work.n >= mid && work.mles.size() % n == 0 - // clang-format on - ); - - s25t::element one_m_r = 0x1_s25; - s25o::sub(one_m_r, one_m_r, r); - - memmg::managed_array mles_p{num_mles * mid, memr::get_device_resource()}; - - auto f = [ - // clang-format off - mles_p = mles_p.data(), - mles = work.mles.data(), - n = n, - num_mles = num_mles, - r = r, - one_m_r = one_m_r - // clang-format on - ] __device__ - __host__(unsigned mid, unsigned i) noexcept { - for (unsigned mle_index = 0; mle_index < num_mles; ++mle_index) { - auto val = mles[i + mle_index * n]; - s25o::mul(val, val, one_m_r); - if (mid + i < n) { - s25o::muladd(val, r, mles[mid + i + mle_index * n], val); - } - mles_p[i + mle_index * mid] = val; - } - }; - auto fut = algi::for_each(f, mid); - - // complete - co_await std::move(fut); - - work.n = mid; - --work.num_variables; - work.mles = std::move(mles_p); -} -} // namespace sxt::prfsk diff --git a/sxt/proof/sumcheck/gpu_driver.h b/sxt/proof/sumcheck/gpu_driver.h deleted file mode 100644 index 891acb58..00000000 --- a/sxt/proof/sumcheck/gpu_driver.h +++ /dev/null @@ -1,43 +0,0 @@ -/** Proofs GPU - Space and Time's cryptographic proof algorithms on the CPU and GPU. - * - * Copyright 2024-present Space and Time Labs, Inc. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#pragma once - -#include "sxt/memory/management/managed_array_fwd.h" -#include "sxt/proof/sumcheck/driver.h" - -namespace sxt::prfsk { -//-------------------------------------------------------------------------------------------------- -// gpu_driver -//-------------------------------------------------------------------------------------------------- -class gpu_driver final : public driver { -public: - xena::future> - make_workspace(basct::cspan mles, - memmg::managed_array>&& product_table_dev, - memmg::managed_array&& product_terms_dev, unsigned n) const noexcept; - - // driver - xena::future> - make_workspace(basct::cspan mles, - basct::cspan> product_table, - basct::cspan product_terms, unsigned n) const noexcept override; - - xena::future<> sum(basct::span polynomial, workspace& ws) const noexcept override; - - xena::future<> fold(workspace& ws, const s25t::element& r) const noexcept override; -}; -} // namespace sxt::prfsk diff --git a/sxt/proof/sumcheck/gpu_driver.t.cc b/sxt/proof/sumcheck/gpu_driver.t.cc deleted file mode 100644 index cf0acbba..00000000 --- a/sxt/proof/sumcheck/gpu_driver.t.cc +++ /dev/null @@ -1,28 +0,0 @@ -/** Proofs GPU - Space and Time's cryptographic proof algorithms on the CPU and GPU. - * - * Copyright 2024-present Space and Time Labs, Inc. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#include "sxt/proof/sumcheck/gpu_driver.h" - -#include "sxt/base/test/unit_test.h" -#include "sxt/proof/sumcheck/driver_test.h" - -using namespace sxt; -using namespace sxt::prfsk; - -TEST_CASE("we can perform the primitive operations for sumcheck proofs") { - gpu_driver drv; - exercise_driver(drv); -} diff --git a/sxt/proof/sumcheck/mle_utility.cc b/sxt/proof/sumcheck/mle_utility.cc deleted file mode 100644 index 8bcf33f3..00000000 --- a/sxt/proof/sumcheck/mle_utility.cc +++ /dev/null @@ -1,98 +0,0 @@ -/** Proofs GPU - Space and Time's cryptographic proof algorithms on the CPU and GPU. - * - * Copyright 2025-present Space and Time Labs, Inc. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#include "sxt/proof/sumcheck/mle_utility.h" - -#include -#include - -#include "sxt/base/container/span_utility.h" -#include "sxt/base/device/memory_utility.h" -#include "sxt/base/device/property.h" -#include "sxt/base/device/stream.h" -#include "sxt/base/error/assert.h" -#include "sxt/base/num/ceil_log2.h" -#include "sxt/base/num/divide_up.h" -#include "sxt/memory/management/managed_array.h" -#include "sxt/scalar25/type/element.h" - -namespace sxt::prfsk { -//-------------------------------------------------------------------------------------------------- -// copy_partial_mles -//-------------------------------------------------------------------------------------------------- -void copy_partial_mles(memmg::managed_array& partial_mles, basdv::stream& stream, - basct::cspan mles, unsigned n, unsigned a, - unsigned b) noexcept { - auto num_variables = std::max(basn::ceil_log2(n), 1); - auto mid = 1u << (num_variables - 1u); - auto num_mles = mles.size() / n; - auto part1_size = b - a; - SXT_DEBUG_ASSERT(a < b && b <= n); - auto ap = std::min(mid + a, n); - auto bp = std::min(mid + b, n); - auto part2_size = bp - ap; - - // resize array - auto partial_length = part1_size + part2_size; - partial_mles.resize(partial_length * num_mles); - - // copy data - for (unsigned mle_index = 0; mle_index < num_mles; ++mle_index) { - // first part - auto src = mles.subspan(n * mle_index + a, part1_size); - auto dst = basct::subspan(partial_mles, partial_length * mle_index, part1_size); - basdv::async_copy_host_to_device(dst, src, stream); - - // second part - src = mles.subspan(n * mle_index + ap, part2_size); - dst = basct::subspan(partial_mles, partial_length * mle_index + part1_size, part2_size); - if (!src.empty()) { - basdv::async_copy_host_to_device(dst, src, stream); - } - } -} - -//-------------------------------------------------------------------------------------------------- -// copy_folded_mles -//-------------------------------------------------------------------------------------------------- -void copy_folded_mles(basct::span host_mles, basdv::stream& stream, - basct::cspan device_mles, unsigned np, unsigned a, - unsigned b) noexcept { - auto num_mles = host_mles.size() / np; - auto slice_n = device_mles.size() / num_mles; - auto slice_np = b - a; - SXT_DEBUG_ASSERT( - // clang-format off - host_mles.size() == num_mles * np && - device_mles.size() == num_mles * slice_n && - b <= np - // clang-format on - ); - for (unsigned mle_index = 0; mle_index < num_mles; ++mle_index) { - auto src = device_mles.subspan(mle_index * slice_n, slice_np); - auto dst = host_mles.subspan(mle_index * np + a, slice_np); - basdv::async_copy_device_to_host(dst, src, stream); - } -} - -//-------------------------------------------------------------------------------------------------- -// get_gpu_memory_fraction -//-------------------------------------------------------------------------------------------------- -double get_gpu_memory_fraction(basct::cspan mles) noexcept { - auto total_memory = static_cast(basdv::get_total_device_memory()); - return static_cast(mles.size() * sizeof(s25t::element)) / total_memory; -} -} // namespace sxt::prfsk diff --git a/sxt/proof/sumcheck/mle_utility.h b/sxt/proof/sumcheck/mle_utility.h deleted file mode 100644 index 9f3285a8..00000000 --- a/sxt/proof/sumcheck/mle_utility.h +++ /dev/null @@ -1,48 +0,0 @@ -/** Proofs GPU - Space and Time's cryptographic proof algorithms on the CPU and GPU. - * - * Copyright 2025-present Space and Time Labs, Inc. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#pragma once - -#include "sxt/base/container/span.h" -#include "sxt/memory/management/managed_array_fwd.h" - -namespace sxt::basdv { -class stream; -} -namespace sxt::s25t { -class element; -} - -namespace sxt::prfsk { -//-------------------------------------------------------------------------------------------------- -// copy_partial_mles -//-------------------------------------------------------------------------------------------------- -void copy_partial_mles(memmg::managed_array& partial_mles, basdv::stream& stream, - basct::cspan mles, unsigned n, unsigned a, - unsigned b) noexcept; - -//-------------------------------------------------------------------------------------------------- -// copy_folded_mles -//-------------------------------------------------------------------------------------------------- -void copy_folded_mles(basct::span host_mles, basdv::stream& stream, - basct::cspan device_mles, unsigned np, unsigned a, - unsigned b) noexcept; - -//-------------------------------------------------------------------------------------------------- -// get_gpu_memory_fraction -//-------------------------------------------------------------------------------------------------- -double get_gpu_memory_fraction(basct::cspan mles) noexcept; -} // namespace sxt::prfsk diff --git a/sxt/proof/sumcheck/mle_utility.t.cc b/sxt/proof/sumcheck/mle_utility.t.cc deleted file mode 100644 index f930943d..00000000 --- a/sxt/proof/sumcheck/mle_utility.t.cc +++ /dev/null @@ -1,94 +0,0 @@ -/** Proofs GPU - Space and Time's cryptographic proof algorithms on the CPU and GPU. - * - * Copyright 2025-present Space and Time Labs, Inc. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#include "sxt/proof/sumcheck/mle_utility.h" - -#include - -#include "sxt/base/device/stream.h" -#include "sxt/base/device/synchronization.h" -#include "sxt/base/test/unit_test.h" -#include "sxt/memory/management/managed_array.h" -#include "sxt/memory/resource/managed_device_resource.h" -#include "sxt/scalar25/type/element.h" -#include "sxt/scalar25/type/literal.h" - -using namespace sxt; -using namespace sxt::prfsk; -using s25t::operator""_s25; - -TEST_CASE("we can copy a slice of mles to device memory") { - std::pmr::vector mles{memr::get_managed_device_resource()}; - memmg::managed_array partial_mles{memr::get_managed_device_resource()}; - - basdv::stream stream; - - SECTION("we can copy an mle with a single element") { - mles = {0x123_s25}; - copy_partial_mles(partial_mles, stream, mles, 1, 0, 1); - basdv::synchronize_stream(stream); - memmg::managed_array expected = {0x123_s25}; - REQUIRE(partial_mles == expected); - } - - SECTION("we can copy a slice of MLEs") { - mles = {0x1_s25, 0x2_s25, 0x3_s25, 0x4_s25, 0x5_s25, 0x6_s25}; - copy_partial_mles(partial_mles, stream, mles, 3, 0, 1); - basdv::synchronize_stream(stream); - memmg::managed_array expected = {0x1_s25, 0x3_s25, 0x4_s25, 0x6_s25}; - REQUIRE(partial_mles == expected); - } -} - -TEST_CASE("we can copy partially folded MLEs to the host") { - std::pmr::vector device_mles{memr::get_managed_device_resource()}; - std::vector host_mles; - - basdv::stream stream; - - SECTION("we can copy a single element") { - device_mles = {0x123_s25}; - host_mles.resize(1); - copy_folded_mles(host_mles, stream, device_mles, 1, 0, 1); - basdv::synchronize_stream(stream); - std::vector expected = {0x123_s25}; - REQUIRE(host_mles == expected); - } - - SECTION("we can copy partially folded MLEs") { - device_mles = {0x123_s25, 0x456_s25}; - host_mles.resize(4); - copy_folded_mles(host_mles, stream, device_mles, 2, 0, 1); - basdv::synchronize_stream(stream); - std::vector expected = {0x123_s25, 0x0_s25, 0x456_s25, 0x0_s25}; - REQUIRE(host_mles == expected); - } -} - -TEST_CASE("we can query the fraction of device memory taken by MLEs") { - std::vector mles; - - SECTION("we handle the zero case") { REQUIRE(get_gpu_memory_fraction(mles) == 0.0); } - - SECTION("the fractions doubles if the length of mles doubles") { - mles.resize(1); - auto f1 = get_gpu_memory_fraction(mles); - REQUIRE(f1 > 0); - mles.resize(2); - auto f2 = get_gpu_memory_fraction(mles); - REQUIRE(f2 == Catch::Approx(2 * f1)); - } -} diff --git a/sxt/proof/sumcheck/reference_transcript.cc b/sxt/proof/sumcheck/reference_transcript.cc deleted file mode 100644 index fcbeb6b5..00000000 --- a/sxt/proof/sumcheck/reference_transcript.cc +++ /dev/null @@ -1,46 +0,0 @@ -/** Proofs GPU - Space and Time's cryptographic proof algorithms on the CPU and GPU. - * - * Copyright 2025-present Space and Time Labs, Inc. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#include "sxt/proof/sumcheck/reference_transcript.h" - -#include "sxt/proof/transcript/transcript_utility.h" -#include "sxt/scalar25/type/element.h" - -namespace sxt::prfsk { -//-------------------------------------------------------------------------------------------------- -// constructor -//-------------------------------------------------------------------------------------------------- -reference_transcript::reference_transcript(prft::transcript& transcript) noexcept - : transcript_{transcript} {} - -//-------------------------------------------------------------------------------------------------- -// init -//-------------------------------------------------------------------------------------------------- -void reference_transcript::init(size_t num_variables, size_t round_degree) noexcept { - prft::set_domain(transcript_, "sumcheck proof v1"); - prft::append_value(transcript_, "n", num_variables); - prft::append_value(transcript_, "k", round_degree); -} - -//-------------------------------------------------------------------------------------------------- -// round_challenge -//-------------------------------------------------------------------------------------------------- -void reference_transcript::round_challenge(s25t::element& r, - basct::cspan polynomial) noexcept { - prft::append_values(transcript_, "P", polynomial); - prft::challenge_value(r, transcript_, "R"); -} -} // namespace sxt::prfsk diff --git a/sxt/proof/sumcheck/reference_transcript.h b/sxt/proof/sumcheck/reference_transcript.h deleted file mode 100644 index 3ee41ed0..00000000 --- a/sxt/proof/sumcheck/reference_transcript.h +++ /dev/null @@ -1,37 +0,0 @@ -/** Proofs GPU - Space and Time's cryptographic proof algorithms on the CPU and GPU. - * - * Copyright 2025-present Space and Time Labs, Inc. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#pragma once - -#include "sxt/proof/sumcheck/sumcheck_transcript.h" -#include "sxt/proof/transcript/transcript.h" - -namespace sxt::prfsk { -//-------------------------------------------------------------------------------------------------- -// reference_transcript -//-------------------------------------------------------------------------------------------------- -class reference_transcript final : public sumcheck_transcript { -public: - explicit reference_transcript(prft::transcript& transcript) noexcept; - - void init(size_t num_variables, size_t round_degree) noexcept override; - - void round_challenge(s25t::element& r, basct::cspan polynomial) noexcept override; - -private: - prft::transcript& transcript_; -}; -} // namespace sxt::prfsk diff --git a/sxt/proof/sumcheck/reference_transcript.t.cc b/sxt/proof/sumcheck/reference_transcript.t.cc deleted file mode 100644 index 890312c5..00000000 --- a/sxt/proof/sumcheck/reference_transcript.t.cc +++ /dev/null @@ -1,58 +0,0 @@ -/** Proofs GPU - Space and Time's cryptographic proof algorithms on the CPU and GPU. - * - * Copyright 2025-present Space and Time Labs, Inc. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#include "sxt/proof/sumcheck/reference_transcript.h" - -#include - -#include "sxt/base/test/unit_test.h" -#include "sxt/proof/transcript/transcript.h" -#include "sxt/scalar25/type/literal.h" - -using namespace sxt; -using namespace sxt::prfsk; -using s25t::operator""_s25; - -TEST_CASE("we provide an implementation of sumcheck transcript") { - prft::transcript base_transcript{"abc"}; - reference_transcript transcript{base_transcript}; - std::vector p = {0x123_s25}; - s25t::element r, rp; - - SECTION("we don't draw the same challenge from a transcript") { - transcript.round_challenge(r, p); - transcript.round_challenge(rp, p); - REQUIRE(r != rp); - - prft::transcript base_transcript_p{"abc"}; - reference_transcript transcript_p{base_transcript_p}; - p[0] = 0x456_s25; - transcript_p.round_challenge(rp, p); - REQUIRE(r != rp); - } - - SECTION("init_transcript produces different results based on parameters") { - transcript.init(1, 2); - transcript.round_challenge(r, p); - - prft::transcript base_transcript_p{"abc"}; - reference_transcript transcript_p{base_transcript_p}; - transcript.init(2, 1); - transcript.round_challenge(rp, p); - - REQUIRE(r != rp); - } -} diff --git a/sxt/proof/sumcheck/sum_gpu.cc b/sxt/proof/sumcheck/sum_gpu.cc deleted file mode 100644 index 8066fd24..00000000 --- a/sxt/proof/sumcheck/sum_gpu.cc +++ /dev/null @@ -1,231 +0,0 @@ -/** Proofs GPU - Space and Time's cryptographic proof algorithms on the CPU and GPU. - * - * Copyright 2025-present Space and Time Labs, Inc. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#include "sxt/proof/sumcheck/sum_gpu.h" - -#include - -#include "sxt/algorithm/reduction/kernel_fit.h" -#include "sxt/algorithm/reduction/thread_reduction.h" -#include "sxt/base/device/memory_utility.h" -#include "sxt/base/device/state.h" -#include "sxt/base/device/stream.h" -#include "sxt/base/iterator/split.h" -#include "sxt/base/num/ceil_log2.h" -#include "sxt/base/num/constexpr_switch.h" -#include "sxt/execution/async/coroutine.h" -#include "sxt/execution/async/future.h" -#include "sxt/execution/device/for_each.h" -#include "sxt/execution/kernel/kernel_dims.h" -#include "sxt/execution/kernel/launch.h" -#include "sxt/memory/management/managed_array.h" -#include "sxt/memory/resource/async_device_resource.h" -#include "sxt/memory/resource/device_resource.h" -#include "sxt/proof/sumcheck/constant.h" -#include "sxt/proof/sumcheck/device_cache.h" -#include "sxt/proof/sumcheck/mle_utility.h" -#include "sxt/proof/sumcheck/polynomial_mapper.h" -#include "sxt/proof/sumcheck/reduction_gpu.h" -#include "sxt/scalar25/operation/add.h" -#include "sxt/scalar25/operation/mul.h" -#include "sxt/scalar25/type/element.h" - -namespace sxt::prfsk { -//-------------------------------------------------------------------------------------------------- -// polynomial_reducer -//-------------------------------------------------------------------------------------------------- -namespace { -template struct polynomial_reducer { - using value_type = std::array; - - CUDA_CALLABLE static void accumulate_inplace(value_type& res, const value_type& e) noexcept { - for (unsigned i = 0; i < res.size(); ++i) { - s25o::add(res[i], res[i], e[i]); - } - } -}; -} // namespace - -//-------------------------------------------------------------------------------------------------- -// partial_sum_kernel_impl -//-------------------------------------------------------------------------------------------------- -template -__device__ static void partial_sum_kernel_impl(s25t::element* __restrict__ shared_data, - const s25t::element* __restrict__ mles, - const unsigned* __restrict__ product_terms, - unsigned split, unsigned n) noexcept { - using Mapper = polynomial_mapper; - using Reducer = polynomial_reducer; - using T = Mapper::value_type; - Mapper mapper{ - .mles = mles, - .product_terms = product_terms, - .split = split, - .n = n, - }; - auto index = blockIdx.x * (BlockSize * 2) + threadIdx.x; - auto step = BlockSize * 2 * gridDim.x; - algr::thread_reduce(reinterpret_cast(shared_data), mapper, split, step, - threadIdx.x, index); -} - -//-------------------------------------------------------------------------------------------------- -// partial_sum_kernel -//-------------------------------------------------------------------------------------------------- -template -__global__ static void -partial_sum_kernel(s25t::element* __restrict__ out, const s25t::element* __restrict__ mles, - const std::pair* __restrict__ product_table, - const unsigned* __restrict__ product_terms, unsigned num_coefficients, - unsigned split, unsigned n) noexcept { - auto product_index = blockIdx.y; - auto num_terms = product_table[product_index].second; - auto thread_index = threadIdx.x; - auto output_step = gridDim.x * gridDim.y; - - // shared data for reduction - __shared__ s25t::element shared_data[2 * BlockSize * (max_degree_v + 1u)]; - - // adjust pointers - out += blockIdx.x; - out += gridDim.x * product_index; - for (unsigned i = 0; i < product_index; ++i) { - product_terms += product_table[i].second; - } - - // sum - basn::constexpr_switch<1, max_degree_v + 1u>( - num_terms, [&](std::integral_constant) noexcept { - partial_sum_kernel_impl(shared_data, mles, product_terms, split, n); - }); - - // write out result - auto mult = product_table[product_index].first; - for (unsigned i = thread_index; i < num_coefficients; i += BlockSize) { - auto output_index = output_step * i; - if (i < num_terms + 1u) { - s25o::mul(out[output_index], mult, shared_data[i]); - } else { - out[output_index] = s25t::element{}; - } - } -} - -//-------------------------------------------------------------------------------------------------- -// partial_sum -//-------------------------------------------------------------------------------------------------- -static xena::future<> partial_sum(basct::span p, basdv::stream& stream, - basct::cspan mles, - basct::cspan> product_table, - basct::cspan product_terms, unsigned split, - unsigned n) noexcept { - auto num_coefficients = p.size(); - auto num_products = product_table.size(); - auto dims = algr::fit_reduction_kernel(split); - memr::async_device_resource resource{stream}; - - // partials - memmg::managed_array partials{num_coefficients * dims.num_blocks * num_products, - &resource}; - xenk::launch_kernel(dims.block_size, [&]( - std::integral_constant) noexcept { - partial_sum_kernel<<>>( - partials.data(), mles.data(), product_table.data(), product_terms.data(), num_coefficients, - split, n); - }); - - // reduce partials - co_await reduce_sums(p, stream, partials); -} - -//-------------------------------------------------------------------------------------------------- -// sum_gpu -//-------------------------------------------------------------------------------------------------- -xena::future<> sum_gpu(basct::span p, device_cache& cache, - const basit::split_options& options, basct::cspan mles, - unsigned n) noexcept { - auto num_variables = std::max(basn::ceil_log2(n), 1); - auto mid = 1u << (num_variables - 1u); - auto num_mles = mles.size() / n; - auto num_coefficients = p.size(); - - // split - auto [chunk_first, chunk_last] = basit::split(basit::index_range{0, mid}, options); - - // sum - size_t counter = 0; - co_await xendv::concurrent_for_each( - chunk_first, chunk_last, [&](basit::index_range rng) noexcept -> xena::future<> { - basdv::stream stream; - memr::async_device_resource resource{stream}; - - // copy partial mles to device - memmg::managed_array partial_mles{&resource}; - copy_partial_mles(partial_mles, stream, mles, n, rng.a(), rng.b()); - auto split = rng.b() - rng.a(); - auto np = partial_mles.size() / num_mles; - - // lookup problem descriptor - basct::cspan> product_table; - basct::cspan product_terms; - cache.lookup(product_table, product_terms, stream); - - // compute - memmg::managed_array partial_p(num_coefficients); - co_await partial_sum(partial_p, stream, partial_mles, product_table, product_terms, split, - np); - - // fill in the result - if (counter == 0) { - for (unsigned i = 0; i < num_coefficients; ++i) { - p[i] = partial_p[i]; - } - } else { - for (unsigned i = 0; i < num_coefficients; ++i) { - s25o::add(p[i], p[i], partial_p[i]); - } - } - ++counter; - }); -} - -xena::future<> sum_gpu(basct::span p, device_cache& cache, - basct::cspan mles, unsigned n) noexcept { - basit::split_options options{ - .min_chunk_size = 100'000u, - .max_chunk_size = 200'000u, - .split_factor = basdv::get_num_devices(), - }; - co_await sum_gpu(p, cache, options, mles, n); -} - -xena::future<> sum_gpu(basct::span p, basct::cspan mles, - basct::cspan> product_table, - basct::cspan product_terms, unsigned n) noexcept { - auto num_variables = std::max(basn::ceil_log2(n), 1); - auto mid = 1u << (num_variables - 1u); - SXT_DEBUG_ASSERT( - // clang-format off - basdv::is_host_pointer(p.data()) && - basdv::is_active_device_pointer(mles.data()) && - basdv::is_active_device_pointer(product_table.data()) && - basdv::is_active_device_pointer(product_terms.data()) - // clang-format on - ); - basdv::stream stream; - co_await partial_sum(p, stream, mles, product_table, product_terms, mid, n); -} -} // namespace sxt::prfsk diff --git a/sxt/proof/sumcheck/sum_gpu.h b/sxt/proof/sumcheck/sum_gpu.h deleted file mode 100644 index a693c7da..00000000 --- a/sxt/proof/sumcheck/sum_gpu.h +++ /dev/null @@ -1,58 +0,0 @@ -/** Proofs GPU - Space and Time's cryptographic proof algorithms on the CPU and GPU. - * - * Copyright 2025-present Space and Time Labs, Inc. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#pragma once - -#include - -#include "sxt/base/container/span.h" -#include "sxt/base/device/property.h" -#include "sxt/execution/async/future_fwd.h" - -namespace sxt::s25t { -class element; -} - -namespace sxt::basit { -struct split_options; -} - -namespace sxt::prfsk { -class device_cache; - -//-------------------------------------------------------------------------------------------------- -// sum_options -//-------------------------------------------------------------------------------------------------- -struct sum_options { - unsigned min_chunk_size = 100'000u; - unsigned max_chunk_size = 250'000u; - unsigned split_factor = unsigned(basdv::get_num_devices()); -}; - -//-------------------------------------------------------------------------------------------------- -// sum_gpu -//-------------------------------------------------------------------------------------------------- -xena::future<> sum_gpu(basct::span p, device_cache& cache, - const basit::split_options& options, basct::cspan mles, - unsigned n) noexcept; - -xena::future<> sum_gpu(basct::span p, device_cache& cache, - basct::cspan mles, unsigned n) noexcept; - -xena::future<> sum_gpu(basct::span p, basct::cspan mles, - basct::cspan> product_table, - basct::cspan product_terms, unsigned n) noexcept; -} // namespace sxt::prfsk diff --git a/sxt/proof/sumcheck/sum_gpu.t.cc b/sxt/proof/sumcheck/sum_gpu.t.cc deleted file mode 100644 index 35807fd5..00000000 --- a/sxt/proof/sumcheck/sum_gpu.t.cc +++ /dev/null @@ -1,138 +0,0 @@ -/** Proofs GPU - Space and Time's cryptographic proof algorithms on the CPU and GPU. - * - * Copyright 2025-present Space and Time Labs, Inc. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#include "sxt/proof/sumcheck/sum_gpu.h" - -#include - -#include "sxt/base/iterator/split.h" -#include "sxt/base/test/unit_test.h" -#include "sxt/execution/async/future.h" -#include "sxt/execution/schedule/scheduler.h" -#include "sxt/proof/sumcheck/device_cache.h" -#include "sxt/scalar25/operation/overload.h" -#include "sxt/scalar25/type/element.h" -#include "sxt/scalar25/type/literal.h" - -using namespace sxt; -using namespace sxt::prfsk; -using s25t::operator""_s25; - -TEST_CASE("we can sum MLEs") { - std::vector> product_table; - std::vector product_terms; - std::vector mles; - std::vector p(2); - - SECTION("we can sum an MLE with a single term and n=1") { - product_table = {{0x1_s25, 1}}; - product_terms = {0}; - device_cache cache{product_table, product_terms}; - mles = {0x123_s25}; - auto fut = sum_gpu(p, cache, mles, 1); - xens::get_scheduler().run(); - REQUIRE(fut.ready()); - REQUIRE(p[0] == mles[0]); - REQUIRE(p[1] == -mles[0]); - } - - SECTION("we can sum an MLE with a single term, n=1, and a non-unity multiplier") { - product_table = {{0x2_s25, 1}}; - product_terms = {0}; - device_cache cache{product_table, product_terms}; - mles = {0x123_s25}; - auto fut = sum_gpu(p, cache, mles, 1); - xens::get_scheduler().run(); - REQUIRE(fut.ready()); - REQUIRE(p[0] == product_table[0].first * mles[0]); - REQUIRE(p[1] == -product_table[0].first * mles[0]); - } - - SECTION("we can sum an MLE with a single term and n=2") { - product_table = {{0x1_s25, 1}}; - product_terms = {0}; - device_cache cache{product_table, product_terms}; - mles = {0x123_s25, 0x456_s25}; - auto fut = sum_gpu(p, cache, mles, 2); - xens::get_scheduler().run(); - REQUIRE(fut.ready()); - REQUIRE(p[0] == mles[0]); - REQUIRE(p[1] == mles[1] - mles[0]); - } - - SECTION("we can sum an MLE with multiple terms and n=1") { - p.resize(3); - product_table = {{0x1_s25, 2}}; - product_terms = {0, 1}; - device_cache cache{product_table, product_terms}; - mles = {0x123_s25, 0x456_s25}; - auto fut = sum_gpu(p, cache, mles, 1); - xens::get_scheduler().run(); - REQUIRE(fut.ready()); - REQUIRE(p[0] == mles[0] * mles[1]); - REQUIRE(p[1] == -mles[0] * mles[1] - mles[1] * mles[0]); - REQUIRE(p[2] == mles[0] * mles[1]); - } - - SECTION("we can sum multiple mles") { - product_table = { - {0x1_s25, 1}, - {0x1_s25, 1}, - }; - product_terms = {0, 1}; - device_cache cache{product_table, product_terms}; - mles = {0x123_s25, 0x456_s25}; - auto fut = sum_gpu(p, cache, mles, 1); - xens::get_scheduler().run(); - REQUIRE(fut.ready()); - REQUIRE(p[0] == mles[0] + mles[1]); - REQUIRE(p[1] == -mles[0] - mles[1]); - } - - SECTION("we can chunk sums with n=4") { - product_table = {{0x1_s25, 1}}; - product_terms = {0}; - device_cache cache{product_table, product_terms}; - mles = {0x123_s25, 0x456_s25, 0x789_s25, 0x91011_s25}; - basit::split_options options{ - .min_chunk_size = 1, - .max_chunk_size = 1, - .split_factor = 2, - }; - auto fut = sum_gpu(p, cache, options, mles, 4); - xens::get_scheduler().run(); - REQUIRE(fut.ready()); - REQUIRE(p[0] == mles[0] + mles[1]); - REQUIRE(p[1] == (mles[2] - mles[0]) + (mles[3] - mles[1])); - } - - SECTION("we can chunk sums with n=4") { - product_table = {{0x1_s25, 1}}; - product_terms = {0}; - device_cache cache{product_table, product_terms}; - mles = {0x2_s25, 0x4_s25, 0x7_s25, 0x9_s25}; - basit::split_options options{ - .min_chunk_size = 16, - .max_chunk_size = 16, - .split_factor = 2, - }; - auto fut = sum_gpu(p, cache, options, mles, 4); - xens::get_scheduler().run(); - REQUIRE(fut.ready()); - REQUIRE(p[0] == mles[0] + mles[1]); - REQUIRE(p[1] == (mles[2] - mles[0]) + (mles[3] - mles[1])); - } -} diff --git a/sxt/proof/sumcheck/sumcheck_random.cc b/sxt/proof/sumcheck/sumcheck_random.cc deleted file mode 100644 index d3f164e3..00000000 --- a/sxt/proof/sumcheck/sumcheck_random.cc +++ /dev/null @@ -1,77 +0,0 @@ -/** Proofs GPU - Space and Time's cryptographic proof algorithms on the CPU and GPU. - * - * Copyright 2025-present Space and Time Labs, Inc. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#include "sxt/proof/sumcheck/sumcheck_random.h" - -#include - -#include "sxt/base/error/assert.h" -#include "sxt/base/num/fast_random_number_generator.h" -#include "sxt/scalar25/random/element.h" -#include "sxt/scalar25/type/element.h" - -namespace sxt::prfsk { -//-------------------------------------------------------------------------------------------------- -// generate_random_sumcheck_problem -//-------------------------------------------------------------------------------------------------- -void generate_random_sumcheck_problem( - std::vector& mles, - std::vector>& product_table, - std::vector& product_terms, unsigned& n, basn::fast_random_number_generator& rng, - const random_sumcheck_descriptor& descriptor) noexcept { - std::mt19937 rng_p{rng()}; - - // n - SXT_RELEASE_ASSERT(descriptor.min_length <= descriptor.max_length); - std::uniform_int_distribution n_dist{descriptor.min_length, descriptor.max_length}; - n = n_dist(rng_p); - - // num_mles - SXT_RELEASE_ASSERT(descriptor.min_num_mles <= descriptor.max_num_mles); - std::uniform_int_distribution num_mles_dist{descriptor.min_num_mles, - descriptor.max_num_mles}; - auto num_mles = num_mles_dist(rng_p); - - // num_products - SXT_RELEASE_ASSERT(descriptor.min_num_products <= descriptor.max_num_products); - std::uniform_int_distribution num_products_dist{descriptor.min_num_products, - descriptor.max_num_products}; - auto num_products = num_products_dist(rng_p); - - // mles - mles.resize(n * num_mles); - s25rn::generate_random_elements(mles, rng); - - // product_table - unsigned num_terms = 0; - product_table.resize(num_products); - SXT_RELEASE_ASSERT(descriptor.min_product_length <= descriptor.max_product_length); - std::uniform_int_distribution product_length_dist{descriptor.min_product_length, - descriptor.max_product_length}; - for (auto& [s, len] : product_table) { - s25rn::generate_random_element(s, rng); - len = product_length_dist(rng_p); - num_terms += len; - } - - // product_terms - product_terms.resize(num_terms); - std::uniform_int_distribution mle_dist{0, num_mles - 1}; - for (auto& term : product_terms) { - term = mle_dist(rng_p); - } -} -} // namespace sxt::prfsk diff --git a/sxt/proof/sumcheck/sumcheck_random.h b/sxt/proof/sumcheck/sumcheck_random.h deleted file mode 100644 index f58cd0f8..00000000 --- a/sxt/proof/sumcheck/sumcheck_random.h +++ /dev/null @@ -1,57 +0,0 @@ -/** Proofs GPU - Space and Time's cryptographic proof algorithms on the CPU and GPU. - * - * Copyright 2025-present Space and Time Labs, Inc. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#pragma once - -#include -#include - -#include "sxt/proof/sumcheck/constant.h" - -namespace sxt::s25t { -class element; -} -namespace sxt::basn { -class fast_random_number_generator; -} - -namespace sxt::prfsk { -//-------------------------------------------------------------------------------------------------- -// random_sumcheck_descriptor -//-------------------------------------------------------------------------------------------------- -struct random_sumcheck_descriptor { - unsigned min_length = 1; - unsigned max_length = 10; - - unsigned min_num_products = 1; - unsigned max_num_products = 5; - - unsigned min_product_length = 2; - unsigned max_product_length = max_degree_v; - - unsigned min_num_mles = 1; - unsigned max_num_mles = 5; -}; - -//-------------------------------------------------------------------------------------------------- -// generate_random_sumcheck_problem -//-------------------------------------------------------------------------------------------------- -void generate_random_sumcheck_problem( - std::vector& mles, - std::vector>& product_table, - std::vector& product_terms, unsigned& n, basn::fast_random_number_generator& rng, - const random_sumcheck_descriptor& descriptor) noexcept; -} // namespace sxt::prfsk diff --git a/sxt/proof/sumcheck/sumcheck_transcript.cc b/sxt/proof/sumcheck/sumcheck_transcript.cc deleted file mode 100644 index 36b1ecd8..00000000 --- a/sxt/proof/sumcheck/sumcheck_transcript.cc +++ /dev/null @@ -1,17 +0,0 @@ -/** Proofs GPU - Space and Time's cryptographic proof algorithms on the CPU and GPU. - * - * Copyright 2025-present Space and Time Labs, Inc. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#include "sxt/proof/sumcheck/sumcheck_transcript.h" diff --git a/sxt/proof/sumcheck/sumcheck_transcript.h b/sxt/proof/sumcheck/sumcheck_transcript.h deleted file mode 100644 index 2d690ded..00000000 --- a/sxt/proof/sumcheck/sumcheck_transcript.h +++ /dev/null @@ -1,40 +0,0 @@ -/** Proofs GPU - Space and Time's cryptographic proof algorithms on the CPU and GPU. - * - * Copyright 2025-present Space and Time Labs, Inc. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#pragma once - -#include - -#include "sxt/base/container/span.h" - -namespace sxt::s25t { -class element; -} - -namespace sxt::prfsk { -//-------------------------------------------------------------------------------------------------- -// sumcheck_transcript -//-------------------------------------------------------------------------------------------------- -class sumcheck_transcript { -public: - virtual ~sumcheck_transcript() noexcept = default; - - virtual void init(size_t num_variables, size_t round_degree) noexcept = 0; - - virtual void round_challenge(s25t::element& r, - basct::cspan polynomial) noexcept = 0; -}; -} // namespace sxt::prfsk diff --git a/sxt/proof/sumcheck/workspace.cc b/sxt/proof/sumcheck/workspace.cc deleted file mode 100644 index d356b4af..00000000 --- a/sxt/proof/sumcheck/workspace.cc +++ /dev/null @@ -1,17 +0,0 @@ -/** Proofs GPU - Space and Time's cryptographic proof algorithms on the CPU and GPU. - * - * Copyright 2024-present Space and Time Labs, Inc. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#include "sxt/proof/sumcheck/workspace.h" diff --git a/sxt/proof/sumcheck/workspace.h b/sxt/proof/sumcheck/workspace.h deleted file mode 100644 index edacbd2b..00000000 --- a/sxt/proof/sumcheck/workspace.h +++ /dev/null @@ -1,27 +0,0 @@ -/** Proofs GPU - Space and Time's cryptographic proof algorithms on the CPU and GPU. - * - * Copyright 2024-present Space and Time Labs, Inc. - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ -#pragma once - -namespace sxt::prfsk { -//-------------------------------------------------------------------------------------------------- -// workspace -//-------------------------------------------------------------------------------------------------- -class workspace { -public: - virtual ~workspace() noexcept = default; -}; -} // namespace sxt::prfsk diff --git a/sxt/proof/sumcheck2/BUILD b/sxt/proof/sumcheck2/BUILD index 41937f96..47dfd00f 100644 --- a/sxt/proof/sumcheck2/BUILD +++ b/sxt/proof/sumcheck2/BUILD @@ -192,6 +192,9 @@ sxt_cc_component( name = "polynomial_mapper", test_deps = [ "//sxt/base/test:unit_test", + "//sxt/scalar25/operation:overload", + "//sxt/scalar25/realization:field", + "//sxt/scalar25/type:literal", ], deps = [ ":polynomial_utility", diff --git a/sxt/proof/sumcheck2/polynomial_mapper.t.cc b/sxt/proof/sumcheck2/polynomial_mapper.t.cc index 57b83724..11d26d27 100644 --- a/sxt/proof/sumcheck2/polynomial_mapper.t.cc +++ b/sxt/proof/sumcheck2/polynomial_mapper.t.cc @@ -16,6 +16,48 @@ */ #include "sxt/proof/sumcheck2/polynomial_mapper.h" +#include + #include "sxt/base/test/unit_test.h" +#include "sxt/scalar25/operation/overload.h" +#include "sxt/scalar25/realization/field.h" +#include "sxt/scalar25/type/literal.h" + +using namespace sxt; +using namespace sxt::prfsk2; +using s25t::operator""_s25; + +using T = s25t::element; + +TEST_CASE("we can map indexes to expanded polynomials") { + std::vector mles; + std::vector product_terms; + + SECTION("we can map a single element mle") { + mles = {0x123_s25}; + product_terms = {0}; + polynomial_mapper<1, T> m{ + .mles = mles.data(), + .product_terms = product_terms.data(), + .split = 1, + .n = 1, + }; + auto p = m.map_index(0); + REQUIRE(p[0] == mles[0]); + REQUIRE(p[1] == -mles[0]); + } -TEST_CASE("todo") {} + SECTION("we can map an mle with two elements") { + mles = {0x123_s25, 0x456_s25}; + product_terms = {0}; + polynomial_mapper<1, T> m{ + .mles = mles.data(), + .product_terms = product_terms.data(), + .split = 1, + .n = 2, + }; + auto p = m.map_index(0); + REQUIRE(p[0] == mles[0]); + REQUIRE(p[1] == mles[1] - mles[0]); + } +}