/* * 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, uint8_t harq_pid, uint8_t ulsch_id, uint8_t C, 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->block_length; // 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