nrLDPC_decoder_CL.c 21 KB
Newer Older
frtabu's avatar
frtabu committed
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19
/*
 * 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
20
 */
frtabu's avatar
frtabu committed
21 22
 /*! \file nrLDPC_decoder_CL.c
* \brief ldpc decoder, openCL implementaion
frtabu's avatar
frtabu committed
23
* \author Francois TABURET
frtabu's avatar
frtabu committed
24 25
* \date 2021
* \version 1.0
frtabu's avatar
frtabu committed
26
* \note initial implem - translation of cuda version
frtabu's avatar
frtabu committed
27
*/
28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44


#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
frtabu's avatar
frtabu committed
45
/* uses HW  component id for log messages ( --log_config.hw_log_level <warning| info|debug|trace>) */
46 47
#include <stdio.h>
#include <unistd.h>
48
#include <sys/stat.h>
49 50 51 52 53 54
#include <CL/opencl.h>
#include "PHY/CODING/nrLDPC_decoder/nrLDPC_types.h"
#include "PHY/CODING/nrLDPC_decoder/nrLDPCdecoder_defs.h"
#include "assertions.h"
#include "common/utils/LOG/log.h"

frtabu's avatar
frtabu committed
55 56 57 58
#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);

frtabu's avatar
frtabu committed
59
#include "../nrLDPC_decoder_LYC/bgs/BG1_compact_in_C.h"
60

frtabu's avatar
frtabu committed
61 62 63 64
typedef struct{
  cl_uint max_CU;
  cl_uint max_WID;
  size_t  *max_WIS;
65 66
} ocldev_t;

67 68 69 70 71 72 73
typedef struct {
  cl_kernel cnp_kernel_1st;
  cl_kernel cnp_kernel;
  cl_kernel vnp_kernel_normal;
  cl_kernel pack_decoded_bit;
} oclkernels_t;

74 75 76 77
typedef struct{
  cl_uint num_devices;
  cl_device_id devices[MAX_OCLDEV];
  ocldev_t ocldev[MAX_OCLDEV];
frtabu's avatar
frtabu committed
78 79
  cl_context context;
  cl_program program;
frtabu's avatar
frtabu committed
80 81
  cl_mem dev_h_compact1;
  cl_mem dev_h_compact2;
82 83 84 85 86
  cl_mem dev_const_llr;
  cl_mem dev_llr;
  cl_mem dev_dt;
  cl_mem dev_tmp;  
  oclkernels_t kernels[MAX_OCLDEV];
87 88
  cl_command_queue queue[MAX_OCLDEV];
} oclruntime_t;
89

frtabu's avatar
frtabu committed
90
typedef struct{
91
  oclruntime_t runtime[MAX_OCLRUNTIME];
frtabu's avatar
frtabu committed
92
} ocl_t;
93 94 95



frtabu's avatar
frtabu committed
96
ocl_t ocl;
97

frtabu's avatar
frtabu committed
98 99 100 101

void set_compact_BG(int Zc,short BG){
	cl_uint rt;
	int row,col;
102 103 104 105 106 107 108 109
	if(BG == 1){
		row = 46;
		col = 68;
	}
	else{
		row = 42;
		col = 52;
	}
frtabu's avatar
frtabu committed
110 111 112 113 114 115 116 117 118 119 120 121 122 123 124 125 126
	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}
	};
127
	
frtabu's avatar
frtabu committed
128 129 130 131 132 133 134 135 136 137 138 139 140 141 142 143 144 145 146 147 148 149 150 151 152 153 154 155 156 157 158 159 160 161 162 163 164 165 166 167 168 169 170 171 172 173 174 175 176 177 178 179 180 181 182 183 184
	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;
185
}
186 187
void cl_error_callback(const char* errinfo, const void* private_info, size_t cb, void* user_data) {
  oclruntime_t *runtime = (oclruntime_t *)user_data;
frtabu's avatar
frtabu committed
188 189
  LOG_E(HW,"OpenCL accelerator error  %s\n", errinfo );
}
190

frtabu's avatar
frtabu committed
191 192 193 194 195 196 197 198 199 200 201 202
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;
}
203

204 205 206 207 208 209 210 211 212 213 214 215 216 217 218 219 220
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);
    }

}

221
size_t load_source(char **source_str, char *filename) {
222
    FILE *fp;
223
    struct stat st ;
224
    size_t source_size;
225
    char *src= NULL;
226
    
227 228 229 230 231 232 233 234 235 236 237 238 239 240 241
  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;
242 243
}

244 245


246 247 248
/* from here: entry points in decoder shared lib */
int ldpc_autoinit(void) {   // called by the library loader 
  cl_platform_id platforms[10];
frtabu's avatar
frtabu committed
249
  cl_uint         num_platforms_found;
250
  int             context_ok=0;
frtabu's avatar
frtabu committed
251
  cl_uint rt = clGetPlatformIDs( sizeof(platforms)/sizeof(cl_platform_id), platforms, &num_platforms_found );
252 253 254
  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++) {
frtabu's avatar
frtabu committed
255 256 257 258 259 260 261
	  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 );	  
262
	  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));
263
	  AssertFatal(rt == CL_SUCCESS, "clGetDeviceIDs error %d\n" , (int)rt);
