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

fix(nrLDPC_decoder_CL): Apply clang-format

Browse Source
This commit is contained in:
Romain Beurdouche
2025-08-08 14:48:41 +00:00
parent 37bf1deb7d
commit a2c03f26e9
1 changed files with 416 additions and 330 deletions
Download Patch File Download Diff File Expand all files Collapse all files
openair1/PHY/CODING/nrLDPC_decoder/nrLDPC_decoder_CL.c
+416 -330
View File
@@ -18,22 +18,21 @@
* For more information about the OpenAirInterface (OAI) Software Alliance: * For more information about the OpenAirInterface (OAI) Software Alliance:
* contact@openairinterface.org * contact@openairinterface.org
*/ */
/*! \file nrLDPC_decoder_CL.c /*! \file nrLDPC_decoder_CL.c
* \brief ldpc decoder, openCL implementaion * \brief ldpc decoder, openCL implementaion
* \author Francois TABURET * \author Francois TABURET
* \date 2021 * \date 2021
* \version 1.0 * \version 1.0
* \note initial implem - translation of cuda version * \note initial implem - translation of cuda version
*/ */
#define MAX_ITERATION 2 #define MAX_ITERATION 2
#define MC 1 #define MC 1
#define MAX_OCLDEV 10 #define MAX_OCLDEV 10
#define MAX_OCLRUNTIME 5 #define MAX_OCLRUNTIME 5
typedef struct{ typedef struct {
char x; char x;
char y; char y;
short value; short value;
@@ -51,16 +50,16 @@ typedef struct{
#include "assertions.h" #include "assertions.h"
#include "common/utils/LOG/log.h" #include "common/utils/LOG/log.h"
#define CLSETKERNELARG(A,B,C,D) \ #define CLSETKERNELARG(A, B, C, D) \
rt=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); 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" #include "../nrLDPC_decoder_LYC/bgs/BG1_compact_in_C.h"
typedef struct{ typedef struct {
cl_uint max_CU; cl_uint max_CU;
cl_uint max_WID; cl_uint max_WID;
size_t *max_WIS; size_t *max_WIS;
} ocldev_t; } ocldev_t;
typedef struct { typedef struct {
@@ -70,7 +69,7 @@ typedef struct {
cl_kernel pack_decoded_bit; cl_kernel pack_decoded_bit;
} oclkernels_t; } oclkernels_t;
typedef struct{ typedef struct {
cl_uint num_devices; cl_uint num_devices;
cl_device_id devices[MAX_OCLDEV]; cl_device_id devices[MAX_OCLDEV];
ocldev_t ocldev[MAX_OCLDEV]; ocldev_t ocldev[MAX_OCLDEV];
@@ -81,250 +80,305 @@ typedef struct{
cl_mem dev_const_llr; cl_mem dev_const_llr;
cl_mem dev_llr; cl_mem dev_llr;
cl_mem dev_dt; cl_mem dev_dt;
cl_mem dev_tmp; cl_mem dev_tmp;
oclkernels_t kernels[MAX_OCLDEV]; oclkernels_t kernels[MAX_OCLDEV];
cl_command_queue queue[MAX_OCLDEV]; cl_command_queue queue[MAX_OCLDEV];
} oclruntime_t; } oclruntime_t;
typedef struct{ typedef struct {
oclruntime_t runtime[MAX_OCLRUNTIME]; oclruntime_t runtime[MAX_OCLRUNTIME];
} ocl_t; } ocl_t;
ocl_t ocl; 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){ for (int i = 0; lift_set[i][0] != 0; i++) {
cl_uint rt; for (int j = 0; lift_set[i][j] != 0; j++) {
int row,col; if (Zc == lift_set[i][j]) {
if(BG == 1){ lift_index = i;
row = 46; break;
col = 68; }
} }
else{ }
row = 42; printf("\nZc = %d BG = %d\n", Zc, BG);
col = 52; 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);
int compact_row = 30; AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_h_compact1 for platform %i \n", (int)rt, 0);
int compact_col = 19; ocl.runtime[0].dev_h_compact2 =
if(BG==2){compact_row = 10, compact_col = 23;} clCreateBuffer(ocl.runtime[0].context, CL_MEM_READ_ONLY | CL_MEM_HOST_WRITE_ONLY, memorySize_h_compact2, NULL, (cl_int *)&rt);
int memorySize_h_compact1 = row * compact_col * sizeof(h_element); AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_h_compact2 for platform %i \n", (int)rt, 0);
int memorySize_h_compact2 = compact_row * col * sizeof(h_element); h_element *h1;
int lift_index = 0; h_element *h2;
short lift_set[][9] = { switch (lift_index) {
{2,4,8,16,32,64,128,256}, case 0:
{3,6,12,24,48,96,192,384}, h1 = host_h_compact1_I0;
{5,10,20,40,80,160,320}, h2 = host_h_compact2_I0;
{7,14,28,56,112,224}, break;
{9,18,36,72,144,288}, case 1:
{11,22,44,88,176,352}, h1 = host_h_compact1_I1;
{13,26,52,104,208}, h2 = host_h_compact2_I1;
{15,30,60,120,240}, break;
{0} case 2:
}; h1 = host_h_compact1_I2;
h2 = host_h_compact2_I2;
for(int i = 0; lift_set[i][0] != 0; i++){ break;
for(int j = 0; lift_set[i][j] != 0; j++){ case 3:
if(Zc == lift_set[i][j]){ h1 = host_h_compact1_I3;
lift_index = i; h2 = host_h_compact2_I3;
break; break;
} case 4:
} h1 = host_h_compact1_I4;
} h2 = host_h_compact2_I4;
printf("\nZc = %d BG = %d\n",Zc,BG); break;
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); case 5:
AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_h_compact1 for platform %i \n" , (int)rt, 0); h1 = host_h_compact1_I5;
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); h2 = host_h_compact2_I5;
AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_h_compact2 for platform %i \n" , (int)rt, 0); break;
h_element *h1; case 6:
h_element *h2; h1 = host_h_compact1_I6;
switch(lift_index){ h2 = host_h_compact2_I6;
case 0: break;
h1 = host_h_compact1_I0; case 7:
h2 = host_h_compact2_I0; h1 = host_h_compact1_I7;
break; h2 = host_h_compact2_I7;
case 1: break;
h1 = host_h_compact1_I1; default:
h2 = host_h_compact2_I1; AssertFatal(0, "Invalid lift_index value %i\n", lift_index);
break; break;
case 2: }
h1 = host_h_compact1_I2; rt = clEnqueueWriteBuffer(ocl.runtime[0].queue[0],
h2 = host_h_compact2_I2; ocl.runtime[0].dev_h_compact1,
break; CL_TRUE,
case 3: 0,
h1 = host_h_compact1_I3; memorySize_h_compact1,
h2 = host_h_compact2_I3; h1,
break; 0,
case 4: NULL,
h1 = host_h_compact1_I4; NULL);
h2 = host_h_compact2_I4; AssertFatal(rt == CL_SUCCESS, "Error %d moving h_compact1 memory to pltf %i dev %i\n", (int)rt, 0, 0);
break; rt = clEnqueueWriteBuffer(ocl.runtime[0].queue[0],
case 5: ocl.runtime[0].dev_h_compact2,
h1 = host_h_compact1_I5; CL_TRUE,
h2 = host_h_compact2_I5; 0,
break; memorySize_h_compact2,
case 6: h2,
h1 = host_h_compact1_I6; 0,
h2 = host_h_compact2_I6; NULL,
break; NULL);
case 7: AssertFatal(rt == CL_SUCCESS, "Error %d moving h_compact2 memory to pltf %i dev %i\n", (int)rt, 0, 0);
h1 = host_h_compact1_I7; // return 0;
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; 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) { char *clutil_getstrdev(int intdev)
static char retstring[255]=""; {
char *retptr=retstring; static char retstring[255] = "";
retptr+=sprintf(retptr,"0x%08x: ",(uint32_t)intdev); char *retptr = retstring;
retptr += sprintf(retptr, "0x%08x: ", (uint32_t)intdev);
if (intdev & CL_DEVICE_TYPE_CPU) if (intdev & CL_DEVICE_TYPE_CPU)
retptr+=sprintf(retptr,"%s","cpu "); retptr += sprintf(retptr, "%s", "cpu ");
if (intdev & CL_DEVICE_TYPE_GPU) if (intdev & CL_DEVICE_TYPE_GPU)
retptr+=sprintf(retptr,"%s","gpu "); retptr += sprintf(retptr, "%s", "gpu ");
if (intdev & CL_DEVICE_TYPE_ACCELERATOR) if (intdev & CL_DEVICE_TYPE_ACCELERATOR)
retptr+=sprintf(retptr,"%s","acc "); retptr += sprintf(retptr, "%s", "acc ");
return retstring; return retstring;
} }
void get_CompilErr(cl_program program, int pltf) { void get_CompilErr(cl_program program, int pltf)
{
// Determine the size of the log // Determine the size of the log
size_t log_size; size_t log_size;
for(int i=0; i<ocl.runtime[pltf].num_devices;i++) { 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); clGetProgramBuildInfo(program, ocl.runtime[pltf].devices[i], CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
// Allocate memory for the log // Allocate memory for the log
char *log = (char *) malloc(log_size); char *log = (char *)malloc(log_size);
// Get the log // Get the log
clGetProgramBuildInfo(program, ocl.runtime[pltf].devices[i], CL_PROGRAM_BUILD_LOG, log_size, log, NULL); clGetProgramBuildInfo(program, ocl.runtime[pltf].devices[i], CL_PROGRAM_BUILD_LOG, log_size, log, NULL);
// Print the log // Print the log
printf("%s\n", log); printf("%s\n", log);
free(log); free(log);
} }
} }
size_t load_source(char **source_str, char *filename) { size_t load_source(char **source_str, char *filename)
FILE *fp; {
struct stat st ; FILE *fp;
size_t source_size; struct stat st;
char *src= NULL; size_t source_size;
char *src = NULL;
if (filename == NULL) { if (filename == NULL) {
src = "nrLDPC_decoder_kernels_CL.clc"; src = "nrLDPC_decoder_kernels_CL.clc";
} else { } else {
src = filename; src = filename;
} }
fp = fopen(src, "r"); fp = fopen(src, "r");
AssertFatal(fp,"failed to open cl source %s: %s\n",src,strerror(errno)); AssertFatal(fp, "failed to open cl source %s: %s\n", src, strerror(errno));
fstat(fileno(fp), &st); fstat(fileno(fp), &st);
source_size = st.st_size; source_size = st.st_size;
*source_str = (char*)malloc(source_size); *source_str = (char *)malloc(source_size);
source_size = fread( *source_str, 1, source_size, fp); source_size = fread(*source_str, 1, source_size, fp);
fclose( fp ); fclose(fp);
LOG_I(HW,"Loaded kernel sources from %s %u bytes\n", (filename==NULL)?"embedded cl code":src,(unsigned int)source_size ); LOG_I(HW, "Loaded kernel sources from %s %u bytes\n", (filename == NULL) ? "embedded cl code" : src, (unsigned int)source_size);
return source_size; return source_size;
} }
/* from here: entry points in decoder shared lib */ /* from here: entry points in decoder shared lib */
int ldpc_autoinit(void) { // called by the library loader int ldpc_autoinit(void)
{ // called by the library loader
cl_platform_id platforms[10]; cl_platform_id platforms[10];
cl_uint num_platforms_found; cl_uint num_platforms_found;
int context_ok=0; int context_ok = 0;
cl_uint rt = clGetPlatformIDs( sizeof(platforms)/sizeof(cl_platform_id), platforms, &num_platforms_found ); 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(rt == CL_SUCCESS, "clGetPlatformIDs error %d\n", (int)rt);
AssertFatal( num_platforms_found>0 , "clGetPlatformIDs: no cl compatible platform found\n"); AssertFatal(num_platforms_found > 0, "clGetPlatformIDs: no cl compatible platform found\n");
for (int i=0 ; i<(int)num_platforms_found ; i++) { for (int i = 0; i < (int)num_platforms_found; i++) {
char stringval[255]; char stringval[255];
rt = clGetPlatformInfo(platforms[i],CL_PLATFORM_PROFILE, sizeof(stringval),stringval,NULL); rt = clGetPlatformInfo(platforms[i], CL_PLATFORM_PROFILE, sizeof(stringval), stringval, NULL);
AssertFatal(rt == CL_SUCCESS, "clGetPlatformInfo PROFILE error %d\n" , (int)rt); AssertFatal(rt == CL_SUCCESS, "clGetPlatformInfo PROFILE error %d\n", (int)rt);
LOG_I(HW,"Platform %i, OpenCL profile %s\n", i,stringval ); LOG_I(HW, "Platform %i, OpenCL profile %s\n", i, stringval);
rt = clGetPlatformInfo(platforms[i],CL_PLATFORM_VERSION, sizeof(stringval),stringval,NULL); rt = clGetPlatformInfo(platforms[i], CL_PLATFORM_VERSION, sizeof(stringval), stringval, NULL);
AssertFatal(rt == CL_SUCCESS, "clGetPlatformInfo VERSION error %d\n" , (int)rt); AssertFatal(rt == CL_SUCCESS, "clGetPlatformInfo VERSION error %d\n", (int)rt);
LOG_I(HW,"Platform %i, OpenCL version %s\n", i,stringval ); 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)); rt = clGetDeviceIDs(platforms[i],
AssertFatal(rt == CL_SUCCESS, "clGetDeviceIDs error %d\n" , (int)rt); CL_DEVICE_TYPE_ALL,
int devok=0; sizeof(ocl.runtime[i].devices) / sizeof(cl_device_id),
for (int j=0; j<ocl.runtime[i].num_devices; j++) { ocl.runtime[i].devices,
cl_bool abool; &(ocl.runtime[i].num_devices));
rt = clGetDeviceInfo(ocl.runtime[i].devices[j],CL_DEVICE_AVAILABLE, sizeof(abool),&abool,NULL); AssertFatal(rt == CL_SUCCESS, "clGetDeviceIDs error %d\n", (int)rt);
AssertFatal(rt == CL_SUCCESS, "clGetDeviceInfo DEVICE_AVAILABLE error %d\n" , (int)rt); int devok = 0;
LOG_I(HW,"Device %i is %s available\n", j, (abool==CL_TRUE?"":"not")); for (int j = 0; j < ocl.runtime[i].num_devices; j++) {
cl_device_type devtype; cl_bool abool;
rt = clGetDeviceInfo(ocl.runtime[i].devices[j],CL_DEVICE_TYPE, sizeof(cl_device_type),&devtype,NULL); rt = clGetDeviceInfo(ocl.runtime[i].devices[j], CL_DEVICE_AVAILABLE, sizeof(abool), &abool, NULL);
AssertFatal(rt == CL_SUCCESS, "clGetDeviceInfo DEVICE_TYPE error %d\n" , (int)rt); AssertFatal(rt == CL_SUCCESS, "clGetDeviceInfo DEVICE_AVAILABLE error %d\n", (int)rt);
LOG_I(HW,"Device %i, type %d = %s\n", j,(int)devtype, clutil_getstrdev(devtype)); LOG_I(HW, "Device %i is %s available\n", j, (abool == CL_TRUE ? "" : "not"));
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); cl_device_type devtype;
AssertFatal(rt == CL_SUCCESS, "clGetDeviceInfo MAX_COMPUTE_UNITS error %d\n" , (int)rt); rt = clGetDeviceInfo(ocl.runtime[i].devices[j], CL_DEVICE_TYPE, sizeof(cl_device_type), &devtype, NULL);
LOG_I(HW,"Device %i, number of Compute Units: %d\n", j,ocl.runtime[i].ocldev[j].max_CU); AssertFatal(rt == CL_SUCCESS, "clGetDeviceInfo DEVICE_TYPE error %d\n", (int)rt);
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); LOG_I(HW, "Device %i, type %d = %s\n", j, (int)devtype, clutil_getstrdev(devtype));
AssertFatal(rt == CL_SUCCESS, "clGetDeviceInfo MAX_WORK_ITEM_DIMENSIONS error %d\n" , (int)rt); rt = clGetDeviceInfo(ocl.runtime[i].devices[j],
LOG_I(HW,"Device %i, max Work Items dimension: %d\n", j,ocl.runtime[i].ocldev[j].max_WID); CL_DEVICE_MAX_COMPUTE_UNITS,
ocl.runtime[i].ocldev[j].max_WIS = (size_t *)malloc(ocl.runtime[i].ocldev[j].max_WID * sizeof(size_t)); sizeof(ocl.runtime[i].ocldev[j].max_CU),
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); &(ocl.runtime[i].ocldev[j].max_CU),
AssertFatal(rt == CL_SUCCESS, "clGetDeviceInfo MAX_WORK_ITEM_SIZES error %d\n" , (int)rt); NULL);
for(int k=0; k<ocl.runtime[i].ocldev[j].max_WID;k++) AssertFatal(rt == CL_SUCCESS, "clGetDeviceInfo MAX_COMPUTE_UNITS error %d\n", (int)rt);
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]); LOG_I(HW, "Device %i, number of Compute Units: %d\n", j, ocl.runtime[i].ocldev[j].max_CU);
devok++; 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].dev_const_llr =
ocl.runtime[i].context = clCreateContext(NULL, ocl.runtime[i].num_devices, ocl.runtime[i].devices, cl_error_callback, &(ocl.runtime[i]), (cl_int *)&rt); 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 context for platform %i\n" , (int)rt, i); AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_const_llr for platform %i \n", (int)rt, i);
for(int dev=0; dev<ocl.runtime[i].num_devices; dev++) { ocl.runtime[i].dev_llr =
ocl.runtime[i].queue[dev] = clCreateCommandQueueWithProperties(ocl.runtime[i].context,ocl.runtime[i].devices[dev] , 0, (cl_int *)&rt); 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 command queue for platform %i device %i\n" , (int)rt, i,dev); AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_llr for platform %i \n", (int)rt, i);
} ocl.runtime[i].dev_dt =
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); 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_const_llr for platform %i \n" , (int)rt, i); AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_dt 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); ocl.runtime[i].dev_tmp =
AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_llr for platform %i \n" , (int)rt, i); clCreateBuffer(ocl.runtime[i].context, CL_MEM_READ_ONLY | CL_MEM_HOST_WRITE_ONLY, 68 * 384, NULL, (cl_int *)&rt);
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_tmp for platform %i \n", (int)rt, i);
AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_dt for platform %i \n" , (int)rt, i); char *source_str;
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); size_t source_size = load_source(&source_str, "nrLDPC_decoder_kernels_CL.clc");
AssertFatal(rt == CL_SUCCESS, "Error %d creating buffer dev_tmp for platform %i \n" , (int)rt, i); cl_program program = clCreateProgramWithSource(ocl.runtime[i].context,
char *source_str; 1,
size_t source_size=load_source(&source_str,"nrLDPC_decoder_kernels_CL.clc"); (const char **)&source_str,
cl_program program = clCreateProgramWithSource(ocl.runtime[i].context, 1, (const size_t *)&source_size,
(const char **)&source_str, (const size_t *)&source_size, (cl_int *)&rt); (cl_int *)&rt);
AssertFatal(rt == CL_SUCCESS, "Error %d creating program for platform %i \n" , (int)rt, i); 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); rt = clBuildProgram(program, ocl.runtime[i].num_devices, ocl.runtime[i].devices, NULL /* compile options */, NULL, NULL);
if (rt == CL_BUILD_PROGRAM_FAILURE) { if (rt == CL_BUILD_PROGRAM_FAILURE) {
get_CompilErr(program,i); 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(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"); AssertFatal(context_ok > 0, "No openCL device available to accelerate ldpc\n");
return 0; return 0;
} }
int32_t LDPCshutdown() int32_t LDPCshutdown()
@@ -344,119 +398,151 @@ int32_t LDPCdecoder(t_nrLDPC_dec_params *p_decParams,
t_nrLDPC_time_stats *time_decoder, t_nrLDPC_time_stats *time_decoder,
decode_abort_t *ab) decode_abort_t *ab)
{ {
uint16_t Zc = p_decParams->Z; uint16_t Zc = p_decParams->Z;
uint8_t BG = p_decParams->BG; uint8_t BG = p_decParams->BG;
// uint8_t numMaxIter = p_decParams->numMaxIter; // uint8_t numMaxIter = p_decParams->numMaxIter;
int block_length = p_decParams->Kprime; int block_length = p_decParams->Kprime;
// e_nrLDPC_outMode outMode = p_decParams->outMode; // e_nrLDPC_outMode outMode = p_decParams->outMode;
uint8_t row,col; uint8_t row, col;
if(BG == 1){ if (BG == 1) {
row = 46; row = 46;
col = 68; col = 68;
} } else {
else{ row = 42;
row = 42; col = 52;
col = 52; }
}
// alloc memory // alloc memory
// unsigned char *hard_decision = (unsigned char*)p_out; // unsigned char *hard_decision = (unsigned char*)p_out;
// gpu // gpu
int memorySize_llr = col * Zc * sizeof(char) * MC; int memorySize_llr = col * Zc * sizeof(char) * MC;
// cudaCheck( cudaMemcpyToSymbol(dev_const_llr, p_llr, memorySize_llr_cuda) ); // cudaCheck( cudaMemcpyToSymbol(dev_const_llr, p_llr, memorySize_llr_cuda) );
// cudaCheck( cudaMemcpyToSymbol(dev_llr, p_llr, memorySize_llr_cuda) ); // cudaCheck( cudaMemcpyToSymbol(dev_llr, p_llr, memorySize_llr_cuda) );
set_compact_BG(p_decParams->Z, p_decParams->BG); set_compact_BG(p_decParams->Z, p_decParams->BG);
int rt = int rt =
clEnqueueWriteBuffer(ocl.runtime[0].queue[0], ocl.runtime[0].dev_const_llr, CL_TRUE, 0, memorySize_llr, p_llr, 0, NULL, NULL); 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); 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, rt = clEnqueueWriteBuffer(ocl.runtime[0].queue[0], ocl.runtime[0].dev_llr, CL_TRUE, 0, memorySize_llr, p_llr, 0, NULL, NULL);
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
// Define CUDA kernel dimension // int blockSizeX = Zc;
// int blockSizeX = Zc; // dim3 dimGridKernel1(row, MC, 1); // dim of the thread blocks
// dim3 dimGridKernel1(row, MC, 1); // dim of the thread blocks // dim3 dimBlockKernel1(blockSizeX, 1, 1);
// dim3 dimBlockKernel1(blockSizeX, 1, 1);
// dim3 dimGridKernel2(col, MC, 1); // dim3 dimGridKernel2(col, MC, 1);
// dim3 dimBlockKernel2(blockSizeX, 1, 1); // dim3 dimBlockKernel2(blockSizeX, 1, 1);
// cudaDeviceSynchronize(); // cudaDeviceSynchronize();
// lauch kernel // lauch kernel
size_t global_item_sizek0[2] = {row*Zc,MC}; // Process the entire lists 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 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 size_t local_item_sizek[2] = {128, 1}; // Divide work items into groups of 128
for(int ii = 0; ii < MAX_ITERATION; ii++){ for (int ii = 0; ii < MAX_ITERATION; ii++) {
// first kernel // first kernel
if(ii == 0){ if (ii == 0) {
// ldpc_cnp_kernel_1st_iter // ldpc_cnp_kernel_1st_iter
// <<<dimGridKernel1, dimBlockKernel1>>> // <<<dimGridKernel1, dimBlockKernel1>>>
// ( BG, row, col, Zc); // ( 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, 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, 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, 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, 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, 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, 5, sizeof(int), (void *)&(col));
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel_1st, 6, sizeof(int), (void *)&(Zc)); 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, rt = clEnqueueNDRangeKernel(ocl.runtime[0].queue[0],
global_item_sizek0, local_item_sizek, 0, NULL, NULL); ocl.runtime[0].kernels[0].cnp_kernel_1st,
AssertFatal(rt == CL_SUCCESS, "Error %d enqueing cnp_kernel_1st \n" , (int)rt); 2,
}else{ NULL,
// ldpc_cnp_kernel global_item_sizek0,
// <<<dimGridKernel1, dimBlockKernel1>>> local_item_sizek,
// ( BG, row, col, Zc); 0,
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 0, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_llr)); NULL,
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 1, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_dt)); NULL);
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 2, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_h_compact1)); AssertFatal(rt == CL_SUCCESS, "Error %d enqueing cnp_kernel_1st \n", (int)rt);
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 3, sizeof(int), (void *)&(BG)); } else {
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 4, sizeof(int), (void *)&(row)); // ldpc_cnp_kernel
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 5, sizeof(int), (void *)&(col)); // <<<dimGridKernel1, dimBlockKernel1>>>
CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 6, sizeof(int), (void *)&(Zc)); // ( BG, row, col, Zc);
rt = clEnqueueNDRangeKernel(ocl.runtime[0].queue[0], ocl.runtime[0].kernels[0].cnp_kernel, 2, NULL, CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 0, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_llr));
global_item_sizek0, local_item_sizek, 0, NULL, NULL); CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 1, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_dt));
AssertFatal(rt == CL_SUCCESS, "Error %d enqueing cnp_kernel\n" , (int)rt); 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));
// second kernel CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 4, sizeof(int), (void *)&(row));
// ldpc_vnp_kernel_normal CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 5, sizeof(int), (void *)&(col));
// <<<dimGridKernel2, dimBlockKernel2>>> CLSETKERNELARG(ocl.runtime[0].kernels[0].cnp_kernel, 6, sizeof(int), (void *)&(Zc));
// // (dev_llr, dev_const_llr,BG, row, col, Zc); rt = clEnqueueNDRangeKernel(ocl.runtime[0].queue[0],
// (BG, row, col, Zc); ocl.runtime[0].kernels[0].cnp_kernel,
CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 0, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_llr)); 2,
CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 1, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_dt)); NULL,
CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 2, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_const_llr)); global_item_sizek0,
CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 3, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_h_compact2)); local_item_sizek,
CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 4, sizeof(int), (void *)&(BG)); 0,
CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 5, sizeof(int), (void *)&(row)); NULL,
CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 6, sizeof(int), (void *)&(col)); NULL);
CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 7, sizeof(int), (void *)&(Zc)); AssertFatal(rt == CL_SUCCESS, "Error %d enqueing cnp_kernel\n", (int)rt);
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); // second kernel
AssertFatal(rt == CL_SUCCESS, "Error %d enqueing vnp_kernel_normal \n" , (int)rt); // ldpc_vnp_kernel_normal
} // <<<dimGridKernel2, dimBlockKernel2>>>
// // (dev_llr, dev_const_llr,BG, row, col, Zc);
// int pack = (block_length/128)+1; // (BG, row, col, Zc);
// dim3 pack_block(pack, MC, 1); CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 0, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_llr));
// pack_decoded_bit<<<pack_block,128>>>( col, Zc); 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].pack_decoded_bit, 0, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_llr)); 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].pack_decoded_bit, 1, sizeof(cl_mem), (void *)&(ocl.runtime[0].dev_tmp)); 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].pack_decoded_bit, 2, sizeof(int), (void *)&(col)); CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 4, sizeof(int), (void *)&(BG));
CLSETKERNELARG(ocl.runtime[0].kernels[0].pack_decoded_bit, 3, sizeof(int), (void *)&(Zc)); 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));
// Execute the OpenCL kernel on the list CLSETKERNELARG(ocl.runtime[0].kernels[0].vnp_kernel_normal, 7, sizeof(int), (void *)&(Zc));
size_t global_item_size[2] = {block_length,MC}; // Process the entire lists rt = clEnqueueNDRangeKernel(ocl.runtime[0].queue[0],
size_t local_item_size[2] = {128,1}; // Divide work items into groups of 128 ocl.runtime[0].kernels[0].vnp_kernel_normal,
rt = clEnqueueNDRangeKernel(ocl.runtime[0].queue[0], ocl.runtime[0].kernels[0].pack_decoded_bit, 2, NULL, 2,
global_item_size, local_item_size, 0, NULL, NULL); NULL,
AssertFatal(rt == CL_SUCCESS, "Error %d enqueing pack_decoded_bit \n" , (int)rt); global_item_sizek1,
// cudaCheck( cudaMemcpyFromSymbol((void*)hard_decision, (const void*)dev_tmp, (block_length/8)*sizeof(unsigned char)) ); local_item_sizek,
// cudaDeviceSynchronize(); 0,
NULL,
rt = clEnqueueReadBuffer(ocl.runtime[0].queue[0], ocl.runtime[0].dev_tmp, CL_TRUE, 0, NULL);
(block_length/8)*sizeof(unsigned char) , p_llr, 0, NULL, NULL); AssertFatal(rt == CL_SUCCESS, "Error %d enqueing vnp_kernel_normal \n", (int)rt);
AssertFatal(rt == CL_SUCCESS, "Error %d moving p_llr data to pltf %i dev %i\n" , (int)rt, 0,0); }
return MAX_ITERATION;
// 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 #endif // NRLDPC_KERNEL_SOURCE