maki/openairinterface5g
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

feat(ldpc_cl): Remove ldpc_cl

Browse Source
This commit is contained in:
Romain Beurdouche
2025-08-11 08:38:06 +00:00
parent a2c03f26e9
commit b9bed5838e
4 changed files with 0 additions and 897 deletions
Download Patch File Download Diff File Expand all files Collapse all files
openair1/PHY/CODING/CMakeLists.txt
-13
View File
@@ -32,19 +32,6 @@ add_dependencies(nr_ulschsim ldpc ldpc_orig)
add_dependencies(nr_dlsim ldpc ldpc_orig)
add_dependencies(nr_dlschsim ldpc ldpc_orig)
add_custom_target(nrLDPC_decoder_kernels_CL
COMMAND gcc nrLDPC_decoder/nrLDPC_decoder_CL.c -dD -DNRLDPC_KERNEL_SOURCE -E -o ${CMAKE_CURRENT_BINARY_DIR}/nrLDPC_decoder_kernels_CL.clc
SOURCES nrLDPC_decoder/nrLDPC_decoder_CL.c
WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}
)
add_library(ldpc_cl MODULE
nrLDPC_decoder/nrLDPC_decoder_CL.c
nrLDPC_encoder/ldpc_encoder_optim8segmulti.c
)
set_target_properties(ldpc_cl PROPERTIES LIBRARY_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR})
target_link_libraries(ldpc_cl PRIVATE ldpc_segment OpenCL)
add_dependencies(ldpc_cl nrLDPC_decoder_kernels_CL)
##############################################
# Base CUDA setting
##############################################
openair1/PHY/CODING/DOC/LDPCImplementation.md
-44
View File
@@ -107,40 +107,6 @@ loading `libldpc_orig.so` instead of `libldpc.so`:
./nr-softmodem -O libconfig:gnb.band78.tm1.106PRB.usrpx300.conf:dbgl5 --loader.ldpc.shlibversion _orig
```
loading `libldpc_cl.so` instead of `libldpc.so`:
`make ldpc_cl`
This command creates the `libldpc_cl.so` shared library. To perform this build successfully, only the OpenCL header `/usr/include/CL/opencl.h` and library `/usr/lib/x86_64-linux-gnu/libOpenCL.so` are required, they implement OpenCL API support which is not hardware dependent.
```
Scanning dependencies of target nrLDPC_decoder_kernels_CL
Built target nrLDPC_decoder_kernels_CL
Scanning dependencies of target ldpc_cl
Building C object CMakeFiles/ldpc_cl.dir/usr/local/oai/oai-develop/openairinterface5g/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c.o
In file included from /usr/include/CL/cl.h:32,
from /usr/include/CL/opencl.h:38,
from /usr/local/oai/oai-develop/openairinterface5g/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c:49:
/usr/include/CL/cl_version.h:34:9: note: #pragma message: cl_version.h: CL_TARGET_OPENCL_VERSION is not defined. Defaulting to 220 (OpenCL 2.2)
#pragma message("cl_version.h: CL_TARGET_OPENCL_VERSION is not defined. Defaulting to 220 (OpenCL 2.2)")
^~~~~~~
Building C object CMakeFiles/ldpc_cl.dir/usr/local/oai/oai-develop/openairinterface5g/openair1/PHY/CODING/nrLDPC_encoder/ldpc_encoder_optim8segmulti.c.o
Linking C shared module libldpc_cl.so
Built target ldpc_cl
```
At runtime, to successfully use hardware acceleration via OpenCL, you need to install vendor specific packages which deliver the required drivers and tools to make use of their GPU (Nvidia, Intel...) , fpga (Xilinx, Intel) or CPU (Intel, AMD, ARM...) through OpenCL.
```
./nr-softmodem -O libconfig:gnb.band78.sa.fr1.106PRB.usrpb210.conf:dbgl5 --rfsim --rfsimulator.serveraddr server --log_config.gtpu_log_level info --loader.ldpc.shlibversion _cl
```
```
./nr-uesoftmodem -r 106 --numerology 1 --band 78 -C 3619200000 --rfsim -O libconfig:/usr/local/oai/conf/nrue_sim.conf:dbgl5 --loader.ldpc.shlibversion _cl --log_config.hw_log_level inf
```
A mechanism to select ldpc implementation is also available in the `ldpctest` phy simulator via the `-v` option, which can be used to specify the version of the ldpc shared library to be used.
#### Examples of ldpc shared lib selection when running ldpctest:
@@ -151,14 +117,6 @@ Loading libldpc_cuda.so, the cuda implementation of the ldpc decoder:
$ ./ldpctest -v _cuda
```
Loading libldpc_cl.so, the opencl implementation of the ldpc decoder:
`make ldpc_cl`
```
$ ./ldpctest -v _cl
```
### LDPC libraries
The prototypes for the functions of the interface are defined in [nrLDPC_defs.h](file://nrLDPC_defs.h) as typedefs.
Libraries implementing the LDPC algorithms must be named `libldpc<_version>.so`, they must implement four functions with the names and types:
@@ -169,6 +127,4 @@ Libraries implementing the LDPC algorithms must be named `libldpc<_version>.so`,
`libldpc_cuda.so` has been tested with the `ldpctest` executable, usage from the softmodem's has to be tested.
`libldpc_cl.so` is under development.
[oai Wikis home](https://gitlab.eurecom.fr/oai/openairinterface5g/wikis/home)
openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c
-548
View File
@@ -1,548 +0,0 @@
/*
* Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
* contributor license agreements. See the NOTICE file distributed with
* this work for additional information regarding copyright ownership.
* The OpenAirInterface Software Alliance licenses this file to You under
* the OAI Public License, Version 1.0 (the "License"); you may not use this file
* except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.openairinterface.org/?page_id=698
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*-------------------------------------------------------------------------------
* For more information about the OpenAirInterface (OAI) Software Alliance:
* contact@openairinterface.org
*/
/*! \file nrLDPC_decoder_CL.c
* \brief ldpc decoder, openCL implementaion
* \author Francois TABURET
* \date 2021
* \version 1.0
* \note initial implem - translation of cuda version
*/
#define MAX_ITERATION 2
#define MC 1
#define MAX_OCLDEV 10
#define MAX_OCLRUNTIME 5
typedef struct {
char x;
char y;
short value;
} h_element;
#ifdef NRLDPC_KERNEL_SOURCE
#include "nrLDPC_decoder_kernels_CL.c"
#else
/* uses HW component id for log messages ( --log_config.hw_log_level <warning| info|debug|trace>) */
#include <stdio.h>
#include <unistd.h>
#include <sys/stat.h>
#include <CL/opencl.h>
#include "openair1/PHY/CODING/nrLDPC_extern.h"
#include "assertions.h"
#include "common/utils/LOG/log.h"
#define CLSETKERNELARG(A, B, C, D) \
rt = clSetKernelArg(A, B, C, D); \
AssertFatal(rt == CL_SUCCESS, "Error %d setting kernel argument index %d\n", (int)rt, B);
#include "../nrLDPC_decoder_LYC/bgs/BG1_compact_in_C.h"
typedef struct {
cl_uint max_CU;
cl_uint max_WID;
size_t *max_WIS;
} ocldev_t;
typedef struct {
cl_kernel cnp_kernel_1st;
cl_kernel cnp_kernel;
cl_kernel vnp_kernel_normal;
cl_kernel pack_decoded_bit;
} oclkernels_t;
typedef struct {
cl_uint num_devices;
cl_device_id devices[MAX_OCLDEV];
ocldev_t ocldev[MAX_OCLDEV];
cl_context context;
cl_program program;
cl_mem dev_h_compact1;
cl_mem dev_h_compact2;
cl_mem dev_const_llr;
cl_mem dev_llr;
cl_mem dev_dt;
cl_mem dev_tmp;
oclkernels_t kernels[MAX_OCLDEV];
cl_command_queue queue[MAX_OCLDEV];
} oclruntime_t;
typedef struct {
oclruntime_t runtime[MAX_OCLRUNTIME];
} ocl_t;
ocl_t ocl;
void set_compact_BG(int Zc, short BG)
{
cl_uint rt;
int row, col;
if (BG == 1) {
row = 46;
col = 68;
} else {
row = 42;
col = 52;
}
int compact_row = 30;
int compact_col = 19;
if (BG == 2) {
compact_row = 10, compact_col = 23;
}
int memorySize_h_compact1 = row * compact_col * sizeof(h_element);
int memorySize_h_compact2 = compact_row * col * sizeof(h_element);
int lift_index = 0;
short lift_set[][9] = {{2, 4, 8, 16, 32, 64, 128, 256},
{3, 6, 12, 24, 48, 96, 192, 384},
{5, 10, 20, 40, 80, 160, 320},
{7, 14, 28, 56, 112, 224},
{9, 18, 36, 72, 144, 288},
{11, 22, 44, 88, 176, 352},
{13, 26, 52, 104, 208},
{15, 30, 60, 120, 240},
{0}};
for (int i = 0; lift_set[i][0] != 0; i++) {
for (int j = 0; lift_set[i][j] != 0; j++) {
if (Zc == lift_set[i][j]) {
lift_index = i;
break;
}
}
}
printf("\nZc = %d BG = %d\n", Zc, BG);
ocl.runtime[0].dev_h_compact1 =
clCreateBuffer(ocl.runtime[0].context, CL_MEM_READ_ONLY | CL_MEM_HOST_WRITE_ONLY, memorySize_h_compact1, NULL, (cl_int *)&rt);
AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_h_compact1 for platform %i \n", (int)rt, 0);
ocl.runtime[0].dev_h_compact2 =
clCreateBuffer(ocl.runtime[0].context, CL_MEM_READ_ONLY | CL_MEM_HOST_WRITE_ONLY, memorySize_h_compact2, NULL, (cl_int *)&rt);
AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_h_compact2 for platform %i \n", (int)rt, 0);
h_element *h1;
h_element *h2;
switch (lift_index) {
case 0:
h1 = host_h_compact1_I0;
h2 = host_h_compact2_I0;
break;
case 1:
h1 = host_h_compact1_I1;
h2 = host_h_compact2_I1;
break;
case 2:
h1 = host_h_compact1_I2;
h2 = host_h_compact2_I2;
break;
case 3:
h1 = host_h_compact1_I3;
h2 = host_h_compact2_I3;
break;
case 4:
h1 = host_h_compact1_I4;
h2 = host_h_compact2_I4;
break;
case 5:
h1 = host_h_compact1_I5;
h2 = host_h_compact2_I5;
break;
case 6:
h1 = host_h_compact1_I6;
h2 = host_h_compact2_I6;
break;
case 7:
h1 = host_h_compact1_I7;
h2 = host_h_compact2_I7;
break;
default:
AssertFatal(0, "Invalid lift_index value %i\n", lift_index);
break;
}
rt = clEnqueueWriteBuffer(ocl.runtime[0].queue[0],
ocl.runtime[0].dev_h_compact1,
CL_TRUE,
0,
memorySize_h_compact1,
h1,
0,
NULL,
NULL);
AssertFatal(rt == CL_SUCCESS, "Error %d moving h_compact1 memory to pltf %i dev %i\n", (int)rt, 0, 0);
rt = clEnqueueWriteBuffer(ocl.runtime[0].queue[0],
ocl.runtime[0].dev_h_compact2,
CL_TRUE,
0,
memorySize_h_compact2,
h2,
0,
NULL,
NULL);
AssertFatal(rt == CL_SUCCESS, "Error %d moving h_compact2 memory to pltf %i dev %i\n", (int)rt, 0, 0);
// return 0;
}
void cl_error_callback(const char *errinfo, const void *private_info, size_t cb, void *user_data)
{
oclruntime_t *runtime = (oclruntime_t *)user_data;
LOG_E(HW, "OpenCL accelerator error %s\n", errinfo);
}
char *clutil_getstrdev(int intdev)
{
static char retstring[255] = "";
char *retptr = retstring;
retptr += sprintf(retptr, "0x%08x: ", (uint32_t)intdev);
if (intdev & CL_DEVICE_TYPE_CPU)
retptr += sprintf(retptr, "%s", "cpu ");
if (intdev & CL_DEVICE_TYPE_GPU)
retptr += sprintf(retptr, "%s", "gpu ");
if (intdev & CL_DEVICE_TYPE_ACCELERATOR)
retptr += sprintf(retptr, "%s", "acc ");
return retstring;
}
void get_CompilErr(cl_program program, int pltf)
{
// Determine the size of the log
size_t log_size;
for (int i = 0; i < ocl.runtime[pltf].num_devices; i++) {
clGetProgramBuildInfo(program, ocl.runtime[pltf].devices[i], CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
// Allocate memory for the log
char *log = (char *)malloc(log_size);
// Get the log
clGetProgramBuildInfo(program, ocl.runtime[pltf].devices[i], CL_PROGRAM_BUILD_LOG, log_size, log, NULL);
// Print the log
printf("%s\n", log);
free(log);
}
}
size_t load_source(char **source_str, char *filename)
{
FILE *fp;
struct stat st;
size_t source_size;
char *src = NULL;
if (filename == NULL) {
src = "nrLDPC_decoder_kernels_CL.clc";
} else {
src = filename;
}
fp = fopen(src, "r");
AssertFatal(fp, "failed to open cl source %s: %s\n", src, strerror(errno));
fstat(fileno(fp), &st);
source_size = st.st_size;
*source_str = (char *)malloc(source_size);
source_size = fread(*source_str, 1, source_size, fp);
fclose(fp);
LOG_I(HW, "Loaded kernel sources from %s %u bytes\n", (filename == NULL) ? "embedded cl code" : src, (unsigned int)source_size);
return source_size;
}
/* from here: entry points in decoder shared lib */
int ldpc_autoinit(void)
{ // called by the library loader
cl_platform_id platforms[10];
cl_uint num_platforms_found;
int context_ok = 0;
cl_uint rt = clGetPlatformIDs(sizeof(platforms) / sizeof(cl_platform_id), platforms, &num_platforms_found);
AssertFatal(rt == CL_SUCCESS, "clGetPlatformIDs error %d\n", (int)rt);
AssertFatal(num_platforms_found > 0, "clGetPlatformIDs: no cl compatible platform found\n");
for (int i = 0; i < (int)num_platforms_found; i++) {
char stringval[255];
rt = clGetPlatformInfo(platforms[i], CL_PLATFORM_PROFILE, sizeof(stringval), stringval, NULL);
AssertFatal(rt == CL_SUCCESS, "clGetPlatformInfo PROFILE error %d\n", (int)rt);
LOG_I(HW, "Platform %i, OpenCL profile %s\n", i, stringval);
rt = clGetPlatformInfo(platforms[i], CL_PLATFORM_VERSION, sizeof(stringval), stringval, NULL);
AssertFatal(rt == CL_SUCCESS, "clGetPlatformInfo VERSION error %d\n", (int)rt);
LOG_I(HW, "Platform %i, OpenCL version %s\n", i, stringval);
rt = clGetDeviceIDs(platforms[i],
CL_DEVICE_TYPE_ALL,
sizeof(ocl.runtime[i].devices) / sizeof(cl_device_id),
ocl.runtime[i].devices,
&(ocl.runtime[i].num_devices));
AssertFatal(rt == CL_SUCCESS, "clGetDeviceIDs error %d\n", (int)rt);
int devok = 0;
for (int j = 0; j < ocl.runtime[i].num_devices; j++) {
cl_bool abool;
rt = clGetDeviceInfo(ocl.runtime[i].devices[j], CL_DEVICE_AVAILABLE, sizeof(abool), &abool, NULL);
AssertFatal(rt == CL_SUCCESS, "clGetDeviceInfo DEVICE_AVAILABLE error %d\n", (int)rt);
LOG_I(HW, "Device %i is %s available\n", j, (abool == CL_TRUE ? "" : "not"));
cl_device_type devtype;
rt = clGetDeviceInfo(ocl.runtime[i].devices[j], CL_DEVICE_TYPE, sizeof(cl_device_type), &devtype, NULL);
AssertFatal(rt == CL_SUCCESS, "clGetDeviceInfo DEVICE_TYPE error %d\n", (int)rt);
LOG_I(HW, "Device %i, type %d = %s\n", j, (int)devtype, clutil_getstrdev(devtype));
rt = clGetDeviceInfo(ocl.runtime[i].devices[j],
CL_DEVICE_MAX_COMPUTE_UNITS,
sizeof(ocl.runtime[i].ocldev[j].max_CU),
&(ocl.runtime[i].ocldev[j].max_CU),
NULL);
AssertFatal(rt == CL_SUCCESS, "clGetDeviceInfo MAX_COMPUTE_UNITS error %d\n", (int)rt);
LOG_I(HW, "Device %i, number of Compute Units: %d\n", j, ocl.runtime[i].ocldev[j].max_CU);
rt = clGetDeviceInfo(ocl.runtime[i].devices[j],
CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS,
sizeof(ocl.runtime[i].ocldev[j].max_WID),
&(ocl.runtime[i].ocldev[j].max_WID),
NULL);
AssertFatal(rt == CL_SUCCESS, "clGetDeviceInfo MAX_WORK_ITEM_DIMENSIONS error %d\n", (int)rt);
LOG_I(HW, "Device %i, max Work Items dimension: %d\n", j, ocl.runtime[i].ocldev[j].max_WID);
ocl.runtime[i].ocldev[j].max_WIS = (size_t *)malloc(ocl.runtime[i].ocldev[j].max_WID * sizeof(size_t));
rt = clGetDeviceInfo(ocl.runtime[i].devices[j],
CL_DEVICE_MAX_WORK_ITEM_SIZES,
sizeof(ocl.runtime[i].ocldev[j].max_WID) * sizeof(size_t),
ocl.runtime[i].ocldev[j].max_WIS,
NULL);
AssertFatal(rt == CL_SUCCESS, "clGetDeviceInfo MAX_WORK_ITEM_SIZES error %d\n", (int)rt);
for (int k = 0; k < ocl.runtime[i].ocldev[j].max_WID; k++)
LOG_I(HW, "Device %i, max Work Items size for dimension: %d %u\n", j, k, (uint32_t)ocl.runtime[i].ocldev[j].max_WIS[k]);
devok++;
}
if (devok > 0) {
ocl.runtime[i].context = clCreateContext(NULL,
ocl.runtime[i].num_devices,
ocl.runtime[i].devices,
cl_error_callback,
&(ocl.runtime[i]),
(cl_int *)&rt);
AssertFatal(rt == CL_SUCCESS, "Error %d creating context for platform %i\n", (int)rt, i);
for (int dev = 0; dev < ocl.runtime[i].num_devices; dev++) {
ocl.runtime[i].queue[dev] =
clCreateCommandQueueWithProperties(ocl.runtime[i].context, ocl.runtime[i].devices[dev], 0, (cl_int *)&rt);
AssertFatal(rt == CL_SUCCESS, "Error %d creating command queue for platform %i device %i\n", (int)rt, i, dev);
}
ocl.runtime[i].dev_const_llr =
clCreateBuffer(ocl.runtime[i].context, CL_MEM_READ_ONLY | CL_MEM_HOST_WRITE_ONLY, 68 * 384, NULL, (cl_int *)&rt);
AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_const_llr for platform %i \n", (int)rt, i);
ocl.runtime[i].dev_llr =
clCreateBuffer(ocl.runtime[i].context, CL_MEM_READ_WRITE | CL_MEM_HOST_WRITE_ONLY, 68 * 384, NULL, (cl_int *)&rt);
AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_llr for platform %i \n", (int)rt, i);
ocl.runtime[i].dev_dt =
clCreateBuffer(ocl.runtime[i].context, CL_MEM_READ_WRITE | CL_MEM_HOST_NO_ACCESS, 46 * 68 * 384, NULL, (cl_int *)&rt);
AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_dt for platform %i \n", (int)rt, i);
ocl.runtime[i].dev_tmp =
clCreateBuffer(ocl.runtime[i].context, CL_MEM_READ_ONLY | CL_MEM_HOST_WRITE_ONLY, 68 * 384, NULL, (cl_int *)&rt);
AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_tmp for platform %i \n", (int)rt, i);
char *source_str;
size_t source_size = load_source(&source_str, "nrLDPC_decoder_kernels_CL.clc");
cl_program program = clCreateProgramWithSource(ocl.runtime[i].context,
1,
(const char **)&source_str,
(const size_t *)&source_size,
(cl_int *)&rt);
AssertFatal(rt == CL_SUCCESS, "Error %d creating program for platform %i \n", (int)rt, i);
rt = clBuildProgram(program, ocl.runtime[i].num_devices, ocl.runtime[i].devices, NULL /* compile options */, NULL, NULL);
if (rt == CL_BUILD_PROGRAM_FAILURE) {
get_CompilErr(program, i);
}
AssertFatal(rt == CL_SUCCESS, "Error %d buildding program for platform %i \n", rt, i);
for (int dev = 0; dev < ocl.runtime[i].num_devices; dev++) {
ocl.runtime[i].kernels[dev].cnp_kernel_1st = clCreateKernel(program, "ldpc_cnp_kernel_1st_iter", (cl_int *)&rt);
AssertFatal(rt == CL_SUCCESS,
"Error %d creating kernel %s platform %i, dev %i\n",
(int)rt,
"ldpc_cnp_kernel_1st_iter",
i,
dev);
ocl.runtime[i].kernels[dev].cnp_kernel = clCreateKernel(program, "ldpc_cnp_kernel", (cl_int *)&rt);
AssertFatal(rt == CL_SUCCESS, "Error %d creating kernel %s platform %i, dev %i\n", (int)rt, "ldpc_cnp_kernel", i, dev);
ocl.runtime[i].kernels[dev].vnp_kernel_normal = clCreateKernel(program, "ldpc_vnp_kernel_normal", (cl_int *)&rt);
AssertFatal(rt == CL_SUCCESS,
"Error %d creating kernel %s platform %i, dev %i\n",
(int)rt,
"ldpc_vnp_kernel_normal",
i,
dev);
ocl.runtime[i].kernels[dev].pack_decoded_bit = clCreateKernel(program, "pack_decoded_bit", (cl_int *)&rt);
AssertFatal(rt == CL_SUCCESS, "Error %d creating kernel %s platform %i, dev %i\n", (int)rt, "pack_decoded_bit", i, dev);
}
context_ok++;
}
devok = 0;
}
AssertFatal(context_ok > 0, "No openCL device available to accelerate ldpc\n");
return 0;
}
int32_t LDPCshutdown()
{
return 0;
}
int32_t LDPCinit()
{
// init_LLR_DMA(p_decParams, p_llr, p_out);
return 0;
}
int32_t LDPCdecoder(t_nrLDPC_dec_params *p_decParams,
int8_t *p_llr,
int8_t *p_out,
t_nrLDPC_time_stats *time_decoder,
decode_abort_t *ab)
{
uint16_t Zc = p_decParams->Z;
uint8_t BG = p_decParams->BG;
// uint8_t numMaxIter = p_decParams->numMaxIter;
int block_length = p_decParams->Kprime;
// e_nrLDPC_outMode outMode = p_decParams->outMode;
uint8_t row, col;
if (BG == 1) {
row = 46;
col = 68;
} else {
row = 42;
col = 52;
}
// alloc memory
// unsigned char *hard_decision = (unsigned char*)p_out;
// gpu
int memorySize_llr = col * Zc * sizeof(char) * MC;
// cudaCheck( cudaMemcpyToSymbol(dev_const_llr, p_llr, memorySize_llr_cuda) );
// cudaCheck( cudaMemcpyToSymbol(dev_llr, p_llr, memorySize_llr_cuda) );
set_compact_BG(p_decParams->Z, p_decParams->BG);
int rt =
clEnqueueWriteBuffer(ocl.runtime[0].queue[0], ocl.runtime[0].dev_const_llr, CL_TRUE, 0, memorySize_llr, p_llr, 0, NULL, NULL);
AssertFatal(rt == CL_SUCCESS, "Error %d moving p_llr data to read only memory in pltf %i dev %i\n", (int)rt, 0, 0);
rt = clEnqueueWriteBuffer(ocl.runtime[0].queue[0], ocl.runtime[0].dev_llr, CL_TRUE, 0, memorySize_llr, p_llr, 0, NULL, NULL);
AssertFatal(rt == CL_SUCCESS, "Error %d moving p_llr data to read-write memory in pltf %i dev %i\n", (int)rt, 0, 0);
// Define CUDA kernel dimension
// int blockSizeX = Zc;
// dim3 dimGridKernel1(row, MC, 1); // dim of the thread blocks
// dim3 dimBlockKernel1(blockSizeX, 1, 1);
// dim3 dimGridKernel2(col, MC, 1);
// dim3 dimBlockKernel2(blockSizeX, 1, 1);
// cudaDeviceSynchronize();
// lauch kernel
size_t global_item_sizek0[2] = {row * Zc, MC}; // Process the entire lists
size_t global_item_sizek1[2] = {col * Zc, MC}; // Process the entire lists
size_t local_item_sizek[2] = {128, 1}; // Divide work items into groups of 128
for (int ii = 0; ii < MAX_ITERATION; ii++) {
// first kernel
if (ii == 0) {
// ldpc_cnp_kernel_1st_iter
// <<<dimGridKernel1, dimBlockKernel1>>>
// ( BG, row, col, Zc);
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel_1st, 0, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_llr));
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel_1st, 1, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_dt));
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel_1st, 2, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_h_compact1));
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel_1st, 3, sizeof(int), (void *)&(BG));
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel_1st, 4, sizeof(int), (void *)&(row));
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel_1st, 5, sizeof(int), (void *)&(col));
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel_1st, 6, sizeof(int), (void *)&(Zc));
rt = clEnqueueNDRangeKernel(ocl.runtime[0].queue[0],
ocl.runtime[0].kernels[0].cnp_kernel_1st,
2,
NULL,
global_item_sizek0,
local_item_sizek,
0,
NULL,
NULL);
AssertFatal(rt == CL_SUCCESS, "Error %d enqueing cnp_kernel_1st \n", (int)rt);
} else {
// ldpc_cnp_kernel
// <<<dimGridKernel1, dimBlockKernel1>>>
// ( BG, row, col, Zc);
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 0, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_llr));
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 1, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_dt));
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 2, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_h_compact1));
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 3, sizeof(int), (void *)&(BG));
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 4, sizeof(int), (void *)&(row));
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 5, sizeof(int), (void *)&(col));
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 6, sizeof(int), (void *)&(Zc));
rt = clEnqueueNDRangeKernel(ocl.runtime[0].queue[0],
ocl.runtime[0].kernels[0].cnp_kernel,
2,
NULL,
global_item_sizek0,
local_item_sizek,
0,
NULL,
NULL);
AssertFatal(rt == CL_SUCCESS, "Error %d enqueing cnp_kernel\n", (int)rt);
}
// second kernel
// ldpc_vnp_kernel_normal
// <<<dimGridKernel2, dimBlockKernel2>>>
// // (dev_llr, dev_const_llr,BG, row, col, Zc);
// (BG, row, col, Zc);
CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 0, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_llr));
CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 1, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_dt));
CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 2, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_const_llr));
CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 3, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_h_compact2));
CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 4, sizeof(int), (void *)&(BG));
CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 5, sizeof(int), (void *)&(row));
CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 6, sizeof(int), (void *)&(col));
CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 7, sizeof(int), (void *)&(Zc));
rt = clEnqueueNDRangeKernel(ocl.runtime[0].queue[0],
ocl.runtime[0].kernels[0].vnp_kernel_normal,
2,
NULL,
global_item_sizek1,
local_item_sizek,
0,
NULL,
NULL);
AssertFatal(rt == CL_SUCCESS, "Error %d enqueing vnp_kernel_normal \n", (int)rt);
}
// int pack = (block_length/128)+1;
// dim3 pack_block(pack, MC, 1);
// pack_decoded_bit<<<pack_block,128>>>( col, Zc);
CLSETKERNELARG(ocl.runtime[0].kernels[0].pack_decoded_bit, 0, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_llr));
CLSETKERNELARG(ocl.runtime[0].kernels[0].pack_decoded_bit, 1, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_tmp));
CLSETKERNELARG(ocl.runtime[0].kernels[0].pack_decoded_bit, 2, sizeof(int), (void *)&(col));
CLSETKERNELARG(ocl.runtime[0].kernels[0].pack_decoded_bit, 3, sizeof(int), (void *)&(Zc));
// Execute the OpenCL kernel on the list
size_t global_item_size[2] = {block_length, MC}; // Process the entire lists
size_t local_item_size[2] = {128, 1}; // Divide work items into groups of 128
rt = clEnqueueNDRangeKernel(ocl.runtime[0].queue[0],
ocl.runtime[0].kernels[0].pack_decoded_bit,
2,
NULL,
global_item_size,
local_item_size,
0,
NULL,
NULL);
AssertFatal(rt == CL_SUCCESS, "Error %d enqueing pack_decoded_bit \n", (int)rt);
// cudaCheck( cudaMemcpyFromSymbol((void*)hard_decision, (const void*)dev_tmp, (block_length/8)*sizeof(unsigned char)) );
// cudaDeviceSynchronize();
rt = clEnqueueReadBuffer(ocl.runtime[0].queue[0],
ocl.runtime[0].dev_tmp,
CL_TRUE,
0,
(block_length / 8) * sizeof(unsigned char),
p_llr,
0,
NULL,
NULL);
AssertFatal(rt == CL_SUCCESS, "Error %d moving p_llr data to pltf %i dev %i\n", (int)rt, 0, 0);
return MAX_ITERATION;
}
#endif // NRLDPC_KERNEL_SOURCE
openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_kernels_CL.c
-292
View File
@@ -1,292 +0,0 @@
/*
* Licensed to the OpenAirInterface (OAI) Software Alliance under one or more
* contributor license agreements. See the NOTICE file distributed with
* this work for additional information regarding copyright ownership.
* The OpenAirInterface Software Alliance licenses this file to You under
* the OAI Public License, Version 1.0 (the "License"); you may not use this file
* except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.openairinterface.org/?page_id=698
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*-------------------------------------------------------------------------------
* For more information about the OpenAirInterface (OAI) Software Alliance:
* contact@openairinterface.org
*/
/*! \file PHY/CODING/nrLDPC_decoder_kernels_CL.cl
* \brief kernel functions for ldpc decoder accelerated via openCL
* \author Francois TABURET
* \date 2021
* \version 1.0
* \company Nokia BellLabs France
* \email: francois.taburet@nokia-bell-labs.com
* \note initial implem - translation of cuda version
* \warning
*/
//__global char dev_dt [46*68*384];
//__local char *dev_t;
//__global char dev_llr[68*384];
//__global unsigned char dev_tmp[68*384];
#define INT32_MAX 2147483647
//__constant h_element dev_h_compact1[46*19] = {}; // used in kernel 1
//__constant h_element dev_h_compact2[68*30] = {}; // used in kernel 2
// __device__ __constantant__ h_element dev_h_compact1[46*19]; // used in kernel 1
// __device__ __constantant__ h_element dev_h_compact2[68*30]; // used in kernel 2
// row and col element count
__constant char h_ele_row_bg1_count[46] = {
19, 19, 19, 19, 3, 8, 9, 7, 10, 9,
7, 8, 7, 6, 7, 7, 6, 6, 6, 6,
6, 6, 5, 5, 6, 5, 5, 4, 5, 5,
5, 5, 5, 5, 5, 5, 5, 4, 5, 5,
4, 5, 4, 5, 5, 4};
__constant char h_ele_col_bg1_count[68] = {
30, 28, 7, 11, 9, 4, 8, 12, 8, 7,
12, 10, 12, 11, 10, 7, 10, 10, 13, 7,
8, 11, 12, 5, 6, 6, 1, 1, 1, 1,
1, 1, 1, 1, 1, 1, 1, 1, 1, 1,
1, 1, 1, 1, 1, 1, 1, 1, 1, 1,
1, 1, 1, 1, 1, 1, 1, 1, 1, 1,
1, 1, 1, 1, 1, 1, 1, 1};
__constant char h_ele_row_bg2_count[42] = {
8, 10, 8, 10, 4, 6, 6, 6, 4, 5,
5, 5, 4, 5, 5, 4, 5, 5, 4, 4,
4, 4, 3, 4, 4, 3, 5, 3, 4, 3,
5, 3, 4, 4, 4, 4, 4, 3, 4, 4,
4, 4};
__constant char h_ele_col_bg2_count[52] = {
22, 23, 10, 5, 5, 14, 7, 13, 6, 8,
9, 16, 9, 12, 1, 1, 1, 1, 1, 1,
1, 1, 1, 1, 1, 1, 1, 1, 1, 1,
1, 1, 1, 1, 1, 1, 1, 1, 1, 1,
1, 1, 1, 1, 1, 1, 1, 1, 1, 1,
1, 1};
// Kernel 1
__kernel void ldpc_cnp_kernel_1st_iter( __global char * dev_llr, __global char * dev_dt, __local h_element *dev_h_compact1, int BG, int row, int col, int Zc)
{
// int iMCW = blockIdx.y; // codeword id
// int iBlkRow = blockIdx.x; // block row in h_base
// int iSubRow = threadIdx.x; // row index in sub_block of h_base
// if(blockIdx.x == 0 && threadIdx.x == 1) printf("cnp %d\n", threadIdx.x);
int iMCW = get_group_id(1); // codeword id
int iBlkRow = get_group_id(0); // block row in h_base
int iBlkCol; // block col in h_base
int iSubRow = get_local_id(0);; // row index in sub_block of h_base
int iCol; // overall col index in h_base
int offsetR;
int shift_t;
// For 2-min algorithm.
int Q_sign = 0;
int sq;
int Q, Q_abs;
int R_temp;
int sign = 1;
int rmin1 = INT32_MAX;
int rmin2 = INT32_MAX;
char idx_min = 0;
h_element h_element_t;
int s = (BG==1)? h_ele_row_bg1_count[iBlkRow]:h_ele_row_bg2_count[iBlkRow];
offsetR = (iMCW * row*col*Zc) + iBlkRow * Zc + iSubRow; // row*col*Zc = size of dev_dt
// if(blockIdx.x == 0 && threadIdx.x == 1) printf("s: %d, offset %d\n", s, offsetR);
// The 1st recursion
for(int i = 0; i < s; i++) // loop through all the ZxZ sub-blocks in a row
{
h_element_t = dev_h_compact1[i*row+iBlkRow]; // compact_col == row
iBlkCol = h_element_t.y;
shift_t = h_element_t.value;
shift_t = (iSubRow + shift_t) % Zc;
iCol = (iMCW * col*Zc) + iBlkCol * Zc + shift_t; // col*Zc = size of llr
Q = dev_llr[iCol];
Q_abs = (Q>0)? Q : -Q;
sq = Q < 0;
// if(blockIdx.x == 0 && threadIdx.x == 1) printf("i %d, icol %d, Q: %d\n", i, iCol, Q);
// quick version
sign = sign * (1 - sq * 2);
Q_sign |= sq << i;
if (Q_abs < rmin1){
rmin2 = rmin1;
rmin1 = Q_abs;
idx_min = i;
} else if (Q_abs < rmin2){
rmin2 = Q_abs;
}
}
// if(blockIdx.x == 0 && threadIdx.x == 1)printf("min1 %d, min2 %d, min1_idx %d\n", rmin1, rmin2, idx_min);
// The 2nd recursion
for(int i = 0; i < s; i++){
// v0: Best performance so far. 0.75f is the value of alpha.
sq = 1 - 2 * ((Q_sign >> i) & 0x01);
R_temp = 0.75f * sign * sq * (i != idx_min ? rmin1 : rmin2);
// write results to global memory
h_element_t = dev_h_compact1[i*row+iBlkRow];
int addr_temp = offsetR + h_element_t.y * row * Zc;
dev_dt[addr_temp] = R_temp;
// if(blockIdx.x == 0 && threadIdx.x == 1)printf("R_temp %d, temp_addr %d\n", R_temp, addr_temp);
}
}
// Kernel_1
__kernel void ldpc_cnp_kernel( __global char * dev_llr, __global char * dev_dt, __local h_element *dev_h_compact1, int BG, int row, int col, int Zc)
{
// if(blockIdx.x == 0 && threadIdx.x == 1) printf("cnp\n");
// int iMCW = blockIdx.y;
// int iBlkRow = blockIdx.x; // block row in h_base // block col in h_base
// int iSubRow = threadIdx.x; // row index in sub_block of h_base
int iMCW = get_group_id(1);
int iBlkRow = get_group_id(0); // block row in h_base
int iBlkCol; // block col in h_base
int iSubRow = get_local_id(0);; // row index in sub_block of h_base
int iCol; // overall col index in h_base
int offsetR;
int shift_t;
// For 2-min algorithm.
int Q_sign = 0;
int sq;
int Q, Q_abs;
int R_temp;
int sign = 1;
int rmin1 = INT32_MAX;
int rmin2 = INT32_MAX;
char idx_min = 0;
h_element h_element_t;
int s = (BG==1)? h_ele_row_bg1_count[iBlkRow]: h_ele_row_bg2_count[iBlkRow];
offsetR = (iMCW *row*col*Zc) + iBlkRow * Zc + iSubRow; // row * col * Zc = size of dev_dt
// if(blockIdx.x == 0 && threadIdx.x == 1) printf("s: %d, offset %d\n", s, offsetR);
// The 1st recursion
for(int i = 0; i < s; i++) // loop through all the ZxZ sub-blocks in a row
{
h_element_t = dev_h_compact1[i*row+iBlkRow];
iBlkCol = h_element_t.y;
shift_t = h_element_t.value;
shift_t = (iSubRow + shift_t) % Zc;
iCol = iBlkCol * Zc + shift_t;
R_temp = dev_dt[offsetR + iBlkCol * row * Zc];
Q = dev_llr[iMCW * (col*Zc) + iCol] - R_temp;
Q_abs = (Q>0)? Q : -Q;
// if(blockIdx.x == 0 && threadIdx.x == 1) printf("i %d, icol %d, Q: %d\n", i, iCol, Q);
sq = Q < 0;
sign = sign * (1 - sq * 2);
Q_sign |= sq << i;
if (Q_abs < rmin1){
rmin2 = rmin1;
rmin1 = Q_abs;
idx_min = i;
} else if (Q_abs < rmin2){
rmin2 = Q_abs;
}
}
// if(blockIdx.x == 0 && threadIdx.x == 1)printf("min1 %d, min2 %d, min1_idx %d\n", rmin1, rmin2, idx_min);
// The 2nd recursion
for(int i = 0; i < s; i ++){
sq = 1 - 2 * ((Q_sign >> i) & 0x01);
R_temp = 0.75f * sign * sq * (i != idx_min ? rmin1 : rmin2);
// write results to global memory
h_element_t = dev_h_compact1[i*row+iBlkRow];
int addr_temp = h_element_t.y * row * Zc + offsetR;
dev_dt[addr_temp] = R_temp;
// if(blockIdx.x == 0 && threadIdx.x == 1)printf("R_temp %d, temp_addr %d\n", R_temp, addr_temp);
}
}
// Kernel 2: VNP processing
__kernel void
ldpc_vnp_kernel_normal(__global char * dev_llr, __global char * dev_dt, __global char * dev_const_llr, __local h_element *dev_h_compact2, int BG, int row, int col, int Zc)
{
// int iMCW = blockIdx.y;
// int iBlkCol = blockIdx.x;
// int iSubCol = threadIdx.x;
int iMCW = get_group_id(1);
int iBlkCol = get_group_id(0);
int iBlkRow;
int iSubCol = get_local_id(0);
int iRow;
int iCol;
int shift_t, sf;
int APP;
h_element h_element_t;
// update all the llr values
iCol = iBlkCol * Zc + iSubCol;
APP = dev_const_llr[iMCW *col*Zc + iCol];
int offsetDt = iMCW *row*col*Zc + iBlkCol * row * Zc;
int s = (BG==1)? h_ele_col_bg1_count[iBlkCol]:h_ele_col_bg2_count[iBlkCol];
for(int i = 0; i < s; i++)
{
h_element_t = dev_h_compact2[i*col+iBlkCol];
shift_t = h_element_t.value%Zc;
iBlkRow = h_element_t.x;
sf = iSubCol - shift_t;
sf = (sf + Zc) % Zc;
iRow = iBlkRow * Zc + sf;
APP = APP + dev_dt[offsetDt + iRow];
}
if(APP > SCHAR_MAX) APP = SCHAR_MAX;
if(APP < SCHAR_MIN) APP = SCHAR_MIN;
// write back to device global memory
dev_llr[iMCW *col*Zc + iCol] = APP;
}
__kernel void pack_decoded_bit(__global unsigned char * dev_llr, __global unsigned char * dev_tmp, int col, int Zc)
{
// int iMCW = blockIdx.y;
// int btid = threadIdx.x;
unsigned char tmp[128];
int iMCW = get_group_id(1);
int btid = get_local_id(0);
int tid = iMCW * col*Zc + get_group_id(0)*128 + btid;
tmp[btid] = 0;
if(dev_llr[tid] < 0){
tmp[btid] = 1 << (7-(btid&7));
}
// __syncthreads();
if(btid < 16){
dev_tmp[iMCW * col*Zc + get_group_id(0)*16+btid] = 0;
for(int i = 0; i < 8; i++){
dev_tmp[iMCW * col*Zc + get_group_id(0)*16+btid] += tmp[btid*8+i];
}
}
}