添加gf256的m4rie接口

README.md

@@ -0,0 +1,45 @@
+# cuElim
+
+使用统一内存实现gf2^8域和素域上的矩阵乘法和高斯消元，无需考虑显存大小.
+
+## 使用说明
+
+### 在现有项目中使用
+
+1. 将`include`文件夹中的所有头文件添加到现有项目中.
+
+2. 在需要使用gpu函数的地方引用`cuelim.cuh`
+
+### 使用当前项目
+
+1. 安装依赖
+
+   1. `C++ CUDA CMake ...`
+   2. [`GoogleTest`](https://github.com/google/googletest)
+   3. [`GoogleBenchmark`](https://github.com/google/benchmark)
+
+2. 构建项目
+
+   ```sh
+   mkdir build && cd build
+   cmake ..
+   make -j  # 同时编译多个目标
+   ctest  # 或make test 执行所有测试
+   ```
+
+3. 运行可执行文件
+
+   ```sh
+   ./cuelim  # 执行主程序
+   ./test/target  # 执行特定测试
+   ./benchmark/target  # 执行特定性能测试
+   ```
+
+## 功能简介
+
+- `gf256::MatGF256`：存储GF2^8矩阵，数据结构已经改为与m4rie一致，从低位到高位排列
+- `gf256::ElimResult`：存储高斯消元的结果，包含秩、主元行（进行行交换前的位置）、主元列
+- `gf256::ElimResult gf256::MatGF256::gpu_elim(const gf256::GF256 &gf)`：进行高斯消元
+
+- `gfp::MatGFP`：储存GF65521矩阵，使用32位存储一个元素
+- `gfp::ElimResult gfp::MatGFP::gpu_elim()`：进行高斯消元

benchmark/bench_gfp_mul.cu

@@ -16,5 +16,4 @@ static void bench_gfp(benchmark::State &state)
     }
 }
 
-BENCHMARK(bench_gfp)->Args({10000, 10000, 10000});
-;
+BENCHMARK(bench_gfp)->Args({10000, 10000, 10000});

include/gf256/gf256_header.cuh

@@ -106,7 +106,7 @@ namespace gf256
         GF256(base_t poly)
         {
             assert(irreducible_polynomials_degree_08.count(poly) == 1);
-            this->poly = poly;
+            this->polynomial = poly;
             for (size_t x = 0; x < (1 << gf256_len); x++)
             {
                 mul_table[x][gf256_zero] = gf256_zero;
@@ -146,6 +146,11 @@ namespace gf256
             return inv_table[x];
         }
 
+        base_t poly(void) const
+        {
+            return polynomial;
+        }
+
         inline cudaError_t cpy_to_constant() const
         {
             return cudaMemcpyToSymbol(d_mul_table, mul_table, (1 << gf256_len) * (1 << gf256_len) * sizeof(gf256_t));
@@ -167,13 +172,13 @@ namespace gf256
             {
                 if (temp & (1 << i))
                 {
-                    temp ^= poly << (i - gf256_len);
+                    temp ^= polynomial << (i - gf256_len);
                 }
             }
             return temp;
         }
 
-        base_t poly;
+        base_t polynomial;
         gf256_t inv_table[1 << gf256_num];
         gf256_t mul_table[1 << gf256_num][1 << gf256_num];
     };

include/gfp/gfp_elim.cuh

@@ -134,25 +134,6 @@ namespace gfp
             s_src[b_r][b_c] = s_src[b_r][b_c] ? gfprime - s_src[b_r][b_c] : 0;
         }
         __syncthreads();
