turbo decoder gpu

cmake_targets/CMakeLists.txt

@@ -1,28 +1,53 @@
-#/*
-# * 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
-# */
-
+################################################################################
+#   OpenAirInterface
+#   Copyright(c) 1999 - 2014 Eurecom
+#
+#   OpenAirInterface is free software: you can redistribute it and/or modify
+#   it under the terms of the GNU General Public License as published by
+#   the Free Software Foundation, either version 3 of the License, or
+#   (at your option) any later version.
+#
+#   OpenAirInterface is distributed in the hope that it will be useful,
+#   but WITHOUT ANY WARRANTY; without even the implied warranty of
+#   MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+#   GNU General Public License for more details.
+#
+#   You should have received a copy of the GNU General Public License
+#   along with OpenAirInterface.The full GNU General Public License is
+#  included in this distribution in the file called "COPYING". If not,
+#  see <http://www.gnu.org/licenses/>.
+#
+# Contact Information
+# OpenAirInterface Admin: openair_admin@eurecom.fr
+# OpenAirInterface Tech : openair_tech@eurecom.fr
+# OpenAirInterface Dev  : openair4g-devel@lists.eurecom.fr
+#
+# Address      : Eurecom, Campus SophiaTech, 450 Route des Chappes, CS 50193 - 06904 Biot Sophia Antipolis cedex, FRANCE
 # Author: laurent THOMAS, Lionel GAUTHIER
+###############################################################################
 
 cmake_minimum_required (VERSION 2.8)
 
+##############################################
+# Base CUDA setting
+##############################################
+find_package(CUDA)
+message ("cuda include ${CUDA_INCLUDE_DIRS}")
+message ("cud library ${CUDA_LIBRARY_DIRS}")
+
+add_definitions("-L/usr/local/cuda/lib64")
+
+SET(CUDA_NVCC_FLAGS
+  "${CUDA_NVCC_FLAGS};-arch=sm_52;")
+
+# Disable warnings for CUDA
+SET(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS};-w;-O3")
+SET(CUDA_VERBOSE_BUILD ON)
+SET(CUDA_HOST_COMPILER "/usr/bin/g++")
+
+SET(CUDA_SEPARABLE_COMPILATION ON)
+
+
 #############################################
 # Base directories, compatible with legacy OAI building
 ################################################
@@ -174,13 +199,14 @@ set(CMAKE_EXE_LINKER_FLAGS  "${CMAKE_EXE_LINKER_FLAGS} -Wl,-rpath -Wl,${CMAKE_CU
 # these changes are related to hardcoded path to include .h files
 add_definitions(-DCMAKER)
 set(CMAKE_C_FLAGS_DEBUG "${CMAKE_C_FLAGS} -g -DMALLOC_CHECK_=3")
-set(CMAKE_C_FLAGS_RELWITHDEBINFO "${CMAKE_C_FLAGS} -g -DMALLOC_CHECK_=3 -O2") 
+set(CMAKE_C_FLAGS_RELWITHDEBINFO "${CMAKE_C_FLAGS} -g -DMALLOC_CHECK_=3 -O3") 
 
 
 set(GIT_BRANCH        "UNKNOWN")
 set(GIT_COMMIT_HASH   "UNKNOWN")
 set(GIT_COMMIT_DATE   "UNKNOWN")
 
+
 find_package(Git)
 if(GIT_FOUND)
   message("git found: ${GIT_EXECUTABLE}")
@@ -256,7 +282,15 @@ if (${ENABLE_ITTI})
 endif (${ENABLE_ITTI})
 
 add_boolean_option(RTAI False "Use RTAI")
-
+if (${RTAI})
+  set(DEADLINE_SCHEDULER False)
+  set(CPU_AFFINITY False)
+  add_definitions("-DENABLE_RTAI_CLOCK")
+  add_definitions("-DCONFIG_RTAI_LXRT_INLINE")
+  include_directories ("/usr/realtime/include")
+  include_directories ("/usr/realtime/include/asm")
+  set(RTAI_SOURCE sched_dlsch.c sched_rx_pdsch.c rt_wrapper.c vcd_signal_dumper.c log.c)
+endif (${RTAI})
 
 #############################
 # ASN.1 grammar C code generation & dependancies
@@ -491,7 +525,7 @@ include_directories ("${OPENAIR_TARGETS}/ARCH/EXMIMO/DEFS/")
 #set (option_HWEXMIMOLIB_lib "-l ")
 set(HWLIB_EXMIMO_SOURCE 
   ${OPENAIR_TARGETS}/ARCH/EXMIMO/USERSPACE/LIB/openair0_lib.c
-#  ${OPENAIR_TARGETS}/ARCH/EXMIMO/USERSPACE/LIB/gain_control.c
+  ${OPENAIR_TARGETS}/ARCH/EXMIMO/USERSPACE/LIB/gain_control.c
   )
 add_library(oai_exmimodevif MODULE ${HWLIB_EXMIMO_SOURCE} )
 
@@ -532,8 +566,8 @@ if (${RF_BOARD} STREQUAL "EXMIMO")
   include_directories ("${OPENAIR_TARGETS}/ARCH/EXMIMO/USERSPACE/LIB/")
   include_directories ("${OPENAIR_TARGETS}/ARCH/EXMIMO/DEFS/")
   set(HW_SOURCE ${HW_SOURCE}
-    ${OPENAIR_TARGETS}/ARCH/EXMIMO/USERSPACE/LIB/openair0_lib.c)
-#    ${OPENAIR_TARGETS}/ARCH/EXMIMO/USERSPACE/LIB/gain_control.c)
+    ${OPENAIR_TARGETS}/ARCH/EXMIMO/USERSPACE/LIB/openair0_lib.c
+    ${OPENAIR_TARGETS}/ARCH/EXMIMO/USERSPACE/LIB/gain_control.c)
   set(option_HW_lib "-rdynamic -ldl")
 
 elseif (${RF_BOARD} STREQUAL "OAI_USRP")
