enable prediction-quantization with compaction

- No external gather is needed any more; compaction inside prediction kernel is utilized.
- Also prediction kernel throughput is significantly improved due to less DRAM write.

Author: jtian0

include/kernel/l23r.hh

@@ -24,7 +24,7 @@ void psz_adhoc_scttr(
     T* val, uint32_t* idx, int const n, T* out, float* milliseconds,
     cudaStream_t stream);
 
-template <typename T, bool UsePnEnc = false, typename Eq = uint32_t>
+template <typename T, typename Eq = uint32_t, bool UsePnEnc = false>
 cusz_error_status psz_comp_l23r(
     T* const data, dim3 const len3, double const eb, int const radius,
     Eq* const eq, void* _outlier, float* time_elapsed, cudaStream_t stream);

src/pipeline/compact_cuda.inl include/mem/compact_cu.hhsrc/pipeline/compact_cuda.inl renamed to include/mem/compact_cu.hh

@@ -1,5 +1,5 @@
 /**
- * @file compact_cuda.inl
+ * @file compact_cu.hh
  * @author Jiannan Tian
  * @brief
  * @version 0.4
@@ -32,9 +32,9 @@ struct CompactCudaDram {
   using type = T;
 
   // `h_` for host-accessible
-  T *val, *h_val;
-  uint32_t *idx, *h_idx;
-  uint32_t *num, h_num{0};
+  T *d_val, *h_val;
+  uint32_t *d_idx, *h_idx;
+  uint32_t *d_num, h_num{0};
   size_t reserved_len;
 
   // CompactCudaDram() {}
@@ -48,10 +48,10 @@ struct CompactCudaDram {
 
   CompactCudaDram& malloc()
   {
-    cudaMalloc(&val, sizeof(T) * reserved_len);
-    cudaMalloc(&idx, sizeof(uint32_t) * reserved_len);
-    cudaMalloc(&num, sizeof(uint32_t) * 1);
-    cudaMemset(num, 0x0, sizeof(T) * 1);  // init val
+    cudaMalloc(&d_val, sizeof(T) * reserved_len);
+    cudaMalloc(&d_idx, sizeof(uint32_t) * reserved_len);
+    cudaMalloc(&d_num, sizeof(uint32_t) * 1);
+    cudaMemset(d_num, 0x0, sizeof(T) * 1);  // init d_val
 
     return *this;
   }
@@ -66,7 +66,7 @@ struct CompactCudaDram {
 
   CompactCudaDram& free()
   {
-    cudaFree(idx), cudaFree(val), cudaFree(num);
+    cudaFree(d_idx), cudaFree(d_val), cudaFree(d_num);
     return *this;
   }
 
@@ -79,10 +79,10 @@ struct CompactCudaDram {
   // memcpy
   CompactCudaDram& make_host_accessible(cudaStream_t stream = 0)
   {
-    cudaMemcpyAsync(&h_num, num, 1 * sizeof(uint32_t), d2h, stream);
+    cudaMemcpyAsync(&h_num, d_num, 1 * sizeof(uint32_t), d2h, stream);
     cudaStreamSynchronize(stream);
-    cudaMemcpyAsync(h_val, val, sizeof(T) * (h_num), d2h, stream);
-    cudaMemcpyAsync(h_idx, idx, sizeof(uint32_t) * (h_num), d2h, stream);
+    cudaMemcpyAsync(h_val, d_val, sizeof(T) * (h_num), d2h, stream);
+    cudaMemcpyAsync(h_idx, d_idx, sizeof(uint32_t) * (h_num), d2h, stream);
     cudaStreamSynchronize(stream);
 
     return *this;
@@ -110,6 +110,9 @@ struct CompactCudaDram {
 
   // accessor
   uint32_t num_outliers() { return h_num; }
+  T* val() { return d_val; }
+  uint32_t* idx() { return d_idx; }
+  uint32_t* num() { return d_num; }
 };
 
 #endif /* F712F74C_7488_4445_83EE_EE7F88A64BBA */

src/pipeline/compact_serial.inl include/mem/compact_ser.hhsrc/pipeline/compact_serial.inl renamed to include/mem/compact_ser.hh

