summaryrefslogtreecommitdiff
path: root/modules
diff options
context:
space:
mode:
Diffstat (limited to 'modules')
-rw-r--r--modules/codec/c++/data.hpp4
-rw-r--r--modules/codec/c++/data_raw.hpp114
-rw-r--r--modules/core/c++/common.hpp6
-rw-r--r--modules/core/c++/error.hpp8
-rw-r--r--modules/remote-hip/c++/data.hpp8
-rw-r--r--modules/remote-hip/c++/device.tmpl.hpp2
-rw-r--r--modules/remote-hip/examples/hip_transfer_data.cpp75
7 files changed, 126 insertions, 91 deletions
diff --git a/modules/codec/c++/data.hpp b/modules/codec/c++/data.hpp
index b14d785..33f301f 100644
--- a/modules/codec/c++/data.hpp
+++ b/modules/codec/c++/data.hpp
@@ -675,11 +675,11 @@ class data<schema::Array<T,Dim>, encode::Native> {
static_assert(sizeof...(Dims)==Dim, "Argument size must be equal to the Dimension");
}
- data<T, encode::Native>& at(const std::array<uint64_t, Dim>& ind){
+ constexpr data<T, encode::Native>& at(const std::array<uint64_t, Dim>& ind){
return value_.at(this->get_flat_index(ind));
}
- const data<T, encode::Native>& at(const std::array<uint64_t, Dim>& ind) const {
+ constexpr const data<T, encode::Native>& at(const std::array<uint64_t, Dim>& ind) const {
return value_.at(this->get_flat_index(ind));
}
diff --git a/modules/codec/c++/data_raw.hpp b/modules/codec/c++/data_raw.hpp
index db57af3..2da49b4 100644
--- a/modules/codec/c++/data_raw.hpp
+++ b/modules/codec/c++/data_raw.hpp
@@ -62,54 +62,54 @@ public:
private:
typename native_data_type<Schema>::type value_;
public:
- data():value_{}{}
- data(data<MetaSchema, encode::NativeRaw>):value_{}{}
+ constexpr data():value_{}{}
+ constexpr data(data<MetaSchema, encode::NativeRaw>):value_{}{}
SAW_DEFAULT_COPY(data);
SAW_DEFAULT_MOVE(data);
- data(typename native_data_type<Schema>::type value__):
+ constexpr data(typename native_data_type<Schema>::type value__):
value_{std::move(value__)}{}
- void set(typename native_data_type<Schema>::type val){
+ constexpr void set(typename native_data_type<Schema>::type val){
value_ = val;
}
- typename native_data_type<Schema>::type get() const {return value_;}
+ constexpr typename native_data_type<Schema>::type get() const {return value_;}
- data<Schema, encode::NativeRaw> operator*(const data<Schema, encode::NativeRaw>& rhs)const{
+ constexpr data<Schema, encode::NativeRaw> operator*(const data<Schema, encode::NativeRaw>& rhs)const{
return {get() * rhs.get()};
}
- data<Schema, encode::NativeRaw> operator/(const data<Schema, encode::NativeRaw>& rhs)const{
+ constexpr data<Schema, encode::NativeRaw> operator/(const data<Schema, encode::NativeRaw>& rhs)const{
return {get() / rhs.get()};
}
- data<Schema, encode::NativeRaw> operator+(const data<Schema, encode::NativeRaw>& rhs)const{
+ constexpr data<Schema, encode::NativeRaw> operator+(const data<Schema, encode::NativeRaw>& rhs)const{
return {get() + rhs.get()};
}
- data<Schema, encode::NativeRaw>& operator+=(const data<Schema, encode::NativeRaw>& rhs)const{
+ constexpr data<Schema, encode::NativeRaw>& operator+=(const data<Schema, encode::NativeRaw>& rhs)const{
value_ += rhs.get();
return *this;
}
- data<Schema, encode::NativeRaw>& operator-=(const data<Schema, encode::NativeRaw>& rhs)const{
+ constexpr data<Schema, encode::NativeRaw>& operator-=(const data<Schema, encode::NativeRaw>& rhs)const{
value_ -= rhs.get();
return *this;
}
- data<Schema, encode::NativeRaw> operator-(const data<Schema, encode::NativeRaw>& rhs)const{
+ constexpr data<Schema, encode::NativeRaw> operator-(const data<Schema, encode::NativeRaw>& rhs)const{
return {get() - rhs.get()};
}
template<typename Enc>
- bool operator==(const data<Schema, Enc>& rhs)const{
+ constexpr bool operator==(const data<Schema, Enc>& rhs)const{
return get() == rhs.get();
}
template<typename Enc>
- bool operator<(const data<Schema, Enc>& rhs) const {
+ constexpr bool operator<(const data<Schema, Enc>& rhs) const {
return get() < rhs.get();
}
@@ -117,7 +117,7 @@ public:
* Casts
*/
template<typename Target>
- data<Target, encode::NativeRaw> cast_to() const {
+ constexpr data<Target, encode::NativeRaw> cast_to() const {
auto raw_to = static_cast<typename saw::native_data_type<Target>::type>(value_);
return {raw_to};
}
@@ -132,64 +132,64 @@ public:
private:
ref<typename native_data_type<ReferencedSchema>::type > value_;
public:
- data(ref<typename native_data_type<ReferencedSchema >::type > value__):
+ constexpr data(ref<typename native_data_type<ReferencedSchema >::type > value__):
value_{value__}
{}
SAW_DEFAULT_COPY(data);
SAW_FORBID_MOVE(data);
- void set(typename native_data_type<ReferencedSchema>::type val){
+ constexpr void set(typename native_data_type<ReferencedSchema>::type val){
value_() = val;
}
- typename native_data_type<ReferencedSchema>::type get() const {
+ constexpr typename native_data_type<ReferencedSchema>::type get() const {
return value_();
}
- data<ReferencedSchema, encode::NativeRaw> operator*(const data<Schema, encode::NativeRaw>& rhs)const{
+ constexpr data<ReferencedSchema, encode::NativeRaw> operator*(const data<Schema, encode::NativeRaw>& rhs)const{
return {get() * rhs.get()};
}
- data<ReferencedSchema, encode::NativeRaw> operator/(const data<Schema, encode::NativeRaw>& rhs)const{
+ constexpr data<ReferencedSchema, encode::NativeRaw> operator/(const data<Schema, encode::NativeRaw>& rhs)const{
return {get() / rhs.get()};
}
- data<ReferencedSchema, encode::NativeRaw> operator+(const data<Schema, encode::NativeRaw>& rhs)const{
+ constexpr data<ReferencedSchema, encode::NativeRaw> operator+(const data<Schema, encode::NativeRaw>& rhs)const{
return {get() + rhs.get()};
}
- data<Schema, encode::NativeRaw>& operator+=(const data<Schema, encode::NativeRaw>& rhs)const{
+ constexpr data<Schema, encode::NativeRaw>& operator+=(const data<Schema, encode::NativeRaw>& rhs)const{
value_ += rhs.get();
return *this;
}
- data<Schema, encode::NativeRaw>& operator-=(const data<Schema, encode::NativeRaw>& rhs)const{
+ constexpr data<Schema, encode::NativeRaw>& operator-=(const data<Schema, encode::NativeRaw>& rhs)const{
value_ -= rhs.get();
return *this;
}
- data<ReferencedSchema, encode::NativeRaw> operator-(const data<Schema, encode::NativeRaw>& rhs)const{
+ constexpr data<ReferencedSchema, encode::NativeRaw> operator-(const data<Schema, encode::NativeRaw>& rhs)const{
return {get() - rhs.get()};
}
template<typename Enc>
- bool operator==(const data<Schema, Enc>& rhs)const{
+ constexpr bool operator==(const data<Schema, Enc>& rhs)const{
return get() == rhs.get();
}
template<typename Enc>
- bool operator==(const data<ReferencedSchema, Enc>& rhs)const{
+ constexpr bool operator==(const data<ReferencedSchema, Enc>& rhs)const{
return get() == rhs.get();
}
template<typename Enc>
- bool operator<(const data<Schema, Enc>& rhs) const {
+ constexpr bool operator<(const data<Schema, Enc>& rhs) const {
return get() < rhs.get();
}
template<typename Enc>
- bool operator<(const data<ReferencedSchema, Enc>& rhs) const {
+ constexpr bool operator<(const data<ReferencedSchema, Enc>& rhs) const {
return get() < rhs.get();
}
};
@@ -210,70 +210,70 @@ public:
data(typename saw::native_data_type<typename Schema::InterfaceSchema>::type val__):value_{static_cast<typename saw::native_data_type<typename Schema::StorageSchema>::type>(val__)}{}
- typename saw::native_data_type<typename Schema::InterfaceSchema>::type get() const {
+ constexpr typename saw::native_data_type<typename Schema::InterfaceSchema>::type get() const {
return value_.template cast_to<typename Schema::InterfaceSchema>().get();
}
- data(const saw::data<typename Schema::InterfaceSchema, encode::NativeRaw>& val){
+ constexpr data(const saw::data<typename Schema::InterfaceSchema, encode::NativeRaw>& val){
value_ = val.template cast_to<typename Schema::StorageSchema>();
}
- void set(typename saw::native_data_type<typename Schema::InterfaceSchema>::type val){
+ constexpr void set(typename saw::native_data_type<typename Schema::InterfaceSchema>::type val){
value_.set(static_cast<typename Schema::StorageSchema>(val));
}
- data<Schema, encode::NativeRaw> operator*(const data<Schema, encode::NativeRaw>& rhs) const {
+ constexpr data<Schema, encode::NativeRaw> operator*(const data<Schema, encode::NativeRaw>& rhs) const {
using CalcType = typename native_data_type<typename Schema::InterfaceSchema>::type;
CalcType left = static_cast<CalcType>(value_.get());
CalcType right = static_cast<CalcType>(rhs.get());
return {left * right};
}
- data<typename Schema::InterfaceSchema, encode::NativeRaw> operator*(const data<typename Schema::InterfaceSchema, encode::NativeRaw>& rhs) const {
+ constexpr data<typename Schema::InterfaceSchema, encode::NativeRaw> operator*(const data<typename Schema::InterfaceSchema, encode::NativeRaw>& rhs) const {
using CalcType = typename native_data_type<typename Schema::InterfaceSchema>::type;
CalcType left = static_cast<CalcType>(value_.get());
CalcType right = rhs.get();
return {left * right};
}
- data<typename Schema::InterfaceSchema, encode::NativeRaw> operator/(const data<Schema, encode::NativeRaw>& rhs)const{
+ constexpr data<typename Schema::InterfaceSchema, encode::NativeRaw> operator/(const data<Schema, encode::NativeRaw>& rhs)const{
using CalcType = typename native_data_type<typename Schema::InterfaceSchema>::type;
CalcType left = static_cast<CalcType>(value_.get());
CalcType right = static_cast<CalcType>(rhs.get());
return {left / right};
}
- data<typename Schema::InterfaceSchema, encode::NativeRaw> operator+(const data<Schema, encode::NativeRaw>& rhs)const{
+ constexpr data<typename Schema::InterfaceSchema, encode::NativeRaw> operator+(const data<Schema, encode::NativeRaw>& rhs)const{
using CalcType = typename native_data_type<typename Schema::InterfaceSchema>::type;
CalcType left = static_cast<CalcType>(value_.get());
CalcType right = static_cast<CalcType>(rhs.get());
return {left + right};
}
- data<Schema, encode::NativeRaw>& operator+=(const data<Schema, encode::NativeRaw>& rhs)const{
+ constexpr data<Schema, encode::NativeRaw>& operator+=(const data<Schema, encode::NativeRaw>& rhs)const{
*this = *this + rhs.get();
return *this;
}
- data<typename Schema::InterfaceSchema, encode::NativeRaw> operator-(const data<Schema, encode::NativeRaw>& rhs)const{
+ constexpr data<typename Schema::InterfaceSchema, encode::NativeRaw> operator-(const data<Schema, encode::NativeRaw>& rhs)const{
using CalcType = typename native_data_type<typename Schema::InterfaceSchema>::type;
CalcType left = static_cast<CalcType>(value_.get());
CalcType right = static_cast<CalcType>(rhs.get());
return {left - right};
}
- data<Schema, encode::NativeRaw>& operator-=(const data<Schema, encode::NativeRaw>& rhs) const {
+ constexpr data<Schema, encode::NativeRaw>& operator-=(const data<Schema, encode::NativeRaw>& rhs) const {
*this = *this - rhs.get();
return *this;
}
template<typename Enc>
- bool operator==(const data<Schema, Enc>& rhs)const{
+ constexpr bool operator==(const data<Schema, Enc>& rhs)const{
return get() == rhs.get();
}
template<typename Enc>
- bool operator<(const data<Schema, Enc>& rhs) const {
+ constexpr bool operator<(const data<Schema, Enc>& rhs) const {
return get() < rhs.get();
}
};
@@ -415,14 +415,14 @@ class data<schema::Array<T,Dim>, encode::NativeRaw> {
typename raw_native_array_type_helper<T>::Type* value_;
uint64_t value_size_;
- uint64_t get_full_size() const {
+ constexpr uint64_t get_full_size() const {
uint64_t s = 1;
for(uint64_t iter = 0; iter < Dim; ++iter){
- assert(dims_.at(iter) > 0);
- s *= dims_.at(iter);
+ s *= dims_[iter];
}
+
return s;
}
public:
@@ -438,7 +438,7 @@ class data<schema::Array<T,Dim>, encode::NativeRaw> {
/**
* Purely internal function for working C interfacing
*/
- typename raw_native_array_type_helper<T>::Type* get_raw_data() {
+ constexpr typename raw_native_array_type_helper<T>::Type* get_raw_data() {
if(value_size_ == 0u){
return nullptr;
}
@@ -477,54 +477,53 @@ class data<schema::Array<T,Dim>, encode::NativeRaw> {
}
template<std::integral... Dims>
- data(Dims... size_):
+ constexpr data(Dims... size_):
data{{static_cast<uint64_t>(size_)...}}
{
static_assert(sizeof...(Dims)==Dim, "Argument size must be equal to the Dimension");
}
- data<schema::Ref<T>, encode::NativeRaw> at(const std::array<uint64_t, Dim>& ind){
+ constexpr data<schema::Ref<T>, encode::NativeRaw> at(const std::array<uint64_t, Dim>& ind){
return {value_[this->get_flat_index(ind)]};
}
- const data<schema::Ref<T>, encode::NativeRaw> at(const std::array<uint64_t, Dim>& ind) const {
+ constexpr const data<schema::Ref<T>, encode::NativeRaw> at(const std::array<uint64_t, Dim>& ind) const {
return {value_[this->get_flat_index(ind)]};
}
template<std::integral... Dims>
- data<schema::Ref<T>, encode::NativeRaw> at(Dims... i){
+ constexpr data<schema::Ref<T>, encode::NativeRaw> at(Dims... i){
return {value_[this->get_flat_index(std::array<uint64_t, Dim>{static_cast<uint64_t>(i)...})]};
}
template<std::integral... Dims>
- const data<schema::Ref<T>, encode::NativeRaw> at(Dims... i) const {
+ constexpr const data<schema::Ref<T>, encode::NativeRaw> at(Dims... i) const {
return {value_[this->get_flat_index(std::array<uint64_t, Dim>{static_cast<uint64_t>(i)...})]};
}
template<typename Encoding>
- data<schema::Ref<T>,encode::NativeRaw> at(const data<schema::FixedArray<schema::UInt64,Dim>, Encoding>& i){
+ constexpr data<schema::Ref<T>,encode::NativeRaw> at(const data<schema::FixedArray<schema::UInt64,Dim>, Encoding>& i){
return {value_[this->get_flat_index(i)]};
}
template<typename Encoding>
- const data<schema::Ref<T>,encode::NativeRaw> at(const data<schema::FixedArray<schema::UInt64,Dim>, Encoding>& i)const{
+ constexpr const data<schema::Ref<T>,encode::NativeRaw> at(const data<schema::FixedArray<schema::UInt64,Dim>, Encoding>& i)const{
return {value_[this->get_flat_index(i)]};
}
- std::size_t get_dim_size(uint64_t i) const {
+ constexpr std::size_t get_dim_size(uint64_t i) const {
return dims_.at(i);
}
- uint64_t size() const { return value_size_;}
+ constexpr uint64_t size() const { return value_size_;}
- data<schema::FixedArray<schema::UInt64, Dim>, encode::NativeRaw> get_dims() const {
+ constexpr data<schema::FixedArray<schema::UInt64, Dim>, encode::NativeRaw> get_dims() const {
return {dims_};
}
private:
template<typename U>
- uint64_t get_flat_index(const U& i) const {
- assert(value_size_ == get_full_size());
+ constexpr uint64_t get_flat_index(const U& i) const {
uint64_t s = 0;
uint64_t stride = 1;
@@ -539,10 +538,9 @@ private:
}else{
static_assert(always_false<V>, "Cases exhausted");
}
- }(i.at(iter));
- assert(ind < dims_.at(iter));
+ }(i[iter]);
s += ind * stride;
- stride *= dims_.at(iter);
+ stride *= dims_[iter];
}
return s;
diff --git a/modules/core/c++/common.hpp b/modules/core/c++/common.hpp
index f63c531..de464b8 100644
--- a/modules/core/c++/common.hpp
+++ b/modules/core/c++/common.hpp
@@ -90,7 +90,7 @@ public:
/**
* Main constructor.
*/
- ref(T& ref__):
+ constexpr ref(T& ref__):
ref_{&ref__}
{}
@@ -99,14 +99,14 @@ public:
/**
* Operator retrieving the itself.
*/
- T& operator()(){
+ constexpr T& operator()(){
return *ref_;
}
/**
* Operator retrieving the itself.
*/
- const T& operator()() const {
+ constexpr const T& operator()() const {
return *ref_;
}
};
diff --git a/modules/core/c++/error.hpp b/modules/core/c++/error.hpp
index dab297e..ac45fbc 100644
--- a/modules/core/c++/error.hpp
+++ b/modules/core/c++/error.hpp
@@ -18,7 +18,7 @@ namespace saw {
* critical and recoverable errors. Additional code ids can be provided to the
* constructor if additional distinctions are necessary.
*/
-class error {
+class error final {
public:
/**
* Type alias
@@ -101,7 +101,7 @@ namespace impl {
/**
* Internal registry for all error types. Is generated dynamically.
*/
-class error_registry {
+class error_registry final {
private:
struct error_info {
error_info() = delete;
@@ -301,6 +301,10 @@ public:
}
};
+/**
+ * This tries to catch cases where error starts including itself as a type which can happen in more complicated cases.
+ * So this acts as a type safe guard.
+ */
template <typename T> class error_or<error_or<T>> {
private:
error_or() = delete;
diff --git a/modules/remote-hip/c++/data.hpp b/modules/remote-hip/c++/data.hpp
index 3e7c3ed..5d3635f 100644
--- a/modules/remote-hip/c++/data.hpp
+++ b/modules/remote-hip/c++/data.hpp
@@ -8,16 +8,16 @@ namespace saw {
* Generic wrapper class which stores data on the sycl side.
* Most of the times this will be a root object.
*/
-template<typename Schema>
-class data<Schema, encode::Hip<encode::Native>> {
+template<typename Schema,typename T>
+class data<Schema, encode::Hip<T>> {
private:
- data<Schema, encode::Native>* data_;
+ data<Schema, T>* data_;
public:
data():
data_{nullptr}
{}
- data<Schema, encode::Native>** get_device_data() {
+ data<Schema, T>** get_device_data() {
return &data_;
}
};
diff --git a/modules/remote-hip/c++/device.tmpl.hpp b/modules/remote-hip/c++/device.tmpl.hpp
index 6edf431..0517f67 100644
--- a/modules/remote-hip/c++/device.tmpl.hpp
+++ b/modules/remote-hip/c++/device.tmpl.hpp
@@ -33,7 +33,7 @@ struct hip_copy_to_device<schema::Array<T,Dim>, Encoding> {
static error_or<void> apply(data<Schema, Encoding>& from, data<Schema,Encoding>** to){
typename native_data_type<T>::type* dat{};
hipError_t data_malloc_err = hipMalloc(&dat,sizeof(typename native_data_type<T>::type) * from.size());
- hipError_t data_copy_err = hipMemcpy(&dat, &(from.get_raw_data()),sizeof(typename native_data_type<T>::type) * from.size(), hipMemcpyHostToDevice);
+ hipError_t data_copy_err = hipMemcpy(&dat, (from.get_raw_data()),sizeof(typename native_data_type<T>::type) * from.size(), hipMemcpyHostToDevice);
if(from.size() == 0u){
// Everything is fine. We just don't want to allocate data which doesn't exist.
diff --git a/modules/remote-hip/examples/hip_transfer_data.cpp b/modules/remote-hip/examples/hip_transfer_data.cpp
index 18c82df..3f02d87 100644
--- a/modules/remote-hip/examples/hip_transfer_data.cpp
+++ b/modules/remote-hip/examples/hip_transfer_data.cpp
@@ -1,20 +1,22 @@
+#include <forstio/codec/data_raw.hpp>
+
#include "../c++/remote.hpp"
#include "../c++/transfer.hpp"
#include <iostream>
-__global__ void print_value(saw::data<saw::schema::Int16,saw::encode::Native>* val){
+__global__ void print_value(saw::data<saw::schema::Int16,saw::encode::NativeRaw>* val){
int v = val->get();
printf("Hello world: %d\n", v);
}
-__global__ void print_array_vals(saw::data<saw::schema::Array<saw::schema::Int16>* val){
+__global__ void print_array_vals(saw::data<saw::schema::Array<saw::schema::Int16>, saw::encode::NativeRaw>* val){
uint64_t orig_len = val->size();
long len = (long) orig_len;
printf("Array size: %ld\n", len);
- for(uint64_t i = 0; i < orig_len; +i){
- int v = val->at(i);
+ for(uint64_t i = 0; i < orig_len; ++i){
+ int v = val->at(i).get();
printf("%d ", v);
}
printf("\n");
@@ -35,28 +37,59 @@ saw::error_or<void> real_main(){
}
auto& addr = eo_addr.get_value();
- auto eo_dat_srv = rmt.data_listen<sch::Int16, encode::Native>(*addr);
- if(eo_dat_srv.is_error()){
- return std::move(eo_dat_srv.get_error());
- }
- auto& dat_srv = eo_dat_srv.get_value();
+ {
+ auto eo_dat_srv = rmt.data_listen<sch::Int16, encode::NativeRaw>(*addr);
+ if(eo_dat_srv.is_error()){
+ return std::move(eo_dat_srv.get_error());
+ }
+ auto& dat_srv = eo_dat_srv.get_value();
- data<sch::Int16> val{42};
- id<sch::Int16> id_val{0u};
- auto eo_send = dat_srv->send(val, id_val);
- if(eo_send.is_error()){
- return std::move(eo_send.get_error());
- }
+ data<sch::Int16,encode::NativeRaw> val{42};
+ id<sch::Int16> id_val{0u};
+ auto eo_send = dat_srv->send(val, id_val);
+ if(eo_send.is_error()){
+ return std::move(eo_send.get_error());
+ }
+
+ auto eo_dfind = dat_srv->find(id_val);
+ if(eo_dfind.is_error()){
+ return std::move(eo_dfind.get_error());
+ }
+ auto dfind = eo_dfind.get_value();
- auto eo_dfind = dat_srv->find(id_val);
- if(eo_dfind.is_error()){
- return std::move(eo_dfind.get_error());
+ auto& v = dfind();
+
+ print_value<<<dim3(2),dim3(2),0,hipStreamDefault>>>(*(v.get_device_data()));
}
- auto dfind = eo_dfind.get_value();
- auto& v = dfind();
+ {
+ auto eo_dat_srv = rmt.data_listen<sch::Array<sch::Int16>, encode::NativeRaw>(*addr);
+ if(eo_dat_srv.is_error()){
+ return std::move(eo_dat_srv.get_error());
+ }
+ auto& dat_srv = eo_dat_srv.get_value();
+
+ data<sch::Array<sch::Int16>,encode::NativeRaw> val{4};
+ val.at(0u).set(5);
+ val.at(1u).set(3);
+ val.at(2u).set(-6);
+ val.at(3u).set(1);
+ id<sch::Array<sch::Int16>> id_val{0u};
+ auto eo_send = dat_srv->send(val, id_val);
+ if(eo_send.is_error()){
+ return std::move(eo_send.get_error());
+ }
- print_value<<<dim3(2),dim3(2),0,hipStreamDefault>>>(*(v.get_device_data()));
+ auto eo_dfind = dat_srv->find(id_val);
+ if(eo_dfind.is_error()){
+ return std::move(eo_dfind.get_error());
+ }
+ auto dfind = eo_dfind.get_value();
+
+ auto& v = dfind();
+
+ print_array_vals<<<dim3(2),dim3(2),0,hipStreamDefault>>>(*(v.get_device_data()));
+ }
return make_void();
}
generated by cgit v1.2.3 (git 2.25.1) at 2024-10-16 14:16:47 +0000