-        // if (bx == 0 && by == 0 && tid == 0)
-        // {
-        //     for (int i = 0; i < StepSize; i++)
-        //     {
-        //         for (int j = 0; j < BlockRow; j++)
-        //         {
-        //             printf("%05d ", s_idx[i][j]);
-        //         }
-        //         printf("\n");
-        //     }
-        //     for (int i = 0; i < StepSize; i++)
-        //     {
-        //         for (int j = 0; j < BlockCol; j++)
-        //         {
-        //             printf("%05d ", s_src[i][j]);
-        //         }
-        //         printf("\n");
-        //     }
-        // }
         for (int k = 0; k < rank; k++)
         {
             for (int j = 0; j < BlockRow / THREAD_Y; j++)

include/interface.cuh

@@ -0,0 +1,41 @@
+#ifndef INTERFACE_CUH
+#define INTERFACE_CUH
+
+#include "cuelim.cuh"
+#include <m4rie/m4rie.h>
+
+void mzedread(mzed_t *A, gf256::MatGF256 &mat)
+{
+    assert(A->nrows == mat.nrows && A->ncols == mat.ncols);
+    for (size_t r = 0; r < A->nrows; r++)
+    {
+        for (size_t cn = 0; cn < A->x->width; cn++)
+        {
+            *mat.at_base(r, cn) = A->x->rows[r][cn];
+        }
+    }
+}
+
+void mzedwrite(gf256::MatGF256 &mat, mzed_t *A)
+{
+    assert(A->nrows == mat.nrows && A->ncols == mat.ncols);
+    for (size_t r = 0; r < mat.nrows; r++)
+    {
+        for (size_t cn = 0; cn < mat.width; cn++)
+        {
+            A->x->rows[r][cn] = *mat.at_base(r, cn);
+        }
+    }
+}
+
+size_t gpu_mzed_elim(mzed_t *A)
+{
+    gf256::MatGF256 mat(A->nrows, A->ncols);
+    mzedread(A, mat);
+    gf256::GF256 gf256(A->finite_field->minpoly);
+    gf256::ElimResult res = mat.gpu_elim(gf256);
+    mzedwrite(mat, A);
+    return res.rank;
+}
+
+#endif

src/main.cu

@@ -20,5 +20,5 @@ bool test_gfp_elim(size_t rank, size_t rank_col, size_t nrows, size_t ncols, uin
 
 int main()
 {
-    cout << test_gfp_elim(1234, 2345, 3456, 4567, 41921095) << endl;
+    cout << test_gfp_elim(2000, 20000, 2500, 25000, 41921095) << endl;
 }

test/CMakeLists.txt

@@ -7,6 +7,7 @@ set(TEST_SRC_FILES
     "test_gf256_matrix.cu"
     "test_gf256_elim.cu"
     "test_gfp_mul.cu"
+    "test_gfp_elim.cu"
 )
 
 foreach(SRC ${TEST_SRC_FILES})
@@ -16,13 +17,13 @@ foreach(SRC ${TEST_SRC_FILES})
     gtest_discover_tests(${SRC_NAME})
 endforeach()
 
-# set(TEST_M4RIE_SRC_FILES
-#     "test_m4rie_interface.cu"
-# )
+set(TEST_M4RIE_SRC_FILES
+    "test_interface.cu"
+)
 
-# foreach(SRC ${TEST_M4RIE_SRC_FILES})
-#     get_filename_component(SRC_NAME ${SRC} NAME_WE)
-#     add_executable(${SRC_NAME} ${SRC})
-#     target_link_libraries(${SRC_NAME} GTest::GTest GTest::Main m4ri m4rie)
-#     gtest_discover_tests(${SRC_NAME})
-# endforeach()
+foreach(SRC ${TEST_M4RIE_SRC_FILES})
+    get_filename_component(SRC_NAME ${SRC} NAME_WE)
+    add_executable(${SRC_NAME} ${SRC})
+    target_link_libraries(${SRC_NAME} GTest::GTest GTest::Main m4ri m4rie)
+    gtest_discover_tests(${SRC_NAME})
+endforeach()

test/test_gfp_elim.cu

@@ -0,0 +1,29 @@
+#include <gtest/gtest.h>
+#include "test_header.cuh"
+
+using namespace gfp;
+
+bool test_gfp_elim(size_t rank, size_t rank_col, size_t nrows, size_t ncols, uint_fast32_t seed)
+{
+    MatGFP rdc(rank, ncols);
+    rdc.randomize(rank_col, seed);
+    MatGFP mix(nrows, rank);
+    mix.randomize(seed);
+    MatGFP a = mix * rdc;
+    ElimResult res = a.gpu_elim();
+    MatGFP win(a, 0, 0, res.rank, a.width);
+    return rdc == win;
+}
+
+TEST(TestGFPMul, Small)
+{
+    uint_fast32_t seed = 41921095;
+    EXPECT_TRUE(test_gfp_elim(5, 6, 7, 8, seed));
+}
+
+TEST(TestGFPMul, Mediem)
+{
+    uint_fast32_t seed = 41921095;
+    EXPECT_TRUE(test_gfp_elim(50, 60, 70, 80, seed));
+    EXPECT_TRUE(test_gfp_elim(500, 600, 700, 800, seed));
+}

test/test_interface.cu

@@ -0,0 +1,40 @@
+#include <gtest/gtest.h>
+#include "test_header.cuh"
+#include "interface.cuh"
+
+using namespace gf256;
+
+bool test_gf256_elim_interface(size_t rank, size_t rank_col, size_t nrows, size_t ncols, const GF256 &gf256, uint_fast32_t seed)
+{
+    assert(rank <= nrows && rank <= rank_col && rank_col <= ncols);
+    MatGF256 rdc(rank, ncols);
+    rdc.randomize(rank_col, seed);
+    MatGF256 mix(nrows, rank);
+    mix.randomize(seed);
+    MatGF256 src = gpu_mul(mix, rdc, gf256);
+
+    gf2e *ff_m4rie = gf2e_init(gf256.poly());
+    mzed_t *A_m4rie = mzed_init(ff_m4rie, src.nrows, src.ncols);
+    mzedwrite(src, A_m4rie);
+    mzed_t *A_m4rie_copy = mzed_copy(NULL, A_m4rie);
+
+    base_t rank_interface = gpu_mzed_elim(A_m4rie);
+    rci_t rank_m4rie = mzed_echelonize_newton_john(A_m4rie_copy, 1);
+
+    return (rank_interface == rank_m4rie) && (mzed_cmp(A_m4rie, A_m4rie_copy) == 0);
+}
+
+TEST(TestInterface, Small)
+{
+    uint_fast32_t seed = 41921095;
+    GF256 gf256(0b100011101);
+    EXPECT_TRUE(test_gf256_elim_interface(5, 7, 6, 8, gf256, seed));
+}
+
+TEST(TestInterface, Mediem)
+{
+    uint_fast32_t seed = 41921095;
+    GF256 gf256(0b100011101);
+    EXPECT_TRUE(test_gf256_elim_interface(50, 70, 60, 80, gf256, seed));
+    EXPECT_TRUE(test_gf256_elim_interface(500, 700, 600, 800, gf256, seed));
+}