Skip to content

Commit

Permalink
Blazingly fast, but incorrect computation.
Browse files Browse the repository at this point in the history
  • Loading branch information
root committed Dec 15, 2024
1 parent fb12094 commit ee9ddd4
Show file tree
Hide file tree
Showing 7 changed files with 452 additions and 93 deletions.
2 changes: 1 addition & 1 deletion parallel-crypto3/CMakeLists.txt
View file
Original file line number Diff line number Diff line change
Expand Up @@ -12,7 +12,7 @@ find_package(CM REQUIRED)
include(CMConfig)

set(CMAKE_EXPORT_COMPILE_COMMANDS ON)
#set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -g -O0")
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -g -O0")

# add_compile_options(-v)
# #add_compile_options(--hipsycl-platform=cuda --hipsycl-gpu-arch=sm_75 --hipsycl-targets=cuda:sm_75 --hipsycl-cuda-path=/var/empty/local/cuda/ --hipsycl-nvcxx=/var/empty/nvidia/hpc_sdk/Linux_x86_64/24.9/compilers/bin/nvc++)
Expand Down
21 changes: 21 additions & 0 deletions parallel-crypto3/libs/parallel-math/include/nil/crypto3/math/detail/global_queue.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -83,3 +83,24 @@ struct writeback_scope {
tracked_buffer->set_write_back(!no_write_back && write_back_val);
}
};

template<typename T>
struct uninitialized_allocator : std::allocator<T> {
using value_type = T;

template<typename U, typename... Args>
void construct(U* ptr, Args&&... args) noexcept {
if constexpr (sizeof...(Args) > 0) {
// construct if arguments are passed
::new (static_cast<void*>(ptr)) U(std::forward<Args>(args)...);
} else {
// no-op if no arguments
}
}

// optional: ensure compatibility with std::allocator_traits
template<typename U>
struct rebind {
using other = uninitialized_allocator<U>;
};
};
320 changes: 320 additions & 0 deletions parallel-crypto3/libs/parallel-math/include/nil/crypto3/math/detail/lazy.hpp
View file
Original file line number Diff line number Diff line change
@@ -0,0 +1,320 @@
//---------------------------------------------------------------------------//
// Copyright (c) 2024 Dmitrii Tabalin <[email protected]>
//
// MIT License
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to deal
// in the Software without restriction, including without limitation the rights
// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
// copies of the Software, and to permit persons to whom the Software is
// furnished to do so, subject to the following conditions:
//
// The above copyright notice and this permission notice shall be included in all
// copies or substantial portions of the Software.
//
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
// SOFTWARE.
//---------------------------------------------------------------------------//

#pragma once

#include <vector>
#include <optional>

