ldpctest memory allocations

memory allocations for CUDA targets (harq->b) for gNB transmitter

CMakeLists.txt

@@ -188,6 +188,10 @@ if(CMAKE_C_COMPILER_ID STREQUAL "Clang" OR CMAKE_CXX_COMPILER_ID STREQUAL "Clang
   set(commonOpts "${commonOpts} -Wno-unused-command-line-argument")
 endif()
 
+if (ENABLE_CUDA)
+  set(commonOpts "${commonOpts} -DENABLE_CUDA")
+endif()
+
 set(CMAKE_C_FLAGS
   "${C_FLAGS_PROCESSOR} ${commonOpts} -std=gnu11 -funroll-loops ${CMAKE_C_FLAGS}")
 set(CMAKE_CXX_FLAGS
@@ -987,6 +991,8 @@ set(PHY_SRC_UE
   ${OPENAIR1_DIR}/PHY/NR_TRANSPORT/srs_rx.c
   ${OPENAIR1_DIR}/PHY/NR_TRANSPORT/nr_uci_tools_common.c
   )
+  set(PHY_NR_SRC_CU
+     )
   set(PHY_NR_UE_SRC
   ${OPENAIR1_DIR}/PHY/INIT/nr_parms.c
   ${OPENAIR1_DIR}/PHY/MODULATION/nr_modulation.c
@@ -1064,7 +1070,17 @@ target_include_directories(PHY_UE PRIVATE ${blas_INCLUDE_DIRS} ${lapacke_INCLUDE
 add_library(PHY_NR_COMMON ${PHY_NR_SRC_COMMON})
 target_link_libraries(PHY_NR_COMMON PUBLIC UTIL)
 
-add_library(PHY_NR ${PHY_NR_SRC})
+if (ENABLE_CUDA)
+  find_package(CUDA REQUIRED)
+  if (NOT CUDA_FOUND)
+    message(FATAL_ERROR "no CUDA found")
+  endif()
+  set(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS};-arch=sm_87;")
+  set(CUDA_VERBOS_BUILD_ON)
+  cuda_add_library(PHY_NR ${PHY_NR_SRC} ${PHY_NR_SRC_CU})
+else()
+  add_library(PHY_NR ${PHY_NR_SRC})
+endif()
 target_link_libraries(PHY_NR nr_phy_common nr_common nr_fapi_p5)
 
 add_library(PHY_NR_NO_AVX_256 ${PHY_NR_SRC})
@@ -1980,13 +1996,21 @@ target_link_libraries(smallblocktest PRIVATE
 )
 
 
+#if (ENABLE_CUDA)
+#  find_package(CUDA REQUIRED)
+#  if (NOT CUDA_FOUND)
+#    message(FATAL_ERROR "no CUDA found")
+#  endif()
+#  cuda_add_library(ldpctest)
+#else()
+#endif()
 add_executable(ldpctest
   ${OPENAIR1_DIR}/PHY/CODING/TESTBENCH/ldpctest.c
   ${OPENAIR1_DIR}/PHY/CODING/nrLDPC_load.c
 )
 target_link_libraries(ldpctest PRIVATE
   -Wl,--start-group UTIL SIMU PHY_COMMON PHY_NR_COMMON -Wl,--end-group
-  m pthread dl shlib_loader ${T_LIB}
+  m pthread dl shlib_loader ${T_LIB} ${CUDA_LIBRARIES}
   # link 'check_crc' to make it resolved in the LDPC coding libraries
   # 'check_crc' is not used in ldpctest so it is not linked in the executable by default
   # --whole-archive links 'check_crc' in the executable even though it is note used, see 'man ld'

cmake_targets/build_oai

@@ -135,6 +135,8 @@ Options:
   CC=/usr/bin/clang CXX=/usr/bin/clang++ ./build_oai ... --sanitize-memory
 --sanitize-thread | -fsanitize=thread
   Enable the thread sanitizer on all targets
+--use-cuda
+  Tell build to use NVCC for appropriate .cu files and cuda_runtime functions  
 -h | --help
    Print this help"
 }
@@ -386,6 +388,9 @@ function main() {
         --sanitize-thread | -fsanitize=thread)
             CMAKE_CMD="$CMAKE_CMD -DSANITIZE_THREAD=True"
             shift;;
+	--use-cuda)
+	    CMAKE_CMD="$CMAKE_CMD -DENABLE_CUDA=True"
+	    shift;;
        --trace-asn1c-enc-dec)
             CMAKE_CMD="$CMAKE_CMD -DTRACE_ASN1C_ENC_DEC=ON"
             echo_info "Enabling asn1c internal traces via OAI logging system"

executables/nr-softmodem-common.h

@@ -96,6 +96,7 @@
 #define CONFIG_HLP_TELN          "Start embedded telnet server \n"
 #define CONFIG_HLP_SNR           "Set average SNR in dB (for --siml1 option)\n"
 #define CONFIG_HLP_NOS1          "Disable s1 interface\n"
+#define CONFIG_HLP_USEGPU        "Use GPU for processing offload"
 
 /*--------------------------------------------------------------------------------------------------------------------------------*/
 /*                                            command line parameters for LOG utility                                             */

executables/nr-softmodem.c

@@ -103,6 +103,7 @@ unsigned int mmapped_dma=0;
 uint64_t downlink_frequency[MAX_NUM_CCs][4];
 int32_t uplink_frequency_offset[MAX_NUM_CCs][4];
 char *uecap_file;
+uint32_t use_gpu;
 
 runmode_t mode = normal_txrx;
 

executables/nr-softmodem.h

@@ -27,6 +27,7 @@
   {"U" ,                    CONFIG_HLP_ULBM_PHYTEST,   0,                .u64ptr=&ulsch_slot_bitmap,          .defintval=0,                     TYPE_UINT64, 0},        \
   {"usrp-tx-thread-config", CONFIG_HLP_USRP_THREAD,    0,                .iptr=&usrp_tx_thread,               .defstrval=0,                     TYPE_INT,    0},        \
   {"uecap_file",            CONFIG_HLP_UECAP_FILE,     0,                .strptr=&uecap_file,                 .defstrval="./uecap_ports1.xml",  TYPE_STRING, 0},        \
+  {"use_gpu",            CONFIG_HLP_USE_GPU,     0,                .uptr=&use_gpu,                 .defintval=0,  TYPE_UINT, 0},        \
 }
 // clang-format on
 
@@ -39,7 +40,7 @@ extern uint32_t target_ul_bw;
 extern uint64_t dlsch_slot_bitmap;
 extern uint64_t ulsch_slot_bitmap;
 extern char *uecap_file;
-
+extern uint32_t use_gpu;
 // In nr-gnb.c
 extern void init_gNB();
 extern void stop_gNB(int);

openair1/PHY/CODING/CMakeLists.txt

@@ -13,6 +13,7 @@ target_link_libraries(coding PRIVATE log_headers)
 add_library(ldpc_orig MODULE
             nrLDPC_decoder/nrLDPC_decoder.c
             nrLDPC_encoder/ldpc_encoder.c
+            nrLDPC_encoder/ldpc_encoder32.c
 )
 set_target_properties(ldpc_orig PROPERTIES LIBRARY_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR})
 target_link_libraries(ldpc_orig PRIVATE ldpc_segment ldpc_gen_HEADERS)
@@ -20,6 +21,7 @@ target_link_libraries(ldpc_orig PRIVATE ldpc_segment ldpc_gen_HEADERS)
 add_library(ldpc MODULE
             nrLDPC_decoder/nrLDPC_decoder.c
             nrLDPC_encoder/ldpc_encoder_optim8segmulti.c
+            nrLDPC_encoder/ldpc_encoder32.c
 )
 set_target_properties(ldpc PROPERTIES LIBRARY_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR})
 target_link_libraries(ldpc PRIVATE ldpc_segment ldpc_gen_HEADERS)
@@ -42,22 +44,28 @@ if (ENABLE_LDPC_CUDA)
   if (NOT CUDA_FOUND)
     message(FATAL_ERROR "no CUDA found")
   endif()
-  SET(CUDA_NVCC_FLAG "${CUDA_NVCC_FLAGS};-arch=sm_60;")
+  SET(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS};-arch=sm_87;")
   SET(CUDA_VERBOSE_BUILD ON)
   cuda_add_library(ldpc_cuda MODULE
-                   nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu
-                   nrLDPC_encoder/ldpc_encoder_optim8segmulti.c
+	  #                   nrLDPC_decoder_LYC/nrLDPC_decoder_LYC.cu
+                   nrLDPC_decoder/nrLDPC_decoder.c
+		   nrLDPC_encoder/ldpc_encoder_optim8segmulti.c
+		   nrLDPC_encoder/ldpc_encoder_cuda32.c
+		   nrLDPC_encoder/ldpc_BG1_Zc384_32bit.cu
+		   nrLDPC_encoder/ldpc_input.cu
                    # The slot coding layer cannot be linked with
                    # target_link_libraries like above
                    # because of cuda_add_library
                    # which already uses target_link_libraries
                    nrLDPC_coding/nrLDPC_coding_segment/nrLDPC_coding_segment_decoder.c
+                   nrLDPC_coding/nrLDPC_coding_segment/nrLDPC_coding_segment_encoder_cuda.c
                    nrLDPC_coding/nrLDPC_coding_segment/nrLDPC_coding_segment_encoder.c
                    nrLDPC_coding/nrLDPC_coding_segment/nr_rate_matching.c
   )
+  target_link_libraries(ldpc_cuda ldpc_gen_HEADERS ${T_LIB})
   set_target_properties(ldpc_cuda PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
   set_target_properties(ldpc_cuda PROPERTIES LIBRARY_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR})
-
+  target_compile_definitions(ldpc_cuda PRIVATE ENABLE_CUDA)
   add_dependencies(ldpctest ldpc_cuda)
   add_dependencies(nr-softmodem ldpc_cuda)
   add_dependencies(nr-uesoftmodem ldpc_cuda)

openair1/PHY/CODING/TESTBENCH/ldpctest.c

@@ -34,6 +34,9 @@
 #include "openair1/PHY/defs_nr_common.h"
 #include "coding_unitary_defs.h"
 #include "common/utils/LOG/log.h"
+#ifdef ENABLE_CUDA
+#include <cuda_runtime.h>
+#endif
 
 #define MAX_BLOCK_LENGTH 8448
 
