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

Refactor LDPC API to make common API with various implementations

Browse Source
This commit is contained in:
Laurent THOMAS
2023-11-28 09:37:37 +01:00
committed by Raymond Knopp
parent e945b260f8
commit 99ca87a1c1
22 changed files with 622 additions and 1299 deletions
Download Patch File Download Diff File Expand all files Collapse all files
openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c
+24 -18
View File
@@ -47,8 +47,7 @@ typedef struct{
#include <unistd.h>
#include <sys/stat.h>
#include <CL/opencl.h>
#include "PHY/CODING/nrLDPC_decoder/nrLDPC_types.h"
#include "PHY/CODING/nrLDPC_decoder/nrLDPCdecoder_defs.h"
#include "openair1/PHY/CODING/nrLDPC_extern.h"
#include "assertions.h"
#include "common/utils/LOG/log.h"
@@ -135,9 +134,9 @@ void set_compact_BG(int Zc,short BG){
}
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);
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);
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){
@@ -328,18 +327,23 @@ int ldpc_autoinit(void) { // called by the library loader
return 0;
}
void nrLDPC_initcall(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out) {
set_compact_BG(p_decParams->Z,p_decParams->BG);
// init_LLR_DMA(p_decParams, p_llr, p_out);
int32_t LDPCshutdown()
{
return 0;
}
int32_t nrLDPC_decod(t_nrLDPC_dec_params *p_decParams,
int8_t *p_llr,
int8_t *p_out,
t_nrLDPC_procBuf *p_procBuf,
t_nrLDPC_time_stats *time_decoder,
decode_abort_t *ab)
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_procBuf *p_procBuf,
t_nrLDPC_time_stats *time_decoder,
decode_abort_t *ab)
{
uint16_t Zc = p_decParams->Z;
uint8_t BG = p_decParams->BG;
@@ -362,12 +366,14 @@ int32_t nrLDPC_decod(t_nrLDPC_dec_params *p_decParams,
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) );
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);
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);
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