@@ -1,5 +1,5 @@
 /**
- * @file compact_serial.inl
+ * @file compact_ser.hh
  * @author Jiannan Tian
  * @brief
  * @version 0.4

include/mem/layout_cxx.hh

@@ -12,8 +12,9 @@
 #ifndef DC62DA60_8211_4C93_9541_950ADEFC2820
 #define DC62DA60_8211_4C93_9541_950ADEFC2820
 
+#include "compact_cu.hh"
 #include "layout.h"
-#include "mem/memseg.h"
+#include "memseg.h"
 #include "memseg_cxx.hh"
 
 template <typename T, typename E, typename H>
@@ -34,6 +35,8 @@ class pszmempool_cxx {
   pszmem_cxx<T> *sv;  // sp-val
   pszmem_cxx<M> *si;  // sp-idx
 
+  CompactCudaDram<T> *compact;
+
   pszmem_cxx<B> *_compressed;  // compressed
 
   size_t len;
@@ -52,8 +55,11 @@ class pszmempool_cxx {
   F *hist() { return ht->dptr(); }
   E *ectrl_lrz() { return el->dptr(); }
   E *ectrl_spl() { return es->dptr(); }
-  B *compressed() { return _compressed->dptr(); };
-  B *compressed_h() { return _compressed->hptr(); };
+  B *compressed() { return _compressed->dptr(); }
+  B *compressed_h() { return _compressed->hptr(); }
+  T *compact_val() { return compact->val(); }
+  M *compact_idx() { return compact->idx(); }
+  M compact_num_outliers() { return compact->num_outliers(); }
 };
 
 #define TPL template <typename T, typename E, typename H>
@@ -90,6 +96,8 @@ TPL POOL::pszmempool_cxx(_u4 x, int _radius, _u4 y, _u4 z)
   sv = new pszmem_cxx<T>(x, y, z, "sp-val");
   si = new pszmem_cxx<M>(x, y, z, "sp-idx");
 
+  compact = new CompactCudaDram<T>;
+
   _compressed->control({Malloc, MallocHost});
   oc->control({Malloc, MallocHost});
   ac->control({Malloc, MallocHost});
@@ -98,13 +106,16 @@ TPL POOL::pszmempool_cxx(_u4 x, int _radius, _u4 y, _u4 z)
   sv->control({Malloc, MallocHost});
   si->control({Malloc, MallocHost});
 
+  compact->reserve_space(len / 5).control({Malloc, MallocHost});
+
   el->asaviewof(e);
   es->asaviewof(e);
 }
 
 TPL POOL::~pszmempool_cxx()
 {
   delete ac, delete e, delete oc, delete ht, delete sv, delete si;
+  compact->control({Free, FreeHost});
 }
 
 TPL POOL *POOL::clear_buffer()

src/kernel/detail/l23r.inl

@@ -17,7 +17,7 @@
 #include <type_traits>
 
 #include "cusz/suint.hh"
-#include "pipeline/compact_cuda.inl"
+#include "mem/compact_cu.hh"
 
 namespace psz {
 namespace rolling {

src/kernel/detail/lproto.inl

@@ -18,7 +18,7 @@
 #include <stdexcept>
 
 #include "../../utils/it_cuda.hh"
-#include "pipeline/compact_cuda.inl"
+#include "mem/compact_cu.hh"
 #include "utils/cuda_err.cuh"
 #include "utils/timer.h"
 
@@ -48,9 +48,9 @@ __global__ void c_lorenzo_1d1l(T* in_data, dim3 len3, dim3 stride3, int radius,
     if (check_boundary1()) {  // postquant
         eq[id] = quantizable * static_cast<Eq>(candidate);
         if (not quantizable) {
-            auto dram_idx         = atomicAdd(compact.num, 1);
-            compact.val[dram_idx] = candidate;
-            compact.idx[dram_idx] = id;
+            auto dram_idx         = atomicAdd(compact.d_num, 1);
+            compact.d_val[dram_idx] = candidate;
+            compact.d_idx[dram_idx] = id;
         }
     }
 }
@@ -79,9 +79,9 @@ __global__ void c_lorenzo_2d1l(T* in_data, dim3 len3, dim3 stride3, int radius,
     if (check_boundary2()) {
         eq[id] = quantizable * static_cast<Eq>(candidate);
         if (not quantizable) {
-            auto dram_idx         = atomicAdd(compact.num, 1);
-            compact.val[dram_idx] = candidate;
-            compact.idx[dram_idx] = id;
+            auto dram_idx         = atomicAdd(compact.d_num, 1);
+            compact.d_val[dram_idx] = candidate;
+            compact.d_idx[dram_idx] = id;
         }
     }
 }
@@ -112,9 +112,9 @@ __global__ void c_lorenzo_3d1l(T* in_data, dim3 len3, dim3 stride3, int radius,
     if (check_boundary3()) {
         eq[id] = quantizable * static_cast<Eq>(candidate);
         if (not quantizable) {
-            auto dram_idx         = atomicAdd(compact.num, 1);
-            compact.val[dram_idx] = candidate;
-            compact.idx[dram_idx] = id;
+            auto dram_idx         = atomicAdd(compact.d_num, 1);
+            compact.d_val[dram_idx] = candidate;
+            compact.d_idx[dram_idx] = id;
         }
     }
 }

src/kernel/detail/subroutine.inl

@@ -12,7 +12,7 @@
 #include <stdint.h>
 #include <type_traits>
 #include "cusz/suint.hh"
-#include "pipeline/compact_cuda.inl"
+#include "mem/compact_cu.hh"
 #include "subsub.inl"
 
 namespace psz {

src/kernel/l23r_cu.cu

@@ -14,11 +14,11 @@
 #include "cusz/type.h"
 #include "detail/l23r.inl"
 #include "kernel/l23r.hh"
-#include "pipeline/compact_cuda.inl"
+#include "mem/compact_cu.hh"
 #include "utils/cuda_err.cuh"
 #include "utils/timer.h"
 
-template <typename T, bool UsePnEnc, typename Eq>
+template <typename T, typename Eq, bool UsePnEnc>
 cusz_error_status psz_comp_l23r(
     T* const data, dim3 const len3, double const eb, int const radius,
     Eq* const eq, void* _outlier, float* time_elapsed, cudaStream_t stream)
@@ -72,17 +72,20 @@ cusz_error_status psz_comp_l23r(
   if (d == 1) {
     psz::rolling::c_lorenzo_1d1l<T, false, Eq, T, Tile1D, Seq1D>
         <<<Grid1D, Block1D, 0, stream>>>(
-            data, len3, leap3, radius, ebx2_r, eq, ot->val, ot->idx, ot->num);
+            data, len3, leap3, radius, ebx2_r, eq, ot->val(), ot->idx(),
+            ot->num());
   }
   else if (d == 2) {
     psz::rolling::c_lorenzo_2d1l<T, false, Eq, T>
         <<<Grid2D, Block2D, 0, stream>>>(
-            data, len3, leap3, radius, ebx2_r, eq, ot->val, ot->idx, ot->num);
+            data, len3, leap3, radius, ebx2_r, eq, ot->val(), ot->idx(),
+            ot->num());
   }
   else if (d == 3) {
     psz::rolling::c_lorenzo_3d1l<T, false, Eq, T>
         <<<Grid3D, Block3D, 0, stream>>>(
-            data, len3, leap3, radius, ebx2_r, eq, ot->val, ot->idx, ot->num);
+            data, len3, leap3, radius, ebx2_r, eq, ot->val(), ot->idx(),
+            ot->num());
   }
 
   STOP_CUDAEVENT_RECORDING(stream);
@@ -93,22 +96,22 @@ cusz_error_status psz_comp_l23r(
   return CUSZ_SUCCESS;
 }
 
-template cusz_error_status psz_comp_l23r<float, false>(
+template cusz_error_status psz_comp_l23r<float, uint32_t, false>(
     float* const data, dim3 const len3, double const eb, int const radius,
     uint32_t* const eq, void* _outlier, float* time_elapsed,
     cudaStream_t stream);
 
-template cusz_error_status psz_comp_l23r<float, true>(
+template cusz_error_status psz_comp_l23r<float, uint32_t, true>(
     float* const data, dim3 const len3, double const eb, int const radius,
     uint32_t* const eq, void* _outlier, float* time_elapsed,
     cudaStream_t stream);
 
-template cusz_error_status psz_comp_l23r<double, false>(
+template cusz_error_status psz_comp_l23r<double, uint32_t, false>(
     double* const data, dim3 const len3, double const eb, int const radius,
     uint32_t* const eq, void* _outlier, float* time_elapsed,
     cudaStream_t stream);
 
-template cusz_error_status psz_comp_l23r<double, true>(
+template cusz_error_status psz_comp_l23r<double, uint32_t, true>(
     double* const data, dim3 const len3, double const eb, int const radius,
     uint32_t* const eq, void* _outlier, float* time_elapsed,
     cudaStream_t stream);

src/kernel/lproto_cu.cu

@@ -11,7 +11,7 @@
 
 #include "cusz/type.h"
 #include "kernel/lproto.hh"
-#include "pipeline/compact_cuda.inl"
+#include "mem/compact_cu.hh"
 #include "utils/cuda_err.cuh"
 #include "utils/timer.h"
 

src/pipeline/compressor.inl

@@ -25,11 +25,11 @@
 #include "kernel/hist.hh"
 #include "kernel/histsp.hh"
 #include "kernel/l23.hh"
+#include "kernel/l23r.hh"
 #include "kernel/spv_gpu.hh"
 #include "mem/layout.h"
 #include "mem/layout_cxx.hh"
 #include "mem/memseg_cxx.hh"
-#include "kernel/hist.hh"
 #include "utils/config.hh"
 #include "utils/cuda_err.cuh"
 
@@ -79,12 +79,10 @@ Compressor<C>* Compressor<C>::init(CONFIG* config, bool debug)
 
   codec->init(len, booklen, pardeg, debug);
   mem = new pszmempool_cxx<T, E, H>(x, radius, y, z);
-  
+
   return this;
 }
 
-
-
 template <class C>
 Compressor<C>* Compressor<C>::compress(
     cusz_context* config, T* in, BYTE*& out, size_t& outlen,
@@ -113,39 +111,50 @@ Compressor<C>* Compressor<C>::compress(
   auto sublen = div(data_len, pardeg);
 
   auto update_header = [&]() {
-    header.x = len3.x;
-    header.y = len3.y;
-    header.z = len3.z;
+    header.x = len3.x, header.y = len3.y, header.z = len3.z,
     header.w = 1;  // placeholder
-    header.radius = radius;
+    header.radius = radius, header.eb = eb;
     header.vle_pardeg = pardeg;
-    header.eb = eb;
     header.splen = splen;
     // header.byte_vle = use_fallback_codec ? 8 : 4;
   };
 
   /******************************************************************************/
 