@@ -104,14 +107,16 @@ one_measurement_t test_ldpc(short max_iterations,
                             unsigned char qbits,
                             short Kprime,
                             unsigned int ntrials,
-                            int n_segments)
+                            int n_segments,
+	                    int gen_code,
+			    int use32bit)
 {
   one_measurement_t ret = {0};
   reset_meas(&ret.time_optim);
   reset_meas(&ret.time_decoder);
   // clock initiate
   // time_stats_t time,time_optim,tinput,tprep,tparity,toutput, time_decoder;
-  time_stats_t time, tinput, tprep, tparity, toutput;
+  time_stats_t time, tinput, tinput_memcpy, tprep, tparity, toutput;
   double n_iter_mean = 0;
   double n_iter_std = 0;
   int n_iter_max = 0;
@@ -119,14 +124,22 @@ one_measurement_t test_ldpc(short max_iterations,
   double sigma;
   sigma = 1.0 / sqrt(2 * SNR);
   cpu_meas_enabled = 1;
-  uint8_t *test_input[MAX_NUM_NR_DLSCH_SEGMENTS_PER_LAYER * NR_MAX_NB_LAYERS];
-  uint8_t estimated_output[MAX_NUM_DLSCH_SEGMENTS][Kprime];
+#ifdef ENABLE_CUDA
+  uint8_t **test_input,*test_input_p;
+#else
+  uint8_t *test_input[n_segments * NR_MAX_NB_LAYERS];
+#endif
+  uint8_t *channel_input[n_segments];
+  uint8_t estimated_output[n_segments][Kprime];
   memset(estimated_output, 0, sizeof(estimated_output));
-  uint8_t *channel_input[MAX_NUM_DLSCH_SEGMENTS];
   uint8_t *channel_input_optim;
+
   // double channel_output[68 * 384];
-  double modulated_input[MAX_NUM_DLSCH_SEGMENTS][68 * 384] = {0};
-  int8_t channel_output_fixed[MAX_NUM_DLSCH_SEGMENTS][68 * 384] = {0};
+  double modulated_input[n_segments][68 * 384];
+  memset(modulated_input,0,sizeof(modulated_input));
+  int8_t channel_output_fixed[n_segments][68 * 384];
+  memset(channel_output_fixed,0,sizeof(channel_output_fixed));
+
   short BG = 0, nrows = 0; //,ncols;
   int i1, Kb = 0;
   int R_ind = 0;
@@ -134,7 +147,7 @@ one_measurement_t test_ldpc(short max_iterations,
   int code_rate_vec[8] = {15, 13, 25, 12, 23, 34, 56, 89};
   // double code_rate_actual_vec[8] = {0.2, 0.33333, 0.4, 0.5, 0.66667, 0.73333, 0.81481, 0.88};
 
-  t_nrLDPC_dec_params decParams[MAX_NUM_DLSCH_SEGMENTS] = {0};
+  t_nrLDPC_dec_params decParams[n_segments];
 
   t_nrLDPC_time_stats decoder_profiler = {0};
 
@@ -142,6 +155,7 @@ one_measurement_t test_ldpc(short max_iterations,
 
   reset_meas(&time);
   reset_meas(&tinput);
+  reset_meas(&tinput_memcpy);
   reset_meas(&tprep);
   reset_meas(&tparity);
   reset_meas(&toutput);
@@ -259,18 +273,29 @@ one_measurement_t test_ldpc(short max_iterations,
   printf("To: %d\n", (Kb + nrows - no_punctured_columns) * Zc - removed_bit);
   printf("number of undecoded bits: %d\n", (Kb + nrows - no_punctured_columns - 2) * Zc - removed_bit);
 
+//  ldpc_toCompare.LDPCinit();
   // generate input block
-  for (int j = 0; j < MAX_NUM_DLSCH_SEGMENTS; j++) {
+#ifdef ENABLE_CUDA
+  cudaHostAlloc((void**)&test_input_p,n_segments*sizeof(uint8_t*),cudaHostAllocMapped);
+  test_input=(uint8_t **)test_input_p;
+  printf("test input %p\n",test_input);
+#endif
+  for (int j = 0; j < n_segments; j++) {
+#ifdef ENABLE_CUDA
+    cudaHostAlloc((void**)&test_input[j],((K + 7) & ~7) / 8,cudaHostAllocMapped);
+    printf("test input[%d] %p\n",j,test_input[j]);
+#else
     test_input[j] = malloc16(((K + 7) & ~7) / 8);
     memset(test_input[j], 0, ((K + 7) & ~7) / 8);
+#endif
     channel_input[j] = malloc16(68 * 384);
     memset(channel_input[j], 0, 68 * 384);
   }
-  channel_input_optim = malloc16(68 * 384);
-  memset(channel_input_optim, 0, 68 * 384);
+  channel_input_optim = malloc16(68 * 384 * sizeof(uint32_t));
+  if (use32bit ==0) memset(channel_input_optim, 0, 68 * 384 * sizeof(uint32_t));
 
   // Fill input segments with random values
-  for (int j = 0; j < MAX_NUM_DLSCH_SEGMENTS; j++) {
+  for (int j = 0; j < n_segments; j++) {
     int i = 0;
     for (i = 0; i < ((Kprime + 7) & ~7) / 8; i++)
       test_input[j][i] = (uint8_t)rand();
@@ -282,13 +307,17 @@ one_measurement_t test_ldpc(short max_iterations,
   }
 
   encoder_implemparams_t impp = {.Zc = Zc, .Kb = Kb, .BG = BG, .K = K};
-  impp.gen_code = 1;
-
+  impp.gen_code = gen_code;
+  impp.tparity = &tparity;
+  impp.tinput = &tinput;
+  impp.tinput_memcpy = &tinput_memcpy;
+  impp.toutput = &toutput;
   if (ntrials == 0)
     ldpc_orig.LDPCencoder(test_input, channel_input[0], &impp);
   impp.gen_code = 0;
   decode_abort_t dec_abort;
   init_abort(&dec_abort);
+  uint32_t **output32;
   for (int trial = 0; trial < ntrials; trial++) {
     unsigned int segment_bler = 0;
     //// encoder
@@ -301,26 +330,27 @@ one_measurement_t test_ldpc(short max_iterations,
     impp.n_segments = n_segments;
     start_meas(&ret.time_optim);
     impp.first_seg = 0;
-    ldpc_toCompare.LDPCencoder(test_input, channel_input_optim, &impp);
+    if (use32bit==0) ldpc_toCompare.LDPCencoder(test_input, channel_input_optim, &impp);
+    else output32=ldpc_toCompare.LDPCencoder32(test_input, &impp);
     stop_meas(&ret.time_optim);
 
     if (ntrials == 1)
       for (int j = 0; j < n_segments; j++)
         for (int i = 0; i < K + (nrows - no_punctured_columns) * Zc - removed_bit; i++) {
-          if (channel_input[j][i] != ((channel_input_optim[i] >> j) & 0x1)) {
-            printf("differ in seg %d pos %d (%u,%u)\n", j, i, channel_input[j][i], (channel_input_optim[i] >> j) & 0x1);
-            return ret;
-          }
+          if (((use32bit == 0) && (channel_input[j][i] != ((channel_input_optim[i] >> j) & 0x1))) ||                 ((use32bit == 1) && (channel_input[j][i] != (((output32[0][i] >> j) & 0x1))))) {
+               printf("differ in seg %d pos %d (%u,%u)\n", j, i, channel_input[j][i], (((uint32_t*)channel_input_optim)[i] >> j) & 0x1);
+               return ret;
+            }
         }
-
+    int bit;
     for (int j = 0; j < n_segments; j++) {
       for (int i = 2 * Zc; i < (Kb + nrows - no_punctured_columns) * Zc - removed_bit; i++) {
 #ifdef DEBUG_CODER
         if ((i & 0xf) == 0)
           printf("\ne %d..%d:    ", i, i + 15);
 #endif
-
-        if (((channel_input_optim[i - 2 * Zc] >> j) & 0x1) == 0)
+        bit = (use32bit==0) ? ((channel_input_optim[i - 2 * Zc] >> j) & 0x1) : ((output32[0][i - 2 * Zc] >> j) & 0x1);
+        if (bit == 0)
           modulated_input[j][i] = 1.0; /// sqrt(2);  //QPSK
         else
           modulated_input[j][i] = -1.0; /// sqrt(2);
@@ -345,7 +375,6 @@ one_measurement_t test_ldpc(short max_iterations,
       decParams[j].numMaxIter = max_iterations;
       decParams[j].outMode = nrLDPC_outMode_BIT;
       decParams[j].Kprime = Kprime;
-      ldpc_toCompare.LDPCinit();
     }
     for (int j = 0; j < n_segments; j++) {
       start_meas(&ret.time_decoder);
@@ -388,8 +417,12 @@ one_measurement_t test_ldpc(short max_iterations,
 
   ret.errors_bit_uncoded = ret.errors_bit_uncoded / (double)((Kb + nrows - no_punctured_columns - 2) * Zc - removed_bit);
 
-  for (int j = 0; j < MAX_NUM_DLSCH_SEGMENTS; j++) {
+  for (int j = 0; j < n_segments; j++) {
+#ifdef ENABLE_CUDA
+    cudaFree(test_input[j]);	  
+#else
     free(test_input[j]);
+#endif
     free(channel_input[j]);
   }
   free(channel_input_optim);
@@ -397,6 +430,7 @@ one_measurement_t test_ldpc(short max_iterations,
   print_meas(&time, "ldpc_encoder", NULL, NULL);
   print_meas(&ret.time_optim, "ldpc_encoder_optim", NULL, NULL);
   print_meas(&tinput, "ldpc_encoder_optim(input)", NULL, NULL);
+  print_meas(&tinput_memcpy, "ldpc_encoder_optim(input memcpy)", NULL, NULL);
   print_meas(&tprep, "ldpc_encoder_optim(prep)", NULL, NULL);
   print_meas(&tparity, "ldpc_encoder_optim(parity)", NULL, NULL);
   print_meas(&toutput, "ldpc_encoder_optim(output)", NULL, NULL);
@@ -437,19 +471,20 @@ int main(int argc, char *argv[])
 
   int n_trials = 1;
   double SNR_step = 0.1;
-
+  int gen_code = 1;
   randominit(0);
   int test_uncoded = 0;
   n_iter_stats_t dec_iter[400] = {0};
 
   short BG = 0, Zc;
+  int use32bit=0;
 
   if ((uniqCfg = load_configmodule(argc, argv, CONFIG_ENABLECMDLINEONLY)) == 0) {
     exit_fun("[LDPCTEST] Error, configuration module init failed\n");
   }
   logInit();
 
-  while ((c = getopt(argc, argv, "--:O:q:r:s:S:l:G:n:d:i:t:u:hv:")) != -1) {
+  while ((c = getopt(argc, argv, "--:O:q:r:s:S:l:Gn:d:i:t:u:hv:g:")) != -1) {
     /* ignore long options starting with '--', option '-O' and their arguments that are handled by configmodule */
     /* with this opstring getopt returns 1 for non-option arguments, refer to 'man 3 getopt' */
     if (c == 1 || c == '-' || c == 'O')
@@ -475,6 +510,7 @@ int main(int argc, char *argv[])
 
       case 'G':
         ldpc_version = "_cuda";
+	use32bit=1;
         break;
 
       case 'n':
@@ -503,6 +539,11 @@ int main(int argc, char *argv[])
       case 'v':
         ldpc_version = strdup(optarg);
         break;
+      case 'g':
+	gen_code = atoi(optarg);
+	AssertFatal(gen_code <= 4, "gen_code %d is not allowed\n",gen_code);
+        n_trials=0;	
+	break;
       case 'h':
       default:
         printf("CURRENTLY SUPPORTED CODE RATES: \n");
@@ -559,7 +600,9 @@ int main(int argc, char *argv[])
                                       qbits,
                                       Kprime, // block length bytes
                                       n_trials,
-                                      n_segments);
+                                      n_segments,
+				      gen_code,
+				      use32bit);
 
     decoded_errors[i] = res.errors;
     dec_iter[i] = res.dec_iter;

openair1/PHY/CODING/nrLDPC_coding/nrLDPC_coding_interface.h

@@ -152,6 +152,7 @@ typedef struct nrLDPC_segment_encoding_parameters_s{
   time_stats_t ts_interleave;
   time_stats_t ts_rate_match;
   time_stats_t ts_ldpc_encode;
+  time_stats_t ts_output;
 } nrLDPC_segment_encoding_parameters_t;
 
 /**
@@ -201,6 +202,7 @@ typedef struct nrLDPC_TB_encoding_parameters_s{
   uint32_t C;
   nrLDPC_segment_encoding_parameters_t *segments;
   unsigned char *output;
+  uint8_t **c_dev;
 } nrLDPC_TB_encoding_parameters_t;
 
 /**
@@ -226,10 +228,13 @@ typedef struct nrLDPC_slot_encoding_parameters_s{
   int nb_TBs;
   tpool_t *threadPool;
   time_stats_t *tinput;
+  time_stats_t *tinput_memcpy;
   time_stats_t *tprep;
   time_stats_t *tparity;
   time_stats_t *toutput;
+  time_stats_t *tconcat;
   nrLDPC_TB_encoding_parameters_t *TBs;
+  uint32_t use_gpu;
 } nrLDPC_slot_encoding_parameters_t;
 
 typedef int32_t(nrLDPC_coding_init_t)(void);

openair1/PHY/CODING/nrLDPC_coding/nrLDPC_coding_segment/nrLDPC_coding_segment_decoder.c

@@ -281,6 +281,11 @@ int nrLDPC_prepare_TB_decoding(nrLDPC_slot_decoding_parameters_t *nrLDPC_slot_de
 
 int32_t nrLDPC_coding_init(void)
 {
+  LOG_I(NR_PHY,"Initializing coding library\n");
+#ifdef ENABLE_CUDA
+  LOG_I(NR_PHY,"Calling cuda_support_init()\n");
+#endif
+  cuda_support_init();	
   return 0;
 }
 

openair1/PHY/CODING/nrLDPC_coding/nrLDPC_coding_segment/nrLDPC_coding_segment_encoder.c

@@ -367,9 +367,11 @@ static int nrLDPC_launch_TB_encoding(nrLDPC_slot_encoding_parameters_t *nrLDPC_s
   encoder_implemparams_t common_segment_params = {
     .n_segments = nrLDPC_TB_encoding_parameters->C,
     .tinput = nrLDPC_slot_encoding_parameters->tinput,
+    .tinput_memcpy = nrLDPC_slot_encoding_parameters->tinput_memcpy,
     .tprep = nrLDPC_slot_encoding_parameters->tprep,
     .tparity = nrLDPC_slot_encoding_parameters->tparity,
     .toutput = nrLDPC_slot_encoding_parameters->toutput,
+    .tconcat = nrLDPC_slot_encoding_parameters->tconcat,
     .Kb = nrLDPC_TB_encoding_parameters->Kb,
     .Zc = nrLDPC_TB_encoding_parameters->Z,
     .BG = nrLDPC_TB_encoding_parameters->BG,
@@ -406,16 +408,24 @@ int nrLDPC_coding_encoder(nrLDPC_slot_encoding_parameters_t *nrLDPC_slot_encodin
   for (int dlsch_id = 0; dlsch_id < nrLDPC_slot_encoding_parameters->nb_TBs; dlsch_id++) {
     // Compute number of tasks to encode TB
     nrLDPC_TB_encoding_parameters_t *nrLDPC_TB_encoding_parameters = &nrLDPC_slot_encoding_parameters->TBs[dlsch_id];
-    size_t n_seg = (nrLDPC_TB_encoding_parameters->C / 8 + ((nrLDPC_TB_encoding_parameters->C & 7) == 0 ? 0 : 1));
-    nbTasks += n_seg;
-
-    // Search for maximum E for sizing encoder output f and f2
-    for (int seg_id = 0; seg_id < nrLDPC_TB_encoding_parameters->C; seg_id++) {
-      uint32_t E = nrLDPC_TB_encoding_parameters->segments[seg_id].E;
-      Emax = E > Emax ? E : Emax;
+#if defined(ENABLE_CUDA)     
+    if (nrLDPC_slot_encoding_parameters->use_gpu && nrLDPC_TB_encoding_parameters->C > 8 && nrLDPC_TB_encoding_parameters->Z == 384) {
+        nrLDPC_coding_encoder32(nrLDPC_slot_encoding_parameters,nrLDPC_TB_encoding_parameters);
     }
-  }
+    else {
+#endif
+      size_t n_seg = (nrLDPC_TB_encoding_parameters->C / 8 + ((nrLDPC_TB_encoding_parameters->C & 7) == 0 ? 0 : 1));
+      nbTasks += n_seg;
 
+      // Search for maximum E for sizing encoder output f and f2
+      for (int seg_id = 0; seg_id < nrLDPC_TB_encoding_parameters->C; seg_id++) {
+        uint32_t E = nrLDPC_TB_encoding_parameters->segments[seg_id].E;
+        Emax = E > Emax ? E : Emax;
+      }
+#if defined(ENABLE_CUDA) 
+    }
+#endif
+  }
   // Create f and f2 to old encoding tasks outputs
   uint32_t Emax_ceil_mod = ceil_mod(Emax, 64);
   uint8_t f[nbTasks][Emax_ceil_mod] __attribute__((aligned(64)));
@@ -432,6 +442,10 @@ int nrLDPC_coding_encoder(nrLDPC_slot_encoding_parameters_t *nrLDPC_slot_encodin
     // For easier indexing we store the pointers to sub arrays of f and f2 in pointer arrays
     // Then a function to which we pass the pointer arrays can directly use f_2d[j] ans f2_2d[j]
     nrLDPC_TB_encoding_parameters_t *nrLDPC_TB_encoding_parameters = &nrLDPC_slot_encoding_parameters->TBs[dlsch_id];
+#if defined(ENABLE_CUDA) 
+    if (nrLDPC_slot_encoding_parameters->use_gpu && nrLDPC_TB_encoding_parameters->C > 8 && nrLDPC_TB_encoding_parameters->Z == 384) 
+      continue;
+#endif
     size_t n_seg = (nrLDPC_TB_encoding_parameters->C / 8 + ((nrLDPC_TB_encoding_parameters->C & 7) == 0 ? 0 : 1));
     uint8_t *f_2d[n_seg];
     uint8_t *f2_2d[n_seg];
@@ -449,16 +463,21 @@ int nrLDPC_coding_encoder(nrLDPC_slot_encoding_parameters_t *nrLDPC_slot_encodin
     completed_many_task_ans(&ans, nbTasks - nbEncode);
   }
   // Execute thread pool tasks
-  join_task_ans(&ans);
+  if (nbTasks > 0)  join_task_ans(&ans);
 
   // Write output
+  time_stats_t *tconcat = nrLDPC_slot_encoding_parameters->tconcat;
+  if(tconcat != NULL) start_meas(tconcat);
   nbTasks = 0;
   for (int dlsch_id = 0; dlsch_id < nrLDPC_slot_encoding_parameters->nb_TBs; dlsch_id++) {
     nrLDPC_TB_encoding_parameters_t *nrLDPC_TB_encoding_parameters = &nrLDPC_slot_encoding_parameters->TBs[dlsch_id];
     uint32_t C = nrLDPC_TB_encoding_parameters->C;
+#if defined(ENABLE_CUDA)
+    if (nrLDPC_slot_encoding_parameters->use_gpu && C > 8 && nrLDPC_TB_encoding_parameters->Z == 384) 
+      continue;
+#endif
     size_t n_seg = (C / 8 + ((C & 7) == 0 ? 0 : 1));
 
-    time_stats_t *toutput = nrLDPC_slot_encoding_parameters->toutput;
 
     for (int j = 0; j < n_seg; j++) {
       unsigned int macro_segment = j * 8;
@@ -478,7 +497,6 @@ int nrLDPC_coding_encoder(nrLDPC_slot_encoding_parameters_t *nrLDPC_slot_encodin
         }
       }
 
-      if(toutput != NULL) start_meas(toutput);
  
       uint32_t Eoffset=0;
       for (int s=0; s<macro_segment; s++)
@@ -494,10 +512,10 @@ int nrLDPC_coding_encoder(nrLDPC_slot_encoding_parameters_t *nrLDPC_slot_encodin
                         nrLDPC_TB_encoding_parameters->output,
                         Eoffset);
 
-      if(toutput != NULL) stop_meas(toutput);
     }
     nbTasks += n_seg;
   }
+  if(tconcat != NULL) stop_meas(tconcat);
 
   return 0;
 }

openair1/PHY/CODING/nrLDPC_coding/nrLDPC_coding_segment/nrLDPC_coding_segment_encoder_cuda.c

@@ -0,0 +1,688 @@
+/*
+ * Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
+ * contributor license agreements.  See the NOTICE file distributed with
+ * this work for additional information regarding copyright ownership.
+ * The OpenAirInterface Software Alliance licenses this file to You under
+ * the OAI Public License, Version 1.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.openairinterface.org/?page_id=698
+ *
+ * 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.
+ *-------------------------------------------------------------------------------
+ * For more information about the OpenAirInterface (OAI) Software Alliance:
+ *      contact@openairinterface.org
+ */
+
+/*! \file PHY/CODING/nrLDPC_coding/nrLDPC_coding_segment/nrLDPC_coding_segment_encoder.c
+ * \brief Top-level routines for implementing LDPC encoding of transport channels
+ */
+
+#include "nr_rate_matching.h"
+#include "PHY/defs_gNB.h"
+#include "PHY/CODING/coding_extern.h"
+#include "PHY/CODING/coding_defs.h"
+#include "PHY/CODING/lte_interleaver_inline.h"
+#include "PHY/CODING/nrLDPC_coding/nrLDPC_coding_interface.h"
+#include "PHY/CODING/nrLDPC_extern.h"
+#include "PHY/NR_TRANSPORT/nr_transport_proto.h"
+#include "PHY/NR_TRANSPORT/nr_transport_common_proto.h"
+#include "PHY/NR_TRANSPORT/nr_dlsch.h"
+#include "SCHED_NR/sched_nr.h"
+#include "common/utils/LOG/vcd_signal_dumper.h"
+#include "common/utils/LOG/log.h"
+#include "common/utils/nr/nr_common.h"
+#include <openair2/UTIL/OPT/opt.h>
+
+#include <syscall.h>
+
+//#define DEBUG_LDPC_ENCODING
+//#define DEBUG_LDPC_ENCODING_FREE 1
+
+extern uint32_t **d_host;
+
+static void unpack_output(uint32_t *f,
+                         uint32_t E,
+                         uint32_t *f2,
+                         uint32_t E2,
+			 uint32_t E2_first_segment32,
+                         uint32_t E2_first_segment,
+                         uint32_t nb_segments,
+                         uint8_t *output) {
+
+  uint32_t s;
+ // int s0;
+  uint32_t *fp;
+  int foffset;
+  uint32_t *output_p = (uint32_t *)output;
+  //printf("unpack: E %d, E2 %d, E2_first_segment %d, E2_first_segment32 %d, nb_segments %d\n",E,E2,E2_first_segment,E2_first_segment32,nb_segments);
+
+  uint32_t bit_index = 0;
+#ifdef __AVX2__
+  simde__m256i shift0=simde_mm256_set_epi32(7,6,5,4,3,2,1,0);
+  simde__m256i shift1=simde_mm256_set_epi32(15,14,13,12,11,10,9,8);
+  simde__m256i shift2=simde_mm256_set_epi32(23,22,21,20,19,18,17,16);
+  simde__m256i shift3=simde_mm256_set_epi32(31,30,29,28,27,26,25,24);
+  simde__m256i vmask0=simde_mm256_set_epi32(0x80,0x40,0x20,0x10,0x8,0x4,0x2,0x1);
+  simde__m256i vmask1=simde_mm256_set_epi32(0x8000,0x4000,0x2000,0x1000,0x800,0x400,0x200,0x100);
+  simde__m256i vmask2=simde_mm256_set_epi32(0x800000,0x400000,0x200000,0x100000,0x80000,0x40000,0x20000,0x10000);
+  simde__m256i vmask3=simde_mm256_set_epi32(0x80000000,0x40000000,0x20000000,0x10000000,0x8000000,0x4000000,0x2000000,0x1000000);
+  int s2=0;
+  for (s = 0; s < E2_first_segment ; s++) { 
+    s2 = s&31;	  
+    foffset = (s>>5)*E;
+    fp = f+foffset;
+    int i;
+    if ((bit_index&31) == 0 ) {
+      for (i = 0; i < (E>>5)<<5; i+=32) {
+	simde__m256i f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i],s2);    
+  	simde__m256i cshift = simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift0),vmask0);
+	f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i+8],s2);    
+  	cshift = simde_mm256_or_si256(simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift1),vmask1),cshift);
+	f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i+16],s2);    
+  	cshift = simde_mm256_or_si256(simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift2),vmask2),cshift);
+	f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i+24],s2);    
+  	cshift = simde_mm256_or_si256(simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift3),vmask3),cshift);
+	*(output_p + (bit_index>>5))     = simde_mm256_extract_epi32(cshift,0) | 
+	                                   simde_mm256_extract_epi32(cshift,1) | 
+	                                   simde_mm256_extract_epi32(cshift,2) | 
+	                                   simde_mm256_extract_epi32(cshift,3) | 
+	                                   simde_mm256_extract_epi32(cshift,4) | 
+	                                   simde_mm256_extract_epi32(cshift,5) | 
+	                                   simde_mm256_extract_epi32(cshift,6) | 
+	                                   simde_mm256_extract_epi32(cshift,7);
+	bit_index+=32;
+      }
+      uint32_t Emod32=E&31;
+      if (Emod32 != 0) {
+	simde__m256i f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i],s2);    
+  	simde__m256i cshift = simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift0),vmask0);
+	f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i+8],s2);    
+  	cshift = simde_mm256_or_si256(simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift1),vmask1),cshift);
+	f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i+16],s2);    
+  	cshift = simde_mm256_or_si256(simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift2),vmask2),cshift);
+	f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i+24],s2);    
+  	cshift = simde_mm256_or_si256(simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift3),vmask3),cshift);
+	*(output_p + (bit_index>>5))     = (simde_mm256_extract_epi32(cshift,0) | 
+	                                   simde_mm256_extract_epi32(cshift,1) | 
+	                                   simde_mm256_extract_epi32(cshift,2) | 
+	                                   simde_mm256_extract_epi32(cshift,3) | 
+	                                   simde_mm256_extract_epi32(cshift,4) | 
+	                                   simde_mm256_extract_epi32(cshift,5) | 
+	                                   simde_mm256_extract_epi32(cshift,6) | 
+	                                   simde_mm256_extract_epi32(cshift,7))&((1<<Emod32)-1);
+        bit_index+=Emod32;
+      }
+    }
+    else {
+      for (i = 0; i < (E>>5)<<5; i+=32) {
+	simde__m256i f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i],s2);    
+  	simde__m256i cshift = simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift0),vmask0);
+	f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i+8],s2);    
+  	cshift = simde_mm256_or_si256(simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift1),vmask1),cshift);
+	f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i+16],s2);    
+  	cshift = simde_mm256_or_si256(simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift2),vmask2),cshift);
+	f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i+24],s2);    
+  	cshift = simde_mm256_or_si256(simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift3),vmask3),cshift);
+	uint32_t tmp    = simde_mm256_extract_epi32(cshift,0) | 
+	                  simde_mm256_extract_epi32(cshift,1) | 
+	                  simde_mm256_extract_epi32(cshift,2) | 
+	                  simde_mm256_extract_epi32(cshift,3) | 
+	                  simde_mm256_extract_epi32(cshift,4) | 
+	                  simde_mm256_extract_epi32(cshift,5) | 
+	                  simde_mm256_extract_epi32(cshift,6) | 
+	                  simde_mm256_extract_epi32(cshift,7);
+        *(output_p + (bit_index>>5))     |= (tmp<<(bit_index&31));
+        *(output_p + (bit_index>>5)+1)   |= (tmp>>(32-(bit_index&31)));					   
+	bit_index+=32;
+      }
+      uint32_t Emod32=E&31;
+      if (Emod32 != 0) {
+	simde__m256i f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i],s2);    
+  	simde__m256i cshift = simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift0),vmask0);
+	f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i+8],s2);    
+  	cshift = simde_mm256_or_si256(simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift1),vmask1),cshift);
+	f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i+16],s2);    
+  	cshift = simde_mm256_or_si256(simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift2),vmask2),cshift);
+	f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i+24],s2);    
+  	cshift = simde_mm256_or_si256(simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift3),vmask3),cshift);
+	uint32_t tmp = 	(simde_mm256_extract_epi32(cshift,0) | 
+	                 simde_mm256_extract_epi32(cshift,1) | 
+	                 simde_mm256_extract_epi32(cshift,2) | 
+	                 simde_mm256_extract_epi32(cshift,3) | 
+	                 simde_mm256_extract_epi32(cshift,4) | 
+	                 simde_mm256_extract_epi32(cshift,5) | 
+	                 simde_mm256_extract_epi32(cshift,6) | 
+	                 simde_mm256_extract_epi32(cshift,7))&((1<<Emod32)-1);
+	*(output_p + (bit_index>>5))     |= (tmp<<(bit_index&31));
+        *(output_p + (bit_index>>5)+1)   |= (tmp>>(32-(bit_index&31)));
+        bit_index+=Emod32;
+      }
+    }
+  }
+  for ( ; s < nb_segments ; s++){
+    s2 = s&31;	  
+    foffset = ((s>>5)-E2_first_segment32)*E2;
+    fp = f2+foffset;
+    int i;
+    if ((bit_index&31) == 0 ) {
+      for (i = 0; i < (E2>>5)<<5; i+=32) {
+	simde__m256i f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i],s2);    
+  	simde__m256i cshift = simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift0),vmask0);
+	f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i+8],s2);    
+  	cshift = simde_mm256_or_si256(simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift1),vmask1),cshift);
+	f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i+16],s2);    
+  	cshift = simde_mm256_or_si256(simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift2),vmask2),cshift);
+	f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i+24],s2);    
+  	cshift = simde_mm256_or_si256(simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift3),vmask3),cshift);
+	*(output_p + (bit_index>>5))     = simde_mm256_extract_epi32(cshift,0) | 
+	                                   simde_mm256_extract_epi32(cshift,1) | 
+	                                   simde_mm256_extract_epi32(cshift,2) | 
+	                                   simde_mm256_extract_epi32(cshift,3) | 
+	                                   simde_mm256_extract_epi32(cshift,4) | 
+	                                   simde_mm256_extract_epi32(cshift,5) | 
+	                                   simde_mm256_extract_epi32(cshift,6) | 
+	                                   simde_mm256_extract_epi32(cshift,7);
+	bit_index+=32;
+      }
+      uint32_t E2mod32=E2&31;
+      if (E2mod32 != 0) {
+	simde__m256i f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i],s2);    
+  	simde__m256i cshift = simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift0),vmask0);
+	f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i+8],s2);    
+  	cshift = simde_mm256_or_si256(simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift1),vmask1),cshift);
+	f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i+16],s2);    
+  	cshift = simde_mm256_or_si256(simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift2),vmask2),cshift);
+	f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i+24],s2);    
+  	cshift = simde_mm256_or_si256(simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift3),vmask3),cshift);
+	*(output_p + (bit_index>>5))     = (simde_mm256_extract_epi32(cshift,0) | 
+	                                   simde_mm256_extract_epi32(cshift,1) | 
+	                                   simde_mm256_extract_epi32(cshift,2) | 
+	                                   simde_mm256_extract_epi32(cshift,3) | 
+	                                   simde_mm256_extract_epi32(cshift,4) | 
+	                                   simde_mm256_extract_epi32(cshift,5) | 
+	                                   simde_mm256_extract_epi32(cshift,6) | 
+	                                   simde_mm256_extract_epi32(cshift,7))&((1<<E2mod32)-1);
+        bit_index+=E2mod32;
+      }
+    }
+    else {
+      for (i = 0; i < (E2>>5)<<5; i+=32) {
+	simde__m256i f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i],s2);    
+  	simde__m256i cshift = simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift0),vmask0);
+	f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i+8],s2);    
+  	cshift = simde_mm256_or_si256(simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift1),vmask1),cshift);
+	f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i+16],s2);    
+  	cshift = simde_mm256_or_si256(simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift2),vmask2),cshift);
+	f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i+24],s2);    
+  	cshift = simde_mm256_or_si256(simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift3),vmask3),cshift);
+	uint32_t tmp    = simde_mm256_extract_epi32(cshift,0) | 
+	                  simde_mm256_extract_epi32(cshift,1) | 
+	                  simde_mm256_extract_epi32(cshift,2) | 
+	                  simde_mm256_extract_epi32(cshift,3) | 
+	                  simde_mm256_extract_epi32(cshift,4) | 
+	                  simde_mm256_extract_epi32(cshift,5) | 
+	                  simde_mm256_extract_epi32(cshift,6) | 
+	                  simde_mm256_extract_epi32(cshift,7);
+        *(output_p + (bit_index>>5))     |= (tmp<<(bit_index&31));
+        *(output_p + (bit_index>>5)+1)   |= (tmp>>(32-(bit_index&31)));					   
+	bit_index+=32;
+      }
+      uint32_t E2mod32=E2&31;
+      if (E2mod32 != 0) {
+	simde__m256i f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i],s2);    
+  	simde__m256i cshift = simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift0),vmask0);
+	f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i+8],s2);    
+  	cshift = simde_mm256_or_si256(simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift1),vmask1),cshift);
+	f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i+16],s2);    
+  	cshift = simde_mm256_or_si256(simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift2),vmask2),cshift);
+	f256 = simde_mm256_srli_epi32(*(simde__m256i*)&fp[i+24],s2);    
+  	cshift = simde_mm256_or_si256(simde_mm256_and_si256(simde_mm256_sllv_epi32(f256,shift3),vmask3),cshift);
+	uint32_t tmp = 	(simde_mm256_extract_epi32(cshift,0) | 
+	                 simde_mm256_extract_epi32(cshift,1) | 
+	                 simde_mm256_extract_epi32(cshift,2) | 
+	                 simde_mm256_extract_epi32(cshift,3) | 
+	                 simde_mm256_extract_epi32(cshift,4) | 
+	                 simde_mm256_extract_epi32(cshift,5) | 
+	                 simde_mm256_extract_epi32(cshift,6) | 
+	                 simde_mm256_extract_epi32(cshift,7))&((1<<E2mod32)-1);
+	*(output_p + (bit_index>>5))     |= (tmp<<(bit_index&31));
+        *(output_p + (bit_index>>5)+1)   |= (tmp>>(32-(bit_index&31)));
+        bit_index+=E2mod32;
+      }
+    }
+  }
+#elif defined(__aarch64__)
+  const int32_t ucShift0[32][4] = { {0,1,2,3}, {-1,0,1,2},{-2,-1,0,1}, {-3,-2,-1,0}, {-4,-3,-2,-1}, {-5,-4,-3,-2}, {-6,-5,-4,-3}, {-7,-6,-5,-4}, {-8,-7,-6,-5}, {-9,-8,-7,-6}, {-10,-9,-8,-7}, {-11,-10,-9,-8}, {-12,-11,-10,-9}, {-13,-12,-11,-10}, {-14,-13,-12,-11}, {-15,-14,-13,-12}, {-16,-15,-14,-13}, {-17,-16,-15,-14}, {-18,-17,-16,-15}, {-19,-18,-17,-16}, {-20,-19,-18,-17}, {-21,-20,-19,-18}, {-22,-21,-20,-19}, {-23,-22,-21,-20}, {-24,-23,-22,-21}, {-25,-24,-23,-22}, {-26,-25,-24,-23}, {-27,-26,-25,-24}, {-28,-27,-26,-25}, {-29,-28,-27,-26}, {-30,-29,-28,-27}, {-31,-30,-29,-28}}; 
+
+  const int32_t ucShift1[32][4] = { {4,5,6,7}, {3,4,5,6}, {2,3,4,5}, {1,2,3,4}, {0,1,2,3}, {-1,0,1,2},{-2,-1,0,1}, {-3,-2,-1,0}, {-4,-3,-2,-1}, {-5,-4,-3,-2}, {-6,-5,-4,-3}, {-7,-6,-5,-4}, {-8,-7,-6,-5}, {-9,-8,-7,-6}, {-10,-9,-8,-7}, {-11,-10,-9,-8}, {-12,-11,-10,-9}, {-13,-12,-11,-10}, {-14,-13,-12,-11}, {-15,-14,-13,-12}, {-16,-15,-14,-13}, {-17,-16,-15,-14}, {-18,-17,-16,-15}, {-19,-18,-17,-16}, {-20,-19,-18,-17}, {-21,-20,-19,-18}, {-22,-21,-20,-19}, {-23,-22,-21,-20}, {-24,-23,-22,-21}, {-25,-24,-23,-22}, {-26,-25,-24,-23}, {-27,-26,-25,-24}}; 
+
+  const int32_t ucShift2[32][4] = { {8,9,10,11},{7,8,9,10}, {6,7,8,9}, {5,6,7,8}, {4,5,6,7}, {3,4,5,6}, {2,3,4,5}, {1,2,3,4}, {0,1,2,3}, {-1,0,1,2},{-2,-1,0,1}, {-3,-2,-1,0}, {-4,-3,-2,-1}, {-5,-4,-3,-2}, {-6,-5,-4,-3}, {-7,-6,-5,-4}, {-8,-7,-6,-5}, {-9,-8,-7,-6}, {-10,-9,-8,-7}, {-11,-10,-9,-8}, {-12,-11,-10,-9}, {-13,-12,-11,-10}, {-14,-13,-12,-11}, {-15,-14,-13,-12}, {-16,-15,-14,-13}, {-17,-16,-15,-14}, {-18,-17,-16,-15}, {-19,-18,-17,-16}, {-20,-19,-18,-17}, {-21,-20,-19,-18}, {-22,-21,-20,-19},{-23,-22,-21,-20}}; 
+
+  const int32_t ucShift3[32][4] = { {12,13,14,15}, {11,12,13,14}, {10,11,12,13}, {9,10,11,12}, {8,9,10,11},{7,8,9,10}, {6,7,8,9}, {5,6,7,8}, {4,5,6,7}, {3,4,5,6}, {2,3,4,5}, {1,2,3,4}, {0,1,2,3}, {-1,0,1,2},{-2,-1,0,1}, {-3,-2,-1,0}, {-4,-3,-2,-1}, {-5,-4,-3,-2}, {-6,-5,-4,-3}, {-7,-6,-5,-4}, {-8,-7,-6,-5}, {-9,-8,-7,-6}, {-10,-9,-8,-7}, {-11,-10,-9,-8}, {-12,-11,-10,-9}, {-13,-12,-11,-10}, {-14,-13,-12,-11}, {-15,-14,-13,-12}, {-16,-15,-14,-13}, {-17,-16,-15,-14}, {-18,-17,-16,-15}, {-19,-18,-17,-16}}; 
+
+  const int32_t ucShift4[32][4] = { {16,17,18,19}, {15,16,17,18}, {14,15,16,17}, {13,14,15,16}, {12,13,14,15}, {11,12,13,14}, {10,11,12,13}, {9,10,11,12}, {8,9,10,11},{7,8,9,10}, {6,7,8,9}, {5,6,7,8}, {4,5,6,7}, {3,4,5,6}, {2,3,4,5}, {1,2,3,4}, {0,1,2,3}, {-1,0,1,2},{-2,-1,0,1}, {-3,-2,-1,0}, {-4,-3,-2,-1}, {-5,-4,-3,-2}, {-6,-5,-4,-3}, {-7,-6,-5,-4}, {-8,-7,-6,-5}, {-9,-8,-7,-6}, {-10,-9,-8,-7}, {-11,-10,-9,-8}, {-12,-11,-10,-9}, {-13,-12,-11,-10}, {-14,-13,-12,-11}, {-15,-14,-13,-12}}; 
+
+  const int32_t ucShift5[32][4] = { {20,21,22,23}, {19,20,21,22}, {18,19,20,21}, {17,18,19,20}, {16,17,18,19}, {15,16,17,18}, {14,15,16,17}, {13,14,15,16}, {12,13,14,15}, {11,12,13,14}, {10,11,12,13}, {9,10,11,12}, {8,9,10,11},{7,8,9,10}, {6,7,8,9}, {5,6,7,8}, {4,5,6,7}, {3,4,5,6}, {2,3,4,5}, {1,2,3,4}, {0,1,2,3}, {-1,0,1,2},{-2,-1,0,1}, {-3,-2,-1,0}, {-4,-3,-2,-1}, {-5,-4,-3,-2}, {-6,-5,-4,-3}, {-7,-6,-5,-4}, {-8,-7,-6,-5}, {-9,-8,-7,-6}, {-10,-9,-8,-7}, {-11,-10,-9,-8}}; 
+
+  const int32_t ucShift6[32][4] = { {24,25,26,27}, {23,24,25,26}, {22,23,24,25}, {21,22,23,24}, {20,21,22,23}, {19,20,21,22}, {18,19,20,21}, {17,18,19,20}, {16,17,18,19}, {15,16,17,18}, {14,15,16,17}, {13,14,15,16}, {12,13,14,15}, {11,12,13,14}, {10,11,12,13}, {9,10,11,12}, {8,9,10,11},{7,8,9,10}, {6,7,8,9}, {5,6,7,8}, {4,5,6,7}, {3,4,5,6}, {2,3,4,5}, {1,2,3,4}, {0,1,2,3}, {-1,0,1,2},{-2,-1,0,1}, {-3,-2,-1,0}, {-4,-3,-2,-1}, {-5,-4,-3,-2}, {-6,-5,-4,-3}, {-7,-6,-5,-4}}; 
+
+  const int32_t ucShift7[32][4] = { {28,29,30,31}, {27,28,29,30}, {26,27,28,29}, {25,26,27,28}, {24,25,26,27}, {23,24,25,26}, {22,23,24,25}, {21,22,23,24}, {20,21,22,23}, {19,20,21,22}, {18,19,20,21}, {17,18,19,20}, {16,17,18,19}, {15,16,17,18}, {14,15,16,17}, {13,14,15,16}, {12,13,14,15}, {11,12,13,14}, {10,11,12,13}, {9,10,11,12}, {8,9,10,11},{7,8,9,10}, {6,7,8,9}, {5,6,7,8}, {4,5,6,7}, {3,4,5,6}, {2,3,4,5}, {1,2,3,4}, {0,1,2,3}, {-1,0,1,2},{-2,-1,0,1}, {-3,-2,-1,0}}; 
+  const uint32_t __attribute__ ((aligned (16))) masks0[4] = {0x1,0x2,0x4,0x8};
+  const uint32_t __attribute__ ((aligned (16))) masks1[4] = {0x10,0x20,0x40,0x80};
+  const uint32_t __attribute__ ((aligned (16))) masks2[4] = {0x100,0x200,0x400,0x800};
+  const uint32_t __attribute__ ((aligned (16))) masks3[4] = {0x1000,0x2000,0x4000,0x8000};
+  const uint32_t __attribute__ ((aligned (16))) masks4[4] = {0x10000,0x20000,0x40000,0x80000};
+  const uint32_t __attribute__ ((aligned (16))) masks5[4] = {0x100000,0x200000,0x400000,0x800000};
+  const uint32_t __attribute__ ((aligned (16))) masks6[4] = {0x1000000,0x2000000,0x4000000,0x8000000};
+  const uint32_t __attribute__ ((aligned (16))) masks7[4] = {0x10000000,0x20000000,0x40000000,0x80000000};
+  int32x4_t vshift0[32],vshift1[32],vshift2[32],vshift3[32],vshift4[32],vshift5[32],vshift6[32],vshift7[32];
+  for (int n=0;n<32;n++) {
+	  vshift0[n] = vld1q_s32(ucShift0[n]);
+	  vshift1[n] = vld1q_s32(ucShift1[n]);
+	  vshift2[n] = vld1q_s32(ucShift2[n]);
+	  vshift3[n] = vld1q_s32(ucShift3[n]);
+	  vshift4[n] = vld1q_s32(ucShift4[n]);
+	  vshift5[n] = vld1q_s32(ucShift5[n]);
+	  vshift6[n] = vld1q_s32(ucShift6[n]);
+	  vshift7[n] = vld1q_s32(ucShift7[n]);
+  }
+  uint32x4_t vmask0  = vld1q_u32(masks0);
+  uint32x4_t vmask1  = vld1q_u32(masks1);
+  uint32x4_t vmask2  = vld1q_u32(masks2);
+  uint32x4_t vmask3  = vld1q_u32(masks3);
+  uint32x4_t vmask4  = vld1q_u32(masks4);
+  uint32x4_t vmask5  = vld1q_u32(masks5);
+  uint32x4_t vmask6  = vld1q_u32(masks6);
+  uint32x4_t vmask7  = vld1q_u32(masks7);
+  uint32_t output_tmp=0;
+  int s2=0;
+  for (s = 0; s < E2_first_segment ; s++) {
+    s2 = s&31;	  
+    foffset = (s>>5)*E;
+    fp = f+foffset;
+    int i;
+    if ((bit_index&31) == 0 ) {
+      for (i = 0; i < (E>>5)<<5; i+=32) {
+	uint32x4_t *fp128 = (uint32x4_t*)&fp[i];    
+  	uint32x4_t cshift = vandq_u32(vshlq_u32(fp128[0],vshift0[s2]),vmask0);
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[1],vshift1[s2]),vmask1));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[2],vshift2[s2]),vmask2));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[3],vshift3[s2]),vmask3));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[4],vshift4[s2]),vmask4));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[5],vshift5[s2]),vmask5));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[6],vshift6[s2]),vmask6));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[7],vshift7[s2]),vmask7));
+
+	*(output_p + (bit_index>>5))     = vaddvq_u32(cshift);
+	bit_index+=32;
+      }
+      uint32_t Emod32=E&31;
+      if (Emod32 != 0) {
+        uint32x4_t *fp128 = (uint32x4_t*)&fp[i];    
+  	uint32x4_t cshift = vandq_u32(vshlq_u32(fp128[0],vshift0[s2]),vmask0);
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[1],vshift1[s2]),vmask1));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[2],vshift2[s2]),vmask2));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[3],vshift3[s2]),vmask3));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[4],vshift4[s2]),vmask4));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[5],vshift5[s2]),vmask5));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[6],vshift6[s2]),vmask6));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[7],vshift7[s2]),vmask7));
+
+	*(output_p + (bit_index>>5))     = vaddvq_u32(cshift)&((1<<Emod32)-1);
+        bit_index+=Emod32;
+      }
+    }
+    else {
+      for (i = 0; i < (E>>5)<<5; i+=32) {
+	uint32x4_t *fp128 = (uint32x4_t*)&fp[i];    
+  	uint32x4_t cshift = vandq_u32(vshlq_u32(fp128[0],vshift0[s2]),vmask0);
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[1],vshift1[s2]),vmask1));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[2],vshift2[s2]),vmask2));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[3],vshift3[s2]),vmask3));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[4],vshift4[s2]),vmask4));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[5],vshift5[s2]),vmask5));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[6],vshift6[s2]),vmask6));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[7],vshift7[s2]),vmask7));
+	/*
+        uint32_t tmp = vaddvq_u32(cshift);
+	*(output_p + (bit_index>>5))     |= (tmp<<(bit_index&31));
+	*(output_p + (bit_index>>5)+1)   |= (tmp>>(32-(bit_index&31)));
+	*/
+        uint64_t tmp = (uint64_t)vaddvq_u32(cshift);
+	*(uint64_t*)(output_p + (bit_index>>5))     |= (tmp<<(bit_index&31));
+	bit_index+=32;
+      }
+      uint32_t Emod32=E&31;
+      if (Emod32 != 0) {
+        uint32x4_t *fp128 = (uint32x4_t*)&fp[i];    
+  	uint32x4_t cshift = vandq_u32(vshlq_u32(fp128[0],vshift0[s2]),vmask0);
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[1],vshift1[s2]),vmask1));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[2],vshift2[s2]),vmask2));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[3],vshift3[s2]),vmask3));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[4],vshift4[s2]),vmask4));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[5],vshift5[s2]),vmask5));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[6],vshift6[s2]),vmask6));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[7],vshift7[s2]),vmask7));
+        uint32_t tmp = vaddvq_u32(cshift);
+        tmp&=((1<<Emod32)-1);
+        *(output_p + (bit_index>>5))     |= (tmp<<(bit_index&31));
+	*(output_p + (bit_index>>5)+1)   |= (tmp>>(32-(bit_index&31)));
+        bit_index+=Emod32;
+      }
+    }
+  }
+//  s0 = s;
+  for ( ; s < nb_segments ; s++){
+    s2 = s&31;	  
+    foffset = ((s>>5)-E2_first_segment32)*E2;
+    fp = f2+foffset;
+    int i;
+    if ((bit_index&31) == 0 ) {
+      for (i = 0; i < (E2>>5)<<5; i+=32) {
+	uint32x4_t *fp128 = (uint32x4_t*)&fp[i];    
+  	uint32x4_t cshift = vandq_u32(vshlq_u32(fp128[0],vshift0[s2]),vmask0);
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[1],vshift1[s2]),vmask1));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[2],vshift2[s2]),vmask2));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[3],vshift3[s2]),vmask3));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[4],vshift4[s2]),vmask4));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[5],vshift5[s2]),vmask5));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[6],vshift6[s2]),vmask6));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[7],vshift7[s2]),vmask7));
+	*(output_p + (bit_index>>5))     = vaddvq_u32(cshift);
+	bit_index+=32;
+      }
+      uint32_t E2mod32=E2&31;
+      if (E2mod32 != 0) {
+        uint32x4_t *fp128 = (uint32x4_t*)&fp[i];    
+        uint32x4_t cshift = vandq_u32(vshlq_u32(fp128[0],vshift0[s2]),vmask0);
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[1],vshift1[s2]),vmask1));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[2],vshift2[s2]),vmask2));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[3],vshift3[s2]),vmask3));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[4],vshift4[s2]),vmask4));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[5],vshift5[s2]),vmask5));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[6],vshift6[s2]),vmask6));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[7],vshift7[s2]),vmask7));
+        *(output_p + (bit_index>>5))     = vaddvq_u32(cshift)&((1<<E2mod32)-1);
+        bit_index+=E2mod32;
+      }
+    }
+    else {
+      for (i = 0; i < (E2>>5)<<5; i+=32) {
+	uint32x4_t *fp128 = (uint32x4_t*)&fp[i];    
+        uint32x4_t cshift = vandq_u32(vshlq_u32(fp128[0],vshift0[s2]),vmask0);
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[1],vshift1[s2]),vmask1));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[2],vshift2[s2]),vmask2));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[3],vshift3[s2]),vmask3));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[4],vshift4[s2]),vmask4));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[5],vshift5[s2]),vmask5));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[6],vshift6[s2]),vmask6));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[7],vshift7[s2]),vmask7));
+        uint32_t tmp = vaddvq_u32(cshift);
+	*(output_p + (bit_index>>5))     |= (tmp<<(bit_index&31));
+	*(output_p + (bit_index>>5)+1)   |= (tmp>>(32-(bit_index&31)));
+	bit_index+=32;
+      }
+      uint32_t E2mod32=E2&31;
+      if (E2mod32 != 0) {
+        uint32x4_t *fp128 = (uint32x4_t*)&fp[i];    
+        uint32x4_t cshift = vandq_u32(vshlq_u32(fp128[0],vshift0[s2]),vmask0);
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[1],vshift1[s2]),vmask1));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[2],vshift2[s2]),vmask2));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[3],vshift3[s2]),vmask3));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[4],vshift4[s2]),vmask4));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[5],vshift5[s2]),vmask5));
+  	cshift = vorrq_u32(cshift,vandq_u32(vshlq_u32(fp128[6],vshift6[s2]),vmask6));
+        uint32_t tmp = vaddvq_u32(cshift);
+        tmp&=((1<<E2mod32)-1);
+        *(output_p + (bit_index>>5))     |= (tmp<<(bit_index&31));
+	*(output_p + (bit_index>>5)+1)   |= (tmp>>(32-(bit_index&31)));
+        bit_index+=E2mod32;
+      }
+    }
+  }
+#else // non SIMD version
+  unsigned int segpos,s2;	
+  for (s = 0; s < E2_first_segment ; s++) {
+    foffset = (s>>5)*E;
+    fp = f+foffset;
+    s2=s&31;
+    segpos = (1<<s2);
+  //  printf("E %d s %d: foffset %d, s2 %d, segpos %u\n",E,s,foffset,s2,segpos);
+    for (int i = 0; i < E; i++) {
+      output_p[bit_index>>5]|=((fp[i] & segpos)!=0)<<(bit_index&31); 
+      //printf("bit_index %d, output_p[%d] %x\n",bit_index, bit_index>>5,output_p[bit_index>>5]);
+      bit_index++;
+    }
+  }
+  for ( ; s < nb_segments ; s++){
+    foffset = ((s>>5)-E2_first_segment32)*E2;
+    fp = f2+foffset;
+    s2=s&31;
+    segpos = (1<<s2);
+//    printf("E2 %d s %d: foffset %d, s2 %d, segpos %u\n",E2,s,foffset,s2,segpos);
+    for (int i = 0; i < E2; i++) {
+      output_p[bit_index>>5]|=((fp[i] & segpos)!=0)<<(bit_index&31); 
+      bit_index++;
+    }
+  }
+#endif  
+}
+/**
+ * \typedef ldpc8blocks_args_t
+ * \struct ldpc8blocks_args_s
+ * \brief Arguments of an encoding task
+ * encode up to 8 code blocks
+ * \var nrLDPC_TB_encoding_parameters TB encoding parameters as defined in the coding library interface
+ * \var impp encoder implementation specific parameters for the task
+ * \var f first interleaver output to be filled by the task
+ * \var f2 second interleaver output to be filled by the task
+ * in case of a shift of E in the code blocks group processed by the task
+ */
+
+static void ldpcnblocks(nrLDPC_TB_encoding_parameters_t *nrLDPC_TB_encoding_parameters, encoder_implemparams_t impp)
+{
+ 
+  uint8_t mod_order = nrLDPC_TB_encoding_parameters->Qm;
+  uint16_t nb_rb = nrLDPC_TB_encoding_parameters->nb_rb;
+  uint32_t A = nrLDPC_TB_encoding_parameters->A;
+
+  unsigned int G = nrLDPC_TB_encoding_parameters->G;
+  LOG_D(PHY, "dlsch coding A %d K %d G %d (nb_rb %d, mod_order %d)\n", A, impp.K, G, nb_rb, (int)mod_order);
+
+  // nrLDPC_encoder output is in "d"
+  // let's make this interface happy!
+//  uint32_t d[4][68*384];
+  uint8_t *c[nrLDPC_TB_encoding_parameters->C];
+
+  if (!nrLDPC_TB_encoding_parameters->c_dev)  
+    for (int r = 0; r < nrLDPC_TB_encoding_parameters->C; r++)
+      c[r] = nrLDPC_TB_encoding_parameters->segments[r].c;
+  start_meas(&nrLDPC_TB_encoding_parameters->segments[impp.first_seg].ts_ldpc_encode);
+  LDPCencoder32(nrLDPC_TB_encoding_parameters->c_dev ? nrLDPC_TB_encoding_parameters->c_dev : c, &impp);
+  stop_meas(&nrLDPC_TB_encoding_parameters->segments[impp.first_seg].ts_ldpc_encode);
+  // Compute where to place in output buffer that is concatenation of all segments
+
+#ifdef DEBUG_LDPC_ENCODING
+  LOG_D(PHY, "rvidx in encoding = %d\n", nrLDPC_TB_encoding_parameters->rv_index);
+#endif
+  const uint32_t E = nrLDPC_TB_encoding_parameters->segments[0].E;
+  uint32_t E2=E;
+  uint32_t Emax = E;
+  int n_seg   = nrLDPC_TB_encoding_parameters->C>>5;
+  int n_seg2  = n_seg;
+  if ((nrLDPC_TB_encoding_parameters->C & 31) > 0) n_seg2++;
+  int r_shift = n_seg2; 
+  int r_shift2 = nrLDPC_TB_encoding_parameters->C;
+  for (int s=0;s<nrLDPC_TB_encoding_parameters->C;s++) {
+      //printf("segment %d E %d\n",s,nrLDPC_TB_encoding_parameters->segments[s].E);	  
+      if (nrLDPC_TB_encoding_parameters->segments[s].E != E) {
+	 E2=nrLDPC_TB_encoding_parameters->segments[s].E;
+         if(E2 > Emax)
+           Emax = E2;
+	 r_shift = s>>5;
+	 r_shift2 = s;
+	// printf("r_shift %d, r_shift2 %d\n",r_shift,r_shift2);
+         break;
+      }	 
+  }    
+
+  LOG_D(NR_PHY,
+        "Rate Matching, Code segment %d...%d r_shift %d n_seg2 %d (coded bits (G) %u, E %d, E2 %d Filler bits %d, Filler offset %d mod_order %d, nb_rb "
+          "%d,nrOfLayer %d)...\n",
+        0,
+        impp.n_segments-1,
+	r_shift,
+	n_seg2,
+        G,
+        E,E2,
+        impp.F,
+        impp.K - impp.F - 2 * impp.Zc,
+        mod_order,
+        nb_rb,
+        nrLDPC_TB_encoding_parameters->nb_layers);
+/*
+  printf("Rate Matching, Code segment 0..%d r_shift %d r_shift2 %d n_seg2 %d (coded bits (G) %u, E %d, E2 %d Filler bits %d, Filler offset %d mod_order %d, nb_rb "
+          "%d,nrOfLayer %d)...\n",
+        impp.n_segments-1,
+	r_shift,
+	r_shift2,
+	n_seg2,
+        G,
+        E,E2,
+        impp.F,
+        impp.K - impp.F - 2 * impp.Zc,
+        mod_order,
+        nb_rb,
+        nrLDPC_TB_encoding_parameters->nb_layers);
+*/
+
+  uint32_t Tbslbrm = nrLDPC_TB_encoding_parameters->tbslbrm;
+
+  uint32_t e[E*(r_shift+1)];
+  uint32_t e2[E2*(n_seg2-r_shift)];
+  uint32_t f[E*(r_shift+1)] __attribute__ ((aligned (64)));
+  uint32_t f2[E2*(n_seg2-r_shift)] __attribute__ ((aligned (64)));
+
+  // Interleaver outputs are stored in the output arrays
+  uint8_t *output = nrLDPC_TB_encoding_parameters->output;
+
+  start_meas(&nrLDPC_TB_encoding_parameters->segments[0].ts_rate_match);
+  memset(e,0,sizeof(e));
+  memset(f,0,sizeof(f));
+  if (1/*r_shift < n_seg2*/) { 
+    memset(e2,0,sizeof(e2));
+    memset(f2,0,sizeof(f2));
+  }
+
+  for (int r=0;r<n_seg2;r++) {
+    if (r<=r_shift)	  
+      nr_rate_matching_ldpc32(Tbslbrm,
+                              impp.BG,
+                              impp.Zc,
+                              d_host[r],
+                              e+(r*E),
+                              impp.n_segments,
+                              impp.F,
+                              impp.K - impp.F - 2 * impp.Zc,
+                              nrLDPC_TB_encoding_parameters->rv_index,
+                              E);
+    if (r>=r_shift)	  
+      nr_rate_matching_ldpc32(Tbslbrm,
+                              impp.BG,
+                              impp.Zc,
+                              d_host[r],
+                              e2+((r-r_shift)*E2),
+                              impp.n_segments,
+                              impp.F,
+                              impp.K - impp.F - 2 * impp.Zc,
+                              nrLDPC_TB_encoding_parameters->rv_index,
+                              E2);
+   /* 
+    if (r==(n_seg2-1)) {
+	    for (int i=0;i<16;i++) printf("rm: %x %x\n",d[n_seg2-1][i],e2[((n_seg2-1)*E2)+i]);
+    }
+    */
+  }
+  stop_meas(&nrLDPC_TB_encoding_parameters->segments[0].ts_rate_match);
+  if (impp.K - impp.F - 2 * impp.Zc > E) {
+    LOG_E(PHY,
+          "dlsch coding A %d  Kr %d G %d (nb_rb %d, mod_order %d)\n",
+          A,
+          impp.K,
+          G,
+          nb_rb,
+          (int)mod_order);
+
+    LOG_E(NR_PHY,
+          "Rate Matching, Code segments 0..%d (coded bits (G) %u, E %d, Kr %d, Filler bits %d, Filler offset %d mod_order %d, "
+          "nb_rb %d)...\n",
+          impp.n_segments,
+          G,
+          E,
+          impp.K,
+          impp.F,
+          impp.K - impp.F - 2 * impp.Zc,
+          mod_order,
+          nb_rb);
+  }
+  
+  //printf("interleaving r_shift %d, n_seg2 %d\n",r_shift,n_seg2);
+  start_meas(&nrLDPC_TB_encoding_parameters->segments[0].ts_interleave);
+  
+  for (int r=0;r<=r_shift;r++)
+    nr_interleaving_ldpc32(E,
+                           mod_order,
+                           e+E*r,
+                           f+E*r);
+
+  for (int r=r_shift;r<n_seg2;r++)
+    nr_interleaving_ldpc32(E2,
+                           mod_order,
+                           e2+E2*(r-r_shift),
+                           f2+E2*(r-r_shift));
+/*
+  for (int i=0;i<16;i++) printf("intl (f offset %d): %x %x\n",(n_seg2-1)*E2,e2[((n_seg2-1)*E2)+i],f2[((n_seg2-1)*E2)+i]);
+  printf("-------------------\n");
+  for (int i=E2-16;i<E2;i++) printf("intl (f offset %d): %x %x\n",(n_seg2-1)*E2,e2[((n_seg2-1)*E2)+i],f2[((n_seg2-1)*E2)+i]);
+  */
+  stop_meas(&nrLDPC_TB_encoding_parameters->segments[0].ts_interleave);
+
+  if (impp.tconcat) start_meas(impp.tconcat);
+  unpack_output(f,E,f2,E2,r_shift,r_shift2,nrLDPC_TB_encoding_parameters->C,output);
+  if (impp.tconcat) stop_meas(impp.tconcat);
+
+}
+
+int nrLDPC_coding_encoder32(nrLDPC_slot_encoding_parameters_t *nrLDPC_slot_encoding_parameters, nrLDPC_TB_encoding_parameters_t *nrLDPC_TB_encoding_parameters)
+{
+
+
+    encoder_implemparams_t common_segment_params = {
+      .n_segments = nrLDPC_TB_encoding_parameters->C,
+      .tinput = nrLDPC_slot_encoding_parameters->tinput,
+      .tinput_memcpy = nrLDPC_slot_encoding_parameters->tinput_memcpy,
+      .tprep = nrLDPC_slot_encoding_parameters->tprep,
+      .tparity = nrLDPC_slot_encoding_parameters->tparity,
+      .toutput = nrLDPC_slot_encoding_parameters->toutput,
+      .tconcat = nrLDPC_slot_encoding_parameters->tconcat,
+      .Kb = nrLDPC_TB_encoding_parameters->Kb,
+      .Zc = nrLDPC_TB_encoding_parameters->Z,
+      .BG = nrLDPC_TB_encoding_parameters->BG,
+      .output = nrLDPC_TB_encoding_parameters->output, 
+      .K = nrLDPC_TB_encoding_parameters->K,
+      .F = nrLDPC_TB_encoding_parameters->F,
+    };
+
+    LOG_D(NR_PHY,"Calling ldpcnblocks (C %d, Z %d, K %d)\n",common_segment_params.n_segments,common_segment_params.Zc,common_segment_params.K);
+    ldpcnblocks(nrLDPC_TB_encoding_parameters, common_segment_params);
+
+
+
+  return 0;
+}