264 265
	  int devok=0;
	  for (int j=0; j<ocl.runtime[i].num_devices; j++) {
frtabu's avatar
frtabu committed
266
		cl_bool abool;
267
		rt = clGetDeviceInfo(ocl.runtime[i].devices[j],CL_DEVICE_AVAILABLE, sizeof(abool),&abool,NULL);
frtabu's avatar
frtabu committed
268 269 270
		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;		
271
		rt = clGetDeviceInfo(ocl.runtime[i].devices[j],CL_DEVICE_TYPE, sizeof(cl_device_type),&devtype,NULL);
frtabu's avatar
frtabu committed
272 273
		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));
274
		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);
frtabu's avatar
frtabu committed
275
		AssertFatal(rt == CL_SUCCESS, "clGetDeviceInfo MAX_COMPUTE_UNITS error %d\n" , (int)rt);
276 277
		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);
frtabu's avatar
frtabu committed
278
		AssertFatal(rt == CL_SUCCESS, "clGetDeviceInfo MAX_WORK_ITEM_DIMENSIONS error %d\n" , (int)rt);
279 280 281
		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);
frtabu's avatar
frtabu committed
282
		AssertFatal(rt == CL_SUCCESS, "clGetDeviceInfo MAX_WORK_ITEM_SIZES error %d\n" , (int)rt);
283 284 285 286 287 288 289 290 291 292 293
		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);
        }
294 295 296 297 298 299 300 301
        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);      
302
        char *source_str;
303
        size_t source_size=load_source(&source_str,"nrLDPC_decoder_kernels_CL.clc");      
304 305 306
        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); 
307
        rt = clBuildProgram(program, ocl.runtime[i].num_devices,ocl.runtime[i].devices, NULL /* compile options */, NULL, NULL);  
308
        if (rt == CL_BUILD_PROGRAM_FAILURE) {
309 310 311 312 313 314 315 316 317 318 319 320 321 322
          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); 		  		  		  
        }
323
        context_ok++;
324
      }
325
      devok=0;
326
  }
327
  AssertFatal(context_ok>0, "No openCL device available to accelerate ldpc\n"); 
328 329 330 331 332
  return 0;  
}


void nrLDPC_initcall(t_nrLDPC_dec_params* p_decParams, int8_t* p_llr, int8_t* p_out) {
frtabu's avatar
frtabu committed
333 334
	set_compact_BG(p_decParams->Z,p_decParams->BG);
//	init_LLR_DMA(p_decParams, p_llr,  p_out);
335 336
}

337 338 339 340 341 342
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)
343 344 345
{
    uint16_t Zc          = p_decParams->Z;
    uint8_t  BG         = p_decParams->BG;
frtabu's avatar
frtabu committed
346
//    uint8_t  numMaxIter = p_decParams->numMaxIter;
347
    int block_length    = p_decParams->block_length;
frtabu's avatar
frtabu committed
348
//    e_nrLDPC_outMode outMode = p_decParams->outMode;
349 350 351 352 353 354 355 356 357 358 359
	uint8_t row,col;
	if(BG == 1){
		row = 46;
		col = 68;
	}
	else{
		row = 42;
		col = 52;
	}

//	alloc memory
frtabu's avatar
frtabu committed
360
//	unsigned char *hard_decision = (unsigned char*)p_out;
361
//	gpu
362
	int memorySize_llr = col * Zc * sizeof(char) * MC;
363 364
//	cudaCheck( cudaMemcpyToSymbol(dev_const_llr, p_llr, memorySize_llr_cuda) );
//	cudaCheck( cudaMemcpyToSymbol(dev_llr, p_llr, memorySize_llr_cuda) );
365 366 367 368 369 370
    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); 	
371
// Define CUDA kernel dimension
frtabu's avatar
frtabu committed
372
//	int blockSizeX = Zc;
373 374 375 376 377 378 379 380
//	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 
frtabu's avatar
frtabu committed
381 382 383 384
    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
    
385 386 387
	for(int ii = 0; ii < MAX_ITERATION; ii++){
		// first kernel	
		if(ii == 0){
frtabu's avatar
frtabu committed
388 389 390 391 392 393 394 395 396 397 398 399 400
//			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);  
401
		}else{
frtabu's avatar
frtabu committed
402 403 404 405 406 407 408 409 410 411 412 413 414
//		 	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);  
415 416
		}
		// second kernel
frtabu's avatar
frtabu committed
417 418 419 420 421 422 423 424 425 426 427 428 429 430 431
//		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);  
432
	}
frtabu's avatar
frtabu committed
433 434
	
//   int pack = (block_length/128)+1;
435 436
//	dim3 pack_block(pack, MC, 1);
//	pack_decoded_bit<<<pack_block,128>>>( col, Zc);
frtabu's avatar
frtabu committed
437 438 439 440
	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));
441 442
 
    // Execute the OpenCL kernel on the list
frtabu's avatar
frtabu committed
443 444
    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
445
    rt = clEnqueueNDRangeKernel(ocl.runtime[0].queue[0], ocl.runtime[0].kernels[0].pack_decoded_bit, 2, NULL, 
frtabu's avatar
frtabu committed
446 447
            global_item_size, local_item_size, 0, NULL, NULL);
	AssertFatal(rt == CL_SUCCESS, "Error %d enqueing pack_decoded_bit \n" , (int)rt);             
448 449 450 451 452 453
//	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); 
454 455 456
	return MAX_ITERATION;
	
}
457
#endif //NRLDPC_KERNEL_SOURCE