text: /* ********************************************************************************************* */
  /* * C Template for Kernel Execution                                                           * */
  /* * Author: André Bannwart Perina                                                             * */
  /* ********************************************************************************************* */
  /* * Copyright (c) 2017 André B. Perina                                                        * */
  /* *                                                                                           * */
  /* * Permission is hereby granted, free of charge, to any person obtaining a copy of this      * */
  /* * software and associated documentation files (the "Software"), to deal in the Software     * */
  /* * without restriction, including without limitation the rights to use, copy, modify,        * */
  /* * merge, publish, distribute, sublicense, and/or sell copies of the Software, and to        * */
  /* * permit persons to whom the Software is furnished to do so, subject to the following       * */
  /* * conditions:                                                                               * */
  /* *                                                                                           * */
  /* * The above copyright notice and this permission notice shall be included in all copies     * */
  /* * or substantial portions of the Software.                                                  * */
  /* *                                                                                           * */
  /* * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IMPLIED,       * */
  /* * INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS FOR A PARTICULAR  * */
  /* * PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE * */
  /* * FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR      * */
  /* * OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER    * */
  /* * DEALINGS IN THE SOFTWARE.                                                                 * */
  /* ********************************************************************************************* */
  
  #include <CL/opencl.h>
  #include <errno.h>
  #include <stdbool.h>
  #include <stdio.h>
  #include <string.h>
  #include <sys/time.h>
  
  #include "common.h"
  
  /**
   * @brief Header where pre/postamble macro functions should be located.
   *        Function headers:
   *            PREAMBLE(s, sSz, c, cSz, cC, cCSz, w, wSz, wC, wCSz, loopCount);
   *            POSTAMBLE(s, sSz, c, cSz, cC, cCSz, w, wSz, wC, wCSz, loopCount);
   *            LOOPPREAMBLE(s, sSz, c, cSz, cC, cCSz, w, wSz, wC, wCSz, loopCount, loopFlag);
   *            LOOPPOSTAMBLE(s, sSz, c, cSz, cC, cCSz, w, wSz, wC, wCSz, loopCount, loopFlag);
   *            CLEANUP(s, sSz, c, cSz, cC, cCSz, w, wSz, wC, wCSz, loopCount);
   *        where:
   *            s: variable (short *);
   *            sSz: number of members in variable (unsigned int);
   *            c: variable (short *);
   *            cSz: number of members in variable (unsigned int);
   *            cC: variable (short *);
   *            cCSz: number of members in variable (unsigned int);
   *            w: variable (short *);
   *            wSz: number of members in variable (unsigned int);
   *            wC: variable (short *);
   *            wCSz: number of members in variable (unsigned int);
   *            loopCount: variable (unsigned char);
   *            loopFlag: loop condition variable (bool).
   */
  #include "prepostambles.h"
  
  /**
   * @brief Test if two operands are outside an epsilon range.
   *
   * @param a First operand.
   * @param b Second operand.
   * @param e Epsilon value.
   */
  #define TEST_EPSILON(a, b, e) (((a > b) && (a - b > e)) || ((b >= a) && (b - a > e)))
  
  /**
   * @brief Standard statements for function error handling and printing.
   *
   * @param funcName Function name that failed.
   */
  #define FUNCTION_ERROR_STATEMENTS(funcName) {\
  	rv = EXIT_FAILURE;\
  	PRINT_FAIL();\
  	fprintf(stderr, "Error: %s failed with return code %d.\n", funcName, fRet);\
  }
  
  /**
   * @brief Standard statements for POSIX error handling and printing.
   *
   * @param arg Arbitrary string to the printed at the end of error string.
   */
  #define POSIX_ERROR_STATEMENTS(arg) {\
  	rv = EXIT_FAILURE;\
  	PRINT_FAIL();\
  	fprintf(stderr, "Error: %s: %s\n", strerror(errno), arg);\
  }
  
  int main(void) {
  	/* Return variable */
  	int rv = EXIT_SUCCESS;
  
  	/* OpenCL and aux variables */
  	int i = 0, j = 0;
  	cl_int platformsLen, devicesLen, fRet;
  	cl_platform_id *platforms = NULL;
  	cl_device_id *devices = NULL;
  	cl_context context = NULL;
  	cl_command_queue queueBerlekamp = NULL;
  	FILE *programFile = NULL;
  	long programSz;
  	char *programContent = NULL;
  	cl_int programRet;
  	cl_program program = NULL;
  	cl_kernel kernelBerlekamp = NULL;
  	bool loopFlag = false;
  	bool invalidDataFound = false;
  	struct timeval tThen, tNow, tDelta, tExecTime;
  	timerclear(&tExecTime);
  	cl_uint workDimBerlekamp = 1;
  	size_t globalSizeBerlekamp[1] = {
  		256
  	};
  	size_t localSizeBerlekamp[1] = {
  		256
  	};
  
  	/* Input/output variables */
  	short *s = malloc(8160 * sizeof(short));
  	cl_mem sK = NULL;
  	short *c = malloc(4590 * sizeof(short));
  	short *cC = malloc(4590 * sizeof(short));
  	cl_mem cK = NULL;
  	short *w = malloc(4590 * sizeof(short));
  	short *wC = malloc(4590 * sizeof(short));
  	cl_mem wK = NULL;
  	unsigned char loopCount = 255;
  
  	/* Calling preamble function */
  	PRINT_STEP("Calling preamble function...");
  	PREAMBLE(s, 8160, c, 4590, cC, 4590, w, 4590, wC, 4590, loopCount);
  	PRINT_SUCCESS();
  
  	/* Get platforms IDs */
  	PRINT_STEP("Getting platforms IDs...");
  	fRet = clGetPlatformIDs(0, NULL, &platformsLen);
  	ASSERT_CALL(CL_SUCCESS == fRet, FUNCTION_ERROR_STATEMENTS("clGetPlatformIDs"));
  	platforms = malloc(platformsLen * sizeof(cl_platform_id));
  	fRet = clGetPlatformIDs(platformsLen, platforms, NULL);
  	ASSERT_CALL(CL_SUCCESS == fRet, FUNCTION_ERROR_STATEMENTS("clGetPlatformIDs"));
  	PRINT_SUCCESS();
  
  	/* Get devices IDs for first platform availble */
  	PRINT_STEP("Getting devices IDs for first platform...");
  	fRet = clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_ALL, 0, NULL, &devicesLen);
  	ASSERT_CALL(CL_SUCCESS == fRet, FUNCTION_ERROR_STATEMENTS("clGetDevicesIDs"));
  	devices = malloc(devicesLen * sizeof(cl_device_id));
  	fRet = clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_ALL, devicesLen, devices, NULL);
  	ASSERT_CALL(CL_SUCCESS == fRet, FUNCTION_ERROR_STATEMENTS("clGetDevicesIDs"));
  	PRINT_SUCCESS();
  
  	/* Create context for first available device */
  	PRINT_STEP("Creating context...");
  	context = clCreateContext(NULL, 1, devices, NULL, NULL, &fRet);
  	ASSERT_CALL(CL_SUCCESS == fRet, FUNCTION_ERROR_STATEMENTS("clCreateContext"));
  	PRINT_SUCCESS();
  
  	/* Create command queue for berlekamp kernel */
  	PRINT_STEP("Creating command queue for \"berlekamp\"...");
  	queueBerlekamp = clCreateCommandQueue(context, devices[0], 0, &fRet);
  	ASSERT_CALL(CL_SUCCESS == fRet, FUNCTION_ERROR_STATEMENTS("clCreateCommandQueue"));
  	PRINT_SUCCESS();
  
  	/* Open binary file */
  	PRINT_STEP("Opening program binary...");
  	programFile = fopen("kern.cl", "rb");
  	ASSERT_CALL(programFile, POSIX_ERROR_STATEMENTS("kern.cl"));
  	PRINT_SUCCESS();
  
  	/* Get size and read file */
  	PRINT_STEP("Reading program binary...");
  	fseek(programFile, 0, SEEK_END);
  	programSz = ftell(programFile);
  	fseek(programFile, 0, SEEK_SET);
  	programContent = malloc(programSz);
  	fread(programContent, programSz, 1, programFile);
  	fclose(programFile);
  	programFile = NULL;
  	PRINT_SUCCESS();
  
  	/* Create program from source file */
  	PRINT_STEP("Creating program from source...");
  	program = clCreateProgramWithSource(context, 1, (const char **) &programContent, &programSz, &fRet);
  	ASSERT_CALL(CL_SUCCESS == fRet, FUNCTION_ERROR_STATEMENTS("clCreateProgramWithSource"));
  	PRINT_SUCCESS();
  
  	/* Build program */
  	PRINT_STEP("Building program...");
  	fRet = clBuildProgram(program, 1, devices, NULL, NULL, NULL);
  	ASSERT_CALL(CL_SUCCESS == fRet, FUNCTION_ERROR_STATEMENTS("clBuildProgram"));
  	PRINT_SUCCESS();
  
  	/* Create berlekamp kernel */
  	PRINT_STEP("Creating kernel \"berlekamp\" from program...");
  	kernelBerlekamp = clCreateKernel(program, "berlekamp", &fRet);
  	ASSERT_CALL(CL_SUCCESS == fRet, FUNCTION_ERROR_STATEMENTS("clCreateKernel"));
  	PRINT_SUCCESS();
  
  	/* Create input and output buffers */
  	PRINT_STEP("Creating buffers...");
  	sK = clCreateBuffer(context, CL_MEM_READ_ONLY, 8160 * sizeof(short), NULL, &fRet);
  	ASSERT_CALL(CL_SUCCESS == fRet, FUNCTION_ERROR_STATEMENTS("clCreateBuffer (sK)"));
  	cK = clCreateBuffer(context, CL_MEM_READ_WRITE, 4590 * sizeof(short), NULL, &fRet);
  	ASSERT_CALL(CL_SUCCESS == fRet, FUNCTION_ERROR_STATEMENTS("clCreateBuffer (cK)"));
  	wK = clCreateBuffer(context, CL_MEM_READ_WRITE, 4590 * sizeof(short), NULL, &fRet);
  	ASSERT_CALL(CL_SUCCESS == fRet, FUNCTION_ERROR_STATEMENTS("clCreateBuffer (wK)"));
  	PRINT_SUCCESS();
  
  	/* Set kernel arguments for berlekamp */
  	PRINT_STEP("Setting kernel arguments for \"berlekamp\"...");
  	fRet = clSetKernelArg(kernelBerlekamp, 0, sizeof(cl_mem), &sK);
  	ASSERT_CALL(CL_SUCCESS == fRet, FUNCTION_ERROR_STATEMENTS("clSetKernelArg (sK)"));
  	fRet = clSetKernelArg(kernelBerlekamp, 1, sizeof(cl_mem), &cK);
  	ASSERT_CALL(CL_SUCCESS == fRet, FUNCTION_ERROR_STATEMENTS("clSetKernelArg (cK)"));
  	fRet = clSetKernelArg(kernelBerlekamp, 2, sizeof(cl_mem), &wK);
  	ASSERT_CALL(CL_SUCCESS == fRet, FUNCTION_ERROR_STATEMENTS("clSetKernelArg (wK)"));
  	fRet = clSetKernelArg(kernelBerlekamp, 3, sizeof(unsigned char), &loopCount);
  	ASSERT_CALL(CL_SUCCESS == fRet, FUNCTION_ERROR_STATEMENTS("clSetKernelArg (loopCount)"));
  	PRINT_SUCCESS();
  
  	do {
  		/* Setting input and output buffers */
  		PRINT_STEP("[%d] Setting buffers...", i);
  		fRet = clEnqueueWriteBuffer(queueBerlekamp, sK, CL_TRUE, 0, 8160 * sizeof(short), s, 0, NULL, NULL);
  		ASSERT_CALL(CL_SUCCESS == fRet, FUNCTION_ERROR_STATEMENTS("clEnqueueWriteBuffer (sK)"));
  		fRet = clEnqueueWriteBuffer(queueBerlekamp, cK, CL_TRUE, 0, 4590 * sizeof(short), c, 0, NULL, NULL);
  		ASSERT_CALL(CL_SUCCESS == fRet, FUNCTION_ERROR_STATEMENTS("clEnqueueWriteBuffer (cK)"));
  		fRet = clEnqueueWriteBuffer(queueBerlekamp, wK, CL_TRUE, 0, 4590 * sizeof(short), w, 0, NULL, NULL);
  		ASSERT_CALL(CL_SUCCESS == fRet, FUNCTION_ERROR_STATEMENTS("clEnqueueWriteBuffer (wK)"));
  		fRet = clSetKernelArg(kernelBerlekamp, 3, sizeof(unsigned char), &loopCount);
  		ASSERT_CALL(CL_SUCCESS == fRet, FUNCTION_ERROR_STATEMENTS("clSetKernelArg (loopCount)"));
  		PRINT_SUCCESS();
  
  		PRINT_STEP("[%d] Running kernels...", i);
  		gettimeofday(&tThen, NULL);
  		fRet = clEnqueueNDRangeKernel(queueBerlekamp, kernelBerlekamp, workDimBerlekamp, NULL, globalSizeBerlekamp, localSizeBerlekamp, 0, NULL, NULL);
  		ASSERT_CALL(CL_SUCCESS == fRet, FUNCTION_ERROR_STATEMENTS("clEnqueueNDRangeKernel"));
  		clFinish(queueBerlekamp);
  		gettimeofday(&tNow, NULL);
  		PRINT_SUCCESS();
  
  		/* Get output buffers */
  		PRINT_STEP("[%d] Getting kernels arguments...", i);
  		fRet = clEnqueueReadBuffer(queueBerlekamp, cK, CL_TRUE, 0, 4590 * sizeof(short), c, 0, NULL, NULL);
  		fRet = clEnqueueReadBuffer(queueBerlekamp, wK, CL_TRUE, 0, 4590 * sizeof(short), w, 0, NULL, NULL);
  		ASSERT_CALL(CL_SUCCESS == fRet, FUNCTION_ERROR_STATEMENTS("clEnqueueReadBuffer"));
  		PRINT_SUCCESS();
  
  		timersub(&tNow, &tThen, &tDelta);
  		timeradd(&tExecTime, &tDelta, &tExecTime);
  		i++;
  	} while(loopFlag);
  
  
  	/* Print profiling results */
  	long totalTime = (1000000 * tExecTime.tv_sec) + tExecTime.tv_usec;
  	printf("Elapsed time spent on kernels: %ld us; Average time per iteration: %lf us.\n", totalTime, totalTime / (double) i);
  
  	/* Validate received data */
  	PRINT_STEP("Validating received data...");
  	for(i = 0; i < 4590; i++) {
  		if(cC[i] != c[i]) {
  			if(!invalidDataFound) {
  				PRINT_FAIL();
  				invalidDataFound = true;
  			}
  			printf("Variable c[%d]: expected %hd got %hd.\n", i, cC[i], c[i]);
  		}
  	}
  	for(i = 0; i < 4590; i++) {
  		if(wC[i] != w[i]) {
  			if(!invalidDataFound) {
  				PRINT_FAIL();
  				invalidDataFound = true;
  			}
  			printf("Variable w[%d]: expected %hd got %hd.\n", i, wC[i], w[i]);
  		}
  	}
  	if(!invalidDataFound)
  		PRINT_SUCCESS();
  
  _err:
  
  	/* Dealloc buffers */
  	if(sK)
  		clReleaseMemObject(sK);
  	if(cK)
  		clReleaseMemObject(cK);
  	if(wK)
  		clReleaseMemObject(wK);
  
  	/* Dealloc variables */
  	free(s);
  	free(c);
  	free(cC);
  	free(w);
  	free(wC);
  
  	/* Dealloc kernels */
  	if(kernelBerlekamp)
  		clReleaseKernel(kernelBerlekamp);
  
  	/* Dealloc program */
  	if(program)
  		clReleaseProgram(program);
  	if(programContent)
  		free(programContent);
  	if(programFile)
  		fclose(programFile);
  
  	/* Dealloc queues */
  	if(queueBerlekamp)
  		clReleaseCommandQueue(queueBerlekamp);
  
  	/* Last OpenCL variables */
  	if(context)
  		clReleaseContext(context);
  	if(devices)
  		free(devices);
  	if(platforms)
  		free(platforms);
  
  
  	return rv;
  }
  