From a2c03f26e9999e0bfc560b79993aaba5af96fd96 Mon Sep 17 00:00:00 2001 From: Romain Beurdouche Date: Thu, 7 Aug 2025 09:50:20 +0000 Subject: [PATCH] fix(nrLDPC_decoder_CL): Apply clang-format --- .../CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c | 746 ++++++++++-------- 1 file changed, 416 insertions(+), 330 deletions(-) diff --git a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c index 49e3c80791..c8d83479a7 100644 --- a/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c +++ b/openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c @@ -18,22 +18,21 @@ * 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 -*/ - +/*! \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 MC 1 -#define MAX_OCLDEV 10 +#define MAX_OCLDEV 10 #define MAX_OCLRUNTIME 5 -typedef struct{ +typedef struct { char x; char y; short value; @@ -51,16 +50,16 @@ typedef struct{ #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); +#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{ +typedef struct { cl_uint max_CU; cl_uint max_WID; - size_t *max_WIS; + size_t *max_WIS; } ocldev_t; typedef struct { @@ -70,7 +69,7 @@ typedef struct { cl_kernel pack_decoded_bit; } oclkernels_t; -typedef struct{ +typedef struct { cl_uint num_devices; cl_device_id devices[MAX_OCLDEV]; ocldev_t ocldev[MAX_OCLDEV]; @@ -81,250 +80,305 @@ typedef struct{ cl_mem dev_const_llr; cl_mem dev_llr; cl_mem dev_dt; - cl_mem dev_tmp; + cl_mem dev_tmp; oclkernels_t kernels[MAX_OCLDEV]; cl_command_queue queue[MAX_OCLDEV]; } oclruntime_t; -typedef struct{ +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}}; -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; + 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) { +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 ); + 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); +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 "); + retptr += sprintf(retptr, "%s", "cpu "); if (intdev & CL_DEVICE_TYPE_GPU) - retptr+=sprintf(retptr,"%s","gpu "); + retptr += sprintf(retptr, "%s", "gpu "); if (intdev & CL_DEVICE_TYPE_ACCELERATOR) - retptr+=sprintf(retptr,"%s","acc "); + 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; i0 , "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 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); } - 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; dev0, "No openCL device available to accelerate ldpc\n"); - return 0; + AssertFatal(context_ok > 0, "No openCL device available to accelerate ldpc\n"); + return 0; } int32_t LDPCshutdown() @@ -344,119 +398,151 @@ int32_t LDPCdecoder(t_nrLDPC_dec_params *p_decParams, 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; - } + 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) ); + // 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); + 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(); + // 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 -// <<>> -// ( 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 -// <<>> -// ( 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 -// <<>> -// // (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<<>>( 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; - + // 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 + // <<>> + // ( 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 + // <<>> + // ( 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 + // <<>> + // // (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<<>>( 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 +#endif // NRLDPC_KERNEL_SOURCE