@@ -614,7 +648,6 @@ add_boolean_option(NAS_ADDRESS_FIX False "specific to oaisim: for nasmesh driver
 add_boolean_option(NAS_NETLINK False "???? Must be True to compile nasmesh driver without rtai")
 add_boolean_option(OAISIM False "specific to oaisim")
 add_boolean_option(OAI_NW_DRIVER_USE_NETLINK True "????")
-
 add_boolean_option(USE_MME False "this flag is used only one time in lte-softmodem.c")
 add_list_string_option(PACKAGE_NAME "NotDefined" "As per attribute name")
 add_boolean_option(MESSAGE_CHART_GENERATOR False         "For generating sequence diagrams")
@@ -838,10 +871,10 @@ include_directories("${OPENAIR_DIR}")
 # Utilities Library
 ################
 add_library(HASHTABLE
-  ${OPENAIR_DIR}/common/utils/hashtable/hashtable.c
-  ${OPENAIR_DIR}/common/utils/hashtable/obj_hashtable.c
+  ${OPENAIR_DIR}/common/utils/collection/hashtable/hashtable.c
+  ${OPENAIR_DIR}/common/utils/collection/hashtable/obj_hashtable.c
 )
-include_directories(${OPENAIR_DIR}/common/utils/hashtable)
+include_directories(${OPENAIR_DIR}/common/utils/collection/hashtable)
 
 if (MESSAGE_CHART_GENERATOR)
   add_library(MSC  
@@ -869,16 +902,20 @@ set(UTIL_SRC
   ${OPENAIR2_DIR}/UTIL/OCG/OCG_parse_filename.c
   ${OPENAIR2_DIR}/UTIL/OCG/OCG_parse_XML.c
   ${OPENAIR2_DIR}/UTIL/OCG/OCG_save_XML.c
+  ${OPENAIR2_DIR}/UTIL/OMG/client_traci_OMG.c
   ${OPENAIR2_DIR}/UTIL/OMG/common.c
   ${OPENAIR2_DIR}/UTIL/OMG/grid.c
+  ${OPENAIR2_DIR}/UTIL/OMG/id_manager.c
   ${OPENAIR2_DIR}/UTIL/OMG/job.c
   ${OPENAIR2_DIR}/UTIL/OMG/mobility_parser.c
   ${OPENAIR2_DIR}/UTIL/OMG/omg.c
-  #${OPENAIR2_DIR}/UTIL/OMG/omg_hashtable.c
+  ${OPENAIR2_DIR}/UTIL/OMG/omg_hashtable.c
   ${OPENAIR2_DIR}/UTIL/OMG/rwalk.c
   ${OPENAIR2_DIR}/UTIL/OMG/rwp.c
+  ${OPENAIR2_DIR}/UTIL/OMG/socket_traci_OMG.c
   ${OPENAIR2_DIR}/UTIL/OMG/static.c
   ${OPENAIR2_DIR}/UTIL/OMG/steadystaterwp.c
+  ${OPENAIR2_DIR}/UTIL/OMG/storage_traci_OMG.c
   ${OPENAIR2_DIR}/UTIL/OMG/trace.c
   ${OPENAIR2_DIR}/UTIL/OMG/trace_hashtable.c
   ${OPENAIR2_DIR}/UTIL/OPT/probe.c
@@ -891,15 +928,6 @@ set(UTIL_SRC
   )
 add_library(UTIL ${UTIL_SRC})
 
-#set(OMG_SUMO_SRC
-#  ${OPENAIR2_DIR}/UTIL/OMG/client_traci_OMG.c
-#  ${OPENAIR2_DIR}/UTIL/OMG/id_manager.c
-#  ${OPENAIR2_DIR}/UTIL/OMG/sumo.c
-#  ${OPENAIR2_DIR}/UTIL/OMG/socket_traci_OMG.c
-#  ${OPENAIR2_DIR}/UTIL/OMG/storage_traci_OMG.c
-#  )
-#add_library(OMG_SUMO ${OMG_SUMO_SRC})
-
 set(SECU_OSA_SRC
   ${OPENAIR2_DIR}/UTIL/OSA/osa_key_deriver.c
   ${OPENAIR2_DIR}/UTIL/OSA/osa_rijndael.c
@@ -960,7 +988,6 @@ set(PHY_SRC
   ${OPENAIR1_DIR}/PHY/LTE_TRANSPORT/pucch.c
   ${OPENAIR1_DIR}/PHY/LTE_TRANSPORT/prach.c
   ${OPENAIR1_DIR}/PHY/LTE_TRANSPORT/pmch.c
-  ${OPENAIR1_DIR}/PHY/LTE_TRANSPORT/pch.c
   ${OPENAIR1_DIR}/PHY/LTE_TRANSPORT/group_hopping.c
   ${OPENAIR1_DIR}/PHY/LTE_TRANSPORT/srs_modulation.c
   ${OPENAIR1_DIR}/PHY/LTE_TRANSPORT/drs_modulation.c
@@ -971,8 +998,6 @@ set(PHY_SRC
   ${OPENAIR1_DIR}/PHY/LTE_TRANSPORT/rar_tools.c
   ${OPENAIR1_DIR}/PHY/LTE_TRANSPORT/print_stats.c
   ${OPENAIR1_DIR}/PHY/LTE_TRANSPORT/initial_sync.c
-  ${OPENAIR1_DIR}/PHY/LTE_TRANSPORT/if4_tools.c
-  ${OPENAIR1_DIR}/PHY/LTE_TRANSPORT/if5_tools.c
   ${OPENAIR1_DIR}/PHY/MODULATION/ofdm_mod.c
   ${OPENAIR1_DIR}/PHY/MODULATION/slot_fep.c
   ${OPENAIR1_DIR}/PHY/MODULATION/slot_fep_mbsfn.c
@@ -1531,25 +1556,11 @@ endif()
 
 pkg_search_module(NETTLE nettle)
 if(NOT ${NETTLE_FOUND})
-  message( FATAL_ERROR "PACKAGE nettle not found: some targets will fail. Run build_oai -I again!")
+  message("PACKAGE nettle not found: some targets will fail")
 else()
   include_directories(${NETTLE_INCLUDE_DIRS})
 endif()
 
-message ("NETTLE VERSION_INSTALLED  = ${NETTLE_VERSION}")
-
-string(REGEX REPLACE "([0-9]+).*" "\\1" NETTLE_VERSION_MAJOR ${NETTLE_VERSION})
-string(REGEX REPLACE "[0-9]+\\.([0-9]+).*" "\\1" NETTLE_VERSION_MINOR ${NETTLE_VERSION})
-message ("NETTLE_VERSION_MAJOR = ${NETTLE_VERSION_MAJOR}")
-message ("NETTLE_VERSION_MINOR = ${NETTLE_VERSION_MINOR}")
-
-if ("${NETTLE_VERSION_MAJOR}" STREQUAL "" OR "${NETTLE_VERSION_MINOR}" STREQUAL "")
-  message( FATAL_ERROR "The nettle version not detected properly. Try to run build_oai -I again" )
-endif()
-
-add_definitions("-DNETTLE_VERSION_MAJOR=${NETTLE_VERSION_MAJOR}")
-add_definitions("-DNETTLE_VERSION_MINOR=${NETTLE_VERSION_MINOR}")
-
 pkg_search_module(XPM xpm)
 if(NOT ${XPM_FOUND})
   message("PACKAGE xpm not found: some targets will fail")
@@ -1619,13 +1630,23 @@ add_definitions(-DASN1_MINIMUM_VERSION=924)
 # lte-softmodem is both eNB and UE implementation
 ###################################################
 
-add_executable(lte-softmodem
+###################################################
+# For CUDA library 
+###################################################
+CUDA_ADD_LIBRARY(PUSCH_CU  
+  ${OPENAIR1_DIR}/PHY/CUDA/LTE_TRANSPORT/turbo_rx_gpu.cu
+  #${OPENAIR1_DIR}/PHY/CUDA/INIT/init_cuda.cu
+  #${OPENAIR1_DIR}/PHY/CUDA/LTE_TRANSPORT/rx_pusch.cu
+)
+
+cuda_add_executable(lte-softmodem
   ${rrc_h}
   ${s1ap_h}
   ${OPENAIR_BIN_DIR}/messages_xml.h
+  ${OPENAIR_TARGETS}/RT/USER/sched_dlsch.c
+  ${OPENAIR_TARGETS}/RT/USER/sched_rx_pdsch.c
   ${OPENAIR_TARGETS}/RT/USER/rt_wrapper.c
   ${OPENAIR_TARGETS}/RT/USER/lte-ue.c
-  ${OPENAIR_TARGETS}/RT/USER/lte-enb.c
   ${OPENAIR_TARGETS}/RT/USER/lte-softmodem.c
   ${OPENAIR1_DIR}/SIMULATION/TOOLS/taus.c
   ${OPENAIR_TARGETS}/SIMU/USER/init_lte.c
@@ -1642,11 +1663,28 @@ add_executable(lte-softmodem
   ${T_SOURCE}
   )
 
+  
+  
+###################################################
+# For link cufft 
+###################################################  
+  
+CUDA_ADD_CUFFT_TO_TARGET(lte-softmodem)
+CUDA_ADD_CUFFT_TO_TARGET(PUSCH_CU)
+
+target_link_libraries (lte-softmodem
+#cudadevrt cudart
+cuda boost_system)  
+
 target_link_libraries (lte-softmodem -ldl
   -Wl,--start-group
-  RRC_LIB S1AP_LIB S1AP_ENB GTPV1U SECU_CN SECU_OSA UTIL HASHTABLE SCTP_CLIENT UDP SCHED_LIB PHY LFDS L2 ${MSC_LIB} ${RAL_LIB} ${NAS_UE_LIB} ${ITTI_LIB} ${MIH_LIB} 
+  RRC_LIB S1AP_LIB S1AP_ENB GTPV1U SECU_CN SECU_OSA UTIL HASHTABLE SCTP_CLIENT UDP SCHED_LIB PHY PUSCH_CU LFDS L2 ${MSC_LIB} ${RAL_LIB} ${NAS_UE_LIB} ${ITTI_LIB} ${MIH_LIB} 
   -Wl,--end-group )
 
+target_link_libraries (lte-softmodem ${LIBXML2_LIBRARIES} cudart -L/usr/local/cuda/lib64)
+target_link_libraries (lte-softmodem ${LIBXML2_LIBRARIES} cudadevrt -L/usr/local/cuda/lib64)  
+  
+  
 target_link_libraries (lte-softmodem ${LIBXML2_LIBRARIES})
 target_link_libraries (lte-softmodem pthread m ${CONFIG_LIBRARIES} rt crypt ${CRYPTO_LIBRARIES} ${OPENSSL_LIBRARIES} ${NETTLE_LIBRARIES} sctp ${option_HW_lib} ${option_TP_lib} ${XFORMS_LIBRARIES} ) 
 target_link_libraries (lte-softmodem ${LIB_LMS_LIBRARIES})
@@ -1658,9 +1696,10 @@ add_executable(lte-softmodem-nos1
   ${rrc_h}
   ${s1ap_h}
   ${OPENAIR_BIN_DIR}/messages_xml.h
+  ${OPENAIR_TARGETS}/RT/USER/sched_dlsch.c
+  ${OPENAIR_TARGETS}/RT/USER/sched_rx_pdsch.c
   ${OPENAIR_TARGETS}/RT/USER/rt_wrapper.c
   ${OPENAIR_TARGETS}/RT/USER/lte-ue.c
-  ${OPENAIR_TARGETS}/RT/USER/lte-enb.c
   ${OPENAIR_TARGETS}/RT/USER/lte-softmodem.c
   ${OPENAIR1_DIR}/SIMULATION/TOOLS/taus.c
   ${OPENAIR_TARGETS}/SIMU/USER/init_lte.c
@@ -1756,9 +1795,6 @@ add_executable(oaisim
   ${s1ap_h}
   ${x2ap_h}
   ${OPENAIR_BIN_DIR}/messages_xml.h
-  ${OPENAIR_TARGETS}/RT/USER/lte-ue.c
-  ${OPENAIR_TARGETS}/RT/USER/lte-enb.c
-  ${OPENAIR_TARGETS}/RT/USER/rt_wrapper.c
   ${OPENAIR_TARGETS}/SIMU/USER/channel_sim.c
   ${OPENAIR_TARGETS}/SIMU/USER/init_lte.c
   ${OPENAIR_TARGETS}/SIMU/USER/oaisim_config.c
@@ -1769,6 +1805,7 @@ add_executable(oaisim
   ${OPENAIR_TARGETS}/SIMU/USER/oaisim.c
   ${OPENAIR2_DIR}/RRC/NAS/nas_config.c
   ${OPENAIR2_DIR}/RRC/NAS/rb_config.c
+  ${OPENAIR2_DIR}/UTIL/OMG/sumo.c
   ${OPENAIR3_DIR}/NAS/UE/nas_ue_task.c
   ${GTPU_need_ITTI}
   ${OPENAIR_TARGETS}/COMMON/create_tasks.c
@@ -1801,9 +1838,6 @@ add_executable(oaisim_nos1
   ${s1ap_h}
   ${x2ap_h}
   ${OPENAIR_BIN_DIR}/messages_xml.h
-  ${OPENAIR_TARGETS}/RT/USER/lte-ue.c
-  ${OPENAIR_TARGETS}/RT/USER/lte-enb.c
-  ${OPENAIR_TARGETS}/RT/USER/rt_wrapper.c
   ${OPENAIR_TARGETS}/SIMU/USER/channel_sim.c
   ${OPENAIR_TARGETS}/SIMU/USER/init_lte.c
   ${OPENAIR_TARGETS}/SIMU/USER/oaisim_config.c
@@ -1814,6 +1848,7 @@ add_executable(oaisim_nos1
   ${OPENAIR_TARGETS}/SIMU/USER/oaisim.c
   ${OPENAIR2_DIR}/RRC/NAS/nas_config.c
   ${OPENAIR2_DIR}/RRC/NAS/rb_config.c
+  ${OPENAIR2_DIR}/UTIL/OMG/sumo.c
   ${OPENAIR_TARGETS}/COMMON/create_tasks.c
   ${HW_SOURCE}
   ${TRANSPORT_SOURCE}  
@@ -1898,6 +1933,7 @@ foreach(myExe s1ap
     aes128_cmac_encrypt
     secu_knas_encrypt_eia2)
   add_executable(test_${myExe}
+    ${OPENAIR3_DIR}/TEST/test_util.c
     ${OPENAIR3_DIR}/TEST/test_${myExe}.c
     )
   target_link_libraries (test_${myExe}

openair1/PHY/CUDA/LTE_TRANSPORT/turbo_parm.h

@@ -0,0 +1,64 @@
+/*! \file PHY\CUDA/LTE_TRANSPORT/turbo_parm.h
+* \brief turbo decoder using gpu 
+* \author TerngYin Hsu, JianYa Chu
+* \date 2018
+* \version 0.1
+* \company ISIP LAB/NCTU CS
+* \email: tyhsu@cs.nctu.edu.tw
+* \note
+* \warning  
+*/
+
+
+#include<stdio.h>
+#include<stdlib.h>
+#include<cuda_runtime.h>
+typedef float llr_t;
+typedef struct
+{
+	llr_t *sys_d;
+	llr_t *sys1_d;
+	llr_t *sys2_d;
+	llr_t *ypar1_d;
+    llr_t *ypar2_d;
+	llr_t *ext_d;
+	llr_t *ext2_d;
+	llr_t *alpha_d;
+	int *decode_tmp;
+	llr_t *decode_ext2;
+	llr_t *alpha_pre_1;
+	llr_t *alpha_pre_2;
+	llr_t *beta_pre_1;
+	llr_t *beta_pre_2;
+	unsigned char *decode_h;
+	unsigned char *decode_d;
+}turbo_parm_s;
+
+typedef struct 
+{
+	// for excution time
+	cudaEvent_t e_start; 
+	cudaEvent_t e_stop;
+	float e_time;
+	// for fetch time
+	cudaEvent_t f_start, f_stop;
+	float f_time;
+	// for memcpy
+	cudaEvent_t m_start, m_stop;
+	float m_time;
+	// for turbo kernel excution time
+	cudaEvent_t t_start, t_stop;
+	float t_time;
+	// for decode excution time
+	cudaEvent_t d_start, d_stop;
+	float d_time;
+	// for crc check
+	cudaEvent_t s_check[3];
+	// for algorithm & decode stream
+	cudaStream_t stream[2];
+
+
+}cuda_parm_s;
+
+turbo_parm_s* turbo_parm;
+cuda_parm_s cuda_parm; 

openair1/PHY/CUDA/LTE_TRANSPORT/turbo_rx_gpu.cu

@@ -0,0 +1,726 @@
+/*! \file PHY\CUDA/LTE_TRANSPORT/turbo_rx_gpu.cu
+* \brief turbo decoder using gpu 
+* \author TerngYin Hsu, JianYa Chu
+* \date 2018
+* \version 0.1
+* \company ISIP LAB/NCTU CS  
+* \email: tyhsu@cs.nctu.edu.tw
+* \note
+* \warning
+*/
+
+#include <stdlib.h>
+#include <stdio.h>
+#include <stdint.h>
+#include <cuda_runtime.h>
+#include <cuda.h>
+#include <math.h>
+
+
+#include "turbo_parm.h"
+#include "PHY/CODING/extern_3GPPinterleaver.h"
+//#include "extern_interleaver.h"
+#include "PHY/CODING/defs.h"
+#include "turbo_rx_gpu.h"
+//#include "crc_byte.h"
+//#include "turbo_rx.h"
+//#include "PHY/CODING/extern_table_gpu.h"
+//#include "extern_table_gpu.h"
+#include "PHY/defs.h"
+//typedef int16_t llr_t;
+
+#define CRC24_A 0
+#define CRC24_B 1
+#define CRC16 2
+#define CRC8 3
+
+int intable_h[188][6144];
+int detable_h[188][6144];
+
+void free_ptr()
+{
+    cudaFree(turbo_parm->sys_d);
+	cudaFree(turbo_parm->sys1_d);
+	cudaFree(turbo_parm->sys2_d);
+	cudaFree(turbo_parm->ypar1_d);
+	cudaFree(turbo_parm->ypar2_d);
+	cudaFree(turbo_parm->alpha_d);
+	cudaFree(turbo_parm->alpha_pre_1);
+	cudaFree(turbo_parm->alpha_pre_2);
+	cudaFree(turbo_parm->beta_pre_1);
+	cudaFree(turbo_parm->beta_pre_2);
+	cudaFree(turbo_parm->ext_d);
+	cudaFree(turbo_parm->ext2_d);
+	cudaFree(turbo_parm->decode_ext2);
+	cudaFree(turbo_parm->decode_tmp);
+	cudaFreeHost(turbo_parm->decode_h);
+	free(turbo_parm);
+
+	char i;
+	for(i=0;i<2;i++)
+	{
+		cudaStreamDestroy(cuda_parm.stream[i]);
+	}
+}
+
+__constant__ int alpha_table_0[32];
+__constant__ int alpha_table_1[32];
+__constant__ int beta_table_0[32];
+__constant__ int beta_table_1[32];
+__constant__ float alpha_par_table_0[32];
+__constant__ float alpha_par_table_1[32];
+__constant__ float beta_par_table_0[32];
+__constant__ float beta_par_table_1[32];
+__constant__ int interleaver[6144];
+__constant__ int de_interleaver[6144];
+
+void init_alloc()
+{
+	size_t pitch;
+	
+    cudaDeviceProp deviceprop;
+    cudaGetDeviceProperties(&deviceprop,0);
+
+    cudaSetDeviceFlags(cudaDeviceMapHost);
+
+    if(deviceprop.canMapHostMemory!=1)
+        printf("cudaError:cannot map host to device memory\n");
+
+    cudaError_t result;
+	
+	turbo_parm = (turbo_parm_s*)malloc(sizeof(turbo_parm_s));
+	
+	// allocate CUDA memory 
+	result = cudaMallocPitch((void**)&turbo_parm->sys_d, &pitch, 16*6144*sizeof(llr_t) ,1);
+	if(result!=cudaSuccess)
+		printf("cudaMalloc turbo_parm->sys_d failed, err_num=%d\n",result);
+	
+	result = cudaMallocPitch((void**)&turbo_parm->sys1_d, &pitch, 16*6144*sizeof(llr_t) ,1);
+	if(result!=cudaSuccess)
+		printf("cudaMalloc turbo_parm->sys1_d failed, err_num=%d\n",result);
+	
+	result = cudaMallocPitch((void**)&turbo_parm->sys2_d, &pitch, 16*6144*sizeof(llr_t) ,1);
+	if(result!=cudaSuccess)
+		printf("cudaMalloc turbo_parm->sys2_d failed, err_num=%d\n",result);
+	
+    result = cudaMallocPitch((void**)&turbo_parm->ypar1_d, &pitch, 16*6144*sizeof(llr_t) ,1);
+	if(result!=cudaSuccess)
+		printf("cudaMalloc turbo_parm->ypar1_d failed, err_num=%d\n",result);
+	
+	result = cudaMallocPitch((void**)&turbo_parm->ypar2_d, &pitch, 16*6144*sizeof(llr_t) ,1);
+	if(result!=cudaSuccess)
+		printf("cudaMalloc turbo_parm->ypar2_d failed, err_num=%d\n",result);
+	
+	result =  cudaMallocPitch((void**)&turbo_parm->alpha_d, &pitch, 16*8*(6144+648)*sizeof(llr_t) ,1);
+	if(result!=cudaSuccess)
+		printf("cudaMalloc turbo_parm->alpha failed, err_num=%d\n",result);
+	
+	result =  cudaMallocPitch((void**)&turbo_parm->alpha_pre_1, &pitch, 16*32*162*4*sizeof(llr_t) ,1);
+	if(result!=cudaSuccess)
+		printf("cudaMalloc turbo_parm->alpha_pre_1 failed, err_num=%d\n",result);
+	
+	result =  cudaMallocPitch((void**)&turbo_parm->alpha_pre_2, &pitch, 16*32*162*4*sizeof(llr_t) ,1);
+	if(result!=cudaSuccess)
+		printf("cudaMalloc turbo_parm->alpha_pre_2 failed, err_num=%d\n",result);
+	
+	result = cudaMallocPitch((void**)&turbo_parm->beta_pre_1, &pitch, 16*32*162*4*sizeof(llr_t) ,1);
+	if(result!=cudaSuccess)
+		printf("cudaMalloc turbo_parm->beta_pre1 failed, err_num=%d\n",result);
+	
+	result = cudaMallocPitch((void**)&turbo_parm->beta_pre_2, &pitch, 16*32*162*4*sizeof(llr_t) ,1);
+	if(result!=cudaSuccess)
+		printf("cudaMalloc turbo_parm->beta_pre2 failed, err_num=%d\n",result);
+	
+	result =  cudaMallocPitch((void**)&turbo_parm->ext_d, &pitch, 16*6144*sizeof(llr_t) ,1);
+	if(result!=cudaSuccess)
+		printf("cudaMalloc turbo_parm->ext_d failed, err_num=%d\n",result);
+	
+	result = cudaMallocPitch((void**)&turbo_parm->ext2_d, &pitch, 16*6144*sizeof(llr_t) ,1);
+	if(result!=cudaSuccess)
+		printf("cudaMalloc turbo_parm->ext2_d failed, err_num=%d\n",result);
+
+	result = cudaMallocPitch((void**)&turbo_parm->decode_ext2, &pitch, 3*16*6144*sizeof(llr_t) ,1);
+	if(result!=cudaSuccess)
+		printf("cudaMalloc turbo_parm->ext2_d failed, err_num=%d\n",result);
+
+	result = cudaMallocPitch((void**)&turbo_parm->decode_tmp, &pitch, 3*16*6144*sizeof(int) ,1);
+	if(result!=cudaSuccess)
+		printf("cudaMalloc turbo_parm->ext2_d failed, err_num=%d\n",result);
+
+	result = cudaHostAlloc((void**)&turbo_parm->decode_h,3*16*768*sizeof(unsigned char), cudaHostAllocMapped);
+	if(result!=cudaSuccess)
+		printf("cudaHostAlloc turbo_parm->decode_h filaed, err_num=%d\n",result);
+	
+	// get device pointer
+	result = cudaHostGetDevicePointer(&turbo_parm->decode_d, turbo_parm->decode_h, 0);
+	if(result!=cudaSuccess)
+		printf("cuda get device pinter decode_d failed, err_num=%d\n",result);
+
+	// memset for mem
+	cudaMemset(turbo_parm->ext2_d,0,16*6144*sizeof(llr_t));
+	cudaMemset(turbo_parm->alpha_d,0,16*8*(6144+648)*sizeof(llr_t));
+	cudaMemset(turbo_parm->decode_tmp,0,16*6144*sizeof(llr_t));
+	//memset(turbo_parm->decode_h,0,16*768*sizeof(char));
+
+	// init table for decoder
+	int a_table_0[32]={0,3,4,7,1,2,5,6,8,11,12,15,9,10,13,14,16,19,20,23,17,18,21,22,24,27,28,31,25,26,29,30};
+	cudaMemcpyToSymbol(alpha_table_0,a_table_0,32*sizeof(int));
+	int a_table_1[32]={1,2,5,6,0,3,4,7,9,10,13,14,8,11,12,15,17,18,21,22,16,19,20,23,25,26,29,30,24,27,28,31};
+	cudaMemcpyToSymbol(alpha_table_1,a_table_1,32*sizeof(int));
+	
+	float a_p_table_0[32] = {0.0 ,1.0 ,1.0 ,0.0 ,0.0 ,1.0 ,1.0 ,0.0 ,0.0 ,1.0 ,1.0 ,0.0 ,0.0 ,1.0 ,1.0 ,0.0 ,0.0 ,1.0 ,1.0 ,0.0 ,0.0 ,1.0 ,1.0 ,0.0 ,0.0 ,1.0 ,1.0 ,0.0 ,0.0 ,1.0 ,1.0 ,0.0};
+	cudaMemcpyToSymbol(alpha_par_table_0,a_p_table_0,32*sizeof(llr_t));
+	float a_p_table_1[32] = {1.0 ,0.0 ,0.0 ,1.0 ,1.0 ,0.0 ,0.0 ,1.0 ,1.0 ,0.0 ,0.0 ,1.0 ,1.0 ,0.0 ,0.0 ,1.0 ,1.0 ,0.0 ,0.0 ,1.0 ,1.0 ,0.0 ,0.0 ,1.0 ,1.0 ,0.0 ,0.0 ,1.0 ,1.0 ,0.0 ,0.0 ,1.0};
+	cudaMemcpyToSymbol(alpha_par_table_1,a_p_table_1,32*sizeof(llr_t));
+	
+	int b_table_0[32]={0,4,5,1,2,6,7,3,8,12,13,9,10,14,15,11,16,20,21,17,18,22,23,19,24,28,29,25,26,30,31,27};
+	cudaMemcpyToSymbol(beta_table_0,b_table_0,32*sizeof(int));
+	int b_table_1[32]={4,0,1,5,6,2,3,7,12,8,9,13,14,10,11,15,20,16,17,21,22,18,19,23,28,24,25,29,30,26,27,31};
+	cudaMemcpyToSymbol(beta_table_1,b_table_1,32*sizeof(int));
+	
+	float b_p_table_0[32] = {0.0 ,0.0 ,1.0 ,1.0 ,1.0 ,1.0 ,0.0 ,0.0 , 0.0 ,0.0 ,1.0 ,1.0 ,1.0 ,1.0 ,0.0 ,0.0 , 0.0 ,0.0 ,1.0 ,1.0 ,1.0 ,1.0 ,0.0 ,0.0 , 0.0 ,0.0 ,1.0 ,1.0 ,1.0 ,1.0 ,0.0 ,0.0};
+	cudaMemcpyToSymbol(beta_par_table_0,b_p_table_0,32*sizeof(llr_t));
+	float b_p_table_1[32] = {1.0 ,1.0 ,0.0 ,0.0 ,0.0 ,0.0 ,1.0 ,1.0 , 1.0 ,1.0 ,0.0 ,0.0 ,0.0 ,0.0 ,1.0 ,1.0 , 1.0 ,1.0 ,0.0 ,0.0 ,0.0 ,0.0 ,1.0 ,1.0 , 1.0 ,1.0 ,0.0 ,0.0 ,0.0 ,0.0 ,1.0 ,1.0};
+	cudaMemcpyToSymbol(beta_par_table_1,b_p_table_1,32*sizeof(llr_t));
+	
+	
+	// build de-interleaver table and interleaver table
+	int i, j;
+	unsigned long n;
+	unsigned short f1, f2;
+	for(j=0;j<188;j++)
+	{
+		n = f1f2mat[j].nb_bits;
+		f1 = f1f2mat[j].f1;
+		f2 = f1f2mat[j].f2;
+		for(i=0;i<n;i++)
+		{
+			intable_h[j][i] = (((f1+f2*i)%n)*i)%n;
+			detable_h[j][(((f1+f2*i)%n)*i)%n] = i;
+		}
+	}
+
+	// for crc check and stream create
+	for(i=0;i<2;i++)
+	{
+		cudaStreamCreate(&cuda_parm.stream[i]);
+	}
+
+	for(i=0;i<3;i++)
+	{
+		cudaEventCreate(&cuda_parm.s_check[i]);
+	}
+}
+__device__ void compute_alpha(float* sys, float* sys1, float* sys2,
+							  float* par, 
+						      float* alpha, float* alpha_tmp,
+						      float* alpha_pre_1, float* alpha_pre_2,
+						      int num_per_block, int iteration_cnt, int decoder_id, int n, int codeword_num)
+{
+	int alpha_start = blockIdx.y*(n+gridDim.x*4)*8 + blockIdx.x*(num_per_block+1)*8*4;
+	int index = blockIdx.y*n + blockIdx.x*num_per_block*4 + num_per_block*threadIdx.y;
+	llr_t r0, r1;
+	char i;
+	
+	alpha_tmp[threadIdx.x + 8*threadIdx.y] = 0;
+	if(!(iteration_cnt==0 || (iteration_cnt==1 && decoder_id==2)))
+	{
+		if(!(blockIdx.x==0 && threadIdx.y==0))
+		{
+			if(decoder_id==1)
+				alpha_tmp[threadIdx.x + 8*threadIdx.y] = alpha_pre_1[blockIdx.y*gridDim.x*32 + blockIdx.x*32 + threadIdx.x + 8*threadIdx.y -8];
+			else
+				alpha_tmp[threadIdx.x + 8*threadIdx.y] = alpha_pre_2[blockIdx.y*gridDim.x*32 + blockIdx.x*32 + threadIdx.x + 8*threadIdx.y -8];
+		}
+	}
+	alpha[alpha_start + threadIdx.x + 8*threadIdx.y] = alpha_tmp[threadIdx.x + 8*threadIdx.y];
+	__syncthreads();
+	
+	for(i=0;i<num_per_block;i++)
+	{
+		if(decoder_id==1)
+		{
+			r0 = alpha_par_table_0[threadIdx.x + 8*threadIdx.y]*par[index + i];
+			r1 = sys1[index + i] + alpha_par_table_1[threadIdx.x + 8*threadIdx.y]*par[index + i];
+		}
+		else
+		{
+			r0 = alpha_par_table_0[threadIdx.x + 8*threadIdx.y]*par[index + i];
+			r1 = sys2[index + i] + alpha_par_table_1[threadIdx.x + 8*threadIdx.y]*par[index + i];
+		}
+		alpha[alpha_start + (i+1)*32 + threadIdx.x + 8*threadIdx.y] = fmaxf(alpha_tmp[alpha_table_0[threadIdx.x+8*threadIdx.y]] + r0, alpha_tmp[alpha_table_1[threadIdx.x + 8*threadIdx.y]] + r1);
+		__syncthreads();
+		alpha_tmp[threadIdx.x + 8*threadIdx.y] = alpha[alpha_start + (i+1)*32 + threadIdx.x + 8*threadIdx.y];
+		if(i==num_per_block-1)
+		{
+			if(decoder_id==1)
+				alpha_pre_1[blockIdx.y*gridDim.x*32 + blockIdx.x*32 + threadIdx.x + 8*threadIdx.y] = alpha_tmp[threadIdx.x + 8*threadIdx.y];
+			else
+				alpha_pre_2[blockIdx.y*gridDim.x*32 + blockIdx.x*32 + threadIdx.x + 8*threadIdx.y] = alpha_tmp[threadIdx.x + 8*threadIdx.y];
+		}
+	}
+	
+}
+
+__device__ void compute_beta_ext(float* sys, float* sys1, float* sys2,
+							float* par, 
+							float* alpha, float* beta_now, float* beta_next, 
+							float* beta_pre_1, float* beta_pre_2,
+							float* ext_tmp0, float* ext_tmp1,
+							float* ext, float* ext2, float* decode_ext2,
+							int num_per_block, int iteration_cnt, int decoder_id, int n, int codeword_num)
+{
+	llr_t a, r0, r1, max_0, max_1;
+	int alpha_start = blockIdx.y*(n+gridDim.x*4)*8 + blockIdx.x*(num_per_block+1)*8*4;
+	int index = blockIdx.y*n + blockIdx.x*num_per_block*4 + threadIdx.y*num_per_block;
+	int index2;
+	char i,j;
+	
+	beta_now[threadIdx.x + 8*threadIdx.y] = 0;
+	if(!(iteration_cnt==0 || (iteration_cnt==1 && decoder_id==2)))
+	{
+		if(!(blockIdx.x==gridDim.x-1 && threadIdx.y==3))
+		{
+			if(decoder_id==1)
+				beta_now[threadIdx.x + 8*threadIdx.y] = beta_pre_1[blockIdx.y*gridDim.x*32 + blockIdx.x*32 + threadIdx.x + 8*threadIdx.y +8];
+			else
+				beta_now[threadIdx.x + 8*threadIdx.y] = beta_pre_2[blockIdx.y*gridDim.x*32 + blockIdx.x*32 + threadIdx.x + 8*threadIdx.y +8];
+		}
+	}
+	__syncthreads();
+	
+	for(i=num_per_block-1;i>=0;i--)
+	{
+		if(decoder_id==1)
+		{
+			r0 = beta_par_table_0[threadIdx.x + 8*threadIdx.y]*par[index + i];
+			r1 = sys1[index + i] + beta_par_table_1[threadIdx.x + 8*threadIdx.y]*par[index + i];
+		}
+		else
+		{
+			r0 = beta_par_table_0[threadIdx.x + 8*threadIdx.y]*par[index + i];
+			r1 = sys2[index + i] + beta_par_table_1[threadIdx.x + 8*threadIdx.y]*par[index + i];
+		}
+		a = alpha[alpha_start + 32*i + threadIdx.x + 8*threadIdx.y];
+		beta_next[threadIdx.x + 8*threadIdx.y] = fmaxf(beta_now[beta_table_0[threadIdx.x+8*threadIdx.y]] + r0, beta_now[beta_table_1[threadIdx.x + 8*threadIdx.y]] + r1);
+		if(i==0)
+		{
+			if(decoder_id==1)
+				beta_pre_1[blockIdx.y*gridDim.x*32 + blockIdx.x*32 + threadIdx.x + 8*threadIdx.y] = beta_next[threadIdx.x + 8*threadIdx.y];
+			else
+				beta_pre_2[blockIdx.y*gridDim.x*32 + blockIdx.x*32 + threadIdx.x + 8*threadIdx.y] = beta_next[threadIdx.x + 8*threadIdx.y];
+		}
+		ext_tmp0[((num_per_block-1-i)&7)*32 + threadIdx.x + 8*threadIdx.y] = a + r0 + beta_now[beta_table_0[threadIdx.x + 8*threadIdx.y]];
+		ext_tmp1[((num_per_block-1-i)&7)*32 + threadIdx.x + 8*threadIdx.y] = a + r1 + beta_now[beta_table_1[threadIdx.x + 8*threadIdx.y]];
+		__syncthreads();
+		beta_now[threadIdx.x + 8*threadIdx.y] = beta_next[threadIdx.x + 8*threadIdx.y];
+		
+		if(((num_per_block-1-i)&7)==7)
+		{
+			max_0 = ext_tmp0[(7-threadIdx.x)*32 + threadIdx.x + 8*threadIdx.y];
+			max_1 = ext_tmp1[(7-threadIdx.x)*32 + threadIdx.x + 8*threadIdx.y];
+			for(j=1;j<8;j++)
+			{
+				index2 = (threadIdx.x + j)&7;
+				max_0 = fmaxf(max_0, ext_tmp0[(7-threadIdx.x)*32 + index2 + 8*threadIdx.y]);
+				max_1 = fmaxf(max_1, ext_tmp1[(7-threadIdx.x)*32 + index2 + 8*threadIdx.y]);
+			}
+			index2 = blockIdx.x*num_per_block*4 + threadIdx.y*num_per_block + i + threadIdx.x;
+			if(decoder_id==1)
+			{
+				ext[blockIdx.y*n + index2] = max_1 - max_0 - sys1[blockIdx.y*n + index2] + sys[blockIdx.y*n + index2];
+				sys2[blockIdx.y*n + de_interleaver[index2]] = ext[blockIdx.y*n + index2];
+			}
+			else
+			{
+				ext2[blockIdx.y*n + index2] = max_1 - max_0;
+				if(iteration_cnt >= 3)
+				{
+					decode_ext2[(iteration_cnt-3)*gridDim.y*n + blockIdx.y*n + interleaver[index2]] = max_1 - max_0;
+				}
+				sys1[blockIdx.y*n + interleaver[index2]] = max_1 - max_0 - ext[blockIdx.y*n + interleaver[index2]] + sys[blockIdx.y*n + interleaver[index2]];
+			}
+		}
+		else if(i==0 && (num_per_block&7)!=0 && threadIdx.x < (num_per_block&7))
+		{
+			max_0 = ext_tmp0[((num_per_block&7)-1-threadIdx.x)*32 + threadIdx.x + 8*threadIdx.y];
+			max_1 = ext_tmp1[((num_per_block&7)-1-threadIdx.x)*32 + threadIdx.x + 8*threadIdx.y];
+			for(j=1;j<8;j++)
+			{
+				index2 = (threadIdx.x + j)&7;
+				max_0 = fmaxf(max_0, ext_tmp0[((num_per_block&7)-1-threadIdx.x)*32 + index2 + 8*threadIdx.y]);
+				max_1 = fmaxf(max_1, ext_tmp1[((num_per_block&7)-1-threadIdx.x)*32 + index2 + 8*threadIdx.y]);
+			}
+			index2 = blockIdx.x*num_per_block*4 + threadIdx.y*num_per_block + i + threadIdx.x;
+			if(decoder_id==1)
+			{
+				ext[blockIdx.y*n + index2] = max_1 - max_0 - sys1[blockIdx.y*n + index2] + sys[blockIdx.y*n + index2];
+				sys2[blockIdx.y*n + de_interleaver[index2]] = ext[blockIdx.y*n + index2];
+			}
+			else
+			{
+				ext2[blockIdx.y*n + index2] = max_1 - max_0;
+				if(iteration_cnt >= 3)
+				{
+					decode_ext2[(iteration_cnt-3)*gridDim.y*n + blockIdx.y*n + interleaver[index2]] = max_1 - max_0;
+				}
+				sys1[blockIdx.y*n + interleaver[index2]] = max_1 - max_0 - ext[blockIdx.y*n + interleaver[index2]] + sys[blockIdx.y*n + interleaver[index2]];
+			}
+		}
+	}
+}
+
+__global__ void log(float* sys, float* sys1, float* sys2,
+					float* par,
+					float* alpha, 
+					float* alpha_pre_1, float* alpha_pre_2, float* beta_pre_1, float* beta_pre_2,
+					float* ext, float* ext2, float* decode_ext2,
+					int num_per_block, int iteration_cnt, int decoder_id, int n, int codeword_num
+					)
+{
+	__shared__ llr_t alpha_tmp[32];
+	__shared__ llr_t beta_tmp[32];
+	__shared__ llr_t ext_tmp0[32*8];
+	__shared__ llr_t ext_tmp1[32*8];
+	
+	compute_alpha(sys, sys1, sys2,
+				   par,
+				   alpha, alpha_tmp,
+				   alpha_pre_1, alpha_pre_2,
+				   num_per_block, iteration_cnt, decoder_id, n , codeword_num
+				   );
+	__syncthreads();			  
+	
+	compute_beta_ext(sys, sys1, sys2,
+					 par,
+					 alpha, alpha_tmp,beta_tmp,
+					 beta_pre_1, beta_pre_2,
+					 ext_tmp0, ext_tmp1,
+					 ext, ext2, decode_ext2,
+					 num_per_block, iteration_cnt, decoder_id, n, codeword_num
+					);
+		 
+}
+
+// for decoding
+__global__ void decode(llr_t* decode_ext2, unsigned char* decode_d, int* decode_tmp, int n, int decode_len, int iteration_cnt)
+{
+	int i, j;
+	for(i=threadIdx.x; i<n; i+=256)
+	{
+		decode_tmp[iteration_cnt*gridDim.x*n + blockIdx.x*n + i] = 0;
+		if(decode_ext2[iteration_cnt*gridDim.x*n + blockIdx.x*n + i] > 0)
+		{
+			decode_tmp[iteration_cnt*gridDim.x*n + blockIdx.x*n + i] = 1 << (7-(i&7));
+		}
+	}
+	__syncthreads();
+	for(i=threadIdx.x; i<decode_len; i+=256)
+	{
+		decode_d[iteration_cnt*gridDim.x*decode_len + blockIdx.x*decode_len + i] = 0;
+		for(j=0;j<8;j++)
+		{
+			decode_d[iteration_cnt*gridDim.x*decode_len + blockIdx.x*decode_len + i] += decode_tmp[iteration_cnt*gridDim.x*n + blockIdx.x*n + i*8 + j];
+		}
+	}
+}
+
+//#define TIME_EST
+
+unsigned char phy_threegpplte_turbo_decoder_gpu(short **y,
+        unsigned char **decoded_bytes,
+		unsigned int codeword_num,
+        unsigned short n,
+        unsigned short f1,
+        unsigned short f2,
+        unsigned char max_iterations,
+        unsigned char crc_type,
+        unsigned char *f_tmp,
+		unsigned char* ret)
+{
+	unsigned int i,j,iind,k;
+#ifdef TIME_EST
+	cudaEventCreate(&cuda_parm.e_start);
+	cudaEventCreate(&cuda_parm.e_stop);
+	cuda_parm.e_time = 0;
+	cudaEventRecord(cuda_parm.e_start, cuda_parm.stream[0]);
+#endif
+
+    llr_t sys_h[n*codeword_num], ypar1_h[n*codeword_num], ypar2_h[n*codeword_num];
+	
+	unsigned char iteration_cnt=0;
+	unsigned int crc,oldcrc,crc_len;
+	uint8_t temp;
+	unsigned char F;
+	
+	if (crc_type > 3) {
+		printf("Illegal crc length!\n");
+		return 255;
+	}
+	
+	for (iind=0; f1f2mat[iind].nb_bits!=n && iind <188; iind++);
+	
+	if ( iind == 188 ) {
+		printf("Illegal frame length!\n");
+		return 255;
+	}
+	
+	switch (crc_type) {
+		case CRC24_A:
+		case CRC24_B:
+			crc_len=3;
+			break;
+
+		case CRC16:
+			crc_len=2;
+			break;
+
+		case CRC8:
+			crc_len=1;
+			break;
+
+		default:
+			crc_len=3;
+	}
+	
+	// fetch data for each codeword
+#ifdef TIME_EST
+	cudaEventCreate(&cuda_parm.f_start);
+	cudaEventCreate(&cuda_parm.f_stop);
+	cuda_parm.f_time = 0;
+	cudaEventRecord(cuda_parm.f_start,cuda_parm.stream[0]);
+#endif
+	short* yp;
+	for(i=0;i<codeword_num;i++)
+	{
+		yp = y[i];
+		for(j=0;j<n;j++)
+		{
+			sys_h[j+n*i] = *yp;
+			ypar1_h[j+n*i] = *(yp+1);
+			ypar2_h[j+n*i] = *(yp+2);
+			yp+=3;
+		}
+	}
+#ifdef TIME_EST
+	cudaEventRecord(cuda_parm.f_stop,cuda_parm.stream[0]);
+	cudaEventSynchronize(cuda_parm.f_stop);
+	cuda_parm.f_time = 0;
+	cudaEventElapsedTime(&cuda_parm.f_time, cuda_parm.f_start, cuda_parm.f_stop);
+	printf("Fetech data time = %f ms\n",cuda_parm.f_time);
+#endif
+
+	// for kernel  memcpy
+#ifdef TIME_EST
+	cudaEventCreate(&cuda_parm.m_start);
+	cudaEventCreate(&cuda_parm.m_stop);
+	cudaEventRecord(cuda_parm.m_start, cuda_parm.stream[0]);
+#endif
+	cudaMemcpyAsync(turbo_parm->sys_d,sys_h,codeword_num*n*sizeof(llr_t),cudaMemcpyHostToDevice, cuda_parm.stream[0]);
+	cudaMemcpyAsync(turbo_parm->ypar1_d,ypar1_h,codeword_num*n*sizeof(llr_t),cudaMemcpyHostToDevice, cuda_parm.stream[0]);
+	cudaMemcpyAsync(turbo_parm->ypar2_d,ypar2_h,codeword_num*n*sizeof(llr_t),cudaMemcpyHostToDevice, cuda_parm.stream[0]);
+	cudaMemcpyToSymbolAsync(interleaver, intable_h[iind],n*sizeof(int), 0, cudaMemcpyHostToDevice,  cuda_parm.stream[0]);
+	cudaMemcpyToSymbolAsync(de_interleaver, detable_h[iind],n*sizeof(int), 0, cudaMemcpyHostToDevice,  cuda_parm.stream[0]);
+#ifdef TIME_EST
+	cudaEventRecord(cuda_parm.m_stop, cuda_parm.stream[0]);
+	cudaEventSynchronize(cuda_parm.m_stop);	
+	cuda_parm.m_time=0;
+	cudaEventElapsedTime(&cuda_parm.m_time, cuda_parm.m_start, cuda_parm.m_stop);
+	printf("Memcpy Time For Kernel = %f ms\n",cuda_parm.m_time);
+#endif
+
+#ifdef TIME_EST	
+	cuda_parm.t_time = 0;
+	cudaEventCreate(&cuda_parm.t_start);
+	cudaEventCreate(&cuda_parm.t_stop);
+	cudaEventRecord(cuda_parm.t_start,cuda_parm.stream[0]);
+#endif
+	// decide block and thread
+	int blocknum=648;
+	while(blocknum!=8)
+	{
+		if(n%blocknum==0 && n/blocknum>=16)
+		{
+			break;
+		}
+		blocknum-=4;
+	}
+	
+	dim3 threadnum(8,4);
+	size_t s_size = 0;
+	int num_per_block = n / blocknum;
+	blocknum = blocknum/4;
+	dim3 bb(blocknum, codeword_num);
+	
+	llr_t ext2[n*codeword_num];
+	memset(ext2,0,sizeof(ext2));
+	llr_t tmp[n*codeword_num];
+	
+	// for crc check
+	char check=0;
+	int z;
+	// log map algorithm
+	log<<<bb, threadnum, s_size, cuda_parm.stream[0]>>>(turbo_parm->sys_d, turbo_parm->sys_d, turbo_parm->sys2_d,
+											  turbo_parm->ypar1_d,
+											  turbo_parm->alpha_d,
+											  turbo_parm->alpha_pre_1, turbo_parm->alpha_pre_2, turbo_parm->beta_pre_1, turbo_parm->beta_pre_2,
+											  turbo_parm->ext_d, turbo_parm->ext2_d, turbo_parm->decode_ext2,
+											  num_per_block, iteration_cnt, 1, n, codeword_num
+											);
+	
+	while(iteration_cnt++ < max_iterations)
+	{
+		
+		log<<<bb, threadnum, s_size, cuda_parm.stream[0]>>>(turbo_parm->sys_d, turbo_parm->sys1_d, turbo_parm->sys2_d,
+												  turbo_parm->ypar2_d,
+												  turbo_parm->alpha_d,
+												  turbo_parm->alpha_pre_1, turbo_parm->alpha_pre_2, turbo_parm->beta_pre_1, turbo_parm->beta_pre_2,
+												  turbo_parm->ext_d, turbo_parm->ext2_d, turbo_parm->decode_ext2,
+												  num_per_block, iteration_cnt, 2, n, codeword_num
+												);
+		
+		if(iteration_cnt>=3)
+		{
+			cudaEventRecord(cuda_parm.s_check[iteration_cnt-3], cuda_parm.stream[0]);	
+		}
+
+		if(iteration_cnt < max_iterations)
+		{
+			log<<<bb, threadnum, s_size, cuda_parm.stream[0]>>>(turbo_parm->sys_d, turbo_parm->sys1_d, turbo_parm->sys2_d,
+													  turbo_parm->ypar1_d,
+													  turbo_parm->alpha_d,
+													  turbo_parm->alpha_pre_1, turbo_parm->alpha_pre_2, turbo_parm->beta_pre_1, turbo_parm->beta_pre_2,
+													  turbo_parm->ext_d, turbo_parm->ext2_d, turbo_parm->decode_ext2,
+													  num_per_block, iteration_cnt, 1, n, codeword_num
+													);
+		}
+	}
+	cudaDeviceSynchronize();
+	
+#ifdef TIME_EST
+	cudaEventRecord(cuda_parm.t_stop, cuda_parm.stream[0]);
+	cudaEventSynchronize(cuda_parm.t_stop);
+	cudaEventElapsedTime(&cuda_parm.t_time, cuda_parm.t_start, cuda_parm.t_stop);
+	
+	printf("Time For turbo algorithm Kernel = %f ms\n",cuda_parm.t_time);	
+#endif	
+
+	int decode_len = n >> 3;
+	for(i=0;i<=2;i++)
+	{
+		check = 0;
+		// wait for turbo ext
+		cudaStreamWaitEvent(cuda_parm.stream[1], cuda_parm.s_check[i], 0);
+#ifdef TIME_EST
+	cuda_parm.d_time = 0;
+	cudaEventCreate(&cuda_parm.d_start);
+	cudaEventCreate(&cuda_parm.d_stop);
+	cudaEventRecord(cuda_parm.d_start, cuda_parm.stream[1]);
+#endif
+		// decode
+		decode<<<codeword_num, 256, 0, cuda_parm.stream[1]>>>(turbo_parm->decode_ext2, turbo_parm->decode_d, turbo_parm->decode_tmp, n, decode_len, i);
+		cudaStreamSynchronize(cuda_parm.stream[1]);
+
+		// crc check
+		for(j=0;j<codeword_num;j++)
+		{
+			F = f_tmp[1];
+			if(j==0)
+			{
+				F = f_tmp[0];
+			}
+			oldcrc = *((unsigned int *)(&turbo_parm->decode_h[i*codeword_num*decode_len + j*decode_len + decode_len-crc_len]));
+			switch(crc_type)
+			{
+				case CRC24_A:
+					oldcrc&=0x00ffffff;
+					crc = crc24a(&turbo_parm->decode_h[i*codeword_num*decode_len + j*decode_len + ( F>>3 )], n-24-F)>>8;
+					temp=((uint8_t *)&crc)[2];
+					((uint8_t *)&crc)[2] = ((uint8_t *)&crc)[0];
+					((uint8_t *)&crc)[0] = temp;
+					break;
+				case CRC24_B:
+					oldcrc&=0x00ffffff;
+					crc = crc24b(&turbo_parm->decode_h[i*codeword_num*decode_len + j*decode_len],
+								 n-24)>>8;
+					temp=((uint8_t *)&crc)[2];
+					((uint8_t *)&crc)[2] = ((uint8_t *)&crc)[0];
+					((uint8_t *)&crc)[0] = temp;
+					break;
+
+				case CRC16:
+					oldcrc&=0x0000ffff;
+					crc = crc16(&turbo_parm->decode_h[i*codeword_num*decode_len + j*decode_len],
+								n-16)>>16;
+					break;
+
+				case CRC8:
+					oldcrc&=0x000000ff;
+					crc = crc8(&turbo_parm->decode_h[i*codeword_num*decode_len + j*decode_len],
+							   n-8)>>24;
+					break;
+
+				default:
+					printf("FATAL: 3gpplte_turbo_decoder_sse.c: Unknown CRC\n");
+					return(255);
+					break;
+			}
+
+			if ((crc == oldcrc) && (crc!=0)) {
+				ret[j] = i+3-1;
+			}
+			else
+			{
+				check = 1;
+			}			
+		}
+
+#ifdef TIME_EST
+	cudaEventRecord(cuda_parm.d_stop, cuda_parm.stream[1]);
+	cudaEventSynchronize(cuda_parm.d_stop);
+	cudaEventElapsedTime(&cuda_parm.d_time, cuda_parm.d_start, cuda_parm.d_stop);
+	printf("Time For decode & crc check = %f ms\n",cuda_parm.d_time);
+#endif	
+		if(check==0)
+		{
+			for(j=0;j<codeword_num;j++)
+			{
+				for(k=0;k<decode_len;k++)
+				{
+					decoded_bytes[j][k] = turbo_parm->decode_h[i*codeword_num*decode_len + j*decode_len + k];
+				}
+			}
+			
+#ifdef TIME_EST
+	cudaEventRecord(cuda_parm.e_stop, cuda_parm.stream[0]);
+	cudaEventSynchronize(cuda_parm.e_stop);
+	cudaEventElapsedTime(&cuda_parm.e_time, cuda_parm.e_start, cuda_parm.e_stop);
+	printf("Time For CUDA = %f ms\n",cuda_parm.e_time);
+#endif
+			return i+3-1;
+		}
+	}
+
+	// crc check fail
+	for(i=0;i<codeword_num;i++)
+	{
+		for(j=0;j<decode_len;j++)
+		{
+			decoded_bytes[i][j] = turbo_parm->decode_h[2*codeword_num*decode_len + i*decode_len + j];
+		}
+		ret[i] = 5;
+		//return 5;
+	}
+	
+#ifdef TIME_EST
+	cudaEventRecord(cuda_parm.e_stop, cuda_parm.stream[0]);
+	cudaEventSynchronize(cuda_parm.e_stop);
+	cudaEventElapsedTime(&cuda_parm.e_time, cuda_parm.e_start, cuda_parm.e_stop);
+	printf("Time For CUDA = %f ms\n",cuda_parm.e_time);
+#endif	
+	for(i=0;i<2;i++)
+	{
+		cudaStreamSynchronize(cuda_parm.stream[i]);
+	}
+    return 5;
+}

openair1/PHY/CUDA/LTE_TRANSPORT/turbo_rx_gpu.h

@@ -0,0 +1,53 @@
+/*! \file PHY\CUDA/LTE_TRANSPORT/turbo_rx_gpu.h
+* \brief turbo decoder using gpu 
+* \author TerngYin Hsu, JianYa Chu
+* \date 2018
+* \version 0.1
+* \company ISIP LAB/NCTU CS  
+* \email: tyhsu@cs.nctu.edu.tw
+* \note
+* \warning
+*/
+
+
+#ifndef __TURBO_RX_GPU__H__
+#define __TURBO_RX_GPU__H__
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <stdint.h>
+
+#define CRC24_A 0
+#define CRC24_B 1
+#define CRC16 2
+#define CRC8 3
+
+typedef char Binary;
+typedef float llr_t;
+typedef short channel_t;
+
+#ifdef __cplusplus
+extern "C"
+#endif
+unsigned char phy_threegpplte_turbo_decoder_gpu(short **y,
+        unsigned char **decoded_bytes,
+		unsigned int codeword_num,
+        unsigned short n,
+        unsigned short f1,
+        unsigned short f2,
+        unsigned char max_iterations,
+        unsigned char crc_type,
+        unsigned char* f_tmp,
+		unsigned char* ret);
+
+#ifdef __cplusplus
+extern "C"
+#endif
+void free_ptr(void);
+
+#ifdef __cplusplus
+extern "C"
+#endif
+void init_alloc(void);
+
+#endif