Exemplo n.º 1
0
int main(int argc, char ** argv) {
	if(argc != 3){
		printf("wrong option... \n");
		return 0;
	}
    char *needleData = argv[1];
    int needleLen = strlen(needleData);
    int cccharsPerItem;
    int ggcharsPerItem;
    int characterSetSize = 128;//ASCii set
    int *skipTable; // skipTable
    char *needle = &needleData[0];//!!!Pass it in a wrong way
    skipTable = QSPrecomputation(needle, characterSetSize); 

    char *cckernelFile = "./kernel/QS_CPU_LocalCounter_OpenCL_Kernel.cl";//the kernel for CPU
    char *ggkernelFile = "./kernel/QS_GPU_LocalMem_LocalCounter_OpenCL_Kernel.cl";//the kernel for GPU
    char *cckernelSrc = LoadKernelSrcFromFile(cckernelFile);
    char *ggkernelSrc = LoadKernelSrcFromFile(ggkernelFile);
    char *fileName = argv[2];
    
    //Load the haystacks from file
    FILE *filePtr;
    filePtr = fopen(fileName, "r");
    if(!filePtr) {
        fprintf(stderr, "can not open the text file!\n");
    }
    fseek(filePtr, 0 , SEEK_END);
    int fileSize = ftell(filePtr);
    int ggfileSize = fileSize*0.5;
    int ccfileSize = fileSize - ggfileSize;
    rewind(filePtr);
    //Split the file then I need overlaping!
    char *gghaystackData = (char*)calloc(ggfileSize, sizeof(char));
    char *cchaystackData = (char*)calloc(ccfileSize, sizeof(char));
    int textLength;
    textLength = fread(gghaystackData, sizeof(char), ggfileSize, filePtr);
    textLength += fread(cchaystackData, sizeof(char), ccfileSize, filePtr);
    if(textLength != fileSize) {
        fprintf(stderr, "reading error");
    }
    fclose(filePtr);
    //~ char *cchaystackData = LoadHaystackDataFromFile(fileName);
    int cchaystackLen = strlen(cchaystackData);
    int gghaystackLen = strlen(gghaystackData);


    cl_int err;
    cl_platform_id platform;
    cl_device_id cdevice;
    cl_device_id ggdevice;

    //Find a platform
    err = clGetPlatformIDs(1, &platform, NULL);
    if(err != CL_SUCCESS) {
        printf("cant find a platform! \n");
    }
    
    //Set up devices
    err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_CPU, 1, &cdevice, NULL);
    if(err != CL_SUCCESS) {
        printf("cant find CPU! \n");
    }
    err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &ggdevice, NULL);
    if(err != CL_SUCCESS) {
        printf("cant find GPU! \n");
    }
    
    //Create context
    cl_context ccontext = clCreateContext(NULL, 1, &cdevice, NULL, NULL, &err);
    if(err != CL_SUCCESS) {
        printf("cant create a ccontext! \n");
    }
    cl_context ggcontext = clCreateContext(NULL, 1, &ggdevice, NULL, NULL, &err);
    if(err != CL_SUCCESS) {
        printf("cant create a ggcontext! \n");
    }
    
    //Create command queues
    cl_command_queue cqueue = clCreateCommandQueue(ccontext, cdevice, 0, &err);
    if(err != CL_SUCCESS) {
        printf("cant create a  cqueue! \n");
    }
    cl_command_queue ggqueue = clCreateCommandQueue(ggcontext, ggdevice, 0, &err);
    if(err != CL_SUCCESS) {
        printf("cant create a  ggqueue! \n");
    }

    
    //Create the programm object
    cl_program cprogram = clCreateProgramWithSource(ccontext, 1, (const char**)&cckernelSrc, NULL, &err);
    if(err != CL_SUCCESS) {
        printf("cant build the program! \n");
    }
    cl_program ggprogram = clCreateProgramWithSource(ggcontext, 1, (const char**)&ggkernelSrc, NULL, &err);
    if(err != CL_SUCCESS) {
        printf("cant build the program! \n");
    }

    
    //Build the programm executable
    err = clBuildProgram(cprogram, 0, NULL, NULL, NULL, NULL);
    if(err != CL_SUCCESS) {
        printf("cant build the cprogramm exe! \n");
    }
    err = clBuildProgram(ggprogram, 0, NULL, NULL, NULL, NULL);
    if(err != CL_SUCCESS) {
        printf("cant build the ggprogramm exe! \n");
    }

    
    //Create the kernel
    cl_kernel ckernel = clCreateKernel(cprogram, "QSMatch", &err);
    if(err != CL_SUCCESS) {
        printf("cant create the ckernel! \n");
    }
    cl_kernel ggkernel = clCreateKernel(ggprogram, "QSMatch", &err);
    if(err != CL_SUCCESS) {
        printf("cant create the ggkernel! \n");
    }

    
    //Create the haystack buffer
    cl_mem cchaystackBuffer = clCreateBuffer(ccontext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, cchaystackLen, cchaystackData, &err);
    if(err != CL_SUCCESS) {
        printf("couldn't create the cchaystackBuffer \n");
    }
    cl_mem gghaystackBuffer = clCreateBuffer(ggcontext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, gghaystackLen, gghaystackData, &err);
    if(err != CL_SUCCESS) {
        printf("couldn't create the gghaystackBuffer \n");
    }

	//Determine global size and local size
    size_t cclocalSize = 64;
    size_t gglocalSize = 256;
    size_t ccGlobalSize = 64*4;
    size_t ggGlobalSize = 256*5*8;
    int gpuGroupNum = (int)ggGlobalSize/gglocalSize;
    cccharsPerItem = cchaystackLen/ccGlobalSize + 1;//Add 1 it important otherwise some patterns will be lost
    ggcharsPerItem = cchaystackLen/ggGlobalSize + 1;
  
    //Create the results buffer
    int ccres[4] = {0};
    int ggres[gpuGroupNum];
    int k;
    for(k=0;k<gpuGroupNum;k++){
		ggres[k] = 0;
	}
    
    cl_mem ccresBuffer = clCreateBuffer(ccontext, CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(ccres), ccres, &err);
    if(err != CL_SUCCESS){
		printf("couldn't create the ccresBuffer");
	}
	cl_mem ggresBuffer = clCreateBuffer(ggcontext, CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(ggres), ggres, &err);
    if(err != CL_SUCCESS){
		printf("couldn't create the ggresBuffer");
	}
	
    //Create the needleBuffer
    cl_mem ccneedleBuffer = clCreateBuffer(ccontext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(char)*needleLen, needleData, &err);
    if(err != CL_SUCCESS){
		printf("couldn't create the ccneedleBuffer \n");
	}
	cl_mem ggneedleBuffer = clCreateBuffer(ggcontext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(char)*needleLen, needleData, &err);
    if(err != CL_SUCCESS){
		printf("couldn't create the ggneedleBuffer \n");
	}
	
    //Create the skipTableBuffer
    cl_mem ccskipTableBuffer = clCreateBuffer(ccontext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int)*characterSetSize, skipTable, &err);
    if(err != CL_SUCCESS){
		printf("couldn't create the ccskipTableBuffer \n");
	}
	cl_mem ggskipTableBuffer = clCreateBuffer(ggcontext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int)*characterSetSize, skipTable, &err);
    if(err != CL_SUCCESS){
		printf("couldn't create the ggskipTableBuffer \n");
	}
	
    //Set the kernel arguments for CPU
    err = clSetKernelArg(ckernel, 0, sizeof(cl_mem), &cchaystackBuffer);
    err |= clSetKernelArg(ckernel, 1, sizeof(cl_mem), &ccneedleBuffer);
    err |= clSetKernelArg(ckernel, 2, sizeof(cl_mem), &ccskipTableBuffer);
    err |= clSetKernelArg(ckernel, 3, sizeof(int)*1, NULL);
    err |= clSetKernelArg(ckernel, 4, sizeof(needleLen), &needleLen);
    err |= clSetKernelArg(ckernel, 5, sizeof(cccharsPerItem), &cccharsPerItem);
    err |= clSetKernelArg(ckernel, 6, sizeof(cl_mem), &ccresBuffer);
    if(err != CL_SUCCESS) {
        printf("couldn't set the ckernel arguments \n");
    }
    
    //Set the kernel arguments for GPU
    err = clSetKernelArg(ggkernel, 0, sizeof(cl_mem), &gghaystackBuffer);
    err |= clSetKernelArg(ggkernel, 1, sizeof(cl_mem), &ggneedleBuffer);
    err |= clSetKernelArg(ggkernel, 2, sizeof(char)*needleLen, NULL);
    err |= clSetKernelArg(ggkernel, 3, sizeof(cl_mem), &ggskipTableBuffer);
    err |= clSetKernelArg(ggkernel, 4, sizeof(int)*128, NULL);
    err |= clSetKernelArg(ggkernel, 5, sizeof(int)*1, NULL);
    err |= clSetKernelArg(ggkernel, 6, sizeof(needleLen), &needleLen);
    err |= clSetKernelArg(ggkernel, 7, sizeof(ggcharsPerItem), &ggcharsPerItem);
    err |= clSetKernelArg(ggkernel, 8, sizeof(cl_mem), &ggresBuffer);
    if(err != CL_SUCCESS) {
        printf("couldn't set the ggkernel arguments \n");
    }
    
    //Execut the kernel
    //On CPU
    err = clEnqueueNDRangeKernel(cqueue, ckernel, 1, NULL, &ccGlobalSize, &cclocalSize, 0, NULL, NULL);
    if(err != CL_SUCCESS) {
        printf("ckernel could not be executed... \n");
    } 
    else {
        printf("ckernel has been executed successfully! \n");
    }
      
    //~ //On GPU
    err = clEnqueueNDRangeKernel(ggqueue, ggkernel, 1, NULL, &ggGlobalSize, &gglocalSize, 0, NULL, NULL);
    if(err != CL_SUCCESS) {
        printf("ggkernel could not be executed... \n");
    } 
    else {
        printf("ggkernel has been executed successfully! \n");
    }
    
    //Finish the queue lists
    clFinish(ggqueue);
    clFinish(cqueue);
    

    //Copy the results
    clEnqueueReadBuffer(cqueue, ccresBuffer, CL_TRUE, 0, sizeof(ccres), ccres, 0, NULL, NULL);
    clEnqueueReadBuffer(ggqueue, ggresBuffer, CL_TRUE, 0, sizeof(ggres), ggres, 0, NULL, NULL);
    PrintTheResults(ccres, ggres, gpuGroupNum);
    
    //Clean up
    free(skipTable);
    free(cckernelSrc);
    free(ggkernelSrc);
    free(cchaystackData);
    free(gghaystackData);
    clReleaseMemObject(cchaystackBuffer);
    clReleaseMemObject(ccneedleBuffer);
    clReleaseMemObject(ccskipTableBuffer);
    clReleaseMemObject(ccresBuffer);
    clReleaseMemObject(gghaystackBuffer);
    clReleaseMemObject(ggneedleBuffer);
    clReleaseMemObject(ggskipTableBuffer);
    clReleaseMemObject(ggresBuffer);
    clReleaseKernel(ckernel);
    clReleaseKernel(ggkernel);
    clReleaseProgram(cprogram);
    clReleaseProgram(ggprogram);
    clReleaseCommandQueue(cqueue);
    clReleaseCommandQueue(ggqueue);
    clReleaseContext(ccontext);
    clReleaseContext(ggcontext);

    return 0;
}
Exemplo n.º 2
0
int main(int argc, char * argv[]) {
	if(argc < 3){
		printf("wrong option... \n");
		return 0;
	}
	int t;// Reusing t can improve the performance somehow
	int needleNum = argc - 2;//Calculate the number of needle
    string needles[needleNum]; 
    for(t=0;t<needleNum;t++){
		needles[t] = argv[t+1];
	} 
	int cpuNeedleNum = needleNum*0.5;
	int gpuNeedleNum = needleNum - cpuNeedleNum;
	if(cpuNeedleNum + gpuNeedleNum == needleNum){
		printf("the needles number is right \n");
	}
	string cpuNeedles[cpuNeedleNum];
	string gpuNeedles[gpuNeedleNum];
	for(t=0;t<cpuNeedleNum;t++){
		cpuNeedles[t] = needles[t];
	}
	for(t=0;t<gpuNeedleNum;t++){
		gpuNeedles[t] = needles[t+cpuNeedleNum];
	}
	
    //Concatenating the needles into one
    char *cpuNeedlesData = ConcatenateNeedlesIntoOneDataBlock(cpuNeedles, cpuNeedleNum);
    printf("cpuNeedlesData = %s \n", cpuNeedlesData);
    char *gpuNeedlesData = ConcatenateNeedlesIntoOneDataBlock(gpuNeedles, gpuNeedleNum);
    printf("gpuNeedlesData = %s \n", gpuNeedlesData);
    int cpuLenOfneedlesData = strlen(cpuNeedlesData);
    int gpuLenOfneedlesData = strlen(gpuNeedlesData);
    int cpuShortestNeedleLen = findShortestNeedle(cpuNeedles,cpuNeedleNum);
    int gpuShortestNeedleLen = findShortestNeedle(gpuNeedles,gpuNeedleNum);
    int cpuLongestNeedleLen = findLongestNeedle(cpuNeedles,cpuNeedleNum);
    int gpuLongestNeedleLen = findLongestNeedle(gpuNeedles,gpuNeedleNum);
   
    int characterSetSize = 128;//ASCii set
    //~ int characterSetSize = 256; //UTF-8
    int *cpuSkipTable; //CPU skipTable
    int *gpuSkipTable; //GPU skipTable
    cpuSkipTable = SetHorspoolPrecomputation(cpuNeedles, characterSetSize, cpuShortestNeedleLen, cpuNeedleNum);
    gpuSkipTable = SetHorspoolPrecomputation(gpuNeedles, characterSetSize, gpuShortestNeedleLen, gpuNeedleNum);
    
    //Calculating the lastPosOfEachNeedle array and lenOfEachNeedle array
    int cpuLastPosOfEachNeedle[cpuNeedleNum];
    int gpuLastPosOfEachNeedle[gpuNeedleNum];
    int cpuLenOfEachNeedle[cpuNeedleNum]; 
    int gpuLenOfEachNeedle[gpuNeedleNum]; 
    for(t=0;t<cpuNeedleNum;t++){
		cpuLenOfEachNeedle[t] = strlen(cpuNeedles[t]); 
		cpuLastPosOfEachNeedle[t] = strlen(cpuNeedles[t]) - 1;
	}
	for(t=0;t<gpuNeedleNum;t++){
		gpuLenOfEachNeedle[t] = strlen(gpuNeedles[t]); 
		gpuLastPosOfEachNeedle[t] = strlen(gpuNeedles[t]) - 1;
	}

    char *cpuKernelFile = "./kernel/SH_CPU_Kernel.cl";//The kernel file's name
    char *gpuKernelFile = "./kernel/SH_GPU_LocalMem_OpenCL_Kernel.cl";//The kernel file's name
    char *cpuKernelSrc = LoadKernelSrcFromFile(cpuKernelFile);
    char *gpuKernelSrc = LoadKernelSrcFromFile(gpuKernelFile);
    char *textFile = argv[argc-1];// The text file's name
    char *haystackData = LoadHaystackDataFromFile(textFile);
    int haystackLen = strlen(haystackData);

    cl_int err;
    cl_platform_id platform;
    cl_device_id cpuDevice;
    cl_device_id gpuDevice;

    //Find a platform
    err = clGetPlatformIDs(1, &platform, NULL);
    if(err != CL_SUCCESS) {
        printf("cant find a platform! \n");
    }
    
    //Set up devices
    err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_CPU, 1, &cpuDevice, NULL);
    if(err != CL_SUCCESS) {
        printf("cant find a cpu device! \n");
    }
     err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &gpuDevice, NULL);
    if(err != CL_SUCCESS) {
        printf("cant find a gpu device! \n");
    }
    //Create context
    cl_context cpuContext = clCreateContext(NULL, 1, &cpuDevice, NULL, NULL, &err);
    if(err != CL_SUCCESS) {
        printf("cant create a cpu context! \n");
    }
    cl_context gpuContext = clCreateContext(NULL, 1, &gpuDevice, NULL, NULL, &err);
    if(err != CL_SUCCESS) {
        printf("cant create a gpu context! \n");
    }
    
    //Create command queue
    cl_command_queue cpuQueue = clCreateCommandQueue(cpuContext, cpuDevice, 0, &err);
    if(err != CL_SUCCESS) {
        printf("cant create a cpu queue! \n");
    }
    cl_command_queue gpuQueue = clCreateCommandQueue(gpuContext, gpuDevice, 0, &err);
    if(err != CL_SUCCESS) {
        printf("cant create a gpu queue! \n");
    }
    
    //Create the programm object
    cl_program cpuProgram = clCreateProgramWithSource(cpuContext, 1, (const char**)&cpuKernelSrc, NULL, &err);
    if(err != CL_SUCCESS) {
        printf("cant build the cpu program! \n");
    }
    cl_program gpuProgram = clCreateProgramWithSource(gpuContext, 1, (const char**)&gpuKernelSrc, NULL, &err);
    if(err != CL_SUCCESS) {
        printf("cant build the gpu program! \n");
    }
    
    //Build the programm executable
    err = clBuildProgram(cpuProgram, 0, NULL, NULL, NULL, NULL);
    if(err != CL_SUCCESS) {
        printf("cant build the cpu programm exe! \n");
    }
    err = clBuildProgram(gpuProgram, 0, NULL, NULL, NULL, NULL);
    if(err != CL_SUCCESS) {
        printf("cant build the gpu programm exe! \n");
    }
    
    //Create the kernel
    cl_kernel cpuKernel = clCreateKernel(cpuProgram, "SetHorspoolMatch", &err);
    if(err != CL_SUCCESS) {
        printf("cant create the cpu kernel! \n");
    }
    cl_kernel gpuKernel = clCreateKernel(gpuProgram, "SetHorspoolMatch", &err);
    if(err != CL_SUCCESS) {
        printf("cant create the gpu kernel! \n");
    }
    
    //Create the haystack buffer 
    //CPU and GPU program share the haystack buffer
    cl_mem cpuHaystackBuffer = clCreateBuffer(cpuContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, haystackLen, haystackData, &err);
    if(err != CL_SUCCESS) {
        printf("couldn't create the cpuHaystackData buffer \n");
    }
    cl_mem gpuHaystackBuffer = clCreateBuffer(cpuContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, haystackLen, haystackData, &err);
    if(err != CL_SUCCESS) {
        printf("couldn't create the gpuHaystackData buffer \n");
    }
    
    size_t cclocalSize;
    size_t gglocalSize;
    cclocalSize = 64;
    gglocalSize = 256;
    size_t cpuGlobalSize = 64*4;//CPU
    size_t gpuGlobalSize = 256*24*8;//GPU
    //~ size_t gpuGlobalSize = 256*5*8;//GPU
    //Group Number
    int cpuGroupNum = (int)cpuGlobalSize/cclocalSize;
    int gpuGroupNum = (int)gpuGlobalSize/gglocalSize;
    //CPU GPU share the charsPerItem
    int cpuCharsPerItem = haystackLen/cpuGlobalSize + 1;
    int gpuCharsPerItem = haystackLen/gpuGlobalSize + 1;
    
    //Create the results buffer
    int cpuRes[cpuNeedleNum*cpuGroupNum];
    int gpuRes[gpuNeedleNum*gpuGroupNum];
    for(t=0;t<cpuNeedleNum*cpuGroupNum;t++){
		cpuRes[t] = 0;
	}
	for(t=0;t<gpuNeedleNum*gpuGroupNum;t++){
		gpuRes[t] = 0;
	}
    cl_mem cpuResBuffer = clCreateBuffer(cpuContext, CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(cpuRes), cpuRes, &err);//Should be write only
    if(err != CL_SUCCESS) {
        printf("cant create the cpuResBuffer! \n");
    }
    cl_mem gpuResBuffer = clCreateBuffer(gpuContext, CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(gpuRes), gpuRes, &err);//Should be write only
    if(err != CL_SUCCESS) {
        printf("cant create the gpuResBuffer! \n");
    }
        
    //Create the needlesBuffer 
    cl_mem cpuNeedlesBuffer = clCreateBuffer(cpuContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(char)*cpuLenOfneedlesData, cpuNeedlesData, &err);//!!cpuNeedlesLen
    if(err != CL_SUCCESS) {
        printf("cant create the cpuNeedlesBuffer! \n");
    }
    cl_mem gpuNeedlesBuffer = clCreateBuffer(gpuContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(char)*gpuLenOfneedlesData, gpuNeedlesData, &err);
    if(err != CL_SUCCESS) {
        printf("cant create the gpuNeedlesBuffer! \n");
    }
    
    //Create the skipTableBuffer
    //The skipTables of CPU and CPU are different
    cl_mem cpuSkipTableBuffer = clCreateBuffer(cpuContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int)*characterSetSize, cpuSkipTable, &err);
    if(err != CL_SUCCESS) {
        printf("cant create the cpuSkipTableBuffer! \n");
    }
    cl_mem gpuSkipTableBuffer = clCreateBuffer(gpuContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int)*characterSetSize, gpuSkipTable, &err);
    if(err != CL_SUCCESS) {
        printf("cant create the gpuSkipTableBuffer! \n");
    }
    
    //Create the lastPosOfEachNeedleBuffer
    cl_mem cpuLastPosOfEachNeedleBuffer = clCreateBuffer(cpuContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int)*cpuNeedleNum, cpuLastPosOfEachNeedle, &err);//!!lastPosOfEachNeedle
    if(err != CL_SUCCESS) {
        printf("cant create the cpuLastPosOfEachNeedleBuffer! \n");
    }
    cl_mem gpuLastPosOfEachNeedleBuffer = clCreateBuffer(gpuContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int)*gpuNeedleNum, gpuLastPosOfEachNeedle, &err);//!!gpuNeedleNum
    if(err != CL_SUCCESS) {
        printf("cant create the gpuLastPosOfEachNeedleBuffer! \n");
    }
    
    //Create the lenOfEachNeedleBuffer
    cl_mem cpuLenOfEachNeedleBuffer = clCreateBuffer(cpuContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int)*cpuNeedleNum, cpuLenOfEachNeedle, &err);
    if(err != CL_SUCCESS) {
        printf("cant create the cpuLenOfEachNeedleBuffer! \n");
    }
    cl_mem gpuLenOfEachNeedleBuffer = clCreateBuffer(gpuContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int)*gpuNeedleNum, gpuLenOfEachNeedle, &err);
    if(err != CL_SUCCESS) {
        printf("cant create the gpuLenOfEachNeedleBuffer! \n");
    }
    
    //Set the kernel arguments for CPU
    err = clSetKernelArg(cpuKernel, 0, sizeof(cl_mem), &cpuHaystackBuffer);
    err |= clSetKernelArg(cpuKernel, 1, sizeof(cl_mem), &cpuNeedlesBuffer);
    err |= clSetKernelArg(cpuKernel, 2, sizeof(cl_mem), &cpuSkipTableBuffer);
    err |= clSetKernelArg(cpuKernel, 3, sizeof(cpuLongestNeedleLen), &cpuLongestNeedleLen);
    err |= clSetKernelArg(cpuKernel, 4, sizeof(cpuCharsPerItem), &cpuCharsPerItem);
    err |= clSetKernelArg(cpuKernel, 5, sizeof(cpuNeedleNum), &cpuNeedleNum);
    err |= clSetKernelArg(cpuKernel, 6, sizeof(int)*needleNum, NULL);//localCounter
    err |= clSetKernelArg(cpuKernel, 7, sizeof(cl_mem), &cpuLastPosOfEachNeedleBuffer);
    err |= clSetKernelArg(cpuKernel, 8, sizeof(cl_mem), &cpuLenOfEachNeedleBuffer);
    err |= clSetKernelArg(cpuKernel, 9, sizeof(cl_mem), &cpuResBuffer);
    if(err != CL_SUCCESS) {
        printf("couldn't set the cpu cpuKernel arguments \n");
    }
    
    //Set the kernel arguments for GPU
    err = clSetKernelArg(gpuKernel, 0, sizeof(cl_mem), &gpuHaystackBuffer);
    err |= clSetKernelArg(gpuKernel, 1, sizeof(cl_mem), &gpuNeedlesBuffer);
    err |= clSetKernelArg(gpuKernel, 2, sizeof(int), &gpuLenOfneedlesData);
    err |= clSetKernelArg(gpuKernel, 3, sizeof(char)*gpuLenOfneedlesData, NULL);//Local needlesData
    err |= clSetKernelArg(gpuKernel, 4, sizeof(cl_mem), &gpuSkipTableBuffer);
    err |= clSetKernelArg(gpuKernel, 5, sizeof(int)*128, NULL);//Local tempTable
    err |= clSetKernelArg(gpuKernel, 6, sizeof(int), &gpuLongestNeedleLen);
    err |= clSetKernelArg(gpuKernel, 7, sizeof(int), &gpuCharsPerItem);
    err |= clSetKernelArg(gpuKernel, 8, sizeof(int), &gpuNeedleNum);
    err |= clSetKernelArg(gpuKernel, 9, sizeof(int)*gpuNeedleNum, NULL);
    err |= clSetKernelArg(gpuKernel, 10, sizeof(cl_mem), &gpuLastPosOfEachNeedleBuffer);
    //~ err |= clSetKernelArg(gpuKernel, 11, sizeof(int)*gpuNeedleNum, NULL);//localLPOEN
    err |= clSetKernelArg(gpuKernel, 11, sizeof(cl_mem), &gpuLenOfEachNeedleBuffer);
    //~ err |= clSetKernelArg(gpuKernel, 13, sizeof(cl_mem), NULL);//localLEN
    err |= clSetKernelArg(gpuKernel, 12, sizeof(cl_mem), &gpuResBuffer);
    if(err != CL_SUCCESS) {
        printf("couldn't set the gpu kernel arguments \n");
    }
    
    //Execut the kernel
    //Exe on CPU
    err = clEnqueueNDRangeKernel(cpuQueue, cpuKernel, 1, NULL, &cpuGlobalSize, &cclocalSize, 0, NULL, NULL);
    if(err != CL_SUCCESS) {
        printf("CPU kernel could not be executed... \n");
    } else {
        printf("CPU kernel has been executed successfully! \n");
    }
    //Exe on GPU
    err = clEnqueueNDRangeKernel(gpuQueue, gpuKernel, 1, NULL, &gpuGlobalSize, &gglocalSize, 0, NULL, NULL);
    if(err != CL_SUCCESS) {
        printf("gPU kernel could not be executed... \n");
    } else {
        printf("gPU kernel has been executed successfully! \n");
    }
  
    //Finish the queue list
    clFinish(gpuQueue);
    clFinish(cpuQueue);
    
    //Copy the results
    clEnqueueReadBuffer(cpuQueue, cpuResBuffer, CL_TRUE, 0, sizeof(cpuRes), cpuRes, 0, NULL, NULL);
    clEnqueueReadBuffer(gpuQueue, gpuResBuffer, CL_TRUE, 0, sizeof(gpuRes), gpuRes, 0, NULL, NULL);
    PrintTheResults(cpuRes, cpuNeedles, cpuNeedleNum, cpuGroupNum);
    PrintTheResults(gpuRes, gpuNeedles, gpuNeedleNum, gpuGroupNum);

    //Clean up
    free(cpuSkipTable); 
    free(gpuSkipTable); 
    free(cpuKernelSrc);
    free(gpuKernelSrc);
    free(haystackData);
    free(cpuNeedlesData);
    free(gpuNeedlesData);
    clReleaseMemObject(cpuHaystackBuffer);
    clReleaseMemObject(gpuHaystackBuffer);
    clReleaseMemObject(cpuNeedlesBuffer);
    clReleaseMemObject(gpuNeedlesBuffer);
    clReleaseMemObject(cpuSkipTableBuffer);
    clReleaseMemObject(gpuSkipTableBuffer);
    clReleaseMemObject(cpuLenOfEachNeedleBuffer);
    clReleaseMemObject(gpuLenOfEachNeedleBuffer);
    clReleaseMemObject(cpuLastPosOfEachNeedleBuffer);
    clReleaseMemObject(gpuLastPosOfEachNeedleBuffer);
    clReleaseMemObject(cpuResBuffer);
    clReleaseMemObject(gpuResBuffer);
    clReleaseKernel(cpuKernel);
    clReleaseKernel(gpuKernel);
    clReleaseProgram(cpuProgram);
    clReleaseProgram(gpuProgram);
    clReleaseCommandQueue(cpuQueue);
    clReleaseCommandQueue(gpuQueue);
    clReleaseContext(cpuContext);
    clReleaseContext(gpuContext);

    return 0;
}