namespace nil {
namespace crypto3 {
namespace math {

template<typename T>
class lazy_element {
std::optional<T> value;
public:
typedef typename T::field_type field_type;

lazy_element() : value(std::nullopt) {}
lazy_element(const T& value) : value(value) {}
lazy_element(T&& value) : value(std::move(value)) {}
lazy_element(const lazy_element& other) : value(other.value) {}

T& get() {
if (!value) {
value = T::zero();
}
return *value;
}

operator T() const {
return get();
}

const T& get() const {
if (!value) {
return T::zero();
}
return *value;
}

// Arithmetic operations
lazy_element& operator+=(const T& other) {
get() += other;
return *this;
}

lazy_element& operator+=(const lazy_element& other) {
get() += other.get();
return *this;
}

lazy_element operator+(const T& other) const {
lazy_element result(*this);
result += other;
return result;
}

lazy_element operator+(const lazy_element& other) const {
lazy_element result(*this);
result += other;
return result;
}

lazy_element operator-(const T& other) const {
lazy_element result(*this);
result.get() -= other;
return result;
}

lazy_element operator-(const lazy_element& other) const {
lazy_element result(*this);
result.get() -= other.get();
return result;
}

lazy_element& operator-=(const T& other) {
get() -= other;
return *this;
}

lazy_element& operator-=(const lazy_element& other) {
get() -= other.get();
return *this;
}

lazy_element operator*(const T& other) const {
lazy_element result;
result.get() = get() * other;
return result;
}

lazy_element operator*(const lazy_element& other) const {
lazy_element result;
result.get() = get() * other.get();
return result;
}

lazy_element& operator*=(const T& other) {
get() *= other;
return *this;
}

lazy_element& operator*=(const lazy_element& other) {
get() *= other.get();
return *this;
}
};

template <typename T, typename Allocator = void>
class lazy_vector {
typedef std::vector<lazy_element<T>> container_type;
container_type val;
public:
typedef typename container_type::value_type value_type;
typedef typename container_type::allocator_type allocator_type;
typedef typename container_type::reference reference;
typedef typename container_type::const_reference const_reference;
typedef typename container_type::size_type size_type;
typedef typename container_type::difference_type difference_type;
typedef typename container_type::pointer pointer;
typedef typename container_type::const_pointer const_pointer;
typedef typename container_type::iterator iterator;
typedef typename container_type::const_iterator const_iterator;
typedef typename container_type::reverse_iterator reverse_iterator;
typedef typename container_type::const_reverse_iterator const_reverse_iterator;

lazy_vector(size_t size) : val(size) {}
lazy_vector() : val() {}

size_type size() const {
return val.size();
}

container_type& get_storage() {
return val;
}

iterator begin() BOOST_NOEXCEPT {
return val.begin();
}

const_iterator begin() const BOOST_NOEXCEPT {
return val.begin();
}
iterator end() BOOST_NOEXCEPT {
return val.end();
}
const_iterator end() const BOOST_NOEXCEPT {
return val.end();
}

reverse_iterator rbegin() BOOST_NOEXCEPT {
return val.rbegin();
}

const_reverse_iterator rbegin() const BOOST_NOEXCEPT {
return val.rbegin();
}

reverse_iterator rend() BOOST_NOEXCEPT {
return reverse_iterator(begin());
}

const_reverse_iterator rend() const BOOST_NOEXCEPT {
return const_reverse_iterator(begin());
}

const_iterator cbegin() const BOOST_NOEXCEPT {
return begin();
}

const_iterator cend() const BOOST_NOEXCEPT {
return end();
}

const_reverse_iterator crbegin() const BOOST_NOEXCEPT {
return rbegin();
}

const_reverse_iterator crend() const BOOST_NOEXCEPT {
return rend();
}

size_type capacity() const BOOST_NOEXCEPT {
return val.capacity();
}
bool empty() const BOOST_NOEXCEPT {
return val.empty();
}
size_type max_size() const BOOST_NOEXCEPT {
return val.max_size();
}
void reserve(size_type _n) {
return val.reserve(_n);
}
void shrink_to_fit() BOOST_NOEXCEPT {
return val.shrink_to_fit();
}

reference operator[](size_type _n) BOOST_NOEXCEPT {
return val[_n];
}

const_reference operator[](size_type _n) const BOOST_NOEXCEPT {
return val[_n];
}

reference at(size_type _n) {
return val.at(_n);
}

const_reference at(size_type _n) const {
return val.at(_n);
}

reference front() BOOST_NOEXCEPT {
return val.front();
}
const_reference front() const BOOST_NOEXCEPT {
return val.front();
}
reference back() BOOST_NOEXCEPT {
return val.back();
}
const_reference back() const BOOST_NOEXCEPT {
return val.back();
}

value_type* data() BOOST_NOEXCEPT {
return val.data();
}

const value_type* data() const BOOST_NOEXCEPT {
return val.data();
}

void push_back(const_reference _x) {
val.push_back(_x);
}

void push_back(value_type&& _x) {
val.emplace_back(_x);
}

template<class... Args>
reference emplace_back(Args&&... _args) {
return val.template emplace_back<>(_args...);
}

void pop_back() {
val.pop_back();
}

void resize(size_type _new_size) {
val.resize(_new_size);
}

iterator insert(const_iterator _position, const_reference _x) {
return val.insert(_position, _x);
}

iterator insert(const_iterator _position, value_type&& _x) {
return val.insert(_position, _x);
}
template<class... Args>
iterator emplace(const_iterator _position, Args&&... _args) {
return val.template emplace<>(_position, _args...);
}

iterator insert(const_iterator _position, size_type _n, const_reference _x) {
return val.insert(_position, _n, _x);
}

template<class InputIterator>
iterator insert(const_iterator _position, InputIterator _first, InputIterator _last) {
return val.insert(_position, _first, _last);
}

iterator insert(const_iterator _position, std::initializer_list<value_type> _il) {
return insert(_position, _il.begin(), _il.end());
}

iterator erase(const_iterator _position) {
return val.erase(_position);
}

iterator erase(const_iterator _first, const_iterator _last) {
return val.erase(_first, _last);
}

void clear() BOOST_NOEXCEPT {
val.clear();
}

};
} // namespace math
} // namespace crypto3
} // namespace nil
19 changes: 11 additions & 8 deletions parallel-crypto3/libs/parallel-math/include/nil/crypto3/math/domains/basic_radix2_domain.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -103,15 +103,18 @@ namespace nil {
}
if (zero_from_size != 0 && zero_from_size < a.size()) {
// resize the vector on the GPU
value_type zero = value_type::zero();
GLOBAL_QUEUE.submit(
[&](sycl::handler& cgh) {
auto a_acc = a.val_buf->template get_access<sycl::access::mode::read_write>(cgh);
cgh.parallel_for(sycl::range<1>(a.size() - zero_from_size), [=](sycl::id<1> index) {
a_acc[zero_from_size + index] = zero;
if constexpr (std::is_same_v<Allocator, uninitialized_allocator<value_type>>) {
value_type zero = value_type::zero();
GLOBAL_QUEUE.submit(
[&](sycl::handler& cgh) {
auto a_acc = a.val_buf->template get_access<sycl::access::mode::read_write>(cgh);
cgh.parallel_for(sycl::range<1>(a.size() - zero_from_size), [=](sycl::id<1> index) {
new (&a_acc[zero_from_size + index]) value_type(zero);
});
});
});
GLOBAL_QUEUE.wait();
GLOBAL_QUEUE.wait();
}
// if default allocator we have already prefilled with zeros
}

detail::basic_radix2_fft_cached<FieldType, polynomial_dfs_type, Allocator>(a, fft_cache_buffer->first);
Expand Down
Loading

0 comments on commit ee9ddd4

Please sign in to comment.