use_gpu/ENABLE_CUDA/USE_CUDA + sims

CMakeLists.txt

@@ -230,6 +230,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
@@ -1084,15 +1088,30 @@ target_link_libraries(PHY_UE PRIVATE asn1_lte_rrc_hdrs asn1_nr_rrc_hdrs UTIL)
 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_80 -g -ggdb") # A100
+  #SET(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS};-arch=sm_90") # GH200
+  #SET(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS};-arch=sm_87") # Jetson
+  #SET(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS};-arch=sm_89") # L40S
+  SET(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS};-arch=sm_121") # GB10
+  set(CUDA_VERBOS_BUILD_ON)
+  cuda_add_library(PHY_NR ${PHY_NR_SRC} ${PHY_NR_SRC_CU})
+  cuda_add_library(PHY_NR_UE ${PHY_NR_UE_SRC} ${PHY_NR_UE_SRC_CU})
+else()
+  add_library(PHY_NR ${PHY_NR_SRC})
+  add_library(PHY_NR_UE ${PHY_NR_UE_SRC})
+endif()
 target_link_libraries(PHY_NR nr_phy_common nr_common nr_fapi_p5 polar smallblock)
 
 add_library(PHY_NR_NO_AVX_256 ${PHY_NR_SRC})
 target_link_libraries(PHY_NR_NO_AVX_256 nr_phy_common nr_common)
 target_compile_definitions(PHY_NR_NO_AVX_256 PUBLIC USE_128BIT)
 
-add_library(PHY_NR_UE ${PHY_NR_UE_SRC})
-target_link_libraries(PHY_NR_UE PRIVATE asn1_nr_rrc_hdrs nr_phy_common nr_common UTIL polar smallblock)
+target_link_libraries(PHY_NR_UE asn1_nr_rrc_hdrs nr_phy_common nr_common UTIL polar smallblock)
 
 add_library(PHY_RU ${PHY_SRC_RU})
 target_link_libraries(PHY_RU PRIVATE asn1_lte_rrc_hdrs UTIL)
@@ -1988,12 +2007,19 @@ add_executable(ldpctest
 )
 target_link_libraries(ldpctest PRIVATE
   -Wl,--start-group UTIL SIMU PHY_COMMON PHY_NR_COMMON -Wl,--end-group
-  m pthread dl ${T_LIB}
+  m pthread dl ${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'
   -Wl,--whole-archive crc_byte -Wl,--no-whole-archive
 )
+if (ENABLE_CUDA)
+  # ldpctest.c uses cudaHostAlloc(), so include dir/link libr to make available
+  # TODO don't write it directly into ldpctest, but separate compilation unit?
+  find_package(CUDA REQUIRED)
+  target_link_libraries(ldpctest PRIVATE ${CUDA_LIBRARIES})
+  target_include_directories(ldpctest PRIVATE ${CUDA_INCLUDE_DIRS})
+endif()
 
 add_library(physim_common OBJECT ${OPENAIR1_DIR}/SIMULATION/NR_PHY/nr_unitary_common.c)
 target_link_libraries(physim_common PRIVATE UTIL)

cmake_targets/build_oai

@@ -136,6 +136,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"
 }
@@ -384,6 +386,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.c

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

executables/nr-softmodem.h

@@ -39,6 +39,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();

executables/nr-uesoftmodem.c

@@ -109,6 +109,8 @@ uint64_t        sidelink_frequency[MAX_NUM_CCs][4];
 // UE and OAI config variables
 double            cpuf;
 
+extern uint32_t use_gpu;
+
 int create_tasks_nrue(uint32_t ue_nb) {
   LOG_D(NR_RRC, "%s(ue_nb:%d)\n", __FUNCTION__, ue_nb);
   itti_wait_ready(1);

executables/softmodem-common.c

@@ -46,6 +46,8 @@ char *worker_config=NULL;
 int usrp_tx_thread = 0;
 uint8_t nfapi_mode=0;
 
+uint32_t use_gpu=0;
+
 static struct timespec start;
 
 static softmodem_params_t softmodem_params;

executables/softmodem-common.h

@@ -100,7 +100,7 @@ extern "C"
 #define CONFIG_HLP_TADV                                                                                                      \
   "Set RF board timing_advance to compensate fix delay inside the RF board between Rx and Tx timestamps (RF board internal " \
   "issues)\n"
-
+#define CONFIG_HLP_USE_GPU       "Use GPU for processing offload"
 /*-----------------------------------------------------------------------------------------------------------------------------------------------------*/
 /*                                            command line parameters common to eNodeB and UE                                                          */
 /*   optname                 helpstr                  paramflags      XXXptr                              defXXXval              type         numelt   */
@@ -128,6 +128,7 @@ extern "C"
 #define DEFAULT_PDU_ID      softmodem_params.default_pdu_session_id
 
 extern int usrp_tx_thread;
+extern uint32_t use_gpu;
 // clang-format off
 #define CMDLINE_PARAMS_DESC {  \
   {"rf-config-file",        CONFIG_HLP_RFCFGF,        0,              .strptr=&RF_CONFIG_FILE,                .defstrval=NULL,          TYPE_STRING, 0},  \
@@ -165,6 +166,7 @@ extern int usrp_tx_thread;
   {"E" ,                    CONFIG_HLP_TQFS,          PARAMFLAG_BOOL, .iptr=&softmodem_params.threequarter_fs, .defintval=0,            TYPE_INT,    0}, \
   {"imscope" ,              CONFIG_HLP_IMSCOPE,       PARAMFLAG_BOOL, .uptr=&enable_imscope,                   .defintval=0,            TYPE_UINT,   0}, \
   {"imscope-record" ,       CONFIG_HLP_IMSCOPE_RECORD,PARAMFLAG_BOOL, .uptr=&enable_imscope_record,            .defintval=0,            TYPE_UINT,   0}, \
+  {"use_gpu",               CONFIG_HLP_USE_GPU,        0,                .uptr=&use_gpu,                      .defintval=0,                     TYPE_UINT, 0},           \
   {"default-pdu-id",        NULL,                     0,              .iptr=&DEFAULT_PDU_ID,                   .defintval=-1,           TYPE_INT,    0}, \
 }
 // clang-format on
@@ -210,6 +212,7 @@ extern int usrp_tx_thread;
     { .s5 = { NULL } },                     \
     { .s5 = { NULL } },                     \
     { .s5 = { NULL } },                     \