-  psz_comp_l23<T, E, FP>(
-      in, len3, eb, radius, mem->ectrl_lrz(), mem->outlier_space(), &time_pred,
+  // Below is substituted with prediction+compaction
+  /*
+    psz_comp_l23<T, E>(
+        in, len3, eb, radius, mem->ectrl_lrz(), mem->outlier_space(),
+    &time_pred, stream);
+  */
+  psz_comp_l23r<T, E>(
+      in, len3, eb, radius, mem->ectrl_lrz(), (void*)mem->compact, &time_pred,
       stream);
-  psz::histogram<psz_policy::CUDA, E>(
+  psz::histogram<CUDA, E>(
+      mem->ectrl_lrz(), len, mem->hist(), booklen, &time_hist, stream);
+  psz::histsp<CUDA, E>(
       mem->ectrl_lrz(), len, mem->hist(), booklen, &time_hist, stream);
   codec->build_codebook(mem->hist(), booklen, stream);
   codec->encode(mem->ectrl_lrz(), len, &d_codec_out, &codec_outlen, stream);
-  psz::spv_gather<T, M>(
-      mem->outlier_space(), len, mem->outlier_val(), mem->outlier_idx(),
-      &splen, &time_sp, stream);
+  /*
+     psz::spv_gather<T, M>(
+         mem->outlier_space(), len, mem->outlier_val(), mem->outlier_idx(),
+         &splen, &time_sp, stream);
+  */
 
-  /* debug */ CHECK_CUDA(cudaStreamSynchronize(stream));
+  mem->compact->make_host_accessible(stream);
+
+  splen = mem->compact->num_outliers();
+
+  // /* debug */ CHECK_CUDA(cudaStreamSynchronize(stream));
 
   /******************************************************************************/
 
   update_header();
 
   merge_subfiles(
-      d_codec_out, codec_outlen, mem->outlier_val(), mem->outlier_idx(), splen,
-      stream);
+      d_codec_out, codec_outlen, mem->compact_val(), mem->compact_idx(),
+      mem->compact->num_outliers(), stream);
 
   // output
   outlen = psz_utils::filesize(&header);