openair1/PHY/CODING/nrLDPC_coding/nrLDPC_coding_segment/nr_rate_matching.c

@@ -24,7 +24,8 @@
 #include "common/utils/LOG/log.h"
 
 // #define RM_DEBUG 1
-
+#define USE128BIT
+ 
 static const uint8_t index_k0[2][4] = {{0, 17, 33, 56}, {0, 13, 25, 43}};
 
 void nr_interleaving_ldpc(uint32_t E, uint8_t Qm, uint8_t *e, uint8_t *f)
@@ -359,14 +360,15 @@ void nr_interleaving_ldpc(uint32_t E, uint8_t Qm, uint8_t *e, uint8_t *f)
       f = (uint8_t *)f_512;
 #endif
 #ifdef USE128BIT
-      e0_128 = (simde__m128i *)e0;
-      e1_128 = (simde__m128i *)e1;
-      e2_128 = (simde__m128i *)e2;
-      e3_128 = (simde__m128i *)e3;
-      e4_128 = (simde__m128i *)e4;
-      e5_128 = (simde__m128i *)e5;
-      e6_128 = (simde__m128i *)e6;
-      e7_128 = (simde__m128i *)e7;
+      simde__m128i *e0_128 = (simde__m128i *)e0;
+      simde__m128i *e1_128 = (simde__m128i *)e1;
+      simde__m128i *e2_128 = (simde__m128i *)e2;
+      simde__m128i *e3_128 = (simde__m128i *)e3;
+      simde__m128i *e4_128 = (simde__m128i *)e4;
+      simde__m128i *e5_128 = (simde__m128i *)e5;
+      simde__m128i *e6_128 = (simde__m128i *)e6;
+      simde__m128i *e7_128 = (simde__m128i *)e7;
+      simde__m128i *f_128  = (simde__m128i *)f;
       for (; i < (EQm & ~15); i += 16) {
         simde__m128i e0j = simde_mm_loadu_si128(e0_128++);
         simde__m128i e1j = simde_mm_loadu_si128(e1_128++);
@@ -444,6 +446,225 @@ void nr_interleaving_ldpc(uint32_t E, uint8_t Qm, uint8_t *e, uint8_t *f)
   }
 }
 