+    { .s5 = { NULL } },                     \
 }
 // clang-format on
 

openair1/PHY/CODING/CMakeLists.txt

@@ -33,6 +33,68 @@ add_dependencies(nr_ulschsim ldpc ldpc_orig)
 add_dependencies(nr_dlsim ldpc ldpc_orig)
 add_dependencies(nr_dlschsim ldpc ldpc_orig)
 
+##############################################
+# Base CUDA setting
+##############################################
+
+add_boolean_option(ENABLE_LDPC_CUDA OFF "Build support for CUDA" OFF)
+if (ENABLE_LDPC_CUDA)
+  enable_language(CUDA)
+
+  # it would be possible to auto-detect compute capailities: https://stackoverflow.com/a/68223399
+  #include(FindCUDA/select_compute_arch)
+  #CUDA_DETECT_INSTALLED_GPUS(INSTALLED_GPU_CCS_1)
+
+  # Nvidia recommends to build against all targeted archs: https://forums.developer.nvidia.com/t/261162
+  # this is for cmake 3.16 (OAI minimum), starting at cmake 3.18 there is another syntax, see below
+
+  SET(CUDA_VERBOSE_BUILD ON)
+
+  # version using NVLink-C2C in Grace Hopper
+  add_library(ldpc_cuda MODULE
+                   nrLDPC_decoder/nrLDPC_decoder_cuda.c
+                   nrLDPC_decoder/nrLDPC_decoder.c
+		   nrLDPC_encoder/ldpc_encoder_optim8segmulti.c
+		   nrLDPC_decoder/nrLDPC_decoder_BG1_cuda.cu
+		   nrLDPC_encoder/ldpc_encoder_cuda32.c
+		   nrLDPC_encoder/ldpc_BG1_Zc384_32bit.cu
+		   nrLDPC_encoder/ldpc_input.cu
+		   nrLDPC_encoder/ldpc_input32.c
+                   # 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_decoder.c
+                   nrLDPC_coding/nrLDPC_coding_segment/nrLDPC_coding_segment_decoder_cuda.c
+                   nrLDPC_coding/nrLDPC_coding_segment/nrLDPC_coding_segment_encoder.c
+                   nrLDPC_coding/nrLDPC_coding_segment/nr_rate_matching.c
+		   nrLDPC_coding/nrLDPC_coding_segment/ldpc_deinterleave.cu
+		   nrLDPC_coding/nrLDPC_coding_segment/ldpc_rate_matching.cu
+  )
+  set_target_properties(ldpc_cuda PROPERTIES CUDA_ARCHITECTURES "80;87;90")
+  # we need to find the right include directory, but ti seems that
+  # enable_language(CUDA) does not make this available(?), so also use the CUDA package
+  find_package(CUDA REQUIRED)
+  target_include_directories(ldpc_cuda PRIVATE ${CUDA_INCLUDE_DIRS})
+  
+  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})
+
+  add_dependencies(ldpctest ldpc_cuda)
+  add_dependencies(nr-softmodem ldpc_cuda)
+  add_dependencies(nr-uesoftmodem ldpc_cuda)
+  add_dependencies(nr_ulsim ldpc_cuda)
+  add_dependencies(nr_ulschsim ldpc_cuda)
+  add_dependencies(nr_dlsim ldpc_cuda)
+  add_dependencies(nr_dlschsim ldpc_cuda)
+
+  target_compile_definitions(ldpc_cuda PRIVATE USE_CUDA)
+
+endif()
+
 add_library(crc_byte OBJECT crc_byte.c)
 
 add_subdirectory(nrLDPC_coding)

openair1/PHY/CODING/nrLDPC_coding/nrLDPC_coding_interface.h

@@ -116,6 +116,7 @@ typedef struct nrLDPC_slot_decoding_parameters_s{
   int nb_TBs;
   tpool_t *threadPool;
   nrLDPC_TB_decoding_parameters_t *TBs;
+  int use_gpu;
 } nrLDPC_slot_decoding_parameters_t;
 
 /**
@@ -185,6 +186,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;
 
 /**
@@ -216,6 +218,7 @@ typedef struct nrLDPC_slot_encoding_parameters_s{
   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)(int);

openair1/PHY/CODING/nrLDPC_coding/nrLDPC_coding_segment/ldpc_deinterleave.cu

+// deinterleave_u16.cu
+#include <cuda_runtime.h>
+#include <stdio.h>
+#include <stdint.h>
+
+__global__ void deinterleave_i16_2(int16_t* __restrict__ e,
+                                   const int16_t* __restrict__ f,
+                                   int E1,
+				   int E2,
+				   int r_firstE2)
+{
+    int g = (int)(blockIdx.x * blockDim.x + threadIdx.x);
+    int r = (int)blockIdx.y;
+    int E = (r<r_firstE2) ? E1 : E2;
+    int EQm = E/2;
+    if (g >= EQm) return;
+
+    int r_off = r<r_firstE2 ? r*E1 : ((r_firstE2*E1)+(r-r_firstE2)*E2);
+    const int16_t *in = f +r_off + 2*g;
+
+    int16_t* e0 = e + r_off;
+    int16_t* e1 = e0 + EQm;
+    e0[g] = in[0];
+    e1[g] = in[1];
+}
+
+__global__ void deinterleave_i16_4(int16_t* __restrict__ e,
+                                   const int16_t* __restrict__ f,
+                                   int E1,
+                       		   int E2,
+				   int r_firstE2)
+{
+    int g = (int)(blockIdx.x * blockDim.x + threadIdx.x);
+    int r = (int)blockIdx.y;
+    int E = (r<r_firstE2) ? E1 : E2;
+    int EQm = E/4;
+    if (g >= EQm) return;
+
+    int r_off = r<r_firstE2 ? r*E1 : ((r_firstE2*E1)+(r-r_firstE2)*E2);
+    const int16_t *in = f + r_off + 4*g;
+    int16_t* e0 = e + r_off;
+    int16_t* e1 = e0 + EQm;
+    int16_t* e2 = e1 + EQm;
+    int16_t* e3 = e2 + EQm;
+    e0[g] = in[0];
+    e1[g] = in[1];
+    e2[g] = in[2];
+    e3[g] = in[3];
+}
+
+__global__ void deinterleave_i16_6(int16_t* __restrict__ e,
+                                   const int16_t* __restrict__ f,
+                                   int E1,
+				   int E2,
+				   int r_firstE2)
+{
+    int g = (int)(blockIdx.x * blockDim.x + threadIdx.x);
+    int r = (int)blockIdx.y;
+    int E = (r<r_firstE2) ? E1 : E2;
+    int EQm = E/6;
+    if (g >= EQm) return;
+
+    int r_off = r<r_firstE2 ? r*E1 : (r_firstE2*E1)+(r-r_firstE2)*E2;
+    const int16_t* in = f + r_off + 6*g;
+    int16_t* e0 = e + r_off;
+    int16_t* e1 = e0 + EQm;
+    int16_t* e2 = e1 + EQm;
+    int16_t* e3 = e2 + EQm;
+    int16_t* e4 = e3 + EQm;
+    int16_t* e5 = e4 + EQm;
+    e0[g] = in[0];
+    e1[g] = in[1];
+    e2[g] = in[2];
+    e3[g] = in[3];
+    e4[g] = in[4];
+    e5[g] = in[5];
+}
+
+__global__ void deinterleave_i16_8(int16_t* __restrict__ e,
+                                   const int16_t* __restrict__ f,
+                                   const int E1,
+				   const int E2,
+				   const int r_firstE2)
+{
+    int g = (int)(blockIdx.x * blockDim.x + threadIdx.x);
+    int r = (int)blockIdx.y;
+    int E = (r<r_firstE2) ? E1 : E2;
+    int EQm = E/8;
+    if (g >= EQm) return;
+    int r_off = r<r_firstE2 ? r*E1 : (r_firstE2*E1)+(r-r_firstE2)*E2;
+
+    const int16_t *in = (f + r_off + 8*g);
+    int16_t* e0 = e + r_off;
+    int16_t* e1 = e0 + EQm;
+    int16_t* e2 = e1 + EQm;
+    int16_t* e3 = e2 + EQm;
+    int16_t* e4 = e3 + EQm;
+    int16_t* e5 = e4 + EQm;
+    int16_t* e6 = e5 + EQm;
+    int16_t* e7 = e6 + EQm;
+    e0[g] = in[0];
+    e1[g] = in[1];
+    e2[g] = in[2];
+    e3[g] = in[3];
+    e4[g] = in[4];
+    e5[g] = in[5];
+    e6[g] = in[6];
+    e7[g] = in[7];
+}
+
+// Host launcher
+extern "C" void launch_deinterleave_i16(int Qm, int E1, int E2, int C, int r_firstE2,int16_t* e, const int16_t* f,cudaStream_t *s,int8_t sidx)
+{
+    const int threads = 256;
+    dim3 blocks(((E2/Qm)  + threads - 1) / threads,C);
+
+    switch (Qm) {
+        case 2: deinterleave_i16_2<<<blocks, threads, 0, s[sidx]>>>(e, f, E1,E2,r_firstE2); break;
+        case 4: deinterleave_i16_4<<<blocks, threads, 0, s[sidx]>>>(e, f, E1,E2,r_firstE2); break;
+        case 6: deinterleave_i16_6<<<blocks, threads, 0, s[sidx]>>>(e, f, E1,E2,r_firstE2); break;
+        case 8: deinterleave_i16_8<<<blocks, threads, 0, s[sidx]>>>(e, f, E1,E2,r_firstE2); break;
+        default: /* unsupported */ break;
+    }
+    cudaError_t err=cudaPeekAtLastError();
+  
+    if (err!=cudaSuccess) {
+      printf("cuda error (deinterleave_i16): %s (e %p, f %p, E1 %d, E2 %d, Qm %d, C %d)\n",cudaGetErrorString(err),e,f,E1,E2,Qm,C);
+      exit(-1);
+    }
+
+}

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

@@ -36,6 +36,7 @@
 #include <stdint.h>
 #include <syscall.h>
 #include <time.h>
+#include <stdbool.h>
 // #define gNB_DEBUG_TRACE
 
 #define OAI_LDPC_DECODER_MAX_NUM_LLR 27000 // 26112 // NR_LDPC_NCOL_BG1*NR_LDPC_ZMAX = 68*384
@@ -225,9 +226,16 @@ static void nr_process_decode_segment(void *arg)
   completed_task_ans(rdata->ans);
 }
 