+void nr_interleaving_ldpc32(uint32_t E, uint8_t Qm, uint32_t *e, uint32_t *f)
+{
+  const uint32_t EQm = E / Qm;
+  memset(f, 0, E * sizeof(uint32_t));
+  switch(Qm) {
+    case 2: {
+      uint32_t *e0 = e;
+      uint32_t *e1 = e0 + EQm;
+      int i = 0;
+      for (; i < EQm; i++) {
+        *f++ = *e0++;
+        *f++ = *e1++;
+      }
+    } break;
+    case 4: {
+      uint32_t *e0 = e;
+      uint32_t *e1 = e0 + EQm;
+      uint32_t *e2 = e1 + EQm;
+      uint32_t *e3 = e2 + EQm;
+      int i = 0;
+      for (; i < EQm; i++) {
+        *f++ = *e0++;
+        *f++ = *e1++;
+        *f++ = *e2++;
+        *f++ = *e3++;
+      }
+    } break;
+    case 6: {
+      uint32_t *e0 = e;
+      uint32_t *e1 = e0 + EQm;
+      uint32_t *e2 = e1 + EQm;
+      uint32_t *e3 = e2 + EQm;
+      uint32_t *e4 = e3 + EQm;
+      uint32_t *e5 = e4 + EQm;
+      int i = 0;
+      simde__m128i *e0_128 = (simde__m128i *)e0;
+      simde__m128i *e1_128 = (simde__m128i *)e1;
+      simde__m128i *e2_128 = (simde__m128i *)e2;
+      simde__m128i *e3_128 = (simde__m128i *)e3;
+      simde__m128i *e4_128 = (simde__m128i *)e4;
+      simde__m128i *e5_128 = (simde__m128i *)e5;
+      simde__m128i *f128 = (simde__m128i *)f;
+      for (; i < (EQm & ~3); i += 4) {
+        simde__m128i e0j = simde_mm_loadu_si128(e0_128++);
+        simde__m128i e1j = simde_mm_loadu_si128(e1_128++);
+        simde__m128i e2j = simde_mm_loadu_si128(e2_128++);
+        simde__m128i e3j = simde_mm_loadu_si128(e3_128++);
+        simde__m128i e4j = simde_mm_loadu_si128(e4_128++);
+        simde__m128i e5j = simde_mm_loadu_si128(e5_128++);
+
+        simde__m128i tmp0 = simde_mm_unpacklo_epi32(e0j, e1j);   // e0(i) e1(i) e0(i+1) e1(i+1) 
+        simde__m128i tmp1 = simde_mm_unpacklo_epi32(e2j, e3j);   // e2(i) e3(i) e2(i+1) e3(i+1) 
+        simde__m128i tmp2 = simde_mm_unpacklo_epi32(e4j, e5j);   // e4(i) e5(i) e4(i+1) e5(i+1) 
+								 
+        simde_mm_storeu_si128(f128++,simde_mm_unpacklo_epi64(tmp0, tmp1));                        // e0(i) e1(i) e2(i) e3(i) 
+        simde_mm_storeu_si128(f128++,simde_mm_unpacklo_epi64(tmp2,simde_mm_unpackhi_epi64(tmp0,tmp0))); // e4(i) e5(i) e0(i+1) e1(i+1) 
+	simde_mm_storeu_si128(f128++,simde_mm_unpackhi_epi64(tmp1,tmp2));                         // e2(i+1) e3(i+1) e4(i+1) e5(i+1)
+        tmp0 = simde_mm_unpackhi_epi32(e0j, e1j);   // e0(i+2) e1(i+2) e0(i+3) e1(i+3) 
+        tmp1 = simde_mm_unpackhi_epi32(e2j, e3j);   // e2(i+2) e3(i+2) e2(i+3) e3(i+3) 
+        tmp2 = simde_mm_unpackhi_epi32(e4j, e5j);   // e4(i+2) e5(i+2) e4(i+3) e5(i+3) 
+								 
+        simde_mm_storeu_si128(f128++,simde_mm_unpacklo_epi64(tmp0, tmp1));                        // e0(i+2) e1(i+2) e2(i+2) e3(i+2) 
+        simde_mm_storeu_si128(f128++,simde_mm_unpacklo_epi64(tmp2,simde_mm_unpackhi_epi64(tmp0,tmp0))); // e4(i+2) e5(i+2) e0(i+3) e1(i+3) 
+	simde_mm_storeu_si128(f128++,simde_mm_unpackhi_epi64(tmp1,tmp2));                         // e2(i+3) e3(i+3) e4(i+3) e5(i+3)
+      }
+      e0 = (uint32_t *)e0_128;
+      e1 = (uint32_t *)e1_128;
+      e2 = (uint32_t *)e2_128;
+      e3 = (uint32_t *)e3_128;
+      e4 = (uint32_t *)e4_128;
+      e5 = (uint32_t *)e5_128;
+      f  = (uint32_t *)f128;
+      for (; i < EQm; i++) {
+        *f++ = *e0++;
+        *f++ = *e1++;
+        *f++ = *e2++;
+        *f++ = *e3++;
+        *f++ = *e4++;
+        *f++ = *e5++;
+      }
+    } break;
+    case 8: {
+      uint32_t *e0 = e;
+      uint32_t *e1 = e0 + EQm;
+      uint32_t *e2 = e1 + EQm;
+      uint32_t *e3 = e2 + EQm;
+      uint32_t *e4 = e3 + EQm;
+      uint32_t *e5 = e4 + EQm;
+      uint32_t *e6 = e5 + EQm;
+      uint32_t *e7 = e6 + EQm;
+      
+      int i = 0;
+#if 0 /*def __AVX2__, this needs to be fixed, revert to 128-bit below*/
+      simde__m256i *e0_256 = (simde__m256i *)e0;
+      simde__m256i *e1_256 = (simde__m256i *)e1;
+      simde__m256i *e2_256 = (simde__m256i *)e2;
+      simde__m256i *e3_256 = (simde__m256i *)e3;
+      simde__m256i *e4_256 = (simde__m256i *)e4;
+      simde__m256i *e5_256 = (simde__m256i *)e5;
+      simde__m256i *e6_256 = (simde__m256i *)e6;
+      simde__m256i *e7_256 = (simde__m256i *)e7;
+      simde__m256i *f256   = (simde__m256i *)f;
+
+      for (; i < (EQm & ~7); i += 8) {
+        simde__m256i e0j = simde_mm256_loadu_si256(e0_256++);
+        simde__m256i e1j = simde_mm256_loadu_si256(e1_256++);
+        simde__m256i e2j = simde_mm256_loadu_si256(e2_256++);
+        simde__m256i e3j = simde_mm256_loadu_si256(e3_256++);
+        simde__m256i e4j = simde_mm256_loadu_si256(e4_256++);
+        simde__m256i e5j = simde_mm256_loadu_si256(e5_256++);
+        simde__m256i e6j = simde_mm256_loadu_si256(e6_256++);
+        simde__m256i e7j = simde_mm256_loadu_si256(e7_256++);
+
+        simde__m256i tmp0 = simde_mm256_unpacklo_epi32(e0j, e1j);   // e0(i) e1(i) e0(i+1) e1(i+1) e0(i+4) e1(i+4) e0(i+5) e1(i+5) 
+        simde__m256i tmp1 = simde_mm256_unpacklo_epi32(e2j, e3j);   // e2(i) e3(i) e2(i+1) e3(i+1) e2(i+4) e3(i+4) e2(i+5) e3(i+5)
+        simde__m256i tmp2 = simde_mm256_unpacklo_epi32(e4j, e5j);   // e4(i) e5(i) e4(i+1) e5(i+1) e4(i+4) e5(i+4) e4(i+5) e5(i+5) 
+        simde__m256i tmp3 = simde_mm256_unpacklo_epi32(e6j, e7j);   // e6(i) e7(i) e6(i+1) e7(i+1) e6(i+4) e7(i+4) e6(i+5) e7(i+5)
+								    //
+        simde__m256i tmp4 = simde_mm256_unpacklo_epi64(tmp0, tmp1);   // e0(i) e1(i) e2(i) e3(i) e0(i+4) e1(i+4) e2(i+4) e3(i+4) 
+        simde__m256i tmp5 = simde_mm256_unpacklo_epi64(tmp2, tmp3);   // e4(i) e5(i) e6(i) e7(i) e4(i+4) e5(i+4) e6(i+4) e7(i+4) 
+								 
+        simde_mm256_storeu_si256(f256++,_mm256_permute2x128_si256(tmp4,tmp5,0x20)); // e0(i) e1(i) e2(i) e3(i) e4(i) e5(i) e6(i) e7(i) 
+										    
+        simde__m256i tmp6 = simde_mm256_unpackhi_epi64(tmp0, tmp1);   // e0(i+1) e1(i+1) e2(i+1) e3(i+1) e0(i+5) e1(i+5) e2(i+5) e3(i+5) 
+        simde__m256i tmp7 = simde_mm256_unpackhi_epi64(tmp2, tmp3);   // e4(i+1) e5(i+1) e6(i+1) e7(i+1) e4(i+5) e5(i+5) e6(i+5) e7(i+5) 
+        simde_mm256_storeu_si256(f256++,_mm256_permute2x128_si256(tmp6,tmp7,0x20)); // e0(i+1) e1(i+1) e2(i+1) e3(i+1) e4(i+1) e5(i+1) e6(i+1) e7(i+1) 
+										    //
+        tmp0 = simde_mm256_unpackhi_epi32(e0j, e1j);   // e0(i+2) e1(i+2) e0(i+3) e1(i+3) e0(i+6) e1(i+6) e0(i+7) e1(i+7) 
+        tmp1 = simde_mm256_unpackhi_epi32(e2j, e3j);   // e2(i+2) e3(i+2) e2(i+3) e3(i+3) e2(i+6) e3(i+6) e2(i+7) e3(i+7)
+        tmp2 = simde_mm256_unpackhi_epi32(e4j, e5j);   // e4(i+2) e5(i+2) e4(i+3) e5(i+3) e4(i+6) e5(i+6) e4(i+7) e5(i+7) 
+        tmp3 = simde_mm256_unpackhi_epi32(e6j, e7j);   // e6(i+2) e7(i+2) e6(i+3) e7(i+3) e6(i+6) e7(i+6) e6(i+7) e7(i+7)
+								    //
+        simde__m256i tmp8 = simde_mm256_unpacklo_epi64(tmp0, tmp1);   // e0(i+2) e1(i+2) e2(i+2) e3(i+2) e0(i+6) e1(i+6) e2(i+6) e3(i+6) 
+        simde__m256i tmp9 = simde_mm256_unpacklo_epi64(tmp2, tmp3);   // e4(i+2) e5(i+2) e6(i+2) e7(i+2) e4(i+6) e5(i+6) e6(i+6) e7(i+6) 
+								 
+        simde_mm256_storeu_si256(f256++,_mm256_permute2x128_si256(tmp8,tmp9,0x20)); // e0(i+2) e1(i+2) e2(i+2) e3(i+2) e4(i+2) e5(i+2) e6(i+2) e7(i+2) 
+        simde__m256i tmp10 = simde_mm256_unpackhi_epi64(tmp0, tmp1);   // e0(i+3) e1(i+3) e2(i+3) e3(i+3) e0(i+7) e1(i+7) e2(i+7) e3(i+7) 
+        simde__m256i tmp11 = simde_mm256_unpackhi_epi64(tmp2, tmp3);   // e4(i+3) e5(i+3) e6(i+3) e7(i+3) e4(i+7) e5(i+7) e6(i+7) e7(i+7) 
+        simde_mm256_storeu_si256(f256++,_mm256_permute2x128_si256(tmp10,tmp11,0x20)); // e0(i+3) e1(i+3) e2(i+3) e3(i+3) e4(i+3) e5(i+3) e6(i+3) e7(i+3) 
+        simde_mm256_storeu_si256(f256++,_mm256_permute2x128_si256(tmp4,tmp5,0x31)); // e0(i+4) e1(i+4) e2(i+4) e3(i+4) e4(i+4) e5(i+4) e6(i+4) e7(i+4) 
+        simde_mm256_storeu_si256(f256++,_mm256_permute2x128_si256(tmp6,tmp7,0x31)); // e0(i+5) e1(i+5) e2(i+5) e3(i+5) e4(i+5) e5(i+5) e6(i+5) e7(i+5) 
+        simde_mm256_storeu_si256(f256++,_mm256_permute2x128_si256(tmp8,tmp9,0x31)); // e0(i+6) e1(i+6) e2(i+6) e3(i+6) e4(i+6) e5(i+6) e6(i+6) e7(i+6) 
+        simde_mm256_storeu_si256(f256++,_mm256_permute2x128_si256(tmp10,tmp11,0x31)); // e0(i+7) e1(i+7) e2(i+7) e3(i+7) e4(i+7) e5(i+7) e6(i+7) e7(i+7) 
+      }
+      e0 = (uint32_t *)e0_256;
+      e1 = (uint32_t *)e1_256;
+      e2 = (uint32_t *)e2_256;
+      e3 = (uint32_t *)e3_256;
+      e4 = (uint32_t *)e4_256;
+      e5 = (uint32_t *)e5_256;
+      e6 = (uint32_t *)e6_256;
+      e7 = (uint32_t *)e7_256;
+      f = (uint32_t *)f256;
+#endif
+      simde__m128i *e0_128 = (simde__m128i *)e0;
+      simde__m128i *e1_128 = (simde__m128i *)e1;
+      simde__m128i *e2_128 = (simde__m128i *)e2;
+      simde__m128i *e3_128 = (simde__m128i *)e3;
+      simde__m128i *e4_128 = (simde__m128i *)e4;
+      simde__m128i *e5_128 = (simde__m128i *)e5;
+      simde__m128i *e6_128 = (simde__m128i *)e6;
+      simde__m128i *e7_128 = (simde__m128i *)e7;
+      simde__m128i *f128   = (simde__m128i *)f;
+
+      for (; i < (EQm & ~3); i += 4) {
+        simde__m128i e0j = simde_mm_loadu_si128(e0_128++);
+        simde__m128i e1j = simde_mm_loadu_si128(e1_128++);
+        simde__m128i e2j = simde_mm_loadu_si128(e2_128++);
+        simde__m128i e3j = simde_mm_loadu_si128(e3_128++);
+        simde__m128i e4j = simde_mm_loadu_si128(e4_128++);
+        simde__m128i e5j = simde_mm_loadu_si128(e5_128++);
+        simde__m128i e6j = simde_mm_loadu_si128(e6_128++);
+        simde__m128i e7j = simde_mm_loadu_si128(e7_128++);
+
+        simde__m128i tmp0 = simde_mm_unpacklo_epi32(e0j, e1j);   // e0(i) e1(i) e0(i+1) e1(i+1) 
+        simde__m128i tmp1 = simde_mm_unpacklo_epi32(e2j, e3j);   // e2(i) e3(i) e2(i+1) e3(i+1) 
+        simde__m128i tmp2 = simde_mm_unpacklo_epi32(e4j, e5j);   // e4(i) e5(i) e4(i+1) e5(i+1) 
+        simde__m128i tmp3 = simde_mm_unpacklo_epi32(e6j, e7j);   // e6(i) e7(i) e6(i+1) e7(i+1)
+								 
+        simde_mm_storeu_si128(f128++,simde_mm_unpacklo_epi64(tmp0, tmp1)); // e0(i) e1(i) e2(i) e3(i) 
+        simde_mm_storeu_si128(f128++,simde_mm_unpacklo_epi64(tmp2, tmp3)); // e4(i) e5(i) e6(i) e7(i)
+	simde_mm_storeu_si128(f128++,simde_mm_unpackhi_epi64(tmp0, tmp1));  // e0(i+1) e1(i+1) e2(i+1) e3(i+1) 
+        simde_mm_storeu_si128(f128++,simde_mm_unpackhi_epi64(tmp2, tmp3)); // e4(i+1) e5(i+1) e6(i+1) e7(i+1)
+									   //
+        tmp0 = simde_mm_unpackhi_epi32(e0j, e1j);   // e0(i+2) e1(i+2) e0(i+3) e1(i+3) 
+        tmp1 = simde_mm_unpackhi_epi32(e2j, e3j);   // e2(i+2) e3(i+2) e2(i+3) e3(i+3) 
+        tmp2 = simde_mm_unpackhi_epi32(e4j, e5j);   // e4(i+2) e5(i+2) e4(i+3) e5(i+3) 
+        tmp3 = simde_mm_unpackhi_epi32(e6j, e7j);   // e6(i+2) e7(i+2) e6(i+3) e7(i+3)
+								 
+        simde_mm_storeu_si128(f128++,simde_mm_unpacklo_epi64(tmp0, tmp1)); // e0(i+2) e1(i+2) e2(i+2) e3(i+2) 
+        simde_mm_storeu_si128(f128++,simde_mm_unpacklo_epi64(tmp2, tmp3)); // e4(i+2) e5(i+2) e6(i+2) e7(i+2)
+	simde_mm_storeu_si128(f128++,simde_mm_unpackhi_epi64(tmp0, tmp1)); // e0(i+3) e1(i+3) e2(i+3) e3(i+3) 
+        simde_mm_storeu_si128(f128++,simde_mm_unpackhi_epi64(tmp2, tmp3)); //e4(i+3) e5(i+3) e6(i+3) e7(i+3)
+      }
+      e0 = (uint32_t *)e0_128;
+      e1 = (uint32_t *)e1_128;
+      e2 = (uint32_t *)e2_128;
+      e3 = (uint32_t *)e3_128;
+      e4 = (uint32_t *)e4_128;
+      e5 = (uint32_t *)e5_128;
+      e6 = (uint32_t *)e6_128;
+      e7 = (uint32_t *)e7_128;
+      f = (uint32_t *)f128;
+      
+      for (; i < EQm; i++) {
+        *f++ = *e0++;
+        *f++ = *e1++;
+        *f++ = *e2++;
+        *f++ = *e3++;
+        *f++ = *e4++;
+        *f++ = *e5++;
+        *f++ = *e6++;
+        *f++ = *e7++;
+      }
+    } break;
+    default:
+      AssertFatal(false, "Should be here!\n");
+  }
+}
 void nr_deinterleaving_ldpc(uint32_t E, uint8_t Qm, int16_t *e, int16_t *f)
 {
   switch (Qm) {
@@ -610,6 +831,103 @@ int nr_rate_matching_ldpc(uint32_t Tbslbrm,
   return 0;
 }
 
+int nr_rate_matching_ldpc32(uint32_t Tbslbrm,
+                            uint8_t BG,
+                            uint16_t Z,
+                            uint32_t *d,
+                            uint32_t *e,
+                            uint8_t C,
+                            uint32_t F,
+                            uint32_t Foffset,
+                            uint8_t rvidx,
+                            uint32_t E)
+{
+  if (C == 0) {
+    LOG_E(PHY, "nr_rate_matching: invalid parameter C %d\n", C);
+    return -1;
+  }
+
+  //Bit selection
+  uint32_t N = (BG == 1) ? (66 * Z) : (50 * Z);
+  uint32_t Ncb;
+  if (Tbslbrm == 0)
+    Ncb = N;
+  else {
+    uint32_t Nref = 3 * Tbslbrm / (2 * C); //R_LBRM = 2/3
+    Ncb = min(N, Nref);
+  }
+
+  uint32_t ind = (index_k0[BG - 1][rvidx] * Ncb / N) * Z;
+
+#ifdef RM_DEBUG
+  printf("nr_rate_matching_ldpc: E %u, F %u, Foffset %u, k0 %u, Ncb %u, rvidx %d, Tbslbrm %u\n",
+         E,
+         F,
+         Foffset,
+         ind,
+         Ncb,
+         rvidx,
+         Tbslbrm);
+#endif
+
+  if (Foffset > E) {
+    LOG_E(PHY,
+          "nr_rate_matching: invalid parameters (Foffset %d > E %d) F %d, k0 %d, Ncb %d, rvidx %d, Tbslbrm %d\n",
+          Foffset,
+          E,
+          F,
+          ind,
+          Ncb,
+          rvidx,
+          Tbslbrm);
+    return -1;
+  }
+  if (Foffset > Ncb) {
+    LOG_E(PHY, "nr_rate_matching: invalid parameters (Foffset %d > Ncb %d)\n", Foffset, Ncb);
+    return -1;
+  }
+
+  if (ind >= Foffset && ind < (F + Foffset))
+    ind = F + Foffset;
+
+  uint32_t k = 0;
+  if (ind < Foffset) { // case where we have some bits before the filler and the rest after
+    memcpy((void *)e, (void *)(d + ind), (Foffset - ind)<<2);
+
+    if (E + F <= Ncb - ind) { // E+F doesn't contain all coded bits
+      memcpy((void *)(e + Foffset - ind), (void *)(d + Foffset + F), (E - Foffset + ind)<<2);
+      k = E;
+    } else {
+      memcpy((void *)(e + Foffset - ind), (void *)(d + Foffset + F), (Ncb - Foffset - F)<<2);
+      k = Ncb - F - ind;
+    }
+  } else {
+    if (E <= Ncb - ind) { // E+F doesn't contain all coded bits
+      memcpy((void *)(e), (void *)(d + ind), E<<2);
+      k = E;
+    } else {
+      memcpy((void *)(e), (void *)(d + ind), (Ncb - ind)<<2);
+      k = Ncb - ind;
+    }
+  }
+
+  while (k < E) { // case where we do repetitions (low mcs)
+    for (ind = 0; (ind < Ncb) && (k < E); ind++) {
+#ifdef RM_DEBUG
+      printf("RM_TX k%u Ind: %u (%d)\n", k, ind, d[ind]);
+#endif
+
+      if (ind == Foffset)
+        ind = F + Foffset; // skip filler bits
+
+      e[k++] = d[ind];
+
+    }
+  }
+
+  return 0;
+}
+
 int nr_rate_matching_ldpc_rx(uint32_t Tbslbrm,
                              uint8_t BG,
                              uint16_t Z,

openair1/PHY/CODING/nrLDPC_coding/nrLDPC_coding_segment/nr_rate_matching.h

@@ -33,6 +33,15 @@
  */
 void nr_interleaving_ldpc(uint32_t E, uint8_t Qm, uint8_t *e, uint8_t *f);
 
+/**
+ * \brief interleave a code segment after encoding and rate matching (32bit)
+ * \param E size of the code segment in bits
+ * \param Qm modulation order
+ * \param e input rate matched segment
+ * \param f output interleaved segment
+ */
+
+void nr_interleaving_ldpc32(uint32_t E, uint8_t Qm, uint32_t *e, uint32_t *f);
 /**
  * \brief deinterleave a code segment before RX rate matching and decoding
  * \param E size of the code segment in bits
@@ -66,6 +75,30 @@ int nr_rate_matching_ldpc(uint32_t Tbslbrm,
                           uint8_t rvidx,
                           uint32_t E);
 
+/**
+ * \brief rate match a code segment after encoding (32bit)
+ * \Tbslbrm Transport Block size LBRM
+ * \param BG LDPC base graph number
+ * \param Z segment lifting size
+ * \param d input encoded segment
+ * \param e output rate matched segment
+ * \param C number of segments in the Transport Block
+ * \param F number of filler bits in the segment
+ * \param Foffset offset of the filler bits in the segment
+ * \param rvidx redundancy version index
+ * \param E size of the code segment in bits
+ */
+int nr_rate_matching_ldpc32(uint32_t Tbslbrm,
+                            uint8_t BG,
+                            uint16_t Z,
+                            uint32_t *d,
+                            uint32_t *e,
+                            uint8_t C,
+                            uint32_t F,
+                            uint32_t Foffset,
+                            uint8_t rvidx,
+                            uint32_t E);
+
 /**
  * \brief rate match a code segment before decoding
  * \Tbslbrm Transport Block size LBRM

openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.c

@@ -161,6 +161,7 @@ static inline uint32_t nrLDPC_decoder_core(int8_t* p_llr,
 
 int32_t LDPCinit()
 {
+  cuda_support_init();
   return 0;
 }
 

openair1/PHY/CODING/nrLDPC_defs.h

@@ -55,9 +55,11 @@ typedef struct {
   unsigned int first_seg; // optim8segmulti
   unsigned char gen_code; // orig
   time_stats_t *tinput;
+  time_stats_t *tinput_memcpy;
   time_stats_t *tprep;
   time_stats_t *tparity;
   time_stats_t *toutput;
+  time_stats_t *tconcat;
   /// Size in bits of the code segments
   uint32_t K;
   /// Number of lifting sizes to fit the payload
@@ -91,5 +93,7 @@ typedef int32_t(LDPC_decoderfunc_t)(t_nrLDPC_dec_params *p_decParams,
                                     t_nrLDPC_time_stats *time_stats,
                                     decode_abort_t *ab);
 typedef int32_t(LDPC_encoderfunc_t)(uint8_t **, uint8_t *, encoder_implemparams_t *);
+typedef uint32_t**(LDPC_encoderfunc32_t)(uint8_t **, encoder_implemparams_t *);
 
+void cuda_support_init(void);
 #endif

openair1/PHY/CODING/nrLDPC_encoder/ldpc_BG1_Zc384_32bit.cu

@@ -0,0 +1,214 @@
+#include <stdio.h>
+#include <stdint.h>
+#include <cuda_runtime.h>
+// generated code for Zc=384, byte encoding
+__global__ void ldpc_BG1_Zc384_worker(uint32_t *c[4],uint32_t *d[4]) {
+  uint32_t *c32=c[blockIdx.x];
+  uint32_t *d32=d[blockIdx.x]+8448-2*384;
+
+  int i2 = threadIdx.x;
+  int i1 = blockIdx.y;
+  if (i2 < 384) {
+    c32+=i2;
+    d32+=i2;
+    switch(i1) {
+    case 0:
+      d32[0]=c32[307]^c32[76]^c32[205]^c32[276]^c32[787]^c32[1018]^c32[855]^c32[1586]^c32[1612]^c32[1864]^c32[2673]^c32[2377]^c32[2304]^c32[3360]^c32[3404]^c32[3347]^c32[4021]^c32[3984]^c32[4096]^c32[4824]^c32[4769]^c32[4807]^c32[5707]^c32[5643]^c32[5529]^c32[6475]^c32[6304]^c32[6200]^c32[7229]^c32[7090]^c32[6975]^c32[7968]^c32[7809]^c32[7812]^c32[8557]^c32[8743]^c32[8753]^c32[9233]^c32[9558]^c32[9447]^c32[10341]^c32[10184]^c32[10325]^c32[10969]^c32[10840]^c32[10964]^c32[11735]^c32[11619]^c32[11573]^c32[12394]^c32[12642]^c32[12592]^c32[13170]^c32[13187]^c32[13356]^c32[14066]^c32[14064]^c32[14095]^c32[14772]^c32[14923]^c32[14797]^c32[15690]^c32[15373]^c32[15399]^c32[16474]^c32[16240]^c32[16485];
+
+       break;
+    case 1:
+      d32[384]=c32[307]^c32[308]^c32[77]^c32[206]^c32[277]^c32[787]^c32[788]^c32[1019]^c32[856]^c32[1586]^c32[1587]^c32[1613]^c32[1865]^c32[2673]^c32[2674]^c32[2378]^c32[2305]^c32[3361]^c32[3405]^c32[3348]^c32[4021]^c32[4022]^c32[3985]^c32[4097]^c32[4824]^c32[4825]^c32[4770]^c32[4808]^c32[5708]^c32[5644]^c32[5530]^c32[6476]^c32[6305]^c32[6201]^c32[7229]^c32[7230]^c32[7091]^c32[6976]^c32[7968]^c32[7969]^c32[7810]^c32[7813]^c32[8557]^c32[8558]^c32[8744]^c32[8754]^c32[9233]^c32[9234]^c32[9559]^c32[9448]^c32[10341]^c32[10342]^c32[10185]^c32[10326]^c32[10970]^c32[10841]^c32[10965]^c32[11735]^c32[11736]^c32[11620]^c32[11574]^c32[12394]^c32[12395]^c32[12643]^c32[12593]^c32[13171]^c32[13188]^c32[13357]^c32[14066]^c32[14067]^c32[14065]^c32[14096]^c32[14772]^c32[14773]^c32[14924]^c32[14798]^c32[15690]^c32[15691]^c32[15374]^c32[15400]^c32[16474]^c32[16475]^c32[16241]^c32[16486];
+
+       break;
+    case 2:
+      d32[768]=c32[308]^c32[77]^c32[205]^c32[206]^c32[276]^c32[277]^c32[788]^c32[1018]^c32[1019]^c32[855]^c32[856]^c32[1587]^c32[1613]^c32[1864]^c32[1865]^c32[2674]^c32[2378]^c32[2304]^c32[2305]^c32[3361]^c32[3404]^c32[3405]^c32[3347]^c32[3348]^c32[4022]^c32[3985]^c32[4096]^c32[4097]^c32[4825]^c32[4769]^c32[4770]^c32[4807]^c32[4808]^c32[5708]^c32[5643]^c32[5644]^c32[5529]^c32[5530]^c32[6476]^c32[6304]^c32[6305]^c32[6200]^c32[6201]^c32[7230]^c32[7091]^c32[6975]^c32[6976]^c32[7969]^c32[7809]^c32[7810]^c32[7812]^c32[7813]^c32[8558]^c32[8744]^c32[8753]^c32[8754]^c32[9234]^c32[9559]^c32[9447]^c32[9448]^c32[10342]^c32[10184]^c32[10185]^c32[10325]^c32[10326]^c32[10970]^c32[10840]^c32[10841]^c32[10964]^c32[10965]^c32[11736]^c32[11620]^c32[11573]^c32[11574]^c32[12395]^c32[12643]^c32[12592]^c32[12593]^c32[13171]^c32[13187]^c32[13188]^c32[13356]^c32[13357]^c32[14067]^c32[14064]^c32[14065]^c32[14095]^c32[14096]^c32[14773]^c32[14924]^c32[14797]^c32[14798]^c32[15691]^c32[15373]^c32[15374]^c32[15399]^c32[15400]^c32[16475]^c32[16241]^c32[16485]^c32[16486];
+
+       break;
+    case 3:
+      d32[1152]=c32[308]^c32[77]^c32[206]^c32[276]^c32[277]^c32[788]^c32[1019]^c32[855]^c32[856]^c32[1587]^c32[1613]^c32[1865]^c32[2674]^c32[2378]^c32[2304]^c32[2305]^c32[3361]^c32[3405]^c32[3347]^c32[3348]^c32[4022]^c32[3985]^c32[4097]^c32[4825]^c32[4770]^c32[4807]^c32[4808]^c32[5708]^c32[5644]^c32[5529]^c32[5530]^c32[6476]^c32[6305]^c32[6200]^c32[6201]^c32[7230]^c32[7091]^c32[6976]^c32[7969]^c32[7810]^c32[7812]^c32[7813]^c32[8558]^c32[8744]^c32[8753]^c32[8754]^c32[9234]^c32[9559]^c32[9447]^c32[9448]^c32[10342]^c32[10185]^c32[10325]^c32[10326]^c32[10970]^c32[10841]^c32[10964]^c32[10965]^c32[11736]^c32[11620]^c32[11574]^c32[12395]^c32[12643]^c32[12592]^c32[12593]^c32[13171]^c32[13188]^c32[13356]^c32[13357]^c32[14067]^c32[14065]^c32[14095]^c32[14096]^c32[14773]^c32[14924]^c32[14798]^c32[15691]^c32[15374]^c32[15399]^c32[15400]^c32[16475]^c32[16241]^c32[16485]^c32[16486];
+
+       break;
+    case 4:
+      d32[1536]=c32[332]^c32[949];
+
+       break;
+    case 5:
+      d32[1920]=c32[80]^c32[233]^c32[362]^c32[49]^c32[195]^c32[944]^c32[791]^c32[1012]^c32[782]^c32[1743]^c32[1769]^c32[1637]^c32[2446]^c32[2534]^c32[2461]^c32[2419]^c32[3133]^c32[3177]^c32[3120]^c32[4178]^c32[4141]^c32[3869]^c32[4981]^c32[4926]^c32[4964]^c32[5480]^c32[5416]^c32[5686]^c32[6248]^c32[6461]^c32[6357]^c32[7002]^c32[7247]^c32[7132]^c32[7741]^c32[7966]^c32[7969]^c32[8714]^c32[8516]^c32[8526]^c32[9390]^c32[9331]^c32[9220]^c32[9382]^c32[10114]^c32[10341]^c32[10098]^c32[11126]^c32[10997]^c32[11121]^c32[11892]^c32[11776]^c32[11730]^c32[12551]^c32[12415]^c32[12365]^c32[12529]^c32[13327]^c32[13344]^c32[13129]^c32[13839]^c32[13837]^c32[13868]^c32[14929]^c32[14696]^c32[14954]^c32[15463]^c32[15530]^c32[15556]^c32[16247]^c32[16397]^c32[16258]^c32[16179];
+
+       break;
+    case 6:
+      d32[2304]=c32[278]^c32[4865]^c32[7681]^c32[8799]^c32[10076]^c32[13309]^c32[13842]^c32[15585];
+
+       break;
+    case 7:
+      d32[2688]=c32[9]^c32[830]^c32[3388]^c32[5709]^c32[6434]^c32[10866];
+
+       break;
+    case 8:
+      d32[3072]=c32[374]^c32[94]^c32[143]^c32[247]^c32[272]^c32[375]^c32[376]^c32[343]^c32[62]^c32[63]^c32[307]^c32[854]^c32[958]^c32[1085]^c32[804]^c32[805]^c32[922]^c32[1025]^c32[1026]^c32[947]^c32[1653]^c32[1757]^c32[1679]^c32[1783]^c32[1547]^c32[1650]^c32[1651]^c32[2356]^c32[2460]^c32[2444]^c32[2548]^c32[2371]^c32[2474]^c32[2475]^c32[2469]^c32[3427]^c32[3147]^c32[3087]^c32[3190]^c32[3191]^c32[3414]^c32[3133]^c32[3134]^c32[4088]^c32[4192]^c32[4051]^c32[4155]^c32[4163]^c32[3882]^c32[3883]^c32[4891]^c32[4611]^c32[4836]^c32[4939]^c32[4940]^c32[4874]^c32[4977]^c32[4978]^c32[5390]^c32[5494]^c32[5710]^c32[5429]^c32[5430]^c32[5596]^c32[5699]^c32[5700]^c32[6158]^c32[6262]^c32[6371]^c32[6474]^c32[6475]^c32[6267]^c32[6370]^c32[6371]^c32[6912]^c32[7016]^c32[7157]^c32[7261]^c32[7042]^c32[7145]^c32[7146]^c32[8035]^c32[7755]^c32[7876]^c32[7979]^c32[7980]^c32[7879]^c32[7982]^c32[7983]^c32[8624]^c32[8728]^c32[8810]^c32[8530]^c32[8820]^c32[8539]^c32[8540]^c32[9300]^c32[9404]^c32[9241]^c32[9345]^c32[9514]^c32[9233]^c32[9234]^c32[9234]^c32[10024]^c32[10128]^c32[10251]^c32[10354]^c32[10355]^c32[10008]^c32[10111]^c32[10112]^c32[11036]^c32[10756]^c32[10907]^c32[11010]^c32[11011]^c32[11031]^c32[11134]^c32[11135]^c32[11802]^c32[11522]^c32[11686]^c32[11790]^c32[11640]^c32[11743]^c32[11744]^c32[12461]^c32[12565]^c32[12325]^c32[12429]^c32[12659]^c32[12378]^c32[12379]^c32[12327]^c32[13237]^c32[13341]^c32[13254]^c32[13357]^c32[13358]^c32[13423]^c32[13142]^c32[13143]^c32[14133]^c32[13853]^c32[14131]^c32[13850]^c32[13851]^c32[14162]^c32[13881]^c32[13882]^c32[14839]^c32[14943]^c32[14606]^c32[14710]^c32[14864]^c32[14967]^c32[14968]^c32[14816]^c32[15373]^c32[15477]^c32[15440]^c32[15543]^c32[15544]^c32[15466]^c32[15569]^c32[15570]^c32[16157]^c32[16261]^c32[16307]^c32[16411]^c32[16168]^c32[16271]^c32[16272]^c32[16496];
+
+       break;
+    case 9:
+      d32[3456]=c32[366]^c32[1000]^c32[8001]^c32[8581]^c32[10041]^c32[13359]^c32[13887]^c32[15442];
+
+       break;
+    case 10:
+      d32[3840]=c32[869]^c32[1875]^c32[3346]^c32[5487]^c32[6527]^c32[11106];
+
+       break;
+    case 11:
+      d32[4224]=c32[257]^c32[38]^c32[39]^c32[26]^c32[192]^c32[155]^c32[321]^c32[226]^c32[8]^c32[48]^c32[1121]^c32[902]^c32[903]^c32[968]^c32[1134]^c32[805]^c32[971]^c32[870]^c32[1536]^c32[1701]^c32[1702]^c32[1562]^c32[1728]^c32[1814]^c32[1596]^c32[2623]^c32[2404]^c32[2405]^c32[2327]^c32[2493]^c32[2638]^c32[2420]^c32[3310]^c32[3092]^c32[3354]^c32[3136]^c32[3297]^c32[3079]^c32[3971]^c32[4136]^c32[4137]^c32[3934]^c32[4100]^c32[4046]^c32[4212]^c32[4774]^c32[4939]^c32[4940]^c32[4719]^c32[4885]^c32[4757]^c32[4923]^c32[5657]^c32[5439]^c32[5593]^c32[5759]^c32[5479]^c32[5645]^c32[6425]^c32[6207]^c32[6254]^c32[6420]^c32[6150]^c32[6316]^c32[7179]^c32[6960]^c32[6961]^c32[7040]^c32[7206]^c32[6925]^c32[7091]^c32[7918]^c32[7699]^c32[7700]^c32[7759]^c32[7925]^c32[7762]^c32[7928]^c32[8507]^c32[8672]^c32[8673]^c32[8693]^c32[8475]^c32[8703]^c32[8485]^c32[9567]^c32[9348]^c32[9349]^c32[9508]^c32[9290]^c32[9397]^c32[9563]^c32[9224]^c32[10291]^c32[10072]^c32[10073]^c32[10134]^c32[10300]^c32[10275]^c32[10057]^c32[10919]^c32[11085]^c32[10790]^c32[10956]^c32[10914]^c32[11080]^c32[11685]^c32[11850]^c32[11851]^c32[11569]^c32[11735]^c32[11523]^c32[11689]^c32[12344]^c32[12509]^c32[12510]^c32[12592]^c32[12374]^c32[12542]^c32[12324]^c32[12335]^c32[13120]^c32[13286]^c32[13137]^c32[13303]^c32[13306]^c32[13088]^c32[14016]^c32[14181]^c32[14182]^c32[14014]^c32[14180]^c32[14045]^c32[13827]^c32[14722]^c32[14887]^c32[14888]^c32[14873]^c32[14655]^c32[14747]^c32[14913]^c32[15640]^c32[15421]^c32[15422]^c32[15707]^c32[15489]^c32[15733]^c32[15515]^c32[16424]^c32[16205]^c32[16206]^c32[16190]^c32[16356]^c32[16435]^c32[16217]^c32[16316];
+
+       break;
+    case 12:
+      d32[4608]=c32[77]^c32[954]^c32[7854]^c32[8680]^c32[10034]^c32[13898];
+
+       break;
+    case 13:
+      d32[4992]=c32[293]^c32[294]^c32[63]^c32[192]^c32[263]^c32[313]^c32[773]^c32[774]^c32[1005]^c32[842]^c32[1572]^c32[1573]^c32[1599]^c32[1851]^c32[2659]^c32[2660]^c32[2364]^c32[2675]^c32[2481]^c32[3347]^c32[3391]^c32[3334]^c32[4007]^c32[4008]^c32[3971]^c32[4083]^c32[4810]^c32[4811]^c32[4756]^c32[4794]^c32[5694]^c32[5630]^c32[5516]^c32[5642]^c32[6462]^c32[6291]^c32[6187]^c32[7215]^c32[7216]^c32[7077]^c32[6962]^c32[7954]^c32[7955]^c32[7796]^c32[7799]^c32[8543]^c32[8544]^c32[8730]^c32[8740]^c32[9219]^c32[9220]^c32[9545]^c32[9434]^c32[10327]^c32[10328]^c32[10171]^c32[10312]^c32[10956]^c32[10827]^c32[10951]^c32[11721]^c32[11722]^c32[11606]^c32[11560]^c32[12380]^c32[12381]^c32[12629]^c32[12579]^c32[13157]^c32[13174]^c32[13343]^c32[14052]^c32[14053]^c32[14051]^c32[14082]^c32[14758]^c32[14759]^c32[14910]^c32[14784]^c32[15676]^c32[15677]^c32[15360]^c32[15386]^c32[15475]^c32[16460]^c32[16461]^c32[16227]^c32[16472];
+
+       break;
+    case 14:
+      d32[5376]=c32[142]^c32[9464]^c32[11657]^c32[12377]^c32[13403]^c32[16140];
+
+       break;
+    case 15:
+      d32[5760]=c32[193]^c32[346]^c32[91]^c32[161]^c32[162]^c32[241]^c32[1057]^c32[904]^c32[1124]^c32[1125]^c32[770]^c32[1856]^c32[1882]^c32[1750]^c32[2559]^c32[2647]^c32[2573]^c32[2574]^c32[3246]^c32[3290]^c32[3232]^c32[3233]^c32[3907]^c32[3870]^c32[3982]^c32[4710]^c32[4655]^c32[4692]^c32[4693]^c32[5593]^c32[5529]^c32[5414]^c32[5415]^c32[6361]^c32[6190]^c32[6469]^c32[6470]^c32[7115]^c32[6976]^c32[7245]^c32[7854]^c32[7695]^c32[7697]^c32[7698]^c32[7890]^c32[8827]^c32[8629]^c32[8638]^c32[8639]^c32[9503]^c32[9444]^c32[9332]^c32[9333]^c32[10227]^c32[10070]^c32[10210]^c32[10211]^c32[10302]^c32[10855]^c32[11110]^c32[10849]^c32[10850]^c32[11621]^c32[11889]^c32[11843]^c32[12664]^c32[12528]^c32[12477]^c32[12478]^c32[13056]^c32[13073]^c32[13241]^c32[13242]^c32[13952]^c32[13950]^c32[13980]^c32[13981]^c32[13879]^c32[14658]^c32[14809]^c32[14683]^c32[15576]^c32[15643]^c32[15668]^c32[15669]^c32[16360]^c32[16510]^c32[16370]^c32[16371];
+
+       break;
+    case 16:
+      d32[6144]=c32[364]^c32[133]^c32[262]^c32[333]^c32[844]^c32[1075]^c32[912]^c32[781]^c32[1643]^c32[1669]^c32[1537]^c32[2346]^c32[2434]^c32[2361]^c32[2642]^c32[3417]^c32[3077]^c32[3404]^c32[4078]^c32[4041]^c32[4153]^c32[4881]^c32[4826]^c32[4864]^c32[5380]^c32[5700]^c32[5586]^c32[6148]^c32[6361]^c32[6257]^c32[7286]^c32[7147]^c32[7032]^c32[8025]^c32[7866]^c32[7869]^c32[8614]^c32[8800]^c32[8810]^c32[8505]^c32[9290]^c32[9231]^c32[9504]^c32[10014]^c32[10241]^c32[9998]^c32[11026]^c32[10897]^c32[11021]^c32[11792]^c32[11676]^c32[11630]^c32[12451]^c32[12315]^c32[12649]^c32[13227]^c32[13244]^c32[13413]^c32[14123]^c32[14121]^c32[14152]^c32[14829]^c32[14596]^c32[14854]^c32[15363]^c32[15430]^c32[15456]^c32[15649]^c32[16147]^c32[16297]^c32[16158];
+
+       break;
+    case 17:
+      d32[6528]=c32[260]^c32[11055]^c32[12369]^c32[13414]^c32[16503];
+
+       break;
+    case 18:
+      d32[6912]=c32[898]^c32[9379]^c32[10264]^c32[13956]^c32[14596];
+
+       break;
+    case 19:
+      d32[7296]=c32[145]^c32[981]^c32[5720]^c32[6386]^c32[7877];
+
+       break;
+    case 20:
+      d32[7680]=c32[366]^c32[135]^c32[264]^c32[335]^c32[187]^c32[846]^c32[1077]^c32[914]^c32[1645]^c32[1671]^c32[1539]^c32[2348]^c32[2436]^c32[2363]^c32[2510]^c32[3419]^c32[3079]^c32[3406]^c32[4080]^c32[4043]^c32[4155]^c32[4883]^c32[4828]^c32[4866]^c32[5382]^c32[5702]^c32[5588]^c32[6150]^c32[6363]^c32[6259]^c32[7288]^c32[7149]^c32[7034]^c32[7176]^c32[8027]^c32[7868]^c32[7871]^c32[8616]^c32[8802]^c32[8812]^c32[8789]^c32[9292]^c32[9233]^c32[9506]^c32[10016]^c32[10243]^c32[10000]^c32[11028]^c32[10899]^c32[11023]^c32[11794]^c32[11678]^c32[11632]^c32[12453]^c32[12317]^c32[12651]^c32[13229]^c32[13246]^c32[13415]^c32[14125]^c32[14123]^c32[14154]^c32[14831]^c32[14598]^c32[14856]^c32[15365]^c32[15432]^c32[15458]^c32[16149]^c32[16299]^c32[16160];
+
+       break;
+    case 21:
+      d32[8064]=c32[973]^c32[3942]^c32[12616]^c32[15573]^c32[16225];
+
+       break;
+    case 22:
+      d32[8448]=c32[30]^c32[9227]^c32[10217]^c32[13078];
+
+       break;
+    case 23:
+      d32[8832]=c32[792]^c32[1625]^c32[7741]^c32[13851];
+
+       break;
+    case 24:
+      d32[9216]=c32[157]^c32[310]^c32[55]^c32[126]^c32[298]^c32[1021]^c32[868]^c32[1089]^c32[1820]^c32[1846]^c32[1714]^c32[2523]^c32[2611]^c32[2538]^c32[2462]^c32[3210]^c32[3254]^c32[3197]^c32[3307]^c32[3871]^c32[4218]^c32[3946]^c32[4674]^c32[4619]^c32[4657]^c32[5557]^c32[5493]^c32[5379]^c32[6325]^c32[6154]^c32[6434]^c32[7079]^c32[6940]^c32[7209]^c32[7818]^c32[8043]^c32[8046]^c32[8791]^c32[8593]^c32[8603]^c32[8787]^c32[9467]^c32[9408]^c32[9297]^c32[10191]^c32[10034]^c32[10175]^c32[10819]^c32[11074]^c32[10814]^c32[11585]^c32[11853]^c32[11807]^c32[12628]^c32[12492]^c32[12442]^c32[13404]^c32[13421]^c32[13206]^c32[13916]^c32[13914]^c32[13945]^c32[14622]^c32[14773]^c32[14647]^c32[15540]^c32[15607]^c32[15633]^c32[16324]^c32[16474]^c32[16335];
+
+       break;
+    case 25:
+      d32[9600]=c32[840]^c32[4625]^c32[5759]^c32[11064];
+
+       break;
+    case 26:
+      d32[9984]=c32[71]^c32[1617]^c32[3148]^c32[11656];
+
+       break;
+    case 27:
+      d32[10368]=c32[962]^c32[4802]^c32[6245];
+
+       break;
+    case 28:
+      d32[10752]=c32[222]^c32[3091]^c32[14836]^c32[16402];
+
+       break;
+    case 29:
+      d32[11136]=c32[2]^c32[155]^c32[284]^c32[354]^c32[355]^c32[866]^c32[1097]^c32[933]^c32[934]^c32[1020]^c32[1665]^c32[1691]^c32[1559]^c32[2368]^c32[2456]^c32[2382]^c32[2383]^c32[3439]^c32[3099]^c32[3425]^c32[3426]^c32[4100]^c32[4063]^c32[4175]^c32[4903]^c32[4848]^c32[4885]^c32[4886]^c32[5402]^c32[5722]^c32[5607]^c32[5608]^c32[6170]^c32[6383]^c32[6278]^c32[6279]^c32[6924]^c32[7169]^c32[7054]^c32[8047]^c32[7888]^c32[7890]^c32[7891]^c32[8636]^c32[8822]^c32[8831]^c32[8448]^c32[9312]^c32[9253]^c32[9525]^c32[9526]^c32[10036]^c32[10263]^c32[10019]^c32[10020]^c32[11048]^c32[10919]^c32[11042]^c32[11043]^c32[10757]^c32[11814]^c32[11698]^c32[11652]^c32[12473]^c32[12337]^c32[12670]^c32[12671]^c32[13249]^c32[13266]^c32[13434]^c32[13435]^c32[14145]^c32[14143]^c32[14173]^c32[14174]^c32[13971]^c32[14851]^c32[14618]^c32[14876]^c32[15385]^c32[15452]^c32[15477]^c32[15478]^c32[16169]^c32[16319]^c32[16179]^c32[16180];
+
+       break;
+    case 30:
+      d32[11520]=c32[14]^c32[167]^c32[295]^c32[296]^c32[366]^c32[367]^c32[159]^c32[878]^c32[1108]^c32[1109]^c32[945]^c32[946]^c32[1677]^c32[1703]^c32[1570]^c32[1571]^c32[2380]^c32[2468]^c32[2394]^c32[2395]^c32[3451]^c32[3110]^c32[3111]^c32[3437]^c32[3438]^c32[4112]^c32[4075]^c32[4186]^c32[4187]^c32[4915]^c32[4859]^c32[4860]^c32[4897]^c32[4898]^c32[5414]^c32[5733]^c32[5734]^c32[5619]^c32[5620]^c32[6182]^c32[6394]^c32[6395]^c32[6290]^c32[6291]^c32[6936]^c32[7181]^c32[7065]^c32[7066]^c32[8059]^c32[7899]^c32[7900]^c32[7902]^c32[7903]^c32[7909]^c32[8648]^c32[8450]^c32[8459]^c32[8460]^c32[9324]^c32[9265]^c32[9537]^c32[9538]^c32[10048]^c32[10274]^c32[10275]^c32[10031]^c32[10032]^c32[10244]^c32[11060]^c32[10930]^c32[10931]^c32[11054]^c32[11055]^c32[11826]^c32[11710]^c32[11663]^c32[11664]^c32[12485]^c32[12349]^c32[12298]^c32[12299]^c32[13261]^c32[13277]^c32[13278]^c32[13062]^c32[13063]^c32[14157]^c32[14154]^c32[14155]^c32[14185]^c32[14186]^c32[14863]^c32[14630]^c32[14887]^c32[14888]^c32[15397]^c32[15463]^c32[15464]^c32[15489]^c32[15490]^c32[16181]^c32[16331]^c32[16191]^c32[16192];
+
+       break;
+    case 31:
+      d32[11904]=c32[181]^c32[180]^c32[334]^c32[333]^c32[79]^c32[78]^c32[150]^c32[148]^c32[149]^c32[1045]^c32[1044]^c32[892]^c32[891]^c32[1113]^c32[1111]^c32[1112]^c32[868]^c32[1844]^c32[1843]^c32[1870]^c32[1869]^c32[1738]^c32[1737]^c32[2547]^c32[2546]^c32[2635]^c32[2634]^c32[2562]^c32[2560]^c32[2561]^c32[3234]^c32[3233]^c32[3278]^c32[3277]^c32[3221]^c32[3219]^c32[3220]^c32[3895]^c32[3894]^c32[3858]^c32[3857]^c32[3970]^c32[3969]^c32[4698]^c32[4697]^c32[4643]^c32[4642]^c32[4681]^c32[4679]^c32[4680]^c32[5581]^c32[5580]^c32[5517]^c32[5516]^c32[5403]^c32[5401]^c32[5402]^c32[5591]^c32[6349]^c32[6348]^c32[6178]^c32[6177]^c32[6458]^c32[6456]^c32[6457]^c32[7103]^c32[7102]^c32[6964]^c32[6963]^c32[7233]^c32[7232]^c32[7842]^c32[7841]^c32[7683]^c32[7682]^c32[7686]^c32[7684]^c32[7685]^c32[8815]^c32[8814]^c32[8617]^c32[8616]^c32[8627]^c32[8625]^c32[8626]^c32[9491]^c32[9490]^c32[9432]^c32[9431]^c32[9321]^c32[9319]^c32[9320]^c32[10215]^c32[10214]^c32[10058]^c32[10057]^c32[10199]^c32[10197]^c32[10198]^c32[10843]^c32[10842]^c32[11098]^c32[11097]^c32[10838]^c32[10836]^c32[10837]^c32[11609]^c32[11608]^c32[11877]^c32[11876]^c32[11831]^c32[11830]^c32[12652]^c32[12651]^c32[12516]^c32[12515]^c32[12466]^c32[12464]^c32[12465]^c32[13428]^c32[13427]^c32[13061]^c32[13060]^c32[13230]^c32[13228]^c32[13229]^c32[13940]^c32[13939]^c32[13938]^c32[13937]^c32[13969]^c32[13967]^c32[13968]^c32[14646]^c32[14645]^c32[14797]^c32[14796]^c32[14671]^c32[14670]^c32[15564]^c32[15563]^c32[15631]^c32[15630]^c32[15657]^c32[15655]^c32[15656]^c32[16348]^c32[16347]^c32[16498]^c32[16497]^c32[16359]^c32[16357]^c32[16358];
+
+       break;
+    case 32:
+      d32[12288]=c32[211]^c32[364]^c32[108]^c32[109]^c32[179]^c32[180]^c32[102]^c32[1075]^c32[921]^c32[922]^c32[1142]^c32[1143]^c32[1874]^c32[1900]^c32[1767]^c32[1768]^c32[2577]^c32[2665]^c32[2591]^c32[2592]^c32[3264]^c32[3307]^c32[3308]^c32[3250]^c32[3251]^c32[3925]^c32[3888]^c32[3999]^c32[4000]^c32[4728]^c32[4672]^c32[4673]^c32[4710]^c32[4711]^c32[5611]^c32[5546]^c32[5547]^c32[5432]^c32[5433]^c32[6379]^c32[6207]^c32[6208]^c32[6487]^c32[6488]^c32[7133]^c32[6994]^c32[7262]^c32[7263]^c32[7872]^c32[7712]^c32[7713]^c32[7715]^c32[7716]^c32[8461]^c32[8647]^c32[8656]^c32[8657]^c32[9521]^c32[9462]^c32[9350]^c32[9351]^c32[9417]^c32[10245]^c32[10087]^c32[10088]^c32[10228]^c32[10229]^c32[10873]^c32[11127]^c32[11128]^c32[10867]^c32[10868]^c32[10927]^c32[11639]^c32[11523]^c32[11860]^c32[11861]^c32[12298]^c32[12546]^c32[12495]^c32[12496]^c32[13074]^c32[13090]^c32[13091]^c32[13259]^c32[13260]^c32[13970]^c32[13967]^c32[13968]^c32[13998]^c32[13999]^c32[14676]^c32[14827]^c32[14700]^c32[14701]^c32[15594]^c32[15660]^c32[15661]^c32[15686]^c32[15687]^c32[16378]^c32[16144]^c32[16388]^c32[16389];
+
+       break;
+    case 33:
+      d32[12672]=c32[1091]^c32[1544]^c32[8809]^c32[16233];
+
+       break;
+    case 34:
+      d32[13056]=c32[230]^c32[5524]^c32[11722]^c32[13368];
+
+       break;
+    case 35:
+      d32[13440]=c32[189]^c32[342]^c32[87]^c32[158]^c32[1053]^c32[900]^c32[1121]^c32[1088]^c32[1852]^c32[1878]^c32[1746]^c32[2555]^c32[2643]^c32[2570]^c32[3242]^c32[3286]^c32[3229]^c32[3903]^c32[3866]^c32[3978]^c32[4706]^c32[4651]^c32[4689]^c32[4943]^c32[5589]^c32[5525]^c32[5411]^c32[6357]^c32[6186]^c32[6466]^c32[7111]^c32[6972]^c32[7241]^c32[7850]^c32[7691]^c32[7694]^c32[8823]^c32[8625]^c32[8635]^c32[9499]^c32[9440]^c32[9329]^c32[9218]^c32[10223]^c32[10066]^c32[10207]^c32[10851]^c32[11106]^c32[10846]^c32[11617]^c32[11885]^c32[11839]^c32[12660]^c32[12524]^c32[12474]^c32[13436]^c32[13069]^c32[13238]^c32[13948]^c32[13946]^c32[13977]^c32[14654]^c32[14805]^c32[14679]^c32[15572]^c32[15639]^c32[15665]^c32[16356]^c32[16506]^c32[16367];
+
+       break;
+    case 36:
+      d32[13824]=c32[210]^c32[11065]^c32[11817]^c32[13845];
+
+       break;
+    case 37:
+      d32[14208]=c32[38]^c32[39]^c32[192]^c32[321]^c32[8]^c32[902]^c32[903]^c32[1134]^c32[971]^c32[1037]^c32[1701]^c32[1702]^c32[1728]^c32[1596]^c32[2404]^c32[2405]^c32[2493]^c32[2420]^c32[3092]^c32[3136]^c32[3079]^c32[4136]^c32[4137]^c32[4100]^c32[4212]^c32[4939]^c32[4940]^c32[4885]^c32[4923]^c32[5439]^c32[5759]^c32[5645]^c32[6207]^c32[6420]^c32[6316]^c32[6960]^c32[6961]^c32[7206]^c32[7091]^c32[7699]^c32[7700]^c32[7925]^c32[7928]^c32[8672]^c32[8673]^c32[8475]^c32[8485]^c32[9348]^c32[9349]^c32[9290]^c32[9563]^c32[10072]^c32[10073]^c32[10300]^c32[10057]^c32[10066]^c32[11085]^c32[10956]^c32[11080]^c32[11850]^c32[11851]^c32[11735]^c32[11689]^c32[12509]^c32[12510]^c32[12374]^c32[12324]^c32[13286]^c32[13303]^c32[13088]^c32[14181]^c32[14182]^c32[14180]^c32[13827]^c32[14887]^c32[14888]^c32[14655]^c32[14913]^c32[15421]^c32[15422]^c32[15489]^c32[15515]^c32[16205]^c32[16206]^c32[16356]^c32[16217];
+
+       break;
+    case 38:
+      d32[14592]=c32[185]^c32[7089]^c32[7969]^c32[9430];
+
+       break;
+    case 39:
+      d32[14976]=c32[1026]^c32[2397]^c32[5722]^c32[14889];
+
+       break;
+    case 40:
+      d32[15360]=c32[175]^c32[6181]^c32[13368];
+
+       break;
+    case 41:
+      d32[15744]=c32[820]^c32[2618]^c32[7051]^c32[14112];
+
+       break;
+    case 42:
+      d32[16128]=c32[142]^c32[295]^c32[39]^c32[40]^c32[110]^c32[111]^c32[113]^c32[1006]^c32[852]^c32[853]^c32[1073]^c32[1074]^c32[1805]^c32[1831]^c32[1698]^c32[1699]^c32[2508]^c32[2596]^c32[2522]^c32[2523]^c32[3195]^c32[3238]^c32[3239]^c32[3181]^c32[3182]^c32[3086]^c32[3856]^c32[4203]^c32[3930]^c32[3931]^c32[4659]^c32[4987]^c32[4988]^c32[4641]^c32[4642]^c32[5542]^c32[5477]^c32[5478]^c32[5747]^c32[5748]^c32[6310]^c32[6522]^c32[6523]^c32[6418]^c32[6419]^c32[7064]^c32[6925]^c32[7193]^c32[7194]^c32[7803]^c32[8027]^c32[8028]^c32[8030]^c32[8031]^c32[8776]^c32[8578]^c32[8587]^c32[8588]^c32[9452]^c32[9393]^c32[9281]^c32[9282]^c32[10176]^c32[10018]^c32[10019]^c32[10159]^c32[10160]^c32[10804]^c32[11058]^c32[11059]^c32[10798]^c32[10799]^c32[11570]^c32[11838]^c32[11791]^c32[11792]^c32[12613]^c32[12477]^c32[12426]^c32[12427]^c32[13389]^c32[13405]^c32[13406]^c32[13190]^c32[13191]^c32[13901]^c32[13898]^c32[13899]^c32[13929]^c32[13930]^c32[14607]^c32[14758]^c32[14631]^c32[14632]^c32[15525]^c32[15591]^c32[15592]^c32[15617]^c32[15618]^c32[16309]^c32[16459]^c32[16319]^c32[16320];
+
+       break;
+    case 43:
+      d32[16512]=c32[92]^c32[245]^c32[374]^c32[60]^c32[61]^c32[956]^c32[803]^c32[1023]^c32[1024]^c32[881]^c32[1755]^c32[1781]^c32[1649]^c32[2458]^c32[2546]^c32[2472]^c32[2473]^c32[3145]^c32[3189]^c32[3131]^c32[3132]^c32[4190]^c32[4153]^c32[3881]^c32[4609]^c32[4938]^c32[4975]^c32[4976]^c32[5492]^c32[5428]^c32[5697]^c32[5698]^c32[6260]^c32[6473]^c32[6368]^c32[6369]^c32[7014]^c32[7259]^c32[7144]^c32[7753]^c32[7978]^c32[7980]^c32[7981]^c32[8726]^c32[8528]^c32[8537]^c32[8538]^c32[9402]^c32[9343]^c32[9231]^c32[9232]^c32[10126]^c32[10353]^c32[10109]^c32[10110]^c32[10754]^c32[11009]^c32[11132]^c32[11133]^c32[11520]^c32[11788]^c32[11742]^c32[12563]^c32[12427]^c32[12376]^c32[12377]^c32[12420]^c32[13339]^c32[13356]^c32[13140]^c32[13141]^c32[13851]^c32[13849]^c32[13879]^c32[13880]^c32[13938]^c32[14941]^c32[14708]^c32[14966]^c32[15475]^c32[15542]^c32[15567]^c32[15568]^c32[16259]^c32[16409]^c32[16269]^c32[16270];
+
+       break;
+    case 44:
+      d32[16896]=c32[197]^c32[350]^c32[95]^c32[166]^c32[80]^c32[1061]^c32[908]^c32[1129]^c32[1860]^c32[1886]^c32[1754]^c32[2563]^c32[2651]^c32[2578]^c32[3250]^c32[3294]^c32[3237]^c32[3911]^c32[3874]^c32[3986]^c32[4714]^c32[4659]^c32[4697]^c32[5597]^c32[5533]^c32[5419]^c32[5454]^c32[6365]^c32[6194]^c32[6474]^c32[7119]^c32[6980]^c32[7249]^c32[7075]^c32[7858]^c32[7699]^c32[7702]^c32[8831]^c32[8633]^c32[8643]^c32[9507]^c32[9448]^c32[9337]^c32[10231]^c32[10074]^c32[10215]^c32[10859]^c32[11114]^c32[10854]^c32[11625]^c32[11893]^c32[11847]^c32[12668]^c32[12532]^c32[12482]^c32[13060]^c32[13077]^c32[13246]^c32[13956]^c32[13954]^c32[13985]^c32[14662]^c32[14813]^c32[14687]^c32[15580]^c32[15647]^c32[15673]^c32[16364]^c32[16130]^c32[16375];
+
+       break;
+    case 45:
+      d32[17280]=c32[903]^c32[4757]^c32[7695];
+
+       break;
+    
+    }
+  }
+}
+extern "C" int ldpc_BG1_Zc384_cuda32(uint32_t *c[4],uint32_t *d[4],int n_inputs) { 
+
+ dim3 numblocks(n_inputs,46);
+ ldpc_BG1_Zc384_worker<<<numblocks,384>>>(c,d);
+ cudaError_t err=cudaPeekAtLastError();
+ if (err!=cudaSuccess) {
+    printf("cuda error: %s (c %p, d %p)\n",cudaGetErrorString(err),c,d);
+    exit(-1);
+ }
+ cudaDeviceSynchronize();
+ return(0);
+}

openair1/PHY/CODING/nrLDPC_encoder/ldpc_encode_parity_check.c

@@ -103,9 +103,11 @@
 #include "ldpc_BG2_Zc80_byte.c"
 #include "ldpc_BG2_Zc72_byte.c"
 
-static void encode_parity_check_part_optim(uint8_t *cc, uint8_t *d, short BG,short Zc,short Kb, int simd_size, int ncols)
+
+static void encode_parity_check_part_optim(uint8_t *cc, uint8_t *d, short BG,short Zc,short Kb, int simd_size, int ncols,time_stats_t *tinput_memcpy)
 {
   unsigned char c[2 * 22 * Zc * simd_size] __attribute__((aligned(64))); //double size matrix of c
+  if (tinput_memcpy) start_meas(tinput_memcpy);
   for (int i1 = 0; i1 < ncols; i1++)   {
     memcpy(&c[2 * i1 * Zc], &cc[i1 * Zc], Zc * sizeof(unsigned char));
     memcpy(&c[(2 * i1 + 1) * Zc], &cc[i1 * Zc], Zc * sizeof(unsigned char));
@@ -118,6 +120,7 @@ static void encode_parity_check_part_optim(uint8_t *cc, uint8_t *d, short BG,sho
       memcpy(&c[(2 * ncols * Zc * i1)], &c[i1], (2 * ncols * Zc * sizeof(unsigned char)) - i1);
     }
   }
+  if (tinput_memcpy) stop_meas(tinput_memcpy);
   if (BG == 1) {
     switch (Zc) {
       case 176:

openair1/PHY/CODING/nrLDPC_encoder/ldpc_encode_parity_check_cuda.c

@@ -0,0 +1,94 @@
+/*
+ * Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
+ * contributor license agreements.  See the NOTICE file distributed with
+ * this work for additional information regarding copyright ownership.
+ * The OpenAirInterface Software Alliance licenses this file to You under
+ * the OAI Public License, Version 1.1  (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.openairinterface.org/?page_id=698
+ *
+ * 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.
+ *-------------------------------------------------------------------------------
+ * For more information about the OpenAirInterface (OAI) Software Alliance:
+ *      contact@openairinterface.org
+ */
+
+/*!\file ldpc_encode_parity_check.c
+ * \brief Parity check function used by ldpc encoders
+ * \author Florian Kaltenberger, Raymond Knopp, Kien le Trung (Eurecom)
+ * \email openair_tech@eurecom.fr
+ * \date 27-03-2018
+ * \version 1.0
+ * \note
+ * \warning
+ */
+
+#include <stdlib.h>
+#include <math.h>
+#include <stdio.h>
+#include <string.h>
+#include "assertions.h"
+#include "common/utils/LOG/log.h"
+#include <cuda_runtime.h>
+
+
+
+int ldpc_BG1_Zc384_cuda32(uint32_t **c,uint32_t **d,int n_inputs);
+
+
+void encode_parity_check_part_cuda(uint32_t **c, uint32_t **d, short BG,short Zc,short Kb, int ncols, int n_inputs)
+{
+  
+  if (BG == 1) {
+    switch (Zc) {
+      case 176:
+      case 192:
+      case 208:
+      case 224:
+      case 240:
+      case 256:
+      case 288:
+      case 320:
+      case 352:
+	AssertFatal(1==0,"BG %d Zc %d not supported yet for CUDA\n",BG, Zc);
+        break;
+      case 384:
+	ldpc_BG1_Zc384_cuda32(c, d, n_inputs);
+        break;
+      default:
+        AssertFatal(false, "BG %d Zc %d is not supported yet\n", BG, Zc);
+    }
+  } else if (BG == 2) {
+    switch (Zc) {
+      case 72:
+      case 80:
+      case 88:
+      case 96:
+      case 104:
+      case 112:
+      case 120:
+      case 128:
+      case 144:
+      case 160:
+      case 176:
+      case 192:
+      case 208:
+      case 224:
+      case 240:
+      case 256:
+      case 288:
+      case 320:
+      case 352:
+      case 384:
+      default:
+        AssertFatal(false , "BG %d Zc %d is not supported yet\n", BG, Zc);
+    }
+  } else
+    AssertFatal(false, "BG %d is not supported\n", BG);
+}

openair1/PHY/CODING/nrLDPC_encoder/ldpc_encoder.c

@@ -58,11 +58,10 @@ int LDPCencoder(unsigned char **inputArray, unsigned char *outputArray, encoder_
   int no_punctured_columns, removed_bit;
   int nind=0;
   int indlist[1000];
-  int indlist2[1000];
 
   const short *Gen_shift_values = choose_generator_matrix(BG, Zc);
   if (Gen_shift_values==NULL) {
-    printf("ldpc_encoder_orig: could not find generator matrix\n");
+    printf("ldpc_encoder_orig: could not find generator matrix BG %d Zc %d\n",BG,Zc);
     return(-1);
   }
 
@@ -93,32 +92,35 @@ int LDPCencoder(unsigned char **inputArray, unsigned char *outputArray, encoder_
 
   // parity check part
 
-  if (gen_code>=1)
+  if (gen_code==1 || gen_code==2 || gen_code==3)
   {
     char fname[100];
     sprintf(fname,"ldpc_BG%d_Zc%d_byte.c",BG,Zc);
     FILE *fd=fopen(fname,"w");
     AssertFatal(fd!=NULL,"cannot open %s\n",fname);
-    sprintf(fname,"ldpc_BG%d_Zc%d_16bit.c",BG,Zc);
-    FILE *fd2=fopen(fname,"w");
-    AssertFatal(fd2!=NULL,"cannot open %s\n",fname);
+
 
     int shift;
     char data_type[100];
     char xor_command[100];
     int mask;
-
-
+    char permutex_command[100];
+    int use_permutex=0;
+    char alignr_command[100];
+    int use_alignr=0;
 
 
     fprintf(fd,"#include \"PHY/sse_intrin.h\"\n");
-    fprintf(fd2,"#include \"PHY/sse_intrin.h\"\n");
+
 
     if (gen_code == 1 && (Zc&63)==0) {
       shift=6;
       mask=63;
       strcpy(data_type,"__m512i");
       strcpy(xor_command,"_mm512_xor_si512");
+      strcpy(permutex_command,"_mm512_permutex2var_epi8");
+      strcpy(alignr_command,"_mm512_alignr_epi8");
+      use_alignr=1;
     }
     else if (gen_code == 1 && (Zc&31)==0) {
       shift=5; // AVX2 - 256-bit SIMD
@@ -126,14 +128,15 @@ int LDPCencoder(unsigned char **inputArray, unsigned char *outputArray, encoder_
       strcpy(data_type,"simde__m256i");
       strcpy(xor_command,"simde_mm256_xor_si256");
     }
-    else if ((Zc&15)==0) {
+    else if ((gen_code <=2) && (Zc&15)==0) {
       shift=4; // SSE4 - 128-bit SIMD
       mask=15;
       strcpy(data_type,"simde__m128i");
       strcpy(xor_command,"simde_mm_xor_si128");
-
+      strcpy(alignr_command,"simde_mm_alignr_epi8");
+      use_alignr=1;
     }
-    else if ((Zc&7)==0) {
+    else if ((gen_code <=2) && (Zc&7)==0) {
       shift=3; // MMX  - 64-bit SIMD
       mask=7;
       strcpy(data_type,"simde__m64");
@@ -145,21 +148,30 @@ int LDPCencoder(unsigned char **inputArray, unsigned char *outputArray, encoder_
       strcpy(data_type,"uint8_t");
       strcpy(xor_command,"scalar_xor");
       fprintf(fd,"#define scalar_xor(a,b) ((a)^(b))\n");
-      fprintf(fd2,"#define scalar_xor(a,b) ((a)^(b))\n");
+
     }
     fprintf(fd,"// generated code for Zc=%d, byte encoding\n",Zc);
-    fprintf(fd2,"// generated code for Zc=%d, 16bit encoding\n",Zc);
+    if (use_permutex==1) {
+	 fprintf(fd,"  const uint8_t abshift_bytes[%d][%d] __attribute__((aligned(64)))= {\n",1+mask,1+mask); 
+         for (int i=0;i<=mask;i++) {
+	    fprintf(fd,"{");
+	    for (int j=0;j<=mask;j++) {
+              fprintf(fd,"%d",i+j);
+	      if (j<mask) fprintf(fd,",");
+	      else fprintf(fd,"}");
+	    }
+	    if (i<mask) fprintf(fd,",\n");
+	    else fprintf(fd,"};\n");
+      	}
+	fprintf(fd,"static inline __m512i abshift_get(int r)\n {\n");
+        fprintf(fd,"   return * (const __m512i *) abshift_bytes[r];\n}");
+    }
     fprintf(fd,"static inline void ldpc_BG%d_Zc%d_byte(uint8_t *c,uint8_t *d) {\n",BG,Zc);
-    fprintf(fd2,"static inline void ldpc_BG%d_Zc%d_16bit(uint16_t *c,uint16_t *d) {\n",BG,Zc);
     fprintf(fd,"  %s *csimd=(%s *)c,*dsimd=(%s *)d;\n\n",data_type,data_type,data_type);
-    fprintf(fd2,"  %s *csimd=(%s *)c,*dsimd=(%s *)d;\n\n",data_type,data_type,data_type);
     fprintf(fd,"  %s *c2,*d2;\n\n",data_type);
-    fprintf(fd2,"  %s *c2,*d2;\n\n",data_type);
     fprintf(fd,"  int i2;\n");
-    fprintf(fd2,"  int i2;\n");
     fprintf(fd,"  for (i2=0; i2<%d; i2++) {\n",Zc>>shift);
-    if (shift > 0)
-      fprintf(fd2,"  for (i2=0; i2<%d; i2++) {\n",Zc>>(shift-1));
+
     for (i2=0; i2 < 1; i2++)
     {
       //t=Kb*Zc+i2;
@@ -169,53 +181,122 @@ int LDPCencoder(unsigned char **inputArray, unsigned char *outputArray, encoder_
 
       fprintf(fd,"     c2=&csimd[i2];\n");
       fprintf(fd,"     d2=&dsimd[i2];\n");
-      fprintf(fd2,"     c2=&csimd[i2];\n");
-      fprintf(fd2,"     d2=&dsimd[i2];\n");
 
       for (i1 = 0; i1 < nrows; i1++)
-
       {
         fprintf(fd,"\n//row: %d\n",i1);
-        fprintf(fd2,"\n//row: %d\n",i1);
-        AssertFatal(shift > 0 , "The result of the right shift is undefined because the right operand is negative\n");
+        AssertFatal(shift >= 0 , "The result of the right shift is undefined because the right operand is negative\n");
 	fprintf(fd,"     d2[%d]=",(Zc*i1)>>shift);
-	fprintf(fd2,"     d2[%d]=",(Zc*i1)>>(shift-1));
 
         nind=0;
 
         for (i3=0; i3 < ncols; i3++)
         {
           temp_prime=i1 * ncols + i3;
-
-
 	  for (i4=0; i4 < no_shift_values[temp_prime]; i4++)
-	    {
-	          
+  	  {
 	      var=(int)((i3*Zc + (Gen_shift_values[ pointer_shift_values[temp_prime]+i4 ]+1)%Zc)/Zc);
 	      int index =var*2*Zc + (i3*Zc + (Gen_shift_values[ pointer_shift_values[temp_prime]+i4 ]+1)%Zc) % Zc;
 	      printf("var %d, i3 %d, i4 %d, index %d, shift %d, Zc %d, pointer_shift_values[%d] %d gen_shift_value %d\n",var,i3,i4,index,shift,Zc,temp_prime,pointer_shift_values[temp_prime],Gen_shift_values[pointer_shift_values[temp_prime]]);
-	      indlist[nind] = ((index&mask)*((2*Zc*ncols)>>shift)/* *Kb */)+(index>>shift);
-	      printf("indlist[%d] %d, index&mask %d, index>>shift %d\n",nind,indlist[nind],index&mask,index>>shift);
-	      indlist2[nind++] = ((index&(mask>>1))*((2*Zc*ncols)>>(shift-1))*Kb)+(index>>(shift-1));
-	    }
-	  
-
-        }
+   	      printf("indlist[%d] %d, index&mask %d, index>>shift %d\n",nind,indlist[nind],index&mask,index>>shift);
+	      if (use_permutex==1 || use_alignr==1) {
+                indlist[nind++] = index;
+	      }
+	      else {
+	        indlist[nind] = ((index&mask)*((2*Zc*ncols)>>shift)/* *Kb */)+(index>>shift);
+	      }
+	  }//i4
+        }//i3
 	for (i4=0;i4<nind-1;i4++) {
-	  fprintf(fd,"%s(c2[%d],",xor_command,indlist[i4]);
-	  fprintf(fd2,"%s(c2[%d],",xor_command,indlist2[i4]);
+           if (use_permutex==0 && use_alignr==0) {
+  	      fprintf(fd,"%s(c2[%d],",xor_command,indlist[i4]);
+	   }
+	   else {
+              if ((indlist[i4]&mask) == 0) {
+  	        fprintf(fd,"%s(c2[%d],",xor_command,indlist[i4]>>shift);
+	      }
+	      else if (use_permutex==1) {
+                fprintf(fd,"%s(%s(c2[%d],abshift_get(%d),c2[%d]),",xor_command,permutex_command,(indlist[i4]>>shift),indlist[i4]&mask,(indlist[i4]>>shift)+1);
+	      }
+	      else if (use_alignr==1) {
+                fprintf(fd,"%s(%s(c2[%d],c2[%d],%d),",xor_command,alignr_command,(indlist[i4]>>shift)+1,(indlist[i4]>>shift),indlist[i4]&mask);
+	      }
+	   }
+	} //i4
+	if (use_permutex==0 && use_alignr==0)
+	  fprintf(fd,"c2[%d]",indlist[i4]);
+	else {
+          if ((indlist[i4]&mask) == 0) {
+  	     fprintf(fd,"c2[%d]",indlist[i4]>>shift);
+	  }
+	  else if (use_permutex==1){
+             fprintf(fd,"%s(c2[%d],abshift_get(%d),c2[%d])",permutex_command,(indlist[i4]>>shift),indlist[i4]&mask,(indlist[i4]>>shift)+1);
+	  }
+	  else if (use_alignr==1) {
+             fprintf(fd,"%s(c2[%d],c2[%d],%d)",alignr_command,(indlist[i4]>>shift)+1,(indlist[i4]>>shift),indlist[i4]&mask);
+
+	  }
 	}
-	fprintf(fd,"c2[%d]",indlist[i4]);
-	fprintf(fd2,"c2[%d]",indlist2[i4]);
-	for (i4=0;i4<nind-1;i4++) { fprintf(fd,")"); fprintf(fd2,")"); }
-	fprintf(fd,";\n");
-  fprintf(fd2, ";\n");
-      }
-      fprintf(fd,"  }\n}\n");
-      fprintf(fd2,"  }\n}\n");
-    }
+	for (i4=0;i4<nind-1;i4++) fprintf(fd,")"); 
+  	fprintf(fd,";\n");
+       } // i1 
+       fprintf(fd,"  }\n}\n");
+    } // i2
+    fclose(fd);
+  }
+  else if (gen_code == 4) { // CUDA
+    char fname[100];
+    sprintf(fname,"ldpc_BG%d_Zc%d_32bit.cu",BG,Zc);
+    FILE *fd=fopen(fname,"w");
+    AssertFatal(fd!=NULL,"cannot open %s\n",fname);
+    printf("Writing to %s\n",fname);
+    fprintf(fd,"#include <stdio.h>\n#include <stdint.h>\n#include <cuda_runtime.h>\n");
+
+    fprintf(fd,"// generated code for Zc=%d, byte encoding\n",Zc);
+    fprintf(fd,"__global__ void ldpc_BG%d_Zc%d_worker(uint32_t *c[4],uint32_t *d[4]) {\n",BG,Zc);
+    fprintf(fd,"  uint32_t *c32=c[blockIdx.x];\n  uint32_t *d32=d[blockIdx.x];\n\n");
+    fprintf(fd,"  int i2 = threadIdx.x;\n");
+    fprintf(fd,"  int i1 = blockIdx.y;\n");
+    fprintf(fd,"  if (i2 < %d) {\n",Zc);
+    fprintf(fd,"    c32+=i2;\n");
+    fprintf(fd,"    d32+=i2;\n");
+    fprintf(fd,"    switch(i1) {\n");
+
+    for (int i1=0;i1<nrows;i1++) {
+      nind = 0;
+      fprintf(fd,"    case %d:\n",i1);
+      fprintf(fd,"      d32[%d]=",(Zc*i1));
+      for (i3=0; i3 < ncols; i3++)
+      {
+          temp_prime=i1 * ncols + i3;
+	  for (i4=0; i4 < no_shift_values[temp_prime]; i4++)
+  	  {
+	     var=(int)((i3*Zc + (Gen_shift_values[ pointer_shift_values[temp_prime]+i4 ]+1)%Zc)/Zc);
+  	     int index =var*2*Zc + (i3*Zc + (Gen_shift_values[ pointer_shift_values[temp_prime]+i4 ]+1)%Zc) % Zc;
+	     printf("var %d, i3 %d, i4 %d, index %d (index mod 2Zc) %d, Zc %d, pointer_shift_values[%d] %d gen_shift_value %d offset %d\n",var,i3,i4,index,index%(2*Zc),Zc,temp_prime,pointer_shift_values[temp_prime],Gen_shift_values[pointer_shift_values[temp_prime]],(i3*Zc + (Gen_shift_values[ pointer_shift_values[temp_prime]+i4 ]+1)%Zc) % Zc);
+	     if (index%(2*Zc) >= Zc) printf("***********************\n");
+   	     indlist[nind] = index;
+	     printf("indlist[%d] %d, index %d\n",nind,indlist[nind],index);
+	     nind++;
+   	  } //i4
+      } // i3
+      for (i4=0;i4<nind-1;i4++) {
+         fprintf(fd,"c32[%d]^",indlist[i4]);
+      } //i4
+      fprintf(fd,"c32[%d];\n\n",indlist[i4]);
+      fprintf(fd,"       break;\n");
+    }// i1
+    fprintf(fd,"     }\n");
+    fprintf(fd,"  }\n");
+    fprintf(fd,"}\n");
+
+    fprintf(fd,"extern \"C\" int ldpc_BG%d_Zc%d_cuda32(uint32_t *c[4],uint32_t *d[4],int n_inputs) { \n",BG,Zc);
+    fprintf(fd,"dim3 numblocks(n_inputs,%d);\n",nrows);
+    fprintf(fd,"ldpc_BG%d_Zc%d_worker<<<numblocks,%d>>>(c,d);\n",BG,Zc,Zc);
+    fprintf(fd," cudaDeviceSynchronize();\n");
+    fprintf(fd,"  return(0);\n");
+    fprintf(fd,"}\n");
     fclose(fd);
-    fclose(fd2);
   }
   else if(gen_code==0)
   {
@@ -241,10 +322,12 @@ int LDPCencoder(unsigned char **inputArray, unsigned char *outputArray, encoder_
 
           for (i4 = 0; i4 < no_shift_values[temp_prime]; i4++) {
             channel_temp = channel_temp ^ c[i3 * Zc + Gen_shift_values[pointer_shift_values[temp_prime] + i4]];
+	    //if (i1==0) printf("index %d\n",i3 * Zc + Gen_shift_values[pointer_shift_values[temp_prime] + i4]);
           }
         }
 
         d[i2+i1*Zc]=channel_temp;
+	//if (i1==0) printf("reference: d[%d] %d c[%d] %d\n",i2,d[i2],i2,c[i2]);
         // output[t+i1*Zc]=channel_temp;
       }
     }

openair1/PHY/CODING/nrLDPC_encoder/ldpc_encoder32.c

@@ -0,0 +1,52 @@
+/*
+ * Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
+ * contributor license agreements.  See the NOTICE file distributed with
+ * this work for additional information regarding copyright ownership.
+ * The OpenAirInterface Software Alliance licenses this file to You under
+ * the OAI Public License, Version 1.1  (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.openairinterface.org/?page_id=698
+ *
+ * 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.
+ *-------------------------------------------------------------------------------
+ * For more information about the OpenAirInterface (OAI) Software Alliance:
+ *      contact@openairinterface.org
+ */
+
+/*!\file ldpc_encoder.c
+ * \brief Defines the LDPC encoder
+ * \author Florian Kaltenberger, Raymond Knopp, Kien le Trung (Eurecom)
+ * \email openair_tech@eurecom.fr
+ * \date 27-03-2018
+ * \version 1.0
+ * \note
+ * \warning
+ */
+
+
+
+#include <stdlib.h>
+#include <math.h>
+#include <stdio.h>
+#include <string.h>
+#include "defs.h"
+#include "assertions.h"
+#include "openair1/PHY/CODING/nrLDPC_defs.h"
+#include "openair1/PHY/CODING/nrLDPC_extern.h"
+#include "ldpc_generate_coefficient.c"
+
+
+void cuda_support_init() {
+   return;
+}
+uint32_t **LDPCencoder32(uint8_t **input,encoder_implemparams_t *impp)
+{
+	AssertFatal(1==0,"Should not be getting here\n");
+}
+

openair1/PHY/CODING/nrLDPC_encoder/ldpc_encoder_cuda32.c

@@ -0,0 +1,791 @@
+/*
+ * Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
+ * contributor license agreements.  See the NOTICE file distributed with
+ * this work for additional information regarding copyright ownership.
+ * The OpenAirInterface Software Alliance licenses this file to You under
+ * the OAI Public License, Version 1.1  (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.openairinterface.org/?page_id=698
+ *
+ * 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.
+ *-------------------------------------------------------------------------------
+ * For more information about the OpenAirInterface (OAI) Software Alliance:
+ *      contact@openairinterface.org
+ */
+
+/*! \file ldpc_encoder_optim8segmulti.c
+ * \brief Defines the optimized LDPC encoder
+ * \author Florian Kaltenberger, Raymond Knopp, Kien le Trung (Eurecom)
+ * \email openair_tech@eurecom.fr
+ * \date 27-03-2018
+ * \version 1.0
+ * \note
+ * \warning
+ */
+
+#include <stdlib.h>
+#include <math.h>
+#include <stdio.h>
+#include <string.h>
+#include "assertions.h"
+#include "common/utils/LOG/log.h"
+#include "time_meas.h"
+#include "openair1/PHY/CODING/nrLDPC_defs.h"
+#include "PHY/sse_intrin.h"
+#include "openair1/PHY/CODING/nrLDPC_extern.h"
+
+#include "ldpc_generate_coefficient.c"
+
+#include <cuda_runtime.h>
+
+//#define DEBUG_LDPC 1 
+
+#include "ldpc_encode_parity_check_cuda.c"
+uint32_t *c_dev;
+uint32_t **c_host;
+uint32_t *c_devh[4];
+uint32_t *d_dev;
+uint32_t **d_host;
+uint32_t *d_devh[4];
+uint32_t *cc_dev;
+uint32_t **cc_host;
+uint32_t *cc_devh[4];
+uint32_t *input_dev;
+uint32_t **input_host;
+uint32_t *input_devh[128];
+int managed = 0, concurrent = 0, uva = 0, pageable = 0, pageable_uses_host = 0, register_host = 0;
+
+#define USE_GPU_CIRCCOPY 1
+#define USE_GPU_FOR_INPUT 1
+
+void cuda_support_init() {
+
+
+    int dev = 0;
+    struct cudaDeviceProp prop;
+    cudaGetDeviceProperties(&prop, dev);
+
+
+    cudaDeviceGetAttribute(&managed, cudaDevAttrManagedMemory, dev);
+    cudaDeviceGetAttribute(&concurrent, cudaDevAttrConcurrentManagedAccess, dev);
+    cudaDeviceGetAttribute(&uva, cudaDevAttrUnifiedAddressing, dev);
+    cudaDeviceGetAttribute(&pageable, cudaDevAttrPageableMemoryAccess, dev);
+    cudaDeviceGetAttribute(&pageable_uses_host, cudaDevAttrPageableMemoryAccessUsesHostPageTables, dev);
+    cudaDeviceGetAttribute(&register_host, cudaDevAttrHostRegisterSupported,dev);
+
+    LOG_I(NR_PHY,"Device: %s (cc %d.%d)\n", prop.name, prop.major, prop.minor);
+    LOG_I(NR_PHY,"Unified Virtual Addressing (UVA): %s\n", uva ? "YES" : "NO");
+    LOG_I(NR_PHY,"Managed (Unified) Memory:        %s\n", managed ? "YES" : "NO");
+    LOG_I(NR_PHY,"Concurrent managed access:       %s\n", concurrent ? "YES" : "NO");
+    LOG_I(NR_PHY,"Pageable memory access:          %s\n", pageable ? "YES" : "NO");
+    LOG_I(NR_PHY,"Uses host page tables:           %s\n", pageable_uses_host ? "YES" : "NO");
+    LOG_I(NR_PHY,"Host Register supported:         %s\n", register_host ? "YES" : "NO");
+
+  // initialize input and output memory
+  if (!pageable && !register_host) {
+    cudaError_t err=cudaMalloc((void **)&c_dev,4*sizeof(uint32_t*));
+    AssertFatal(err == cudaSuccess,"CUDA Error (c_dev): %s\n", cudaGetErrorString(err));
+    err=cudaHostAlloc((void **)&c_host,4*sizeof(uint32_t*),cudaHostAllocDefault);
+    AssertFatal(err == cudaSuccess,"CUDA Error (c_host): %s\n", cudaGetErrorString(err));
+    for (int i=0;i<4;i++) {
+      err=cudaMalloc((void**)&c_devh[i],2*22*384*sizeof(uint32_t));
+      AssertFatal(err == cudaSuccess,"CUDA Error (c_devh[%d]): %s\n", i,cudaGetErrorString(err));
+      err=cudaHostAlloc((void**)&c_host[i],2*22*384*sizeof(uint32_t),cudaHostAllocDefault);
+      AssertFatal(err == cudaSuccess,"CUDA Error (chost[%d]): %s\n", i,cudaGetErrorString(err));
+    }
+    err = cudaMemcpy(c_dev,c_devh,4*sizeof(uint32_t*),cudaMemcpyHostToDevice);
+    AssertFatal(err == cudaSuccess,"CUDA Error (memcpy c_devh -> c_dev): %s\n", cudaGetErrorString(err));
+    err=cudaMalloc((void**)&d_dev,4*sizeof(uint32_t*));
+    AssertFatal(err == cudaSuccess,"CUDA Error: %s\n", cudaGetErrorString(err));
+    err=cudaHostAlloc((void **)&d_host,4*sizeof(uint32_t*),cudaHostAllocDefault);
+    AssertFatal(err == cudaSuccess,"CUDA Error (d_host): %s\n", cudaGetErrorString(err));
+    for (int i=0;i<4;i++) {
+      err=cudaMalloc((void**)&d_devh[i],68*384*sizeof(uint32_t));
+      AssertFatal(err == cudaSuccess,"CUDA Error (d_devh[%d]: %s\n", i,cudaGetErrorString(err));
+      err=cudaHostAlloc((void**)&d_host[i],68*384*sizeof(uint32_t),cudaHostAllocDefault);
+      AssertFatal(err == cudaSuccess,"CUDA Error (d_host[%d]): %s\n", i,cudaGetErrorString(err));
+    }
+    err=cudaMemcpy(d_dev,d_devh,4*sizeof(uint32_t*),cudaMemcpyHostToDevice);
+    AssertFatal(err == cudaSuccess,"CUDA Error (memcpy d_devh -> d_dev): %s\n", cudaGetErrorString(err));
+    err=cudaMalloc((void**)&cc_dev,4*sizeof(uint32_t*));
+    AssertFatal(err == cudaSuccess,"CUDA Error: %s\n", cudaGetErrorString(err));
+    err=cudaHostAlloc((void **)&cc_host,4*sizeof(uint32_t*),cudaHostAllocDefault);
+    AssertFatal(err == cudaSuccess,"CUDA Error (cc_host): %s\n", cudaGetErrorString(err));
+    for (int i=0;i<4;i++) {
+      err=cudaMalloc((void**)&cc_devh[i],68*384*sizeof(uint32_t));
+      AssertFatal(err == cudaSuccess,"CUDA Error (cc_devh[%d]: %s\n", i,cudaGetErrorString(err));
+      err=cudaHostAlloc((void**)&cc_host[i],68*384*sizeof(uint32_t),cudaHostAllocDefault);
+      AssertFatal(err == cudaSuccess,"CUDA Error (cc_host[%d]): %s\n", i,cudaGetErrorString(err));
+    }
+    err=cudaMemcpy(cc_dev,cc_devh,4*sizeof(uint32_t*),cudaMemcpyHostToDevice);
+    AssertFatal(err == cudaSuccess,"CUDA Error (memcpy cc_devh -> d_dev): %s\n", cudaGetErrorString(err));
+    err=cudaMalloc((void**)&input_dev,128*sizeof(uint8_t*));
+    AssertFatal(err == cudaSuccess,"CUDA Error: %s\n", cudaGetErrorString(err));
+    err=cudaHostAlloc((void **)&input_host,128*sizeof(uint8_t*),cudaHostAllocDefault);
+    AssertFatal(err == cudaSuccess,"CUDA Error (cc_host): %s\n", cudaGetErrorString(err));
+    for (int i=0;i<128;i++) {
+      err=cudaMalloc((void**)&input_devh[i],(8448/8)*sizeof(uint8_t));
+      AssertFatal(err == cudaSuccess,"CUDA Error (input_devh[%d]: %s\n", i,cudaGetErrorString(err));
+      err=cudaHostAlloc((void**)&input_host[i],(8448/8)*sizeof(uint8_t),cudaHostAllocDefault);
+      AssertFatal(err == cudaSuccess,"CUDA Error (input_host[%d]): %s\n", i,cudaGetErrorString(err));
+    }
+    err=cudaMemcpy(input_dev,input_devh,128*sizeof(uint8_t*),cudaMemcpyHostToDevice);
+    AssertFatal(err == cudaSuccess,"CUDA Error (memcpy cc_devh -> d_dev): %s\n", cudaGetErrorString(err));
+  }
+  else {
+    LOG_I(NR_PHY,"Allocating c,d,cc arrays for CPU/GPU shared-memory\n");
+    cudaError_t err=cudaHostAlloc((void **)&c_host,4*sizeof(uint32_t*),cudaHostAllocMapped|cudaHostAllocPortable);
+    AssertFatal(err == cudaSuccess,"CUDA Error (c_host): %s\n", cudaGetErrorString(err));
+    err = cudaHostGetDevicePointer((void**)&c_dev, c_host, 0);
+    AssertFatal(err == cudaSuccess,"CUDA Error (c_dev): %s\n", cudaGetErrorString(err));
+    LOG_I(NR_PHY,"c_host %p, c_dev %p\n",c_host,c_dev);
+    for (int i=0;i<4;i++) {
+      err=cudaHostAlloc((void**)&c_host[i],2*22*384*sizeof(uint32_t),cudaHostAllocMapped);
+      AssertFatal(err == cudaSuccess,"CUDA Error (c_host[%d]): %s\n", i,cudaGetErrorString(err));
+      err = cudaHostGetDevicePointer((void**)&c_devh[i], c_host[i], 0);
+      AssertFatal(err == cudaSuccess,"CUDA Error (c_devh[%d]): %s\n", i,cudaGetErrorString(err));
+    }
+    err=cudaMemcpy(c_dev,c_devh,4*sizeof(uint32_t*),cudaMemcpyHostToDevice);
+    AssertFatal(err == cudaSuccess,"CUDA Error (memcpy c_devh -> c_dev): %s\n", cudaGetErrorString(err));
+    err=cudaHostAlloc((void **)&d_host,4*sizeof(uint32_t*),cudaHostAllocMapped);
+    AssertFatal(err == cudaSuccess,"CUDA Error (d_host): %s\n", cudaGetErrorString(err));
+    err=cudaHostGetDevicePointer((void**)&d_dev, d_host, 0);
+    AssertFatal(err == cudaSuccess,"CUDA Error cudaHostGetDevicePointer(d_dev): %s\n", cudaGetErrorString(err));
+    LOG_I(NR_PHY,"d_host %p, d_dev %p\n",d_host,d_dev);
+    for (int i=0;i<4;i++) {
+      err=cudaHostAlloc((void**)&d_host[i],68*384*sizeof(uint32_t),cudaHostAllocMapped);
+      AssertFatal(err == cudaSuccess,"CUDA Error (d_host[%d]): %s\n", i,cudaGetErrorString(err));
+      err=cudaHostGetDevicePointer((void**)&d_devh[i], d_host[i], 0);
+      AssertFatal(err == cudaSuccess,"CUDA Error (cudaHostGetDevicePointer) d_devh[%d]: %s\n", i,cudaGetErrorString(err));
+    }
+    err=cudaMemcpy(d_dev,d_devh,4*sizeof(uint32_t*),cudaMemcpyHostToDevice);
+    AssertFatal(err == cudaSuccess,"CUDA Error (memcpy d_devh -> d_dev): %s\n", cudaGetErrorString(err));
+    err=cudaHostAlloc((void **)&cc_host,4*sizeof(uint32_t*),cudaHostAllocMapped);
+    AssertFatal(err == cudaSuccess,"CUDA Error (cc_host): %s\n", cudaGetErrorString(err));
+    err=cudaHostGetDevicePointer((void**)&cc_dev, cc_host, 0);
+    AssertFatal(err == cudaSuccess,"CUDA Error cudaHostGetDevicePointer(cc_dev): %s\n", cudaGetErrorString(err));
+    LOG_I(NR_PHY,"cc_host %p, cc_dev %p\n",cc_host,cc_dev);
+    for (int i=0;i<4;i++) {
+      err=cudaHostAlloc((void**)&cc_host[i],68*384*sizeof(uint32_t),cudaHostAllocMapped);
+      AssertFatal(err == cudaSuccess,"CUDA Error (cc_host[%d]): %s\n", i,cudaGetErrorString(err));
+      err=cudaHostGetDevicePointer((void**)&cc_devh[i], cc_host[i], 0);
+      AssertFatal(err == cudaSuccess,"CUDA Error (cudaHostGetDevicePointer) cc_devh[%d]: %s\n", i,cudaGetErrorString(err));
+    }
+    err=cudaMemcpy(cc_dev,cc_devh,4*sizeof(uint32_t*),cudaMemcpyHostToDevice);
+    AssertFatal(err == cudaSuccess,"CUDA Error (memcpy d_devh -> d_dev): %s\n", cudaGetErrorString(err));
+    err=cudaHostAlloc((void **)&input_host,128*sizeof(uint8_t*),cudaHostAllocMapped);
+    AssertFatal(err == cudaSuccess,"CUDA Error (input_host): %s\n", cudaGetErrorString(err));
+    err=cudaHostGetDevicePointer((void**)&input_dev, input_host, 0);
+    AssertFatal(err == cudaSuccess,"CUDA Error cudaHostGetDevicePointer(cc_host): %s\n", cudaGetErrorString(err));
+    LOG_I(NR_PHY,"input_host %p, input_dev %p\n",input_host,input_dev);
+    /*
+    for (int i=0;i<128;i++) {
+      err=cudaHostAlloc((void**)&input_host[i],(8448/8)*sizeof(uint8_t),cudaHostAllocMapped);
+      AssertFatal(err == cudaSuccess,"CUDA Error (input_host[%d]): %s\n", i,cudaGetErrorString(err));
+      err=cudaHostGetDevicePointer((void**)&input_devh[i], input_host[i], 0);
+      AssertFatal(err == cudaSuccess,"CUDA Error (cudaHostGetDevicePointer) input_devh[%d]: %s\n", i,cudaGetErrorString(err));
+    }
+    err=cudaMemcpy(input_dev,input_devh,128*sizeof(uint8_t*),cudaMemcpyHostToDevice);
+    AssertFatal(err == cudaSuccess,"CUDA Error (memcpy input_devh -> input_dev): %s\n", cudaGetErrorString(err));
+    */
+  }
+}
+
+uint32_t **LDPCencoder32(uint8_t **input, encoder_implemparams_t *impp)
+{
+  //set_log(PHY, 4);
+
+  int Zc = impp->Zc;
+  int Kb = impp->Kb;
+  short block_length = impp->K;
+  short BG = impp->BG;
+  int nrows=0,ncols=0;
+  int rate=3;
+  int no_punctured_columns,removed_bit;
+
+  if(impp->tinput != NULL) start_meas(impp->tinput);
+  //determine number of bits in codeword
+  if (BG==1)
+    {
+      nrows=46; //parity check bits
+      ncols=22; //info bits
+      rate=3;
+    }
+    else if (BG==2)
+    {
+      nrows=42; //parity check bits
+      ncols=10; // info bits
+      rate=5;
+    }
+
+#ifdef DEBUG_LDPC
+  LOG_I(PHY,"ldpc_encoder_cuda32: BG %d, Zc %d, Kb %d, block_length %d, segments %d\n",BG,Zc,Kb,block_length,impp->n_segments);
+  LOG_I(PHY,"ldpc_encoder_cuda32: PDU (seg 0) %x %x %x %x\n",input[0][0],input[0][1],input[0][2],input[0][3]);
+#endif
+
+  AssertFatal(Zc > 0, "no valid Zc found for block length %d\n", block_length);
+
+  int n_inputs = (impp->n_segments/32)+(((impp->n_segments&31) > 0) ? 1: 0);
+//  uint32_t  cc[4][22*Zc]; //padded input, unpacked, max size
+
+  // calculate number of punctured bits
+  no_punctured_columns=(int)((nrows-2)*Zc+block_length-block_length*rate)/Zc;
+  removed_bit=(nrows-no_punctured_columns-2) * Zc+block_length-(int)(block_length*rate);
+  // clear input
+  /*
+  for (int i=0;i<n_inputs;i++) { 
+    if (!pageable && !register_host)
+      cudaMemset(cc_devh[i],0,22*Zc*sizeof(uint32_t));
+    else
+      memset(cc_host[i],0,22*Zc*sizeof(uint32_t));
+  }*/
+
+
+#if 0
+  // unoptimized version of input processing
+  unsigned int temp;
+  for (int i_dword=0; i_dword < block_length; i_dword++) {
+    unsigned int i = i_dword;
+    for (int j = 0; j < impp->n_segments; j++) {
+
+      temp = (input[j][i/8]&(128>>(i&7)))>>(7-(i&7));
+      cc[j>>5][i] |= (temp << (j&31));
+    }
+  }
+#else
+//  uint32_t *ccp[4];
+//  for (int s=0;s<n_inputs;s++) ccp[s]=cc[s];  
+#ifdef USE_GPU_FOR_INPUT
+  if (!pageable && !register_host) {
+    for (int r=0;r<impp->n_segments;r++) {
+        cudaMemcpy(input_devh[r],input[r],block_length>>3,cudaMemcpyHostToDevice);
+    }
+  }
+  ldpc_input(pageable||register_host? input : input_dev,(uint32_t**)cc_dev,block_length,impp->n_segments);
+//  for (int i=0;i<16;i++) printf("i %d: cc[0] %x\n",i,cc[0][i]);
+#else  
+#ifndef __aarch64__
+  simde__m256i andmask = simde_mm256_set_epi32(0x1,0x2,0x4,0x8,0x10,0x20,0x40,0x80);  // every 8 bits -> 8 bytes, pattern repeats.
+  simde__m256i zero256 = simde_mm256_setzero_si256();
+  simde__m256i masks[32];
+  register simde__m256i c256;
+  for (int i=0;i<32;i++) {
+    masks[i] = simde_mm256_set1_epi32(1<<i);
+  }
+
+  int i;
+  for (i=0; i < block_length; i += 8) {
+    unsigned int i8 = i >> 3;
+    for (int j0=0;j0<n_inputs;j0++) {
+      c256 = simde_mm256_and_si256(simde_mm256_cmpeq_epi32(simde_mm256_andnot_si256(simde_mm256_set1_epi32(input[j0<<5][i8]),andmask),zero256),masks[0]);
+      int j2 = 32;
+      if (j0==(n_inputs-1)) j2= impp->n_segments-(j0<<5);
+      for (int j=1; j < j2; j++) {
+        c256 = simde_mm256_or_si256(simde_mm256_and_si256(simde_mm256_cmpeq_epi32(simde_mm256_andnot_si256(simde_mm256_set1_epi32(input[(j0<<5)+j][i8]),andmask),zero256),masks[j]),c256);
+      }
+      ((simde__m256i *)cc_host[j0])[i8] = c256;
+    }
+  }
+#else
+  int i2=0;
+  const int32_t ucShifta[32][4] = { 
+	  {-7,-6,-5,-4}, // 0  
+	  {-6,-5,-4,-3}, // 1
+	  {-5,-4,-3,-2}, // 2 
+	  {-4,-3,-2,-1}, // 3 
+	  {-3,-2,-1,0}, // 4 
+	  {-2,-1,0,1}, // 5
+	  {-1,0,1,2}, // 6
+	  {0,1,2,3}, // 7
+	  {1,2,3,4}, // 8
+	  {2,3,4,5}, // 9
+	  {3,4,5,6}, // 10
+	  {4,5,6,7}, // 11
+	  {5,6,7,8}, // 12
+	  {6,7,8,9}, // 13
+	  {7,8,9,10}, // 14
+	  {8,9,10,11}, // 15
+	  {9,10,11,12}, // 16
+	  {10,11,12,13}, // 17
+	  {11,12,13,14}, // 18
+	  {12,13,14,15}, // 19
+	  {13,14,15,16}, // 20
+	  {14,15,16,17}, // 21
+	  {15,16,17,18}, // 22
+	  {16,17,18,19}, // 23
+	  {17,18,19,20}, // 24
+	  {18,19,20,21}, // 25
+	  {19,20,21,22}, // 26
+	  {20,21,22,23}, // 27
+	  {21,22,23,24}, // 28
+	  {22,23,24,25}, // 29
+	  {23,24,25,26}, // 30
+	  {24,25,26,27}}; // 31
+  const int32_t ucShiftb[32][4] = { 
+	  {-3,-2,-1,0}, // 0 
+	  {-2,-1,0,1}, // 1
+	  {-1,0,1,2}, // 2
+	  {0,1,2,3}, // 3
+	  {1,2,3,4}, // 4
+	  {2,3,4,5}, // 5
+	  {3,4,5,6}, // 6 
+	  {4,5,6,7}, // 7
+	  {5,6,7,8}, // 8
+	  {6,7,8,9}, // 9
+	  {7,8,9,10}, // 10
+	  {8,9,10,11}, // 11
+	  {9,10,11,12}, // 12
+	  {10,11,12,13}, // 13
+	  {11,12,13,14}, // 14
+	  {12,13,14,15}, // 15
+	  {13,14,15,16}, // 16
+	  {14,15,16,17}, // 17
+	  {15,16,17,18}, // 18
+	  {16,17,18,19}, // 19
+	  {17,18,19,20}, // 20
+	  {18,19,20,21}, // 21
+	  {19,20,21,22}, // 22
+	  {20,21,22,23}, // 23
+	  {21,22,23,24}, // 24
+	  {22,23,24,25}, // 25
+	  {23,24,25,26}, // 26
+	  {24,25,26,27}, // 27
+	  {25,26,27,28}, // 28
+	  {26,27,28,29}, // 29
+	  {27,28,29,30}, // 30
+	  {28,29,30,31}}; // 31
+			      
+  const int32_t ucShiftc[32][4] = { 
+	  {-15,-14,-13,-12}, // 0
+	  {-14,-13,-12,-11}, // 1
+	  {-13,-12,-11,-10}, // 2
+          {-12,-11,-10,-9}, // 3
+	  {-11,-10,-9,-8}, // 4
+	  {-10,-9,-8,-7}, // 5
+	  {-9,-8,-7,-6}, // 6
+	  {-8,-7,-6,-5}, // 7
+	  {-7,-6,-5,-4}, // 8  
+	  {-6,-5,-4,-3}, // 9
+	  {-5,-4,-3,-2}, // 10 
+	  {-4,-3,-2,-1}, // 11 
+	  {-3,-2,-1,0}, // 12 
+	  {-2,-1,0,1}, // 13
+	  {-1,0,1,2}, // 14
+	  {0,1,2,3}, // 15
+	  {1,2,3,4}, // 16
+	  {2,3,4,5}, // 17
+	  {3,4,5,6}, // 18
+	  {4,5,6,7}, // 19
+	  {5,6,7,8}, // 20
+	  {6,7,8,9}, // 21
+	  {7,8,9,10}, // 22
+	  {8,9,10,11}, // 23
+	  {9,10,11,12}, // 24
+	  {10,11,12,13}, // 25
+	  {11,12,13,14}, // 26
+	  {12,13,14,15}, // 27
+	  {13,14,15,16}, // 28
+	  {14,15,16,17}, // 29
+	  {15,16,17,18}, // 30
+	  {16,17,18,19}}; // 31
+  const int32_t ucShiftd[32][4] = { 
+	  {-11,-10,-9,-8}, // 0
+	  {-10,-9,-8,-7}, // 1
+	  {-9,-8,-7,-6}, // 2
+	  {-8,-7,-6,-5}, // 3
+	  {-7,-6,-5,-4}, // 4  
+	  {-6,-5,-4,-3}, // 5
+	  {-5,-4,-3,-2}, // 6
+	  {-4,-3,-2,-1}, // 7 
+	  {-3,-2,-1,0}, // 8
+	  {-2,-1,0,1}, // 9
+	  {-1,0,1,2}, // 10
+	  {0,1,2,3}, // 11 
+	  {1,2,3,4}, // 12
+	  {2,3,4,5}, // 13
+	  {3,4,5,6}, // 14 
+	  {4,5,6,7}, // 15
+	  {5,6,7,8}, // 16
+	  {6,7,8,9}, // 17
+	  {7,8,9,10}, // 18
+	  {8,9,10,11}, // 19
+	  {9,10,11,12}, // 20
+	  {10,11,12,13}, // 21
+	  {11,12,13,14}, // 22
+	  {12,13,14,15}, // 23
+	  {13,14,15,16}, // 24
+	  {14,15,16,17}, // 25
+	  {15,16,17,18}, // 26
+	  {16,17,18,19}, // 27
+	  {17,18,19,20}, // 28
+	  {18,19,20,21}, // 29
+	  {19,20,21,22}, // 30
+	  {20,21,22,23}};// 31
+			  
+  const int32_t ucShifte[32][4] = { 
+	  {-23,-22,-21,-20}, // 0 
+	  {-22,-21,-20,-19}, // 1
+	  {-21,-20,-19,-18}, // 2
+	  {-20,-19,-18,-17}, // 3
+	  {-19,-18,-17,-16}, // 4 
+	  {-18,-17,-16,-15}, // 5
+	  {-17,-16,-15,-14}, // 6
+	  {-16,-15,-14,-13}, // 7
+	  {-15,-14,-13,-12}, // 8
+	  {-14,-13,-12,-11}, // 9
+	  {-13,-12,-11,-10}, // 10
+          {-12,-11,-10,-9}, // 11
+	  {-11,-10,-9,-8}, // 12
+	  {-10,-9,-8,-7}, // 13
+	  {-9,-8,-7,-6}, // 14 
+	  {-8,-7,-6,-5}, // 15
+	  {-7,-6,-5,-4}, // 16 
+	  {-6,-5,-4,-3}, // 17
+	  {-5,-4,-3,-2}, // 18 
+	  {-4,-3,-2,-1}, // 19 
+	  {-3,-2,-1,0}, // 20 
+	  {-2,-1,0,1}, // 21
+	  {-1,0,1,2}, // 22
+	  {0,1,2,3}, // 23
+	  {1,2,3,4}, // 24
+	  {2,3,4,5}, // 25
+	  {3,4,5,6}, // 26
+	  {4,5,6,7}, // 27
+	  {5,6,7,8}, // 28
+	  {6,7,8,9}, // 29
+	  {7,8,9,10}, // 30
+	  {8,9,10,11}}; // 31 
+			  
+  const int32_t ucShiftf[32][4] = { 
+	  {-19,-18,-17,-16}, // 0 
+	  {-18,-17,-16,-15}, // 1
+	  {-17,-16,-15,-14}, // 2
+	  {-16,-15,-14,-13}, // 3
+	  {-15,-14,-13,-12}, // 4
+	  {-14,-13,-12,-11}, // 5
+	  {-13,-12,-11,-10}, // 6
+          {-12,-11,-10,-9}, // 7
+	  {-11,-10,-9,-8}, // 8
+	  {-10,-9,-8,-7}, // 9  
+	  {-9,-8,-7,-6}, // 10 
+	  {-8,-7,-6,-5}, // 11
+	  {-7,-6,-5,-4}, // 12 
+	  {-6,-5,-4,-3}, // 13
+	  {-5,-4,-3,-2}, // 14
+	  {-4,-3,-2,-1}, // 15 
+	  {-3,-2,-1,0}, // 16
+	  {-2,-1,0,1}, // 17
+	  {-1,0,1,2}, // 18
+	  {0,1,2,3}, // 19
+	  {1,2,3,4}, // 20
+	  {2,3,4,5}, // 21
+	  {3,4,5,6}, // 22 
+	  {4,5,6,7}, // 23
+	  {5,6,7,8}, // 24
+	  {6,7,8,9}, // 25
+	  {7,8,9,10}, // 26
+	  {8,9,10,11}, // 27
+	  {9,10,11,12}, // 28
+	  {10,11,12,13}, // 29
+	  {11,12,13,14}, // 30
+	  {12,13,14,15}}; // 31
+			  
+  const int32_t ucShiftg[32][4] = { 
+	  {-31,-30,-29,-28},
+	  {-30,-29,-28,-27},
+	  {-29,-28,-27,-26},
+	  {-28,-27,-26,-25},
+	  {-27,-26,-25,-24},
+	  {-26,-25,-24,-23},
+	  {-25,-24,-23,-22},
+	  {-24,-23,-22,-21},
+	  {-23,-22,-21,-20}, // 0 
+	  {-22,-21,-20,-19}, // 1
+	  {-21,-20,-19,-18}, // 2
+	  {-20,-19,-18,-17}, // 3
+	  {-19,-18,-17,-16}, // 4 
+	  {-18,-17,-16,-15}, // 5
+	  {-17,-16,-15,-14}, // 6
+	  {-16,-15,-14,-13}, // 7
+	  {-15,-14,-13,-12}, // 8
+	  {-14,-13,-12,-11}, // 9
+	  {-13,-12,-11,-10}, // 10
+          {-12,-11,-10,-9}, // 11
+	  {-11,-10,-9,-8}, // 12
+	  {-10,-9,-8,-7}, // 13
+	  {-9,-8,-7,-6}, // 14 
+	  {-8,-7,-6,-5}, // 15
+	  {-7,-6,-5,-4}, // 16 
+	  {-6,-5,-4,-3}, // 17
+	  {-5,-4,-3,-2}, // 18 
+	  {-4,-3,-2,-1}, // 19 
+	  {-3,-2,-1,0}, // 20 
+	  {-2,-1,0,1}, // 21
+	  {-1,0,1,2}, // 22
+	  {0,1,2,3}}; // 23
+			  
+  const int32_t ucShifth[32][4] = { 
+	  {-27,-26,-25,-24}, // 0
+	  {-26,-25,-24,-23}, // 1
+	  {-25,-24,-23,-22}, // 2
+	  {-24,-23,-22,-21}, // 3
+	  {-23,-22,-21,-20}, // 4 
+	  {-22,-21,-20,-19}, // 5
+	  {-21,-20,-19,-18}, // 6
+	  {-20,-19,-18,-17}, // 7
+	  {-19,-18,-17,-16}, // 8 
+	  {-18,-17,-16,-15}, // 9
+	  {-17,-16,-15,-14}, // 10
+	  {-16,-15,-14,-13}, // 11
+	  {-15,-14,-13,-12}, // 12
+	  {-14,-13,-12,-11}, // 13
+	  {-13,-12,-11,-10}, // 14
+          {-12,-11,-10,-9}, // 15
+	  {-11,-10,-9,-8}, // 16
+	  {-10,-9,-8,-7}, // 17 
+	  {-9,-8,-7,-6}, // 18 
+	  {-8,-7,-6,-5}, // 19
+	  {-7,-6,-5,-4}, // 20 
+	  {-6,-5,-4,-3}, // 21
+	  {-5,-4,-3,-2}, // 22
+	  {-4,-3,-2,-1}, // 23 
+	  {-3,-2,-1,0}, // 24 
+	  {-2,-1,0,1}, // 25
+	  {-1,0,1,2}, // 26
+	  {0,1,2,3}, // 27
+	  {1,2,3,4}, // 28
+	  {2,3,4,5}, // 29
+	  {3,4,5,6}, // 30
+          {4,5,6,7}}; // 31 
+  const uint32_t __attribute__ ((aligned (16))) masksa[4] = {0x80,0x40,0x20,0x10};
+  const uint32_t __attribute__ ((aligned (16))) masksb[4] = {0x8,0x4,0x2,0x1};
+  const uint32_t __attribute__ ((aligned (16))) masksc[4] = {0x8000,0x4000,0x2000,0x1000};
+  const uint32_t __attribute__ ((aligned (16))) masksd[4] = {0x800,0x400,0x200,0x100};
+  const uint32_t __attribute__ ((aligned (16))) maskse[4] = {0x800000,0x400000,0x200000,0x100000};
+  const uint32_t __attribute__ ((aligned (16))) masksf[4] = {0x80000,0x40000,0x20000,0x10000};
+  const uint32_t __attribute__ ((aligned (16))) masksg[4] = {0x80000000,0x40000000,0x20000000,0x10000000};
+  const uint32_t __attribute__ ((aligned (16))) masksh[4] = {0x8000000,0x4000000,0x2000000,0x1000000};
+  int32x4_t vshifta[32],vshiftb[32],vshiftc[32],vshiftd[32],vshifte[32],vshiftf[32],vshiftg[32],vshifth[32];
+  uint32x4_t vmasksa  = vld1q_u32(masksa);
+  uint32x4_t vmasksb  = vld1q_u32(masksb);
+  uint32x4_t vmasksc  = vld1q_u32(masksc);
+  uint32x4_t vmasksd  = vld1q_u32(masksd);
+  uint32x4_t vmaskse  = vld1q_u32(maskse);
+  uint32x4_t vmasksf  = vld1q_u32(masksf);
+  uint32x4_t vmasksg  = vld1q_u32(masksg);
+  uint32x4_t vmasksh  = vld1q_u32(masksh);
+  uint32x4_t in;
+
+  for (int n=0;n<32;n++) { 
+    vshifta[n] = vld1q_s32(ucShifta[n]);
+    vshiftb[n] = vld1q_s32(ucShiftb[n]);
+    vshiftc[n] = vld1q_s32(ucShiftc[n]);
+    vshiftd[n] = vld1q_s32(ucShiftd[n]);
+    vshifte[n] = vld1q_s32(ucShifte[n]);
+    vshiftf[n] = vld1q_s32(ucShiftf[n]);
+    vshiftg[n] = vld1q_s32(ucShiftg[n]);
+    vshifth[n] = vld1q_s32(ucShifth[n]);
+  }
+  i2=0;
+  int j0=0,j1=0,j2=0,j3=0;
+  if (impp->n_segments <= 32) {
+	  j0=impp->n_segments;
+  }
+  else if (impp->n_segments <= 64) {
+     j0=32;
+     j1=impp->n_segments-32;
+  }
+  else if (impp->n_segments <= 96) {
+     j0=32; j1=32;
+     j2=impp->n_segments-64;
+  }
+  else if (impp->n_segments <= 128) {
+     j0=32; j1=32; j2=32;
+     j3=impp->n_segments-96;
+  }
+  uint32x4_t *ccp,cc0,cc1,cc2,cc3,cc4,cc5,cc6,cc7;
+  for (int i=0; i < (block_length>>5); i++,i2+=8) {
+    in = vdupq_n_u32(((uint32_t*)input[0])[i]);
+    cc0 = vshlq_u32(vandq_u32(in,vmasksa),vshifta[0]);
+    cc1 = vshlq_u32(vandq_u32(in,vmasksb),vshiftb[0]);
+    cc2 = vshlq_u32(vandq_u32(in,vmasksc),vshiftc[0]);
+    cc3 = vshlq_u32(vandq_u32(in,vmasksd),vshiftd[0]);
+    cc4 = vshlq_u32(vandq_u32(in,vmaskse),vshifte[0]);
+    cc5 = vshlq_u32(vandq_u32(in,vmasksf),vshiftf[0]);
+    cc6 = vshlq_u32(vandq_u32(in,vmasksg),vshiftg[0]);
+    cc7 = vshlq_u32(vandq_u32(in,vmasksh),vshifth[0]);
+    for (int j = 1; j < j0; j++) {
+      in = vdupq_n_u32(((uint32_t*)input[j])[i]);
+      cc0 = vorrq_u32(cc0,vshlq_u32(vandq_u32(in,vmasksa),vshifta[j]));
+      cc1 = vorrq_u32(cc1,vshlq_u32(vandq_u32(in,vmasksb),vshiftb[j]));
+      cc2 = vorrq_u32(cc2,vshlq_u32(vandq_u32(in,vmasksc),vshiftc[j]));
+      cc3 = vorrq_u32(cc3,vshlq_u32(vandq_u32(in,vmasksd),vshiftd[j]));
+      cc4 = vorrq_u32(cc4,vshlq_u32(vandq_u32(in,vmaskse),vshifte[j]));
+      cc5 = vorrq_u32(cc5,vshlq_u32(vandq_u32(in,vmasksf),vshiftf[j]));
+      cc6 = vorrq_u32(cc6,vshlq_u32(vandq_u32(in,vmasksg),vshiftg[j]));
+      cc7 = vorrq_u32(cc7,vshlq_u32(vandq_u32(in,vmasksh),vshifth[j]));
+    }
+    ccp=&((uint32x4_t *)cc_host[0])[i2];
+    ccp[0] = cc0;
+    ccp[1] = cc1;
+    ccp[2] = cc2;
+    ccp[3] = cc3;
+    ccp[4] = cc4;
+    ccp[5] = cc5;
+    ccp[6] = cc6;
+    ccp[7] = cc7;
+    if (j1>0) {
+      in = vdupq_n_u32(((uint32_t*)input[32])[i]);
+      cc0 = vshlq_u32(vandq_u32(in,vmasksa),vshifta[0]);
+      cc1 = vshlq_u32(vandq_u32(in,vmasksb),vshiftb[0]);
+      cc2 = vshlq_u32(vandq_u32(in,vmasksc),vshiftc[0]);
+      cc3 = vshlq_u32(vandq_u32(in,vmasksd),vshiftd[0]);
+      cc4 = vshlq_u32(vandq_u32(in,vmaskse),vshifte[0]);
+      cc5 = vshlq_u32(vandq_u32(in,vmasksf),vshiftf[0]);
+      cc6 = vshlq_u32(vandq_u32(in,vmasksg),vshiftg[0]);
+      cc7 = vshlq_u32(vandq_u32(in,vmasksh),vshifth[0]);
+      for (int j = 1; j < j1; j++) {
+        in = vdupq_n_u32(((uint32_t*)input[32+j])[i]);
+        cc0 = vorrq_u32(cc0,vshlq_u32(vandq_u32(in,vmasksa),vshifta[j]));
+        cc1 = vorrq_u32(cc1,vshlq_u32(vandq_u32(in,vmasksb),vshiftb[j]));
+        cc2 = vorrq_u32(cc2,vshlq_u32(vandq_u32(in,vmasksc),vshiftc[j]));
+        cc3 = vorrq_u32(cc3,vshlq_u32(vandq_u32(in,vmasksd),vshiftd[j]));
+        cc4 = vorrq_u32(cc4,vshlq_u32(vandq_u32(in,vmaskse),vshifte[j]));
+        cc5 = vorrq_u32(cc5,vshlq_u32(vandq_u32(in,vmasksf),vshiftf[j]));
+        cc6 = vorrq_u32(cc6,vshlq_u32(vandq_u32(in,vmasksg),vshiftg[j]));
+        cc7 = vorrq_u32(cc7,vshlq_u32(vandq_u32(in,vmasksh),vshifth[j]));
+      }
+      ccp=&((uint32x4_t *)cc_host[1])[i2];
+      ccp[0] = cc0;
+      ccp[1] = cc1;
+      ccp[2] = cc2;
+      ccp[3] = cc3;
+      ccp[4] = cc4;
+      ccp[5] = cc5;
+      ccp[6] = cc6;
+      ccp[7] = cc7;
+    }
+    if (j2>0) {
+      in = vdupq_n_u32(((uint32_t*)input[64])[i]);
+      cc0 = vshlq_u32(vandq_u32(in,vmasksa),vshifta[0]);
+      cc1 = vshlq_u32(vandq_u32(in,vmasksb),vshiftb[0]);
+      cc2 = vshlq_u32(vandq_u32(in,vmasksc),vshiftc[0]);
+      cc3 = vshlq_u32(vandq_u32(in,vmasksd),vshiftd[0]);
+      cc4 = vshlq_u32(vandq_u32(in,vmaskse),vshifte[0]);
+      cc5 = vshlq_u32(vandq_u32(in,vmasksf),vshiftf[0]);
+      cc6 = vshlq_u32(vandq_u32(in,vmasksg),vshiftg[0]);
+      cc7 = vshlq_u32(vandq_u32(in,vmasksh),vshifth[0]);
+      for (int j = 1; j < j2; j++) {
+        in = vdupq_n_u32(((uint32_t*)input[64+j])[i]);
+        cc0 = vorrq_u32(cc0,vshlq_u32(vandq_u32(in,vmasksa),vshifta[j]));
+        cc1 = vorrq_u32(cc1,vshlq_u32(vandq_u32(in,vmasksb),vshiftb[j]));
+        cc2 = vorrq_u32(cc2,vshlq_u32(vandq_u32(in,vmasksc),vshiftc[j]));
+        cc3 = vorrq_u32(cc3,vshlq_u32(vandq_u32(in,vmasksd),vshiftd[j]));
+        cc4 = vorrq_u32(cc4,vshlq_u32(vandq_u32(in,vmaskse),vshifte[j]));
+        cc5 = vorrq_u32(cc5,vshlq_u32(vandq_u32(in,vmasksf),vshiftf[j]));
+        cc6 = vorrq_u32(cc6,vshlq_u32(vandq_u32(in,vmasksg),vshiftg[j]));
+        cc7 = vorrq_u32(cc7,vshlq_u32(vandq_u32(in,vmasksh),vshifth[j]));
+      }
+      ccp=&((uint32x4_t *)cc_host[2])[i2];
+      ccp[0] = cc0;
+      ccp[1] = cc1;
+      ccp[2] = cc2;
+      ccp[3] = cc3;
+      ccp[4] = cc4;
+      ccp[5] = cc5;
+      ccp[6] = cc6;
+      ccp[7] = cc7;
+    }
+    if (j3>0) {
+      in = vdupq_n_u32(((uint32_t*)input[96])[i]);
+      cc0 = vshlq_u32(vandq_u32(in,vmasksa),vshifta[0]);
+      cc1 = vshlq_u32(vandq_u32(in,vmasksb),vshiftb[0]);
+      cc2 = vshlq_u32(vandq_u32(in,vmasksc),vshiftc[0]);
+      cc3 = vshlq_u32(vandq_u32(in,vmasksd),vshiftd[0]);
+      cc4 = vshlq_u32(vandq_u32(in,vmaskse),vshifte[0]);
+      cc5 = vshlq_u32(vandq_u32(in,vmasksf),vshiftf[0]);
+      cc6 = vshlq_u32(vandq_u32(in,vmasksg),vshiftg[0]);
+      cc7 = vshlq_u32(vandq_u32(in,vmasksh),vshifth[0]);
+      for (int j = 1; j < j3; j++) {
+        in = vdupq_n_u32(((uint32_t*)input[96+j])[i]);
+        cc0 = vorrq_u32(cc0,vshlq_u32(vandq_u32(in,vmasksa),vshifta[j]));
+        cc1 = vorrq_u32(cc1,vshlq_u32(vandq_u32(in,vmasksb),vshiftb[j]));
+        cc2 = vorrq_u32(cc2,vshlq_u32(vandq_u32(in,vmasksc),vshiftc[j]));
+        cc3 = vorrq_u32(cc3,vshlq_u32(vandq_u32(in,vmasksd),vshiftd[j]));
+        cc4 = vorrq_u32(cc4,vshlq_u32(vandq_u32(in,vmaskse),vshifte[j]));
+        cc5 = vorrq_u32(cc5,vshlq_u32(vandq_u32(in,vmasksf),vshiftf[j]));
+        cc6 = vorrq_u32(cc6,vshlq_u32(vandq_u32(in,vmasksg),vshiftg[j]));
+        cc7 = vorrq_u32(cc7,vshlq_u32(vandq_u32(in,vmasksh),vshifth[j]));
+      }
+      ccp=&((uint32x4_t *)cc_host[3])[i2];
+      ccp[0] = cc0;
+      ccp[1] = cc1;
+      ccp[2] = cc2;
+      ccp[3] = cc3;
+      ccp[4] = cc4;
+      ccp[5] = cc5;
+      ccp[6] = cc6;
+      ccp[7] = cc7;
+    }
+  }
+#endif
+#endif
+#endif
+  if(impp->tinput != NULL) stop_meas(impp->tinput);
+  if(impp->tinput_memcpy != NULL) start_meas(impp->tinput_memcpy);
+#ifdef USE_GPU_CIRCCOPY
+  circcopy_c(cc_dev,c_dev,n_inputs);
+#else
+#ifdef USE_CUDA_MEMCPY 
+  for (int s=0;s<n_inputs;s++)  {
+    for (int i1 = 0; i1 < ncols; i1++)   {
+      cudaError_t err = cudaMemcpy(&c_devh[s][2 * i1 * Zc], &cc_devh[s][i1 * Zc], Zc * sizeof(uint32_t),cudaMemcpyDeviceToDevice);
+      AssertFatal(err == cudaSuccess, "1st memcpy cc_devh[%d] -> c_devh[%d] %p CUDA Error: %s\n", s, s, c_devh[s],cudaGetErrorString(err)); 			
+      err = cudaMemcpy(&c_devh[s][(2 * i1 + 1) * Zc], &cc_devh[s][i1 * Zc], Zc * sizeof(uint32_t),cudaMemcpyDeviceToDevice);
+      AssertFatal(err == cudaSuccess, "2nd memcpy cc_devh[%d] -> c_devh[%d] %p CUDA Error: %s\n", s, s, c_devh[s],cudaGetErrorString(err)); 			
+    }
+#else
+  for (int s=0;s<n_inputs;s++)  {
+    for (int i1 = 0; i1 < ncols; i1++)   {
+      memcpy(&c_host[s][2 * i1 * Zc], &cc_host[s][i1 * Zc], Zc * sizeof(uint32_t));
+      memcpy(&c_host[s][(2 * i1 + 1) * Zc], &cc_host[s][i1 * Zc], Zc * sizeof(uint32_t));
+    }
+    if (!pageable && !register_host) {
+      cudaError_t err = cudaMemcpy(c_devh[s],c_host[s],sizeof(uint32_t)*2*22*Zc,cudaMemcpyHostToDevice);
+      AssertFatal(err == cudaSuccess, "memcpy c_host[%d] -> c_devh[%d] %p CUDA Error: %s\n", s, s, c_devh[s],cudaGetErrorString(err)); 			
+    }
+  }
+#endif 
+#endif
+  if(impp->tinput != NULL) stop_meas(impp->tinput_memcpy);
+  if (BG==1 && Zc==384)  {
+    //parity check part
+    if(impp->tparity != NULL) start_meas(impp->tparity);
+    encode_parity_check_part_cuda((uint32_t**)c_dev, (uint32_t**)d_dev, BG, Zc, Kb, ncols,n_inputs);
+    if(impp->tparity != NULL) stop_meas(impp->tparity);
+  }
+  else {
+    AssertFatal(1==0,"Only BG1 Zc=384 for now\n");
+  }
+  if(impp->toutput != NULL) start_meas(impp->toutput);
+  for (int s=0;s<n_inputs;s++) {
+	  
+    if (!pageable && !register_host) {
+      cudaMemcpy(d_host[s],&cc_devh[s][2*Zc],sizeof(uint32_t)*(block_length-(2*Zc)),cudaMemcpyDeviceToHost);
+      cudaError_t err = cudaMemcpy(&d_host[s][block_length-(2*Zc)],d_devh[s]+block_length-(2*Zc),sizeof(uint32_t)*((nrows-no_punctured_columns) * Zc-removed_bit),cudaMemcpyDeviceToHost);
+      AssertFatal(err == cudaSuccess, "d_dev[%d] %p CUDA Error: %s\n", s, d_devh[s],cudaGetErrorString(err)); 			
+    }
+    else {
+       memcpy(d_host[s],&cc_host[s][2*Zc],sizeof(uint32_t)*(block_length-(2*Zc)));
+    }
+  }
+  if(impp->toutput != NULL) stop_meas(impp->toutput);
+  return d_host;
+}
+

openair1/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim8segmulti.c

@@ -43,6 +43,11 @@
 #include "ldpc_encode_parity_check.c"
 #include "ldpc_generate_coefficient.c"
 
+
+void ini_input32_luts() {
+	return;
+}
+
 int LDPCencoder(uint8_t **input, uint8_t *output, encoder_implemparams_t *impp)
 {
   //set_log(PHY, 4);
@@ -107,7 +112,6 @@ int LDPCencoder(uint8_t **input, uint8_t *output, encoder_implemparams_t *impp)
   //interleave up to 8 transport-block segements at a time
 
   unsigned int i_byte = 0;
-
 #if defined(__AVX512F__) && defined(__AVX512BW__) && defined(__AVX512VBMI__)
   const __m512i masks5[8] = { _mm512_set1_epi8(0x1), _mm512_set1_epi8(0x2),
                               _mm512_set1_epi8(0x4), _mm512_set1_epi8(0x8),
@@ -193,12 +197,9 @@ int LDPCencoder(uint8_t **input, uint8_t *output, encoder_implemparams_t *impp)
   if(impp->tinput != NULL) stop_meas(impp->tinput);
 
   if ((BG==1 && Zc>=176) || (BG==2 && Zc>=72)) {
-    // extend matrix
-    if(impp->tprep != NULL) start_meas(impp->tprep);
-    if(impp->tprep != NULL) stop_meas(impp->tprep);
     //parity check part
     if(impp->tparity != NULL) start_meas(impp->tparity);
-    encode_parity_check_part_optim(cc, dd, BG, Zc, Kb, simd_size, ncols);
+    encode_parity_check_part_optim(cc, dd, BG, Zc, Kb, simd_size, ncols,impp->tinput_memcpy);
     if(impp->tparity != NULL) stop_meas(impp->tparity);
   }
   else {
@@ -207,8 +208,10 @@ int LDPCencoder(uint8_t **input, uint8_t *output, encoder_implemparams_t *impp)
       return(-1);
     }
   }
+  if(impp->toutput != NULL) start_meas(impp->toutput);
   memcpy(output,&cc[2*Zc],(block_length-(2*Zc)));
   memcpy(output+block_length-(2*Zc),dd,((nrows-no_punctured_columns) * Zc-removed_bit));
+  if(impp->toutput != NULL) stop_meas(impp->toutput);
   return 0;
 }
 

openair1/PHY/CODING/nrLDPC_encoder/ldpc_input.cu

@@ -0,0 +1,175 @@
+#include <stdio.h>
+#include <stdint.h>
+#include <cuda_runtime.h>
+/*
+__device__ const uint32_t masks[4] = {0x80,0x8000,0x800000,0x80000000};
+__global__ void ldpc_input_worker(uint32_t **input,uint32_t *cc[4],int block_length,int nseg) {
+
+  int block_off = blockIdx.x*blockDim.x<<2;	
+  int i2 = threadIdx.x<<2;
+  uint32_t *out=cc[blockIdx.y] + block_off + i2;
+  int nseg0 = (blockIdx.y << 5);
+  int nseg1;
+  if ((nseg0 + 32) <= nseg) nseg1 = nseg0+32;
+  else nseg1 = nseg0 + (nseg&31);
+  int bit_offset = i2+block_off;
+  int uint32_offset = bit_offset>>5;  
+  uint32_t mask = masks[(bit_offset&31)>>3]; 
+  uint32_t mask0 = mask>>(bit_offset&7);bit_offset++;
+  uint32_t mask1 = mask>>(bit_offset&7);bit_offset++;
+  uint32_t mask2 = mask>>(bit_offset&7);bit_offset++;
+  uint32_t mask3 = mask>>(bit_offset&7);
+  uint32_t tmp,jmod;
+  uint32_t otmp0,otmp1,otmp2,otmp3;
+  if (bit_offset  < block_length) {
+      tmp=input[nseg0][uint32_offset];
+      otmp0 = ((tmp&mask0) > 0); 
+      otmp1 = ((tmp&mask1) > 0); 
+      otmp2 = ((tmp&mask2) > 0); 
+      otmp3 = ((tmp&mask3) > 0); 
+      for (int j=nseg0+1;j<nseg1;j++) {
+	 tmp=input[j][uint32_offset];
+	 jmod = j&31;
+         otmp0 |= (((tmp&mask0) > 0)<<jmod); 
+         otmp1 |= (((tmp&mask1) > 0)<<jmod); 
+         otmp2 |= (((tmp&mask2) > 0)<<jmod); 
+         otmp3 |= (((tmp&mask3) > 0)<<jmod); 
+      }	      
+      out[0]=otmp0;
+      out[1]=otmp1;
+      out[2]=otmp2;
+      out[3]=otmp3;
+  } 
+}
+*/
+
+#define ITERATIONS 1
+#define NTHREADS 384 
+__device__ uint32_t masks[32] = {
+	0x80,0x40,0x20,0x10,0x8,0x4,0x2,0x1,
+	0x8000,0x4000,0x2000,0x1000,0x800,0x400,0x200,0x100,
+	0x800000,0x400000,0x200000,0x100000,0x80000,0x40000,0x20000,0x10000,
+	0x80000000,0x40000000,0x20000000,0x10000000,0x8000000,0x4000000,0x2000000,0x1000000};
+__global__ void ldpc_input_worker(uint32_t **input,uint32_t *cc[4],int block_length,int nseg) {
+
+  int block_off = blockIdx.x*blockDim.x*ITERATIONS;	
+  int i2 = threadIdx.x*ITERATIONS;
+  uint32_t *out=cc[blockIdx.y] + block_off + i2;
+  int nseg0 = (blockIdx.y << 5);
+  int nseg1;
+  if ((nseg0 + 32) <= nseg) nseg1 = nseg0+32;
+  else nseg1 = nseg0 + (nseg&31);
+  int bit_offset = i2+block_off;
+  int uint32_offset = bit_offset>>5;  
+#if ITERATIONS==1
+  uint32_t mask0 = masks[bit_offset&31];
+#else
+  uint32_t *mask = &masks[bit_offset&31]; 
+  uint32_t mask0 = mask[0];
+#endif
+#if ITERATIONS==4 || ITERATIONS==8
+  uint32_t mask1 = mask[1];
+  uint32_t mask2 = mask[2];
+  uint32_t mask3 = mask[3];
+#endif
+#if ITERATIONS==8
+  uint32_t mask4 = mask[4];
+  uint32_t mask5 = mask[5];
+  uint32_t mask6 = mask[6];
+  uint32_t mask7 = mask[7];
+#endif
+  uint32_t tmp,jmod;
+  uint32_t otmp0;
+#if ITERATIONS==4 || ITERATIONS==8
+  uint32_t otmp1,otmp2,otmp3;
+#endif
+#if ITERATIONS==8
+  uint32_t otmp4,otmp5,otmp6,otmp7;
+#endif
+  if (bit_offset  < block_length) {
+      tmp=input[nseg0][uint32_offset];
+      otmp0 = ((tmp&mask0) > 0); 
+#if ITERATIONS==4 || ITERATIONS==8
+      otmp1 = ((tmp&mask1) > 0); 
+      otmp2 = ((tmp&mask2) > 0); 
+      otmp3 = ((tmp&mask3) > 0); 
+#endif
+#if ITERATIONS==8
+      otmp4 = ((tmp&mask4) > 0); 
+      otmp5 = ((tmp&mask5) > 0); 
+      otmp6 = ((tmp&mask6) > 0); 
+      otmp7 = ((tmp&mask7) > 0);
+#endif 
+      for (int j=nseg0+1;j<nseg1;j++) {
+	 tmp=input[j][uint32_offset];
+	 jmod = j&31;
+         otmp0 |= (((tmp&mask0) > 0)<<jmod); 
+#if ITERATIONS==4 || ITERATIONS==8
+         otmp1 |= (((tmp&mask1) > 0)<<jmod); 
+         otmp2 |= (((tmp&mask2) > 0)<<jmod); 
+         otmp3 |= (((tmp&mask3) > 0)<<jmod);
+#endif 
+#if ITERATIONS==8
+         otmp4 |= (((tmp&mask4) > 0)<<jmod); 
+         otmp5 |= (((tmp&mask5) > 0)<<jmod); 
+         otmp6 |= (((tmp&mask6) > 0)<<jmod); 
+         otmp7 |= (((tmp&mask7) > 0)<<jmod); 
+#endif
+      }	      
+      out[0]=otmp0;
+#if ITERATIONS==4 || ITERATIONS==8
+      out[1]=otmp1;
+      out[2]=otmp2;
+      out[3]=otmp3;
+#endif
+#if ITERATIONS==8
+      out[4]=otmp4;
+      out[5]=otmp5;
+      out[6]=otmp6;
+      out[7]=otmp7;
+#endif
+  } 
+}
+__global__ void circcopy_c_worker(uint32_t **cc,uint32_t **c) {
+
+  int s = blockIdx.x;	
+  int i1 = blockIdx.y;
+  int i  = threadIdx.x;
+  uint32_t tmp;
+  if (i<384) {
+    tmp = cc[s][(i1*384) + i];	  
+    c[s][2*i1*384 + i] = tmp;
+    c[s][(2*i1+1)*384 + i] = tmp; 
+  }
+}
+
+extern "C" int ldpc_input(uint32_t **input,uint32_t *cc[4],int block_length,int nseg) { 
+
+ int numb = block_length/(NTHREADS*ITERATIONS);
+ if ((block_length%(NTHREADS*ITERATIONS)) > 0) numb++;
+ int ns = nseg>>5;
+ if ((nseg&31)>0) ns++;
+
+ dim3 numblocks(numb,ns);
+ ldpc_input_worker<<<numblocks,NTHREADS>>>(input,cc,block_length,nseg);
+ cudaError_t err=cudaPeekAtLastError();
+ if (err!=cudaSuccess) {
+    printf("cuda error: %s (input %p, cc %p, block_length %d, nseg %d, numb %d, ns %d)\n",cudaGetErrorString(err),input,cc,block_length,nseg,numb,ns);
+    exit(-1);
+ }
+ cudaDeviceSynchronize();
+ return(0);
+}
+
+extern "C" int circcopy_c(uint32_t **cc,uint32_t **c,int n_inputs) { 
+
+ dim3 numblocks(n_inputs,22);
+ circcopy_c_worker<<<numblocks,384>>>(cc,c);
+ cudaError_t err=cudaPeekAtLastError();
+ if (err!=cudaSuccess) {
+    printf("cuda error: %s (cc %p, c %p, n_inputs %d)\n",cudaGetErrorString(err),cc,c,n_inputs);
+    exit(-1);
+ }
+ cudaDeviceSynchronize();
+ return(0);
+}

openair1/PHY/CODING/nrLDPC_extern.h

@@ -29,6 +29,7 @@ typedef struct ldpc_interface_s {
   LDPC_shutdownfunc_t *LDPCshutdown;
   LDPC_decoderfunc_t *LDPCdecoder;
   LDPC_encoderfunc_t *LDPCencoder;
+  LDPC_encoderfunc32_t *LDPCencoder32;
 } ldpc_interface_t;
 
 /* functions to load the LDPC shared lib, implemented in openair1/PHY/CODING/nrLDPC_load.c */
@@ -37,6 +38,7 @@ int free_LDPClib(ldpc_interface_t *ldpc_interface);
 
 LDPC_decoderfunc_t LDPCdecoder;
 LDPC_encoderfunc_t LDPCencoder;
+LDPC_encoderfunc32_t LDPCencoder32;
 
 // inline functions:
 #endif

openair1/PHY/CODING/nrLDPC_load.c

@@ -56,14 +56,17 @@ int load_LDPClib(char *version, ldpc_interface_t *itf)
   loader_shlibfunc_t shlib_fdesc[] = {{.fname = "LDPCinit"},
                                       {.fname = "LDPCshutdown"},
                                       {.fname = "LDPCdecoder"},
-                                      {.fname = "LDPCencoder"}};
+                                      {.fname = "LDPCencoder"},
+                                      {.fname = "LDPCencoder32"}};
   int ret;
   ret = load_module_version_shlib(libname, version, shlib_fdesc, sizeofArray(shlib_fdesc), NULL);
   AssertFatal((ret >= 0), "Error loading ldpc decoder");
+  LOG_I(NR_PHY,"loading LDPC version %s\n",version);
   itf->LDPCinit = (LDPC_initfunc_t *)shlib_fdesc[0].fptr;
   itf->LDPCshutdown = (LDPC_shutdownfunc_t *)shlib_fdesc[1].fptr;
   itf->LDPCdecoder = (LDPC_decoderfunc_t *)shlib_fdesc[2].fptr;
   itf->LDPCencoder = (LDPC_encoderfunc_t *)shlib_fdesc[3].fptr;
+  itf->LDPCencoder32 = (LDPC_encoderfunc32_t *)shlib_fdesc[4].fptr;
 
   AssertFatal(itf->LDPCinit() == 0, "error starting LDPC library %s %s\n", libname, version);
 

openair1/PHY/INIT/nr_init.c

@@ -21,6 +21,7 @@
 
 #include "executables/softmodem-common.h"
 #include "executables/nr-softmodem-common.h"
+#include "executables/nr-softmodem.h"
 #include "common/utils/nr/nr_common.h"
 #include "common/ran_context.h"
 #include "PHY/defs_gNB.h"
@@ -107,6 +108,8 @@ void phy_init_nr_gNB(PHY_VARS_gNB *gNB)
   NR_gNB_COMMON *const common_vars  = &gNB->common_vars;
   NR_gNB_PRACH *const prach_vars   = &gNB->prach_vars;
 
+  gNB->use_gpu = use_gpu;
+
   common_vars->analog_bf = cfg->analog_beamforming_ve.analog_bf_vendor_ext.value;
   LOG_I(PHY, "L1 configured with%s analog beamforming\n", common_vars->analog_bf ? "" : "out");
   if (common_vars->analog_bf) {
@@ -421,7 +424,7 @@ void init_DLSCH_struct(PHY_VARS_gNB *gNB, processingData_L1tx_t *msg)
     LOG_D(PHY, "Allocating Transport Channel Buffers for DLSCH %d/%d\n", i, gNB->max_nb_pdsch);
     msg->dlsch[i] = (NR_gNB_DLSCH_t *)malloc16(num_cw * sizeof(NR_gNB_DLSCH_t));
     for (int j = 0; j < num_cw; j++) {
-      msg->dlsch[i][j] = new_gNB_dlsch(fp, grid_size);
+      msg->dlsch[i][j] = new_gNB_dlsch(fp, grid_size, gNB->use_gpu);
     }
   }
 }
@@ -434,7 +437,7 @@ void reset_DLSCH_struct(const PHY_VARS_gNB *gNB, processingData_L1tx_t *msg)
   int num_cw = NR_MAX_NB_LAYERS > 4? 2:1;
   for (int i = 0; i < gNB->max_nb_pdsch; i++) {
     for (int j = 0; j < num_cw; j++) {
-      free_gNB_dlsch(&msg->dlsch[i][j], grid_size, fp);
+      free_gNB_dlsch(&msg->dlsch[i][j], grid_size, fp, gNB->use_gpu);
     }
     free(msg->dlsch[i]);
   }

openair1/PHY/NR_TRANSPORT/nr_dlsch.c

@@ -778,9 +778,11 @@ void nr_generate_pdsch(processingData_L1tx_t *msgTx, int frame, int slot)
   NR_DL_FRAME_PARMS *frame_parms = &gNB->frame_parms;
   time_stats_t *dlsch_encoding_stats = &gNB->dlsch_encoding_stats;
   time_stats_t *tinput = &gNB->tinput;
+  time_stats_t *tinput_memcpy = &gNB->tinput_memcpy;
   time_stats_t *tprep = &gNB->tprep;
   time_stats_t *tparity = &gNB->tparity;
   time_stats_t *toutput = &gNB->toutput;
+  time_stats_t *tconcat = &gNB->tconcat;
   time_stats_t *dlsch_rate_matching_stats = &gNB->dlsch_rate_matching_stats;
   time_stats_t *dlsch_interleaving_stats = &gNB->dlsch_interleaving_stats;
   time_stats_t *dlsch_segmentation_stats = &gNB->dlsch_segmentation_stats;
@@ -827,9 +829,8 @@ void nr_generate_pdsch(processingData_L1tx_t *msgTx, int frame, int slot)
   }
 
   unsigned char output[size_output >> 3] __attribute__((aligned(64)));
-  bzero(output, sizeof(output));
-
   start_meas(dlsch_encoding_stats);
+  bzero(output, sizeof(output));
   if (nr_dlsch_encoding(gNB,
                         msgTx,
                         frame,
@@ -837,15 +838,18 @@ void nr_generate_pdsch(processingData_L1tx_t *msgTx, int frame, int slot)
                         frame_parms,
                         output,
                         tinput,
+                        tinput_memcpy,
                         tprep,
                         tparity,
                         toutput,
+                        tconcat,
                         dlsch_rate_matching_stats,
                         dlsch_interleaving_stats,
                         dlsch_segmentation_stats)
       == -1) {
     return;
   }
+
   stop_meas(dlsch_encoding_stats);
 
   unsigned char *output_ptr = output;

openair1/PHY/NR_TRANSPORT/nr_dlsch.h

@@ -49,12 +49,14 @@ int nr_dlsch_encoding(PHY_VARS_gNB *gNB,
                       NR_DL_FRAME_PARMS *frame_parms,
                       unsigned char *output,
                       time_stats_t *tinput,
+                      time_stats_t *tinput_memcpy,
                       time_stats_t *tprep,
                       time_stats_t *tparity,
                       time_stats_t *toutput,
                       time_stats_t *dlsch_rate_matching_stats,
                       time_stats_t *dlsch_interleaving_stats,
-                      time_stats_t *dlsch_segmentation_stats);
+                      time_stats_t *dlsch_segmentation_stats,
+                      time_stats_t *dlsch_output_stats);
 
 void dump_pdsch_stats(FILE *fd,PHY_VARS_gNB *gNB);
 

openair1/PHY/NR_TRANSPORT/nr_dlsch_coding.c

@@ -39,10 +39,13 @@
 #include <syscall.h>
 #include <openair2/UTIL/OPT/opt.h>
 
+#ifdef ENABLE_CUDA
+#include <cuda_runtime.h>
+#endif
 // #define DEBUG_DLSCH_CODING
 // #define DEBUG_DLSCH_FREE 1
 
-void free_gNB_dlsch(NR_gNB_DLSCH_t *dlsch, uint16_t N_RB, const NR_DL_FRAME_PARMS *frame_parms)
+void free_gNB_dlsch(NR_gNB_DLSCH_t *dlsch, uint16_t N_RB, const NR_DL_FRAME_PARMS *frame_parms, int use_gpumem)
 {
   int max_layers = (frame_parms->nb_antennas_tx < NR_MAX_NB_LAYERS) ? frame_parms->nb_antennas_tx : NR_MAX_NB_LAYERS;
   uint16_t a_segments = MAX_NUM_NR_DLSCH_SEGMENTS_PER_LAYER * max_layers;
@@ -54,7 +57,12 @@ void free_gNB_dlsch(NR_gNB_DLSCH_t *dlsch, uint16_t N_RB, const NR_DL_FRAME_PARM
 
   NR_DL_gNB_HARQ_t *harq = &dlsch->harq_process;
   if (harq->b) {
-    free16(harq->b, a_segments * 1056);
+#ifdef ENABLE_CUDA	  
+    if (use_gpumem) 
+      cudaFreeHost(harq->b);
+    else	  
+#endif  
+      free16(harq->b, a_segments * 1056);
     harq->b = NULL;
   }
   if (harq->f) {
@@ -68,7 +76,7 @@ void free_gNB_dlsch(NR_gNB_DLSCH_t *dlsch, uint16_t N_RB, const NR_DL_FRAME_PARM
   free(harq->c);
 }
 
-NR_gNB_DLSCH_t new_gNB_dlsch(NR_DL_FRAME_PARMS *frame_parms, uint16_t N_RB)
+NR_gNB_DLSCH_t new_gNB_dlsch(NR_DL_FRAME_PARMS *frame_parms, uint16_t N_RB, int use_gpumem)
 {
   int max_layers = (frame_parms->nb_antennas_tx < NR_MAX_NB_LAYERS) ? frame_parms->nb_antennas_tx : NR_MAX_NB_LAYERS;
   uint16_t a_segments = MAX_NUM_NR_DLSCH_SEGMENTS_PER_LAYER * max_layers; // number of segments to be allocated
@@ -89,16 +97,35 @@ NR_gNB_DLSCH_t new_gNB_dlsch(NR_DL_FRAME_PARMS *frame_parms, uint16_t N_RB)
   bzero(harq->b, dlsch_bytes);
 
   harq->c = (uint8_t **)malloc16(a_segments * sizeof(uint8_t *));
+#ifdef ENABLE_CUDA
+  uint8_t *c_devh[a_segments];
+#endif
   for (int r = 0; r < a_segments; r++) {
     // account for filler in first segment and CRCs for multiple segment case
     // [hna] 8448 is the maximum CB size in NR
     //       68*348 = 68*(maximum size of Zc)
     //       In section 5.3.2 in 38.212, the for loop is up to N + 2*Zc (maximum size of N is 66*Zc, therefore 68*Zc)
-    harq->c[r] = malloc16(8448);
+#ifdef ENABLE_CUDA
+    if (use_gpumem) {
+      cudaHostAlloc((void**)&harq->c[r], 8448/8, cudaHostAllocMapped);
+      cudaError_t err=cudaHostAlloc((void**)&harq->c[r],(8448/8)*sizeof(uint8_t),cudaHostAllocMapped);
+      AssertFatal(err == cudaSuccess,"CUDA Error (harq->c[%d]): %s\n", r,cudaGetErrorString(err));
+      err=cudaHostGetDevicePointer((void**)&c_devh[r], harq->c[r], 0);
+      AssertFatal(err == cudaSuccess,"CUDA Error (cudaHostGetDevicePointer) harq->c_devh[%d]: %s\n", r,cudaGetErrorString(err));
+    }
+    else 
+#endif
+      harq->c[r] = malloc16(8448/8);
     AssertFatal(harq->c[r], "cannot allocate harq->c[%d]\n", r);
-    bzero(harq->c[r], 8448);
+    bzero(harq->c[r], 8448/8);
   }
-
+#ifdef ENABLE_CUDA
+  if (use_gpumem) {
+    cudaError_t err=cudaMalloc((void**)&harq->c_dev,a_segments*sizeof(uint8_t*));
+    err=cudaMemcpy(harq->c_dev,c_devh,a_segments*sizeof(uint8_t*),cudaMemcpyHostToDevice);
+    AssertFatal(err == cudaSuccess,"CUDA Error (memcpy c_devh -> input_dev): %s\n", cudaGetErrorString(err));
+  }
+#endif
   harq->f = malloc16(N_RB * NR_SYMBOLS_PER_SLOT * NR_NB_SC_PER_RB * 8 * NR_MAX_NB_LAYERS);
   AssertFatal(harq->f, "cannot allocate harq->f\n");
   bzero(harq->f, N_RB * NR_SYMBOLS_PER_SLOT * NR_NB_SC_PER_RB * 8 * NR_MAX_NB_LAYERS);
@@ -113,9 +140,11 @@ int nr_dlsch_encoding(PHY_VARS_gNB *gNB,
                       NR_DL_FRAME_PARMS *frame_parms,
                       unsigned char *output,
                       time_stats_t *tinput,
+                      time_stats_t *tinput_memcpy,
                       time_stats_t *tprep,
                       time_stats_t *tparity,
                       time_stats_t *toutput,
+		      time_stats_t *tconcat,
                       time_stats_t *dlsch_rate_matching_stats,
                       time_stats_t *dlsch_interleaving_stats,
                       time_stats_t *dlsch_segmentation_stats)
@@ -129,10 +158,13 @@ int nr_dlsch_encoding(PHY_VARS_gNB *gNB,
                                                        .nb_TBs = msgTx->num_pdsch_slot,
                                                        .threadPool = &gNB->threadPool,
                                                        .tinput = tinput,
+                                                       .tinput_memcpy = tinput_memcpy,
                                                        .tprep = tprep,
                                                        .tparity = tparity,
                                                        .toutput = toutput,
-                                                       .TBs = TBs};
+						       .tconcat = tconcat,
+                                                       .TBs = TBs,
+  						       .use_gpu = gNB->use_gpu};
 
   int num_segments = 0;
 
@@ -171,6 +203,7 @@ int nr_dlsch_encoding(PHY_VARS_gNB *gNB,
       phy_stats->dlsch_stats.current_Qm = rel15->qamModOrder[0];
     }
 
+    start_meas(dlsch_segmentation_stats);
     int max_bytes = MAX_NUM_NR_DLSCH_SEGMENTS_PER_LAYER * rel15->nrOfLayers * 1056;
     int B;
     if (A > NR_MAX_PDSCH_TBS) {
@@ -205,7 +238,6 @@ int nr_dlsch_encoding(PHY_VARS_gNB *gNB,
     TB_parameters->BG = rel15->maintenance_parms_v3.ldpcBaseGraph;
     TB_parameters->Z = harq->Z;
     TB_parameters->A = A;
-    start_meas(dlsch_segmentation_stats);
     TB_parameters->Kb = nr_segmentation(harq->b,
                                         harq->c,
                                         B,
@@ -277,7 +309,9 @@ int nr_dlsch_encoding(PHY_VARS_gNB *gNB,
       reset_meas(&segment_parameters->ts_rate_match);
       reset_meas(&segment_parameters->ts_ldpc_encode);
     }
-
+#ifdef ENABLE_CUDA
+    if (gNB->use_gpu) TB_parameters->c_dev = (uint8_t**)harq->c_dev;
+#endif  
     segments_offset += TB_parameters->C;
 
     /* output and its parts for each dlsch should be aligned on 64 bytes (or 8 * 64 bits)

openair1/PHY/NR_TRANSPORT/nr_transport_proto.h

@@ -106,9 +106,9 @@ void nr_generate_pbch(PHY_VARS_gNB *gNB,
 void nr_init_pbch_interleaver(uint8_t *interleaver);
 uint32_t nr_pbch_extra_byte_generation(int sfn, int n_hf, int ssb_index, int ssb_sc_offset, int Lmax);
 
-NR_gNB_DLSCH_t new_gNB_dlsch(NR_DL_FRAME_PARMS *frame_parms, uint16_t N_RB);
+NR_gNB_DLSCH_t new_gNB_dlsch(NR_DL_FRAME_PARMS *frame_parms, uint16_t N_RB,int use_gpu);
 
-void free_gNB_dlsch(NR_gNB_DLSCH_t *dlsch, uint16_t N_RB, const NR_DL_FRAME_PARMS *frame_parms);
+void free_gNB_dlsch(NR_gNB_DLSCH_t *dlsch, uint16_t N_RB, const NR_DL_FRAME_PARMS *frame_parms,int use_gpu);
 
 /** \brief This function is the top-level entry point to PUSCH demodulation, after frequency-domain transformation and channel estimation.  It performs
     - RB extraction (signal and channel estimates)

openair1/PHY/defs_gNB.h

@@ -62,6 +62,9 @@ typedef struct {
   uint8_t *b;
   /// Pointers to transport block segments
   uint8_t **c;
+#ifdef ENABLE_CUDA
+  uint8_t *c_dev;
+#endif
   /// Frame where current HARQ round was sent
   uint32_t frame;
   /// Interleaver outputs
@@ -504,9 +507,11 @@ typedef struct PHY_VARS_gNB_s {
   time_stats_t dlsch_resource_mapping_stats;
   time_stats_t dlsch_precoding_stats;
   time_stats_t tinput;
+  time_stats_t tinput_memcpy;
   time_stats_t tprep;
   time_stats_t tparity;
   time_stats_t toutput;
+  time_stats_t tconcat;
   
   time_stats_t dlsch_rate_matching_stats;
   time_stats_t dlsch_interleaving_stats;
@@ -556,6 +561,7 @@ typedef struct PHY_VARS_gNB_s {
   int L1_tx_thread_core;
   struct processingData_L1tx *msgDataTx;
   void *scopeData;
+  int use_gpu;
 } PHY_VARS_gNB;
 
 struct puschSymbolReqId {

openair1/SIMULATION/NR_PHY/dlschsim.c

@@ -516,7 +516,7 @@ int main(int argc, char **argv)
         bzero(output, sizeof(output));
 	if (input_fd == NULL) {
     msgDataTx.num_pdsch_slot = 1;
-	  nr_dlsch_encoding(gNB, &msgDataTx, frame, slot, frame_parms, output, NULL, NULL, NULL, NULL, NULL, NULL, NULL);
+	  nr_dlsch_encoding(gNB, &msgDataTx, frame, slot, frame_parms, output, NULL, NULL, NULL, NULL, NULL, NULL, NULL, NULL, NULL);
 	}
 
 	for (SNR = snr0; SNR < snr1 && !stop; SNR += snr_step) {

openair1/SIMULATION/NR_PHY/dlsim.c

@@ -182,6 +182,7 @@ extern void fix_scd(NR_ServingCellConfig_t *scd);// forward declaration
 /* specific dlsim DL preprocessor: uses rbStart/rbSize/mcs/nrOfLayers from command line of dlsim */
 int g_mcsIndex = -1, g_mcsTableIdx = 0, g_rbStart = -1, g_rbSize = -1, g_nrOfLayers = 1, g_pmi = 0;
 
+
 void nr_dlsim_preprocessor(module_id_t module_id, frame_t frame, slot_t slot)
 {
   NR_UE_info_t *UE_info = RC.nrmac[module_id]->UE_info.connected_ue_list[0];
@@ -296,8 +297,9 @@ void validate_input_pmi(nfapi_nr_config_request_scf_t *gNB_config,
               num_antenna_ports, pmi_pdu->num_ant_ports, pmi);
 }
 
-
 configmodule_interface_t *uniqCfg = NULL;
+extern uint32_t use_gpu;
+
 int main(int argc, char **argv)
 {
   stop = false;
@@ -393,7 +395,7 @@ int main(int argc, char **argv)
 
   FILE *scg_fd=NULL;
 
-  while ((c = getopt(argc, argv, "--:O:f:hA:p:f:g:i:n:s:S:t:v:x:y:z:o:H:M:N:F:GR:d:PI:L:a:b:e:m:w:T:U:q:X:Y:Z:")) != -1) {
+  while ((c = getopt(argc, argv, "--:O:f:hA:p:f:g:i:n:s:S:t:v:x:y:z:o:H:M:N:F:GR:d:PQI:L:a:b:e:m:w:T:U:q:X:Y:Z:")) != -1) {
 
     /* ignore long options starting with '--', option '-O' and their arguments that are handled by configmodule */
     /* with this opstring getopt returns 1 for non-option arguments, refer to 'man 3 getopt' */
@@ -513,7 +515,10 @@ int main(int argc, char **argv)
       print_perf=1;
       cpu_meas_enabled = 1;
       break;
-      
+     
+    case 'Q':
+      use_gpu=1;
+      break; 
     case 'I':
       max_ldpc_iterations = atoi(optarg);
       break;
@@ -1036,9 +1041,11 @@ printf("%d\n", slot);
     reset_meas(&gNB->dlsch_encoding_stats);
     reset_meas(&gNB->dci_generation_stats);
     reset_meas(&gNB->tinput);
+    reset_meas(&gNB->tinput_memcpy);
     reset_meas(&gNB->tprep);
     reset_meas(&gNB->tparity);
     reset_meas(&gNB->toutput);
+    reset_meas(&gNB->tconcat);
     reset_meas(&gNB->phase_comp_stats);
 
     uint32_t errors_scrambling[16] = {0};
@@ -1337,12 +1344,19 @@ printf("%d\n", slot);
       printStatIndent2(&gNB->dci_generation_stats, "DCI encoding time");
       printStatIndent2(&gNB->dlsch_encoding_stats,"DLSCH encoding time");
       printStatIndent3(&gNB->dlsch_segmentation_stats,"DLSCH segmentation time");
+      gNB->tinput.trials = gNB->dlsch_encoding_stats.trials;
       printStatIndent3(&gNB->tinput,"DLSCH LDPC input processing time");
-      printStatIndent3(&gNB->tprep,"DLSCH LDPC input preparation time");
+      gNB->tinput_memcpy.trials = gNB->dlsch_encoding_stats.trials;
+      printStatIndent3(&gNB->tinput_memcpy,"DLSCH LDPC input memcpy time");
+      gNB->tparity.trials = gNB->dlsch_encoding_stats.trials;
       printStatIndent3(&gNB->tparity,"DLSCH LDPC parity generation time");
+      gNB->toutput.trials = gNB->dlsch_encoding_stats.trials;
       printStatIndent3(&gNB->toutput,"DLSCH LDPC output generation time");
+      gNB->dlsch_rate_matching_stats.trials = gNB->dlsch_encoding_stats.trials;
       printStatIndent3(&gNB->dlsch_rate_matching_stats,"DLSCH Rate Matching time");
+      gNB->dlsch_interleaving_stats.trials = gNB->dlsch_encoding_stats.trials;
       printStatIndent3(&gNB->dlsch_interleaving_stats,  "DLSCH Interleaving time");
+      printStatIndent3(&gNB->tconcat,  "DLSCH Segment Concatenation time");
       printStatIndent2(&gNB->dlsch_modulation_stats,"DLSCH modulation time");
       printStatIndent2(&gNB->dlsch_scrambling_stats, "DLSCH scrambling time");
       printStatIndent2(&gNB->dlsch_precoding_stats,"DLSCH Mapping/Precoding time");

openair1/SIMULATION/NR_PHY/nr_unitary_common.c

@@ -32,6 +32,7 @@
 
 // Define signal handler to attempt graceful termination
 bool stop = false;
+uint32_t use_gpu=0;
 
 void sigint_handler(int arg)
 {

radio/fhi_72/CMakeLists.txt

@@ -16,8 +16,8 @@ add_library(oran_fhlib_5g MODULE
             oran-init.c
 )
 
-set(E_VERSION 5.1.6)
-set(F_VERSION 6.1.4)
+set(E_VERSION 5.1.5)
+set(F_VERSION 6.1.3)
 
 find_package(xran REQUIRED)
 if(xran_VERSION VERSION_EQUAL E_VERSION)