+#ifdef ENABLE_CUDA
+void nr_process_decode_segment_cuda(nrLDPC_TB_decoding_parameters_t *);
+#endif
 int nrLDPC_prepare_TB_decoding(nrLDPC_slot_decoding_parameters_t *nrLDPC_slot_decoding_parameters,
                                int pusch_id,
-                               thread_info_tm_t *t_info)
+                               thread_info_tm_t *t_info
+#ifdef ENABLE_CUDA
+			       ,int use_gpu
+#endif
+			       )
 {
   nrLDPC_TB_decoding_parameters_t *nrLDPC_TB_decoding_parameters = &nrLDPC_slot_decoding_parameters->TBs[pusch_id];
 
@@ -239,6 +247,14 @@ int nrLDPC_prepare_TB_decoding(nrLDPC_slot_decoding_parameters_t *nrLDPC_slot_de
   decParams.outMode = nrLDPC_outMode_BIT;
   
   for (int r = 0; r < nrLDPC_TB_decoding_parameters->C; r++) {
+#ifdef ENABLE_CUDA
+    if (use_gpu == 1 && decParams.Z >= 128 && decParams.BG == 1 && r==0) {
+    // Call CUDA LDPC decoder for all segments
+      nr_process_decode_segment_cuda(nrLDPC_TB_decoding_parameters);
+      break;
+    }
+    else 
+#endif
     {
       nrLDPC_decoding_parameters_t *rdata = &((nrLDPC_decoding_parameters_t *)t_info->buf)[t_info->len];
       DevAssert(t_info->len < t_info->cap);
@@ -286,19 +302,38 @@ int nrLDPC_prepare_TB_decoding(nrLDPC_slot_decoding_parameters_t *nrLDPC_slot_de
   return nrLDPC_TB_decoding_parameters->C;
 }
 
-int32_t nrLDPC_coding_init(void)
+#ifdef ENABLE_CUDA
+void nrLDPC_coding_init_cuda(int);
+#endif
+
+int32_t nrLDPC_coding_init(int max_num_pxsch)
 {
+  LOG_I(NR_PHY, "Initializing coding library\n");
+#ifdef ENABLE_CUDA
+  LOG_I(NR_PHY, "Calling cuda_support_init()\n");
+  nrLDPC_coding_init_cuda(max_num_pxsch);
+#endif
   return 0;
 }
 
+#ifdef ENABLE_CUDA
+void nrLDPC_coding_shutdown_cuda(void);
+#endif
+
 int32_t nrLDPC_coding_shutdown(void)
 {
+#ifdef ENABLE_CUDA
+  nrLDPC_coding_shutdown_cuda();
+#endif
   return 0;
 }
 
 int32_t nrLDPC_coding_decoder(nrLDPC_slot_decoding_parameters_t *nrLDPC_slot_decoding_parameters)
 {
   int nbSegments = 0;
+#ifdef ENABLE_CUDA
+  int use_gpu = nrLDPC_slot_decoding_parameters->use_gpu;
+#endif
   for (int pusch_id = 0; pusch_id < nrLDPC_slot_decoding_parameters->nb_TBs; pusch_id++) {
     nrLDPC_TB_decoding_parameters_t *nrLDPC_TB_decoding_parameters = &nrLDPC_slot_decoding_parameters->TBs[pusch_id];
     nbSegments += nrLDPC_TB_decoding_parameters->C;
@@ -309,11 +344,27 @@ int32_t nrLDPC_coding_decoder(nrLDPC_slot_decoding_parameters_t *nrLDPC_slot_dec
   thread_info_tm_t t_info = {.buf = (uint8_t *)arr, .len = 0, .cap = nbSegments, .ans = &ans};
 
   for (int pusch_id = 0; pusch_id < nrLDPC_slot_decoding_parameters->nb_TBs; pusch_id++) {
-    (void)nrLDPC_prepare_TB_decoding(nrLDPC_slot_decoding_parameters, pusch_id, &t_info);
+    (void)nrLDPC_prepare_TB_decoding(nrLDPC_slot_decoding_parameters, pusch_id, &t_info
+#ifdef ENABLE_CUDA
+		    ,use_gpu
+#endif
+		    );
   }
 
   // Execute thread pool tasks
-  join_task_ans(t_info.ans);
+#ifdef ENABLE_CUDA
+  bool do_join=false;
+  // check if at least one PUSCH has a Zc<384 or BG=2
+  for (int pusch_id = 0; pusch_id < nrLDPC_slot_decoding_parameters->nb_TBs; pusch_id++) {
+    nrLDPC_TB_decoding_parameters_t *nrLDPC_TB_decoding_parameters = &nrLDPC_slot_decoding_parameters->TBs[pusch_id];
+    if (use_gpu == 0 || nrLDPC_TB_decoding_parameters->Z < 128 ||  nrLDPC_TB_decoding_parameters->BG == 2 ) {
+	do_join=true;    
+	break;
+    }
+  }
+  if (do_join)
+#endif
+    join_task_ans(t_info.ans);
 
   for (int pusch_id = 0; pusch_id < nrLDPC_slot_decoding_parameters->nb_TBs; pusch_id++) {
     nrLDPC_TB_decoding_parameters_t *nrLDPC_TB_decoding_parameters = &nrLDPC_slot_decoding_parameters->TBs[pusch_id];

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

@@ -404,16 +404,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;
+#if defined(USE_CUDA)     
+    if (nrLDPC_slot_encoding_parameters->use_gpu && nrLDPC_TB_encoding_parameters->BG == 1 && 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;
+      // 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(USE_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)));
@@ -430,6 +438,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(USE_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];
@@ -456,6 +468,10 @@ 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++) {
     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(USE_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));
 
 

openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_CUDA_mPassKernel_BG1.h

+/*
+ * 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 nrLDPC_CUDA_mPassKernel.h
+ * \brief Defines the kernels for message passing in CUDA version of LDPC decoder
+ * \author Qizhi Pan, Raymond Knopp
+ * \company EURECOM
+ * \email: qizhi.pan@eurecom.fr, raymond.knopp@eurecom.fr
+ * \date 2025-12-30
+ * \version 1.0
+ * \note
+ * \warning
+ */
+#pragma once
+
+#include <cuda_runtime.h>
+#include <stdint.h>
+#include <stdio.h>
+#include "nrLDPC_types.h"
+#include "nrLDPC_CUDA_public.h"
+#include "nrLDPC_CUDA_shared_param.h"
+
+__device__ __forceinline__ void llrPreProc_Kernel_BG1_int8_Gn_stream(const int8_t *p_llr,
+                                                                     int8_t *p_llrProcBuf,
+                                                                     int8_t *p_cnProcBuf,
+                                                                     uint32_t MsgIdx,
+                                                                     uint32_t lane,
+                                                                     uint32_t colIdx,
+                                                                     uint32_t idxBn,
+                                                                     uint32_t GrpIdx,
+                                                                     uint32_t circShift,
+                                                                     uint32_t Zc,
+                                                                     uint32_t R)
+{
+  {
+    uint32_t *p_cnProcBufBit;
+
+    uint8_t bricksLocal[4];
+    uint8_t *BricksToBeMoved = bricksLocal;
+
+    p_cnProcBufBit = (uint32_t *)(p_cnProcBuf + d_lut_numCnInCnGroups_BG1_R13[GrpIdx] * NR_LDPC_ZMAX * MsgIdx + lane * 4);
+
+    moveBricks_invget_circ((int8_t *)&p_llr[idxBn * Zc], lane * 4, BricksToBeMoved, Zc, circShift);
+
+    *p_cnProcBufBit = *(uint32_t *)BricksToBeMoved;
+  }
+
+  if (colIdx >= 68)
+    return;
+
+  const uint8_t numBn2CnG1 =
+      (R == 13) ? d_lut_numBnInBnGroups_BG1_R13[0]
+                : ((R == 89) ? d_lut_numBnInBnGroups_BG1_R89[0] : d_lut_numBnInBnGroups_BG1_R23[0]); // for R13 is 42
+  const uint32_t startColParity = NR_LDPC_START_COL_PARITY_BG1; // 26 for BG1
+  const uint32_t colG1 = startColParity * Zc;
+
+  const uint32_t *lut_llr2llrProcBufAddr =
+      (R == 13) ? d_llr2llrProcBufAddr_BG1_R13 : ((R == 89) ? d_llr2llrProcBufAddr_BG1_R89 : d_llr2llrProcBufAddr_BG1_R23);
+  const uint32_t *lut_llr2llrProcBufBnPos =
+      (R == 13) ? d_llr2llrProcBufBnPos_BG1_R13 : ((R == 89) ? d_llr2llrProcBufBnPos_BG1_R89 : d_llr2llrProcBufBnPos_BG1_R23);
+
+  if (colIdx < startColParity) {
+    const uint32_t idxBn = lut_llr2llrProcBufAddr[colIdx] + lut_llr2llrProcBufBnPos[colIdx] * NR_LDPC_ZMAX;
+    int32_t *dst = (int32_t *)(&p_llrProcBuf[idxBn] + lane * 4);
+    int32_t *src = (int32_t *)(&p_llr[colIdx * Zc] + lane * 4);
+    *dst = *src;
+  } else {
+    colIdx = colIdx - startColParity;
+    if (numBn2CnG1 > 0 && colIdx < numBn2CnG1) {
+      int32_t *dst = (int32_t *)(&p_llrProcBuf[colIdx * NR_LDPC_ZMAX] + lane * 4);
+      int32_t *src = (int32_t *)(&p_llr[colG1 + colIdx * Zc] + lane * 4);
+      *dst = *src;
+    }
+  }
+}
+
+__device__ void llr2bit_Kernel_BG1_int8(uint32_t R,
+                                        uint8_t *__restrict__ out,
+                                        const int8_t *__restrict__ llrRes,
+                                        uint32_t numLLR,
+                                        uint32_t Zc)
+{
+  uint32_t lane = threadIdx.x;
+  uint32_t outColIdx = (blockIdx.x << 2) + threadIdx.y;
+
+  if (outColIdx >= num_TotalBlocks_llr_llrRes)
+    return;
+
+  // Constants Setup
+  const uint8_t numBn2CnG1 = (R == 13) ? d_lut_numBnInBnGroups_BG1_R13[0]
+                                       : ((R == 89) ? d_lut_numBnInBnGroups_BG1_R89[0] : d_lut_numBnInBnGroups_BG1_R23[0]);
+  const uint32_t startColParity = NR_LDPC_START_COL_PARITY_BG1;
+  const uint32_t *lut_Addr =
+      (R == 13) ? d_llr2llrProcBufAddr_BG1_R13 : ((R == 89) ? d_llr2llrProcBufAddr_BG1_R89 : d_llr2llrProcBufAddr_BG1_R23);
+  const uint32_t *lut_Pos =
+      (R == 13) ? d_llr2llrProcBufBnPos_BG1_R13 : ((R == 89) ? d_llr2llrProcBufBnPos_BG1_R89 : d_llr2llrProcBufBnPos_BG1_R23);
+
+  int32_t raw_llrs;
+
+  // Load Data: Handle Systematic (Scatter/Gather) vs Parity (Linear) mapping
+  if (outColIdx < startColParity) {
+    uint32_t idxBn = lut_Addr[outColIdx] + lut_Pos[outColIdx] * NR_LDPC_ZMAX;
+    raw_llrs = *(const int32_t *)(&llrRes[idxBn] + lane * 4);
+  } else {
+    uint32_t srcParityIdx = outColIdx - startColParity;
+    if (numBn2CnG1 > 0 && outColIdx < numBn2CnG1) {
+      raw_llrs = *(const int32_t *)(llrRes + srcParityIdx * NR_LDPC_ZMAX + lane * 4);
+    } else {
+      raw_llrs = 0;
+    }
+  }
+
+  // Hard Decision: Convert 4 int8 LLRs -> 4 bytes (0 or 1)
+  int8_t *p_val = (int8_t *)&raw_llrs;
+  uint32_t my_word = 0;
+
+#pragma unroll
+  for (int i = 0; i < 4; i++) {
+    // Hard decision: LLR < 0 implies bit 1
+    uint32_t byte_val = (p_val[i] < 0) ? 1 : 0;
+    my_word |= (byte_val << (i * 8));
+  }
+
+  // Store output linearly
+  *(uint32_t *)(&out[outColIdx * Zc + lane * 4]) = my_word;
+}
+
+__device__ void llr2bitPacked_Kernel_BG1_int8(uint32_t R,
+                                              uint8_t *__restrict__ out,
+                                              const int8_t *__restrict__ llrRes,
+                                              uint32_t numLLR,
+                                              uint32_t Zc)
+{
+  uint32_t lane = threadIdx.x;
+  uint32_t outColIdx = (blockIdx.x << 2) + threadIdx.y;
+
+  if (outColIdx >= num_TotalBlocks_llr_llrRes)
+    return;
+
+  const uint8_t numBn2CnG1 = (R == 13) ? d_lut_numBnInBnGroups_BG1_R13[0]
+                                       : ((R == 89) ? d_lut_numBnInBnGroups_BG1_R89[0] : d_lut_numBnInBnGroups_BG1_R23[0]);
+  const uint32_t startColParity = NR_LDPC_START_COL_PARITY_BG1;
+  const uint32_t *lut_Addr =
+      (R == 13) ? d_llr2llrProcBufAddr_BG1_R13 : ((R == 89) ? d_llr2llrProcBufAddr_BG1_R89 : d_llr2llrProcBufAddr_BG1_R23);
+  const uint32_t *lut_Pos =
+      (R == 13) ? d_llr2llrProcBufBnPos_BG1_R13 : ((R == 89) ? d_llr2llrProcBufBnPos_BG1_R89 : d_llr2llrProcBufBnPos_BG1_R23);
+
+  int32_t raw_llrs;
+
+  // Load Data: Handle Systematic vs Parity mapping
+  if (outColIdx < startColParity) {
+    uint32_t idxBn = lut_Addr[outColIdx] + lut_Pos[outColIdx] * NR_LDPC_ZMAX;
+    raw_llrs = *(const int32_t *)(&llrRes[idxBn] + lane * 4);
+  } else {
+    uint32_t srcParityIdx = outColIdx - startColParity;
+    if (numBn2CnG1 > 0 && outColIdx < numBn2CnG1) {
+      raw_llrs = *(const int32_t *)(llrRes + srcParityIdx * NR_LDPC_ZMAX + lane * 4);
+    } else {
+      raw_llrs = 0;
+    }
+  }
+
+  // Extract 4 bits from 4 LLRs (Local Packing)
+  // Assuming MSB order within nibble: [3][2][1][0]
+  uint32_t my_4_bits = 0;
+  int8_t *p_val = (int8_t *)&raw_llrs;
+
+#pragma unroll
+  for (int i = 0; i < 4; i++) {
+    if (p_val[i] < 0) {
+      my_4_bits |= (1 << (3 - i));
+    }
+  }
+
+  // Thread Cooperation: Pair threads to pack 8 bits (1 byte)
+  // Even thread (0,2..) takes low nibble, Odd thread (1,3..) takes high nibble.
+  uint32_t neighbor_bits = __shfl_xor_sync(0xffffffff, my_4_bits, 1);
+
+  if ((lane & 1) == 0) {
+    // Combine: [Odd Thread Bits (High)] | [My Bits (Low)]
+    uint8_t packed_byte = (neighbor_bits & 0xF) | ((my_4_bits & 0xF) << 4);
+
+    // Output addr: lane steps by 4 LLRs, but we output 1 byte per 8 LLRs -> lane >> 1
+    uint32_t outAddr = outColIdx * (Zc >> 3) + (lane >> 1);
+    out[outAddr] = packed_byte;
+  }
+}
\ No newline at end of file

openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_CUDA_public.h

+/*
+ * 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 nrLDPC_CUDA_public.h
+ * \brief Shared functions in CUDA implementation of LDPC decoder 
+ * \author Qizhi Pan, Raymond Knopp
+ * \company EURECOM
+ * \email: qizhi.pan@eurecom.fr, raymond.knopp@eurecom.fr
+ * \date 2025-12-30
+ * \version 1.0
+ * \note 
+ * \warning
+ */
+
+#pragma once
+
+#include <cuda_runtime.h>
+#include <stdint.h>
+#include <stdio.h>
+#include "nrLDPC_types.h"
+
+#include <cuda_runtime.h>
+#include <stdint.h>
+
+
+__device__ __forceinline__ void moveBricks_invput_circ(int8_t *__restrict__ dstBuf,
+                                                       uint32_t dstBuf_Offset,
+                                                       const uint8_t *__restrict__ Four_Bricks,
+                                                       uint32_t Z,
+                                                       uint32_t cshift)
+{
+    uint32_t tmp = cshift + dstBuf_Offset;
+    uint32_t pos = (tmp >= Z) ? tmp - Z : tmp;
+
+    if (pos <= Z - 4) {
+        uint32_t val = *(const uint32_t*)Four_Bricks;
+        memcpy(dstBuf + pos, &val, 4);
+    } 
+    else {
+        uint32_t bytes_at_end = Z - pos; // 1, 2, or 3
+        
+        #pragma unroll
+        for(int i=0; i<4; i++) {
+            if (i < bytes_at_end) {
+                dstBuf[pos + i] = Four_Bricks[i];
+            } else {
+                dstBuf[i - bytes_at_end] = Four_Bricks[i];
+            }
+        }
+    }
+}
+
+__device__ __forceinline__ void moveBricks_forput_circ(int8_t *__restrict__ dstBuf,
+                                                       uint32_t dstBuf_Offset,
+                                                       const uint8_t *__restrict__ Four_Bricks,
+                                                       uint32_t Z,
+                                                       uint32_t cshift)
+{
+    uint32_t tmp = dstBuf_Offset + Z - cshift;
+    uint32_t pos = (tmp >= Z) ? tmp - Z : tmp;
+
+    if (pos <= Z - 4) {
+        uint32_t val = *(const uint32_t*)Four_Bricks;
+        memcpy(dstBuf + pos, &val, 4);
+    } 
+    else {
+        uint32_t bytes_at_end = Z - pos;
+        
+        #pragma unroll
+        for(int i=0; i<4; i++) {
+            if (i < bytes_at_end) {
+                dstBuf[pos + i] = Four_Bricks[i];
+            } else {
+                dstBuf[i - bytes_at_end] = Four_Bricks[i];
+            }
+        }
+    }
+}
+
+__device__ __forceinline__ void moveBricks_invget_circ(const int8_t *__restrict__ dstBuf,
+                                                       uint32_t dstBuf_Offset,
+                                                       uint8_t *__restrict__ Four_Bricks,
+                                                       uint32_t Z,
+                                                       uint32_t cshift)
+{
+    uint32_t tmp = cshift + dstBuf_Offset;
+    uint32_t pos = (tmp >= Z) ? tmp - Z : tmp;
+
+    if (pos <= Z - 4) {
+        uint32_t val;
+        memcpy(&val, dstBuf + pos, 4);
+        *(uint32_t*)Four_Bricks = val;
+    } 
+    else {
+        uint32_t bytes_at_end = Z - pos;
+        
+        #pragma unroll
+        for(int i=0; i<4; i++) {
+            if (i < bytes_at_end) {
+                Four_Bricks[i] = dstBuf[pos + i];
+            } else {
+                Four_Bricks[i] = dstBuf[i - bytes_at_end];
+            }
+        }
+    }
+}
+
+__device__ __forceinline__ uint32_t __vxor4(const uint32_t a, uint32_t b)
+{
+  return a ^ b; 
+}
+
+__device__ __forceinline__ uint32_t __vsign4(const uint32_t a, uint32_t b)
+{
+    uint32_t mask = __vcmplts4(b, 0); 
+    uint32_t bneg = __vneg4(a); 
+    return (mask & bneg) | (~mask & a);
+    //uint32_t is_zero_mask = __vcmpeq4(b, 0);
+    //return result & (~is_zero_mask);
+}
+

openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_CUDA_shared_param.h

+/*
+ * 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 nrLDPC_CUDA_shared_param.h
+ * \brief Shared parameters in CUDA implementation of LDPC decoder 
+ * \author Qizhi Pan, Raymond Knopp
+ * \company EURECOM
+ * \email: qizhi.pan@eurecom.fr, raymond.knopp@eurecom.fr
+ * \date 2025-12-30
+ * \version 1.0
+ * \note 
+ * \warning
+ */
+
+#ifndef NRLDPC_CUDA_SHARED_PARAM_H_
+#define NRLDPC_CUDA_SHARED_PARAM_H_
+
+#include <cuda_runtime.h>
+#define MAX_NUM_DLSCH_SEGMENTS_DL 132
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+#define RowLength 96 //Zc = 384/4 = 96
+
+#define num_TotalBlocks_BG1_R13_Edge 316//based on number of Cn2Bn Msgs
+#define num_TotalBlocks_BG1_R23_Edge 144
+#define num_TotalBlocks_BG1_R89_Edge 79
+#define num_TotalBlocks_llr_llrRes 22 //Only includes systematic bits
+
+#define num_TotalBlocks_cn_BG1_R13_Node 46 //based on number of CNs
+#define num_TotalBlocks_bn_BG1_R13_Node 68 //based on number of BNs
+#define num_TotalBlocks_cn_BG1_R23_Node 13
+#define num_TotalBlocks_bn_BG1_R23_Node 35
+#define num_TotalBlocks_cn_BG1_R89_Node 5
+#define num_TotalBlocks_bn_BG1_R89_Node 27
+
+#define JETSON_ORIN 1 
+
+#if defined(GH200)
+#define NodeEdge_Switch_Cn_R13 32
+#define NodeEdge_Switch_Bn_R13 10
+#define NodeEdge_Switch_Cn_R23 32
+#define NodeEdge_Switch_Bn_R23 12
+#define NodeEdge_Switch_Cn_R89 48
+#define NodeEdge_Switch_Bn_R89 24
+#elif defined(JETSON_ORIN)
+#define NodeEdge_Switch_Cn_R13 3
+#define NodeEdge_Switch_Bn_R13 1
+#define NodeEdge_Switch_Cn_R23 3
+#define NodeEdge_Switch_Bn_R23 1
+#define NodeEdge_Switch_Cn_R89 4
+#define NodeEdge_Switch_Bn_R89 2
+#endif
+
+#ifdef __cplusplus
+}
+#endif
+
+typedef struct KernelLaunchConfig {
+    dim3 grid;
+    dim3 block;
+}KernelLaunchConfig;
+
+
+typedef struct {
+    int idxBn;
+    int idxCn;
+    int preBuf;
+    int circShift;
+    int8_t dd;
+} DumpEntry;
+
+typedef struct {
+    int8_t* p_llr_ptr;     
+    int8_t* p_out_ptr;      
+} ldpc_cuda_bridge_t;
+
+#endif /* NRLDPC_CUDA_SHARED_PARAM_H_ */

openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder.c

@@ -169,6 +169,16 @@ int32_t LDPCshutdown()
   return 0;
 }
 
+#ifndef USE_CUDA
+int32_t LDPCinit_cuda() {
+
+  return 0;
+}
+int32_t LDPCshutdown_cuda() {
+
+  return 0;
+}
+#endif
 int32_t LDPCdecoder(t_nrLDPC_dec_params* p_decParams,
                     int8_t* p_llr,
                     uint8_t* p_out,
@@ -191,6 +201,17 @@ int32_t LDPCdecoder(t_nrLDPC_dec_params* p_decParams,
     return numIter;
 }
 
+#ifndef USE_CUDA
+int32_t LDPCdecoder_cuda(t_nrLDPC_dec_params* p_decParams,
+                         int8_t* p_llr,
+                         uint8_t* p_out,
+                         t_nrLDPC_time_stats* p_profiler,
+                         decode_abort_t* ab)
+{
+
+     return(0);
+}
+#endif
 /**
    \brief PerformsnrLDPC decoding of one code block
    \param p_llr Input LLRs

openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_types.h

@@ -90,6 +90,8 @@ typedef struct nrLDPC_dec_params {
     e_nrLDPC_outMode outMode; /**< Output format */
     int crc_type; /**< Size and type of the parity check bits (16, 24A or 24B) */
     int (*check_crc)(uint8_t* decoded_bytes, uint32_t n, uint8_t crc_type); /**< Parity check function */
+    int n_segments;/**Infomation for cuda streaming*/
+    int LastTrial;/**Infomation for cuda streaming to create and destroy*/
 } t_nrLDPC_dec_params;
 
 /**

openair1/PHY/CODING/nrLDPC_defs.h

@@ -76,6 +76,8 @@ typedef struct {
 } encoder_implemparams_t;
 
 typedef int32_t(LDPC_initfunc_t)(void);
+typedef int32_t(LDPC_initfunc_cuda_t)(int8_t*, int8_t*);
+
 typedef int32_t(LDPC_shutdownfunc_t)(void);
 
 // decoder interface
@@ -93,5 +95,11 @@ 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 *);
+typedef int32_t(LDPC_decoderfunc_cuda_t)(t_nrLDPC_dec_params *p_decParams,
+                                         int8_t *p_llr,
+                                         uint8_t *p_out,
+                                         t_nrLDPC_time_stats *time_stats,
+                                         decode_abort_t *ab);
+void cuda_support_init(void);
 #endif

openair1/PHY/CODING/nrLDPC_encoder/ldpc_encode_parity_check_cuda.c

+/*
+ * 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, cudaStream_t *stream,int sidx);
+
+
+void encode_parity_check_part_cuda(uint32_t **c, uint32_t **d, short BG,short Zc,short Kb, int ncols, int n_inputs, cudaStream_t *stream,int sidx)
+{
+  
+  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, stream,sidx);
+        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_encoder32.c

+/*
+ * 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_input.cu

+#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;
+  } 
+}
+*/
+
+__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 nseg) {
+
+//  int block_off = blockIdx.y*blockDim.x;	
+  int i1 = blockIdx.y;	
+  int i2 = threadIdx.x;
+  int nseg0 = (blockIdx.x << 5);
+  int nseg1;
+  if ((nseg0 + 32) <= nseg) nseg1 = nseg0+32;
+  else nseg1 = nseg0 + (nseg&31);
+  int bit_offset = i2+(i1*384);
+  int uint32_offset = bit_offset>>5;  
+  uint32_t mask0 = masks[bit_offset&31];
+  uint32_t tmp,jmod;
+  uint32_t otmp0;
+  if (bit_offset  < 8448) {
+      tmp=input[nseg0][uint32_offset];
+      otmp0 = ((tmp&mask0) > 0); 
+      for (int j=nseg0+1;j<nseg1;j++) {
+	 tmp=input[j][uint32_offset];
+	 jmod = j&31;
+         otmp0 |= (((tmp&mask0) > 0)<<jmod); 
+      }	      
+      
+      cc[blockIdx.x][(2*i1*384) + i2] = otmp0;
+      cc[blockIdx.x][(2*i1+1)*384 + i2]=otmp0;
+  } 
+}
+extern "C" int ldpc_input(uint32_t **input,uint32_t *cc[4],int nseg,cudaStream_t *stream,int sidx) { 
+
+ int ns = nseg>>5;
+ if ((nseg&31)>0) ns++;
+
+ dim3 numblocks(ns,22);
+ //printf("input %p\n",input);
+ ldpc_input_worker<<<numblocks,384,0,stream[sidx]>>>(input,cc,nseg);
+ cudaError_t err=cudaPeekAtLastError();
+ if (err!=cudaSuccess) {
+    printf("cuda error: %s (input %p, cc %p, nseg %d, ns %d)\n",cudaGetErrorString(err),input,cc,nseg,ns);
+    exit(-1);
+ }
+ return(0);
+}

openair1/PHY/CODING/nrLDPC_extern.h

@@ -24,10 +24,13 @@
 /* ldpc coder/decoder API*/
 typedef struct ldpc_interface_s {
   LDPC_initfunc_t *LDPCinit;
+  LDPC_initfunc_t *LDPCinit_cuda;
   LDPC_shutdownfunc_t *LDPCshutdown;
+  LDPC_shutdownfunc_t *LDPCshutdown_cuda;
   LDPC_decoderfunc_t *LDPCdecoder;
   LDPC_encoderfunc_t *LDPCencoder;
   LDPC_encoderfunc32_t *LDPCencoder32;
+  LDPC_decoderfunc_t *LDPCdecoder_cuda;
 } ldpc_interface_t;
 
 /* functions to load the LDPC shared lib, implemented in openair1/PHY/CODING/nrLDPC_load.c */
@@ -37,6 +40,7 @@ int free_LDPClib(ldpc_interface_t *ldpc_interface);
 LDPC_decoderfunc_t LDPCdecoder;
 LDPC_encoderfunc_t LDPCencoder;
 LDPC_encoderfunc32_t LDPCencoder32;
+LDPC_decoderfunc_t LDPCdecoder_cuda;
 
 // inline functions:
 #endif

openair1/PHY/CODING/nrLDPC_load.c

@@ -54,19 +54,25 @@ int load_LDPClib(char *version, ldpc_interface_t *itf)
   }
   /* function description array, to be used when loading the encoding/decoding shared lib */
   loader_shlibfunc_t shlib_fdesc[] = {{.fname = "LDPCinit"},
+                                      {.fname = "LDPCinit_cuda"},
                                       {.fname = "LDPCshutdown"},
+                                      {.fname = "LDPCshutdown_cuda"},
                                       {.fname = "LDPCdecoder"},
                                       {.fname = "LDPCencoder"},
-                                      {.fname = "LDPCencoder32"}};
+                                      {.fname = "LDPCencoder32"},
+                                      {.fname = "LDPCdecoder_cuda"}};
   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;
+  itf->LDPCinit_cuda = (LDPC_initfunc_t *)shlib_fdesc[1].fptr;
+  itf->LDPCshutdown = (LDPC_shutdownfunc_t *)shlib_fdesc[2].fptr;
+  itf->LDPCshutdown_cuda = (LDPC_shutdownfunc_t *)shlib_fdesc[3].fptr;
+  itf->LDPCdecoder = (LDPC_decoderfunc_t *)shlib_fdesc[4].fptr;
+  itf->LDPCencoder = (LDPC_encoderfunc_t *)shlib_fdesc[5].fptr;
+  itf->LDPCencoder32 = (LDPC_encoderfunc32_t *)shlib_fdesc[6].fptr;
+  itf->LDPCdecoder_cuda = (LDPC_decoderfunc_cuda_t *)shlib_fdesc[7].fptr;
 
   AssertFatal(itf->LDPCinit() == 0, "error starting LDPC library %s %s\n", libname, version);
 
@@ -75,5 +81,6 @@ int load_LDPClib(char *version, ldpc_interface_t *itf)
 
 int free_LDPClib(ldpc_interface_t *ldpc_interface)
 {
+  if (ldpc_interface->LDPCshutdown_cuda) ldpc_interface->LDPCshutdown_cuda();
   return ldpc_interface->LDPCshutdown();
 }

openair1/PHY/NR_TRANSPORT/nr_transport_proto.h

@@ -101,9 +101,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/NR_TRANSPORT/nr_ulsch_decoding.c

@@ -113,7 +113,8 @@ int nr_ulsch_decoding(PHY_VARS_gNB *phy_vars_gNB,
                                                        .slot = nr_tti_rx,
                                                        .nb_TBs = nb_pusch,
                                                        .threadPool = &phy_vars_gNB->threadPool,
-                                                       .TBs = TBs};
+                                                       .TBs = TBs,
+  						       .use_gpu = phy_vars_gNB->use_gpu};
 
   int max_num_segments = 0;
 
@@ -226,7 +227,11 @@ int nr_ulsch_decoding(PHY_VARS_gNB *phy_vars_gNB,
     uint8_t ULSCH_id = ULSCH_ids[pusch_id];
     NR_gNB_ULSCH_t *ulsch = &phy_vars_gNB->ulsch[ULSCH_id];
     NR_UL_gNB_HARQ_t *harq_process = ulsch->harq_process;
-    short *ulsch_llr = phy_vars_gNB->pusch_vars[ULSCH_id].llr;
+#ifdef ENABLE_CUDA
+    int16_t *ulsch_llr = phy_vars_gNB->pusch_vars[ULSCH_id].llr_dev;
+#else
+    int16_t *ulsch_llr = phy_vars_gNB->pusch_vars[ULSCH_id].llr;
+#endif
 
     if (!ulsch_llr) {
       LOG_E(PHY, "ulsch_decoding.c: NULL ulsch_llr pointer\n");

openair1/SIMULATION/NR_PHY/dlschsim.c

@@ -60,6 +60,7 @@ int64_t uplink_frequency_offset[MAX_NUM_CCs][4];
 uint64_t downlink_frequency[MAX_NUM_CCs][4];
 
 double cpuf;
+uint32_t use_gpu=0;
 
 uint8_t const nr_rv_round_map[4] = {0, 2, 3, 1};
 // needed for some functions