Replaced stream write with a logger.

New contructor added for each class.
The old constructors, if called as before, will
default to printing all the messages:
-stdout for default constructor
-file fname for (fname) constructor if fname is not
NULL, stdout or stderr.
This commit is contained in:
Dimitris Karkoulis 2012-06-24 18:23:02 +02:00
parent 69b0ee5170
commit 0e43e89a09
12 changed files with 1664 additions and 919 deletions

View File

@ -20,7 +20,7 @@
* Grenoble, France
*
* Principal authors: D. Karkoulis (karkouli@esrf.fr)
* Last revision: 21/06/2011
* Last revision: 24/06/2011
*
* This program is free software: you can redistribute it and/or modify
* it under the terms of the GNU Lesser General Public License as published
@ -66,16 +66,46 @@ typedef unsigned long lui;
/**
* \brief Overloaded constructor for base class.
*
* Output is set to file "fname"
* Complete logging functionality
*
*/
ocl::ocl(const char *fname){
ocl::ocl(FILE *stream, const char *fname, int safe, int depth, int perf_time, int timestamp, const char *identity):exec_identity(identity)
{
cLog_init(&hLog,stream,fname,0,static_cast<enum_LOGTYPE>(safe),static_cast<enum_LOGDEPTH>(depth),perf_time,timestamp);
if(identity)
cLog_date_text(&hLog,LOGDNONE,"(%s)\n",identity);
else
cLog_date(&hLog,LOGDNONE);
usesStdout = 0;
ContructorInit();
setDocstring("OpenCL base functionality for xrpd1d. \nFeel free to play around but you will not be able to perform integrations"\
"at this level.\nYou may check the OpenCL platforms and devices found in your system at this point. Try print_devices\n"\
"Try any of the derived classes xrpd1d and xrpd2d for complete functionality. \n","ocl_base.hpp");
}
/**
* \brief Overloaded constructor for base class.
*
* Output is set to file "fname"
*
*/
ocl::ocl(const char *fname, const char *identity):exec_identity(identity)
{
cLog_init(&hLog,NULL,fname,0,LOGTFAST,LOGDDEBUG,1,1);
if(identity)
cLog_date_text(&hLog,LOGDNONE,"(%s)\n",identity);
else
cLog_date(&hLog,LOGDNONE);
usesStdout = 0;
ContructorInit();
setDocstring("OpenCL base functionality for xrpd1d. \nFeel free to play around but you will not be able to perform integrations"\
"at this level.\nYou may check the OpenCL platforms and devices found in your system at this point. Try print_devices\n"\
"Try any of the derived classes xrpd1d and xrpd2d for complete functionality. \n","ocl_base.hpp");
stream=fopen(fname,"w");
usesStdout=0;
}
@ -85,11 +115,17 @@ ocl::ocl(const char *fname){
* Output is set to stdout
*
*/
ocl::ocl(){
ocl::ocl():exec_identity(NULL){
cLog_init(&hLog,stdout,NULL,0,LOGTFAST,LOGDDEBUG,1,1);
cLog_date(&hLog,LOGDNONE);
usesStdout = 1;
ContructorInit();
setDocstring("OpenCL base functionality for xrpd1d. \nFeel free to play around but you will not be able to perform integrations"\
"at this level.\nYou may check the OpenCL platforms and devices found in your system at this point. Try print_devices\n"\
"Try any of the derived classes xrpd1d and xrpd2d for complete functionality. \n","ocl_base.hpp");
}
ocl::~ocl(){
@ -97,7 +133,7 @@ ocl::~ocl(){
ocl_tools_destroy(oclconfig);
delete oclconfig;
delete sgs;
if(!usesStdout) fclose(stream);
cLog_fin(&hLog);
delete[] docstr;
}
@ -128,7 +164,7 @@ void ocl::ContructorInit()
reset_time();
oclconfig = new ocl_config_type;
ocl_tools_initialise(oclconfig);
ocl_tools_initialise(oclconfig,&hLog);
sgs = new az_argtype;
sgs->Nx = 0;
@ -142,14 +178,25 @@ void ocl::ContructorInit()
}
void ocl::update_logger(FILE *stream, const char *fname, int safe, int depth, int perf_time, int timestamp)
{
cLog_fin(&hLog);
cLog_init(&hLog,stream,fname,0,static_cast<enum_LOGTYPE>(safe),static_cast<enum_LOGDEPTH>(depth),perf_time,timestamp);
}
/**
* \brief Prints a list of OpenCL capable devices, their platforms and their ids to stream
*/
void ocl::show_devices(int ignoreStream){
//Print all the probed devices
if(ignoreStream) ocl_check_platforms();
else ocl_check_platforms(stream);
if(ignoreStream)
{
oclconfig->hLog->status=0;
ocl_check_platforms(oclconfig->hLog);
oclconfig->hLog->status=1;
}
else ocl_check_platforms(oclconfig->hLog);
return;
}
@ -191,34 +238,38 @@ void ocl::show_device_details(int ignoreStream){
if(hasActiveContext)
{
int dev,plat;
std::string heading;
std::ostringstream heading_stream;
char *heading;
char cast_plat,cast_dev;
FILE *tmp;
get_contexed_Ids(plat,dev);
cast_plat = '0' + (char)plat;
cast_dev = '0' + (char)dev;
heading = '(' + cast_plat + '.' + cast_dev + ')' + ' ';
heading_stream << '(' << cast_plat << '.' << cast_dev << ')' << ' ';
heading = new char [heading_stream.str().length() + 1];
strcpy(heading,heading_stream.str().c_str());
//Force cLogger to print to stdout
if(ignoreStream) hLog.status = 0;
//We need to do the following with fprintf because the stream is FILE and not ofstream.
if(ignoreStream) tmp = stdout;
else tmp = stream;
cLog_extended(&hLog,"%s Platform name: %s\n", heading, oclconfig->platform_info.name);
cLog_extended(&hLog,"%s Platform version: %s\n", heading, oclconfig->platform_info.version);
cLog_extended(&hLog,"%s Platform vendor: %s\n", heading, oclconfig->platform_info.vendor);
cLog_extended(&hLog,"%s Platform extensions: %s\n", heading, oclconfig->platform_info.extensions);
fprintf(tmp,"%s Platform name: %s\n", heading.c_str(), oclconfig->platform_info.name);
fprintf(tmp,"%s Platform version: %s\n", heading.c_str(), oclconfig->platform_info.version);
fprintf(tmp,"%s Platform vendor: %s\n", heading.c_str(), oclconfig->platform_info.vendor);
fprintf(tmp,"%s Platform extensions: %s\n", heading.c_str(), oclconfig->platform_info.extensions);
cLog_extended(&hLog,"\n");
fprintf(tmp,"\n");
cLog_extended(&hLog,"%s Device name: %s\n", heading, oclconfig->device_info.name);
cLog_extended(&hLog,"%s Device type: %s\n", heading, oclconfig->device_info.type);
cLog_extended(&hLog,"%s Device version: %s\n", heading, oclconfig->device_info.version);
cLog_extended(&hLog,"%s Device driver version: %s\n", heading, oclconfig->device_info.driver_version);
cLog_extended(&hLog,"%s Device extensions: %s\n", heading, oclconfig->device_info.extensions);
cLog_extended(&hLog,"%s Device Max Memory: %f (MB)\n", heading, oclconfig->device_info.global_mem/1024.f/1024.f);
fprintf(tmp,"%s Device name: %s\n", heading.c_str(), oclconfig->device_info.name);
fprintf(tmp,"%s Device type: %s\n", heading.c_str(), oclconfig->device_info.type);
fprintf(tmp,"%s Device version: %s\n", heading.c_str(), oclconfig->device_info.version);
fprintf(tmp,"%s Device driver version: %s\n", heading.c_str(), oclconfig->device_info.driver_version);
fprintf(tmp,"%s Device extensions: %s\n", heading.c_str(), oclconfig->device_info.extensions);
fprintf(tmp,"%s Device Max Memory: %ul\n", heading.c_str(), oclconfig->device_info.global_mem);
//Revert cLogger to normal operation
if(ignoreStream) hLog.status = 1;
delete [] heading;
}
return;
}
@ -259,7 +310,7 @@ int ocl::init(const bool useFp64){
//Pick a device and initiate a context. If a context exists destroy it
clean();
if(ocl_init_context(oclconfig,"DEF",(int)useFp64,stream)) return -1;
if(ocl_init_context(oclconfig,"DEF",(int)useFp64)) return -1;
else hasActiveContext=1;
promote_device_details();
@ -281,7 +332,7 @@ int ocl::init(const char *devicetype,const bool useFp64){
//Pick a device and initiate a context. If a context exists destroy it
this->clean();
if(ocl_init_context(oclconfig,devicetype,(int)useFp64,stream)) return -1;
if(ocl_init_context(oclconfig,devicetype,(int)useFp64)) return -1;
else hasActiveContext=1;
promote_device_details();
@ -306,7 +357,7 @@ int ocl::init(const char *devicetype,int platformid,int devid,const bool useFp64
//Pick a device and initiate a context. If a context exists destroy it
clean();
if(ocl_init_context(oclconfig,devicetype,platformid,devid,(int)useFp64,stream)) return -1;
if(ocl_init_context(oclconfig,devicetype,platformid,devid,(int)useFp64)) return -1;
else hasActiveContext=1;
promote_device_details();
@ -327,9 +378,9 @@ int ocl::clean(int preserve_context){
if(!preserve_context)
{
if(hasActiveContext){
ocl_destroy_context(oclconfig->oclcontext);
ocl_destroy_context(oclconfig->oclcontext, &hLog);
hasActiveContext=0;
fprintf(stream,"--released OpenCL context\n");
cLog_debug(&hLog,"--released OpenCL context\n");
return 0;
}
}
@ -346,10 +397,10 @@ return -2;
void ocl::kill_context(){
if(hasActiveContext)
{
ocl_destroy_context(this->oclconfig->oclcontext);
ocl_destroy_context(this->oclconfig->oclcontext, &hLog);
hasActiveContext=0;
fprintf(stream,"Forced destroy context\n");
}else fprintf(stream,"Attempted Forced destroy context ignored\n");
cLog_debug(&hLog,"Forced destroy context\n");
}else cLog_debug(&hLog,"Attempted Forced destroy context ignored\n");
return;
}
@ -463,7 +514,7 @@ void ocl::setDocstring(const char *default_text, const char *filename)
ifstream readme;
std::streamoff len=0;
readme.open(filename,ios::in);
readme.open(filename,ios::in | ios::binary);
//If the file exists:
if(readme){
@ -498,7 +549,7 @@ void ocl::setDocstring(const char *default_text, const char *filename)
}else delete[] bkp;
//Read from file and check we read ALL the data
if( readme.read(docstr,len).gcount() != len) fprintf(stderr,"setDocstring read size mismatch!\n");
if( readme.read(docstr,len).gcount() != len) cLog_critical(&hLog,"setDocstring read size mismatch!\n");
docstr[len] = '\0';
readme.close();
}

View File

@ -20,7 +20,7 @@
* Grenoble, France
*
* Principal authors: D. Karkoulis (karkouli@esrf.fr)
* Last revision: 21/06/2012
* Last revision: 24/06/2012
*
* This program is free software: you can redistribute it and/or modify
* it under the terms of the GNU Lesser General Public License as published
@ -58,10 +58,11 @@
#endif
#include <iostream>
#include <sstream>
#include <CL/opencl.h>
#include "ocl_ckerr.h"
#include "ocl_tools.h"
#include "ocl_tools/ocl_tools.h"
#include "ocl_tools/cLogger/cLogger.h"
/**
* \brief Holds the integration configuration parameters
@ -88,10 +89,13 @@ public:
class ocl{
public:
explicit ocl(const char *fname);
explicit ocl(FILE *stream, const char *fname, int safe, int depth, int perf_time, int timestamp, const char *identity=NULL);
explicit ocl(const char *fname, const char *identity=NULL);
ocl();
virtual ~ocl();
void update_logger(FILE *stream, const char *fname, int safe, int depth, int perf_time, int timestamp);
/*
* Initial configuration: Choose a device and initiate a context. Devicetypes can be GPU,gpu,CPU,cpu,DEF,ACC,ALL.
* Suggested are GPU,CPU. For each setting to work there must be such an OpenCL device and properly installed.
@ -188,6 +192,8 @@ protected:
/**@}*/
FILE *stream;
logger_t hLog;
const char *exec_identity;
/**
* \defgroup guards Status flags/guards

View File

@ -0,0 +1,354 @@
/**
* \file
* \brief Minimal logger
*
* Provides logger functions with basic functionality while trying to remain
* fast and minimal in depedency requirements
*/
/*
* Project: Minimal logger suite. Logging with no depedencies
*
* Copyright (C) 2012 Dimitrios Karkoulis
*
* Principal authors: D. Karkoulis (dimitris.karkoulis@gmail.com)
* Last revision: 24/06/2012
*
* This program is free software: you can redistribute it and/or modify
* it under the terms of the GNU Lesser General Public License as published
* by the Free Software Foundation, either version 3 of the License, or
* (at your option) any later version.
*
* This program is distributed in the hope that it will be useful,
* but WITHOUT ANY WARRANTY; without even the implied warranty of
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
* GNU Lesser General Public License for more details.
*
* You should have received a copy of the GNU General Public License
* and the GNU Lesser General Public License along with this program.
* If not, see <http://www.gnu.org/licenses/>.
*/
#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <string.h>
#include <stdarg.h>
#include "cLogger.h"
#ifdef _MSC_VER
#define STATINL static __inline
#else
#define STATINL static inline
#endif
//#ifdef _WIN32
//#define _BUFLIMIT 1
//#define fprintf(...) \
// do{ \
// if( ((*hLog).safebuf > _BUFLIMIT) && ((*hLog).type != LOGTSAFE)) \
// { \
// fflush((*hLog).stream); \
// (*hLog).safebuf = 0; \
// } \
// fprintf(__VA_ARGS__); \
// (*hLog).safebuf++; \
// }while(0)
//
//
//#define vfprintf(...) \
// do{ \
// if( ((*hLog).safebuf > _BUFLIMIT) && ((*hLog).type != LOGTSAFE)) \
// { \
// fflush((*hLog).stream); \
// (*hLog).safebuf = 0; \
// } \
// vfprintf(__VA_ARGS__); \
// (*hLog).safebuf++; \
// }while(0)
//
//#else
//#define fprintf fprintf
//#define vfprintf vfprintf
//#endif
const char *get_date()
{
static char date[50];
time_t t = time(0);
sprintf(date, "%s" , asctime(localtime(&t)));
return date;
}
STATINL const char *get_timestamp()
{
static char date[13];
time_t t = time(0);
strftime(date, sizeof(date), "(%H:%M:%S)", localtime(&t));
return date;
}
void cLog_init(logger_t *hLog, FILE *stream, const char *fname, int severity, enum_LOGTYPE type, enum_LOGDEPTH depth, int perf, int timestamps)
{
(*hLog).status = 0;
(*hLog).fname = NULL;
(*hLog).severity = 0;
(*hLog).type = 0;
(*hLog).depth = 0;
(*hLog).perf = 0;
(*hLog).timestamps = 0;
if( stream == NULL && fname == NULL ) stream = stdout;
if (stream == stdout || stream == stderr )
{
(*hLog).stream=stdout;
(*hLog).fname = (char *)malloc(6*sizeof(char));
strcpy((*hLog).fname,"NULL");
}else if( stream != NULL && fname == NULL )
{
//Very dangerous. Not Allowed!
fprintf(stderr, "\n"
"/-------------cLog---------------------------------------\\\n"
"| You are trying to use cLog with a stream that is not |\n"
"| NULL, stdout or stderr. This is not allowed when a |\n"
"| a filename is not set and cLog will be disabled. |\n"
"| All messages through cLog will be directed to stdout |\n"
"| or stderr (for critical messages). |\n"
"\--------------------------------------------------------/\n"
"\n"
);
(*hLog).status = 0;
return;
}else
{
(*hLog).stream = fopen(fname,"w");
(*hLog).fname = (char *)malloc(strlen(fname) + 1);
strcpy((*hLog).fname,fname);
}
(*hLog).severity = severity;
(*hLog).type = type;
(*hLog).depth = depth;
(*hLog).perf = perf;
(*hLog).timestamps = timestamps;
(*hLog).status = 1;
return;
}
void cLog_fin(logger_t *hLog)
{
if((*hLog).status)
{
if( (*hLog).stream && (*hLog).stream != stdout && (*hLog).stream != stderr)
fclose((*hLog).stream);
(*hLog).stream = NULL;
if((*hLog).fname)free((*hLog).fname);
(*hLog).status = 0;
}
return;
}
void cLog_date(logger_t *hLog, enum_LOGDEPTH depth)
{
if( ((*hLog).depth >= depth) && ((*hLog).status == 1) )
{
switch((*hLog).type)
{
case LOGTFAST:
fprintf((*hLog).stream,"%s",get_date());
break;
case LOGTSAFE:
fflush((*hLog).stream);
fprintf((*hLog).stream,"%s",get_date());
fflush((*hLog).stream);
break;
}
}
}
void cLog_date_text(logger_t *hLog, enum_LOGDEPTH depth, const char *format, ...)
{
va_list argp;
if( ((*hLog).depth >= depth) && ((*hLog).status == 1) )
{
va_start(argp,format);
switch((*hLog).type)
{
case LOGTFAST:
fprintf((*hLog).stream,"%s",get_date());
vfprintf((*hLog).stream,format,argp);
break;
case LOGTSAFE:
fflush((*hLog).stream);
fprintf((*hLog).stream,"%s",get_date());
vfprintf((*hLog).stream,format,argp);
fflush((*hLog).stream);
break;
}
va_end(argp);
} else if ( (*hLog).status == 0 )
{
va_start(argp,format);
vprintf(format,argp);
va_end(argp);
}
}
void cLog_basic(logger_t *hLog, const char * format, ...)
{
va_list argp;
if( ((*hLog).depth >= LOGDBASIC) && (*hLog).status)
{
if((*hLog).timestamps)fprintf((*hLog).stream,"%s ",get_timestamp());
va_start(argp,format);
switch((*hLog).type)
{
case LOGTFAST:
vfprintf((*hLog).stream,format,argp);
break;
case LOGTSAFE:
fflush((*hLog).stream);
vfprintf((*hLog).stream,format,argp);
fflush((*hLog).stream);
break;
}
va_end(argp);
} else if ( (*hLog).status == 0 )
{
va_start(argp,format);
vprintf(format,argp);
va_end(argp);
}
return;
}
void cLog_extended(logger_t *hLog, const char * format, ...)
{
va_list argp;
if( ((*hLog).depth >= LOGDEXTENDED) && ((*hLog).status == 1) )
{
if((*hLog).timestamps)fprintf((*hLog).stream,"%s ",get_timestamp());
va_start(argp,format);
switch((*hLog).type)
{
case LOGTFAST:
vfprintf((*hLog).stream,format,argp);
break;
case LOGTSAFE:
fflush((*hLog).stream);
vfprintf((*hLog).stream,format,argp);
fflush((*hLog).stream);
break;
}
va_end(argp);
} else if ( (*hLog).status == 0 )
{
va_start(argp,format);
vprintf(format,argp);
va_end(argp);
}
return;
}
void cLog_debug(logger_t *hLog, const char * format, ...)
{
va_list argp;
if((*hLog).depth >= LOGDDEBUG && (*hLog).status)
{
if((*hLog).timestamps)fprintf((*hLog).stream,"%s ",get_timestamp());
va_start(argp,format);
switch((*hLog).type)
{
case LOGTFAST:
vfprintf((*hLog).stream,format,argp);
break;
case LOGTSAFE:
fflush((*hLog).stream);
vfprintf((*hLog).stream,format,argp);
fflush((*hLog).stream);
break;
}
va_end(argp);
} else if ( (*hLog).status == 0 )
{
va_start(argp,format);
vprintf(format,argp);
va_end(argp);
}
return;
}
void cLog_critical(logger_t *hLog, const char * format, ...)
{
va_list argp;
if( ((*hLog).depth >= LOGDONLYERRORS) && ((*hLog).status == 1) )
{
va_start(argp,format);
if((*hLog).timestamps)fprintf((*hLog).stream,"%s ",get_timestamp());
fflush((*hLog).stream);
vfprintf((*hLog).stream,format,argp);
fflush((*hLog).stream);
va_end(argp);
}
if( ((*hLog).stream != stderr) || ((*hLog).status == 0) )
{
va_start(argp,format);
fflush(stderr);
vfprintf(stderr,format,argp);
fflush(stderr);
va_end(argp);
}
return;
}
void cLog_bench(logger_t *hLog, const char * format, ...)
{
va_list argp;
if((*hLog).perf == 1 && (*hLog).depth >= LOGDEXTENDED && (*hLog).status)
{
if((*hLog).timestamps)fprintf((*hLog).stream,"%s ",get_timestamp());
va_start(argp,format);
switch((*hLog).type)
{
case LOGTFAST:
vfprintf((*hLog).stream,format,argp);
break;
case LOGTSAFE:
fflush((*hLog).stream);
vfprintf((*hLog).stream,format,argp);
fflush((*hLog).stream);
break;
}
va_end(argp);
} else if ( (*hLog).status == 0 )
{
va_start(argp,format);
vprintf(format,argp);
va_end(argp);
}
return;
}

View File

@ -0,0 +1,86 @@
/*
* Project: Minimal logger suite. Logging with no depedencies
*
* Copyright (C) 2012 Dimitrios Karkoulis
*
* Principal authors: D. Karkoulis (dimitris.karkoulis@gmail.com)
* Last revision: 24/06/2012
*
* This program is free software: you can redistribute it and/or modify
* it under the terms of the GNU Lesser General Public License as published
* by the Free Software Foundation, either version 3 of the License, or
* (at your option) any later version.
*
* This program is distributed in the hope that it will be useful,
* but WITHOUT ANY WARRANTY; without even the implied warranty of
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
* GNU Lesser General Public License for more details.
*
* You should have received a copy of the GNU General Public License
* and the GNU Lesser General Public License along with this program.
* If not, see <http://www.gnu.org/licenses/>.
*/
#ifndef CLOGGER_H
#define CLOGGER_H
#include <stdarg.h>
#ifdef _WIN32
#ifdef _DLLIMPORT
#define DllInterface __declspec( dllimport )
#else
#define DllInterface __declspec( dllexport )
#endif
#else
#define DllInterface
#endif
typedef struct
{
FILE * stream;
int severity;
int type;
int depth;
int perf;
int timestamps;
int status;
char *fname;
} logger_t;
typedef enum
{
LOGTFAST,
LOGTSAFE
}enum_LOGTYPE;
typedef enum
{
LOGDNONE,
LOGDONLYERRORS,
LOGDBASIC,
LOGDEXTENDED,
LOGDDEBUG
}enum_LOGDEPTH;
#ifdef __cplusplus
extern "C"{
#endif
DllInterface void cLog_basic(logger_t *hLog, const char * format, ...);
DllInterface void cLog_extended(logger_t *hLog, const char* format, ...);
DllInterface void cLog_debug(logger_t *hLog, const char * format, ...);
DllInterface void cLog_critical(logger_t *hLog, const char * format, ...);
DllInterface void cLog_bench(logger_t *hLog, const char* format, ...);
DllInterface void cLog_date(logger_t *hLog, enum_LOGDEPTH depth);
DllInterface void cLog_date_text(logger_t *hLog, enum_LOGDEPTH depth, const char *format, ...);
DllInterface void cLog_init(logger_t *hLog, FILE *stream, const char *fname, int severity, enum_LOGTYPE type, enum_LOGDEPTH depth, int perf, int timestamps);
DllInterface void cLog_fin(logger_t *hLog);
#ifdef __cplusplus
}
#endif
#endif

View File

@ -1,106 +1,106 @@
/*
* Project: Macros for OpenCL API error handling. Requires ocl_tools
*
* Copyright (C) 2011 European Synchrotron Radiation Facility
* Grenoble, France
*
* Principal authors: D. Karkoulis (karkouli@esrf.fr)
* Last revision: 27/05/2011
*
* This program is free software: you can redistribute it and/or modify
* it under the terms of the GNU Lesser General Public License as published
* by the Free Software Foundation, either version 3 of the License, or
* (at your option) any later version.
*
* This program is distributed in the hope that it will be useful,
* but WITHOUT ANY WARRANTY; without even the implied warranty of
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
* GNU Lesser General Public License for more details.
*
* You should have received a copy of the GNU General Public License
* and the GNU Lesser General Public License along with this program.
* If not, see <http://www.gnu.org/licenses/>.
*/
#ifdef _WIN32
#define typeof(_expr) typeid(_expr)
#endif
#define CL_CHECK(_expr) \
do { \
cl_int _err = _expr; \
if (_err == CL_SUCCESS) \
break; \
fprintf(stderr, "OpenCL: '%s' returned %d!\n", #_expr, (int)_err); \
exit(1); \
} while (0)
#define CL_CHECK_ERR(_expr) \
({ \
cl_int _err = CL_INVALID_VALUE; \
typeof(_expr) _ret = _expr; \
if (_err != CL_SUCCESS) { \
fprintf(stderr, "OpenCL: '%s' returned %d!\n", #_expr, (int)_err); \
exit(1); \
} \
_ret; \
})
#define CL_CHECK_PR(_expr) \
do { \
cl_int _err = _expr; \
if (_err == CL_SUCCESS) \
break; \
fprintf(stderr, "OpenCL: '%s:%d' returned %s!\n",__FILE__,__LINE__, ocl_perrc(_err)); \
exit(1); \
} while (0)
#define CL_CHECK_PRN(_expr) \
do { \
cl_int _err = _expr; \
if (_err == CL_SUCCESS) \
break; \
fprintf(stderr, "OpenCL: '%s:%d' returned %s!\n",__FILE__,__LINE__, ocl_perrc(_err)); \
break; \
} while (0)
#define CL_CHECK_PR_RET(_expr) \
do { \
cl_int _err = _expr; \
if (_err == CL_SUCCESS) \
break; \
fprintf(stderr, "OpenCL: '%s:%d' returned %s!\n",__FILE__,__LINE__, ocl_perrc(_err)); \
return -1; \
} while (0)
#define CL_CHECK_ERR_PR(_expr) \
({ \
cl_int err = CL_INVALID_VALUE; \
typeof(_expr) _ret = _expr; \
if (err != CL_SUCCESS) { \
fprintf(stderr, "OpenCL: '%s:%d' returned %s!\n",__FILE__,__LINE__, ocl_perrc(err)); \
exit(1); \
} \
_ret; \
})
#define CL_CHECK_ERR_PRN(_expr) \
({ \
cl_int err = CL_INVALID_VALUE; \
typeof(_expr) _ret = _expr; \
if (err != CL_SUCCESS) { \
fprintf(stderr, "OpenCL: '%s:%d' returned %s!\n",__FILE__,__LINE__, ocl_perrc(err)); \
} \
_ret; \
})
#define CL_CHECK_ERR_PR_RET(_expr) \
({ \
cl_int err = CL_INVALID_VALUE; \
typeof(_expr) _ret = _expr; \
if (err != CL_SUCCESS) { \
fprintf(stderr, "OpenCL: '%s:%d' returned %s!\n",__FILE__,__LINE__, ocl_perrc(err)); \
return -1; \
} \
_ret; \
})
/*
* Project: Macros for OpenCL API error handling. Requires ocl_tools
*
* Copyright (C) 2011 European Synchrotron Radiation Facility
* Grenoble, France
*
* Principal authors: D. Karkoulis (karkouli@esrf.fr)
* Last revision: 27/05/2011
*
* This program is free software: you can redistribute it and/or modify
* it under the terms of the GNU Lesser General Public License as published
* by the Free Software Foundation, either version 3 of the License, or
* (at your option) any later version.
*
* This program is distributed in the hope that it will be useful,
* but WITHOUT ANY WARRANTY; without even the implied warranty of
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
* GNU Lesser General Public License for more details.
*
* You should have received a copy of the GNU General Public License
* and the GNU Lesser General Public License along with this program.
* If not, see <http://www.gnu.org/licenses/>.
*/
#ifdef _WIN32
#define typeof(_expr) typeid(_expr)
#endif
#define CL_CHECK(_expr) \
do { \
cl_int _err = _expr; \
if (_err == CL_SUCCESS) \
break; \
fprintf(stderr, "OpenCL: '%s' returned %d!\n", #_expr, (int)_err); \
exit(1); \
} while (0)
#define CL_CHECK_ERR(_expr) \
({ \
cl_int _err = CL_INVALID_VALUE; \
typeof(_expr) _ret = _expr; \
if (_err != CL_SUCCESS) { \
fprintf(stderr, "OpenCL: '%s' returned %d!\n", #_expr, (int)_err); \
exit(1); \
} \
_ret; \
})
#define CL_CHECK_PR(_expr) \
do { \
cl_int _err = _expr; \
if (_err == CL_SUCCESS) \
break; \
fprintf(stderr, "OpenCL: '%s:%d' returned %s!\n",__FILE__,__LINE__, ocl_perrc(_err)); \
exit(1); \
} while (0)
#define CL_CHECK_PRN(_expr) \
do { \
cl_int _err = _expr; \
if (_err == CL_SUCCESS) \
break; \
fprintf(stderr, "OpenCL: '%s:%d' returned %s!\n",__FILE__,__LINE__, ocl_perrc(_err)); \
break; \
} while (0)
#define CL_CHECK_PR_RET(_expr) \
do { \
cl_int _err = _expr; \
if (_err == CL_SUCCESS) \
break; \
fprintf(stderr, "OpenCL: '%s:%d' returned %s!\n",__FILE__,__LINE__, ocl_perrc(_err)); \
return -1; \
} while (0)
#define CL_CHECK_ERR_PR(_expr) \
({ \
cl_int err = CL_INVALID_VALUE; \
typeof(_expr) _ret = _expr; \
if (err != CL_SUCCESS) { \
fprintf(stderr, "OpenCL: '%s:%d' returned %s!\n",__FILE__,__LINE__, ocl_perrc(err)); \
exit(1); \
} \
_ret; \
})
#define CL_CHECK_ERR_PRN(_expr) \
({ \
cl_int err = CL_INVALID_VALUE; \
typeof(_expr) _ret = _expr; \
if (err != CL_SUCCESS) { \
fprintf(stderr, "OpenCL: '%s:%d' returned %s!\n",__FILE__,__LINE__, ocl_perrc(err)); \
} \
_ret; \
})
#define CL_CHECK_ERR_PR_RET(_expr) \
({ \
cl_int err = CL_INVALID_VALUE; \
typeof(_expr) _ret = _expr; \
if (err != CL_SUCCESS) { \
fprintf(stderr, "OpenCL: '%s:%d' returned %s!\n",__FILE__,__LINE__, ocl_perrc(err)); \
return -1; \
} \
_ret; \
})

View File

@ -0,0 +1,105 @@
/*
* Project: Macros for OpenCL API error handling through cLogger. Requires ocl_tools
*
* Copyright (C) 2012 Dimitris Karkoulis
*
* Principal authors: D. Karkoulis (dimitris.karkoulis@gmail.com)
* Last revision: 24/06/2012
*
* This program is free software: you can redistribute it and/or modify
* it under the terms of the GNU Lesser General Public License as published
* by the Free Software Foundation, either version 3 of the License, or
* (at your option) any later version.
*
* This program is distributed in the hope that it will be useful,
* but WITHOUT ANY WARRANTY; without even the implied warranty of
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
* GNU Lesser General Public License for more details.
*
* You should have received a copy of the GNU General Public License
* and the GNU Lesser General Public License along with this program.
* If not, see <http://www.gnu.org/licenses/>.
*/
#ifdef _WIN32
#define typeof(_expr) typeid(_expr)
#endif
#define CL_CHECK(_expr, hLog) \
do { \
cl_int _err = _expr; \
if (_err == CL_SUCCESS) \
break; \
cLog_critical( hLog, "OpenCL: '%s' returned %d!\n", #_expr, (int)_err); \
exit(1); \
} while (0)
#define CL_CHECK_ERR(_expr, hLog) \
do { \
cl_int _err = CL_INVALID_VALUE; \
typeof(_expr) _ret = _expr; \
if (_err != CL_SUCCESS) { \
cLog_critical( hLog, "OpenCL: '%s' returned %d!\n", #_expr, (int)_err); \
exit(1); \
} \
_ret; \
} while (0)
#define CL_CHECK_PR(_expr, hLog) \
do { \
cl_int _err = _expr; \
if (_err == CL_SUCCESS) \
break; \
cLog_critical( hLog, "OpenCL: '%s:%d' returned %s!\n",__FILE__,__LINE__, ocl_perrc(_err)); \
exit(1); \
} while (0)
#define CL_CHECK_PRN(_expr, hLog) \
do { \
cl_int _err = _expr; \
if (_err == CL_SUCCESS) \
break; \
cLog_critical( hLog, "OpenCL: '%s:%d' returned %s!\n",__FILE__,__LINE__, ocl_perrc(_err)); \
break; \
} while (0)
#define CL_CHECK_PR_RET(_expr, hLog) \
do { \
cl_int _err = _expr; \
if (_err == CL_SUCCESS) \
break; \
cLog_critical( hLog, "OpenCL: '%s:%d' returned %s!\n",__FILE__,__LINE__, ocl_perrc(_err)); \
return -1; \
} while (0)
#define CL_CHECK_ERR_PR(_expr, hLog) \
do { \
cl_int err = CL_INVALID_VALUE; \
typeof(_expr) _ret = _expr; \
if (err != CL_SUCCESS) { \
cLog_critical( hLog, "OpenCL: '%s:%d' returned %s!\n",__FILE__,__LINE__, ocl_perrc(err)); \
exit(1); \
} \
_ret; \
} while (0)
#define CL_CHECK_ERR_PRN(_expr, hLog) \
do ({ \
cl_int err = CL_INVALID_VALUE; \
typeof(_expr) _ret = _expr; \
if (err != CL_SUCCESS) { \
cLog_critical( hLog, "OpenCL: '%s:%d' returned %s!\n",__FILE__,__LINE__, ocl_perrc(err)); \
} \
_ret; \
} while (0)
#define CL_CHECK_ERR_PR_RET(_expr, hLog) \
do { \
cl_int err = CL_INVALID_VALUE; \
typeof(_expr) _ret = _expr; \
if (err != CL_SUCCESS) { \
cLog_critical( hLog, "OpenCL: '%s:%d' returned %s!\n",__FILE__,__LINE__, ocl_perrc(err)); \
return -1; \
} \
_ret; \
} while (0)

File diff suppressed because it is too large Load Diff

View File

@ -1,346 +1,352 @@
/**
* \file
* \brief OpenCL tools header
*
* OpenCL tools for device probe, selection, deletion, error notification
* and vector type conversion. This source is the low-level layer of our
* OpenCL Toolbox (ocl_init_exec.cpp). However, it can be used directly
* as an API
*/
/*
* Project: OpenCL tools for device probe, selection, deletion, error notification
* and vector type conversion. This source is the low-level layer of our
* OpenCL Toolbox (ocl_init_context.cpp). However, it can be used directly
* as an API
*
* Copyright (C) 2011 - 2012 European Synchrotron Radiation Facility
* Grenoble, France
*
* Principal authors: D. Karkoulis (karkouli@esrf.fr)
* Last revision: 21/06/2012
*
* This program is free software: you can redistribute it and/or modify
* it under the terms of the GNU Lesser General Public License as published
* by the Free Software Foundation, either version 3 of the License, or
* (at your option) any later version.
*
* This program is distributed in the hope that it will be useful,
* but WITHOUT ANY WARRANTY; without even the implied warranty of
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
* GNU Lesser General Public License for more details.
*
* You should have received a copy of the GNU General Public License
* and the GNU Lesser General Public License along with this program.
* If not, see <http://www.gnu.org/licenses/>.
*/
/* Header for OpenCL Utilities */
#ifndef OCLTOOLS_H
#define OCLTOOLS_H
#include <CL/opencl.h>
#include "ocl_ckerr.h"
/*This is required for OpenCL callbacks in windows*/
#ifdef _WIN32
#define __call_compat __stdcall
#else
#define __call_compat
#ifndef _CRT_SECURE_NO_WARNINGS
#define _CRT_SECURE_NO_WARNINGS
#endif
#pragma warning(disable : 4996)
#endif
/**
* \brief Maximum number of OpenCL platforms to scan
*/
#define OCL_MAX_PLATFORMS 5
/**
* \brief Maximum number of OpenCL devices-per platform- to scan
*/
#define OCL_MAX_DEVICES 5
/**
* \brief OpenCL tools platform information struct
*
* It can be passed to ocl_platform_info to
* retrieve and save platform information
* for the currect context
*/
typedef struct
{
char *name;
char *vendor;
char *version;
char *extensions;
}ocl_plat_t;
/**
* \brief OpenCL tools platform information struct
*
* It can be passed to ocl_device_info to
* retrieve and save device information
* for the currect context
*/
typedef struct
{
char *name;
char type[4];
char *version;
char *driver_version;
char *extensions;
unsigned long global_mem;
}ocl_dev_t;
/**
* \brief OpenCL tools cl_program structure
*
* Substruct of ocl_configuration_parameters.
* It is used when multiple cl programs must be
* built, i.e. from multiple OpenCL sources.
*/
typedef struct{
cl_program oclprogram;
size_t *kernelstring_lens;
char **kernelstrings;
}ocl_program_type;
/**
* \brief OpenCL tools configuration parameters
*
* OpenCL configuration structure.
* This version supports single device, but multiple
* memory buffers, kernels and sources.
*/
typedef struct ocl_configuration_parameters{
cl_context oclcontext;
cl_device_id ocldevice;
cl_platform_id oclplatform;
cl_command_queue oclcmdqueue;
cl_mem *oclmemref;
//Active device and platform info
int devid;
int platfid;
ocl_plat_t platform_info;
ocl_dev_t device_info;
//If single .cl file:
cl_program oclprogram;
size_t *kernelstring_lens;
char **kernelstrings;
//if multiple .cl files:
ocl_program_type *prgs;
int nprgs;
char compiler_options[1000];
cl_kernel *oclkernels;
int fp64;
size_t work_dim[3];
size_t thread_dim[3];
cl_event t_s[20];
cl_int event_counter;
cl_ulong dev_mem;
cl_int Nbuffers;
cl_int Nkernels;
}ocl_config_type;
/* All production functions return 0 on success, -1 on OpenCL error and -2 on other errors.
when an error is encountered internally, it will print the message to stderr and fallback.
It is important for the user to decide how to continue.*/
/**
* \brief Initialises the internals of an ocl_config_type
*/
void ocl_tools_initialise(ocl_config_type *oclconfig);
/**
* \brief Deallocations inside ocl_config_type
*/
void ocl_tools_destroy(ocl_config_type *oclconfig);
/**
* \brief Simple check all platforms and devices and print information
*/
int ocl_check_platforms(FILE *stream=stdout);
/**
* \brief Simple check for a "device_type" device. Returns the first occurance only
*/
int ocl_find_devicetype(cl_device_type device_type, cl_platform_id &platform, cl_device_id &devid,FILE *stream=stdout);
/**
* \brief Simple check for a "device_type" device. Returns the first occurance that supports double precision only
*/
int ocl_find_devicetype_FP64(cl_device_type device_type, cl_platform_id& platform, cl_device_id& devid);
/**
* \brief Probes OpenCL platforms & devices of a given cl_device_type. Keeps the selected device in oclconfig
*/
int ocl_probe(ocl_config_type *oclconfig,cl_device_type ocldevtype, int usefp64, FILE *stream=stdout);
/**
* \brief Probes OpenCL device of a specific cl_device_type, platform and device number. Keeps the selected device in oclconfig
*/
int ocl_probe(ocl_config_type *oclconfig,cl_device_type ocldevtype,int preset_platform,int preset_device, int usefp64,
FILE *stream=stdout);
/**
* \brief Probes OpenCL device of a specific cl_platform_id and cl_device_id. Keeps the selected device in oclconfig
*/
int ocl_probe(ocl_config_type *oclconfig,cl_platform_id platform,cl_device_id device, int usefp64,
FILE *stream=stdout);
/**
* \brief Create an OpenCL context by device type
*/
/* Needs a string with the type of device: GPU,CPU,ACC,ALL,DEF. Runs ocl_probe and creates the context,
adding it on the appropriate ocl_config_type field*/
int ocl_init_context(ocl_config_type *oclconfig,const char *device_type, int usefp64, FILE *stream=stdout);
/**
* \brief Create an OpenCL context by device type, platform and device number
*/
int ocl_init_context(ocl_config_type *oclconfig,const char *device_type,int preset_platform,int devid, int usefp64,
FILE *stream=stdout);
/**
* \brief Create an OpenCL context by cl_platform_id and cl_device_id
*/
int ocl_init_context(ocl_config_type *oclconfig,cl_platform_id platform,cl_device_id device, int usefp64,
FILE *stream=stdout);
/**
* \brief Destroy an OpenCL context
*/
int ocl_destroy_context(cl_context oclcontext);
/**
* \brief Queries the fp64 capability of an OpenCL device that has been selected by ocl_probe
*/
/* Queries device capabilities to figure if it meets the minimum requirement for double precision*/
/* Returns 0 on successful FP64 evaluation and -1 if only FP32 */
int ocl_eval_FP64(ocl_config_type *oclconfig, FILE *stream);
/**
* \brief Queries the fp64 capability of an OpenCL device directly via the cl_device_id
*/
/* Same as above but directly query a device (as not set in ocl_config_type)
* It is designed to be used while probing for devices so it does not print anything
* neither it sets the fp64 field */
int ocl_eval_FP64(cl_device_id devid);
/**
* \brief Release N buffers referenced by oclconfig
*/
void ocl_relNbuffers_byref(ocl_config_type *oclconfig,int level);
/**
* \brief Release N kernels referenced by oclconfig
*/
void ocl_relNkernels_byref(ocl_config_type *oclconfig,int level);
/**
* \brief OpenCL sources compiler for a .cl file
*/
/* OpenCL Compiler for dynamic kernel creation. It will always report success or failure of the build.*/
int ocl_compiler(ocl_config_type *oclconfig,const char *kernelfilename,int BLOCK_SIZE,const char *optional=NULL, FILE *stream=stdout);
/**
* \brief OpenCL sources compiler for multiple .cl files
*/
int ocl_compiler(ocl_config_type *oclconfig,const char **clList,int clNum,int BLOCK_SIZE,const char *optional, FILE *stream);
/**
* \brief OpenCL sources compiler for cl string
*/
int ocl_compiler(ocl_config_type *oclconfig,unsigned char **clList,unsigned int *clLen,int clNum,int BLOCK_SIZE,const char *optional, FILE *stream);
/**
* \brief Profiler function based on OpenCL events, for display
*/
/* A simple function to get OpenCL profiler information*/
float ocl_get_profT(cl_event *start, cl_event *stop, const char *message,FILE *stream=stdout);
/**
* \brief Profiler function based on OpenCL events, only return value
*/
float ocl_get_profT(cl_event *start, cl_event *stop);
/**
* \brief Convert simple string to cl_device_type
*/
int ocl_string_to_cldevtype(const char *devicetype, cl_device_type &ocldevtype);
/**
* \brief Initialise an ocl_plat_t struct
*/
void ocl_platform_info_init(ocl_plat_t &platinfo);
/**
* \brief Release the memory held by the strings inside an ocl_plat_t struct
*/
void ocl_platform_info_del(ocl_plat_t &platinfo);
/**
* \brief Initialise an ocl_dev_t struct
*/
void ocl_device_info_init(ocl_dev_t &devinfo);
/**
* \brief Release the memory held by the strings inside an ocl_dev_t struct
*/
void ocl_device_info_del(ocl_dev_t &devinfo);
/**
* \brief Populates an ocl_plat_t struct
*/
int ocl_current_platform_info(ocl_config_type *oclconfig);
/**
* \brief Populates an ocl_dev_t struct
*/
int ocl_current_device_info(ocl_config_type *oclconfig);
/**
* \brief Translate error code to error message
*/
/* This function get OpenCL error codes and returns the appropriate string with the error name. It is
REQUIRED by the error handling macros*/
/*inline*/ const char *ocl_perrc(cl_int err);
/**
* \brief OpenCL callback function
*/
/* Opencl error function. Some Opencl functions allow pfn_notify to report errors, by passing it as pointer.
Consult the OpenCL reference card for these functions. */
void __call_compat pfn_notify(const char *errinfo, const void *private_info, size_t cb, void *user_data);
/* Basic function to handle error messages */
void ocl_errmsg(const char *userstring, const char *file, const int line);
/* Vector creation functions. OpenCL allows reintepretation and conversion, but no direct clean vector initialisation like CUDA*/
#ifdef CL_HAS_NAMED_VECTOR_FIELDS
void make_int2(int x,int y, cl_int2 &conv);
cl_int2 make_int2(int x,int y);
void make_uint2(unsigned int x,unsigned int y, cl_uint2 &conv);
cl_uint2 make_uint2(unsigned int x,unsigned int y);
void make_float2(float x,float y, cl_float2 &conv);
cl_float2 make_float2(float x,float y);
void make_double2(double x,double y, cl_double2 &conv);
cl_double2 make_double2(double x,double y);
void make_uint4(unsigned int x,unsigned int y,unsigned int z,unsigned int w,cl_uint4 &conv);
cl_uint4 make_uint4(unsigned int x,unsigned int y,unsigned int z,unsigned int w);
#endif
/**
* \file
* \brief OpenCL tools header
*
* OpenCL tools for device probe, selection, deletion, error notification
* and vector type conversion. This source is the low-level layer of our
* OpenCL Toolbox (ocl_init_exec.cpp). However, it can be used directly
* as an API
*/
/*
* Project: OpenCL tools for device probe, selection, deletion, error notification
* and vector type conversion. This source is the low-level layer of our
* OpenCL Toolbox (ocl_init_context.cpp). However, it can be used directly
* as an API
*
* Copyright (C) 2011 - 2012 European Synchrotron Radiation Facility
* Grenoble, France
*
* Principal authors: D. Karkoulis (karkouli@esrf.fr)
* Last revision: 24/06/2012
*
* This program is free software: you can redistribute it and/or modify
* it under the terms of the GNU Lesser General Public License as published
* by the Free Software Foundation, either version 3 of the License, or
* (at your option) any later version.
*
* This program is distributed in the hope that it will be useful,
* but WITHOUT ANY WARRANTY; without even the implied warranty of
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
* GNU Lesser General Public License for more details.
*
* You should have received a copy of the GNU General Public License
* and the GNU Lesser General Public License along with this program.
* If not, see <http://www.gnu.org/licenses/>.
*/
/* Header for OpenCL Utilities */
#ifndef OCLTOOLS_H
#define OCLTOOLS_H
#include <CL/opencl.h>
#include "ocl_clogger_ckerr.h"
#include "cLogger/cLogger.h"
/*This is required for OpenCL callbacks in windows*/
#ifdef _WIN32
#define __call_compat __stdcall
#ifndef _CRT_SECURE_NO_WARNINGS
#define _CRT_SECURE_NO_WARNINGS
#endif
#pragma warning(disable : 4996)
#else
#define __call_compat
#endif
/**
* \brief Maximum number of OpenCL platforms to scan
*/
#define OCL_MAX_PLATFORMS 5
/**
* \brief Maximum number of OpenCL devices-per platform- to scan
*/
#define OCL_MAX_DEVICES 5
/**
* \brief OpenCL tools platform information struct
*
* It can be passed to ocl_platform_info to
* retrieve and save platform information
* for the currect context
*/
typedef struct
{
char *name;
char *vendor;
char *version;
char *extensions;
}ocl_plat_t;
/**
* \brief OpenCL tools platform information struct
*
* It can be passed to ocl_device_info to
* retrieve and save device information
* for the currect context
*/
typedef struct
{
char *name;
char type[4];
char *version;
char *driver_version;
char *extensions;
unsigned long global_mem;
}ocl_dev_t;
/**
* \brief OpenCL tools cl_program structure
*
* Substruct of ocl_configuration_parameters.
* It is used when multiple cl programs must be
* built, i.e. from multiple OpenCL sources.
*/
typedef struct{
cl_program oclprogram;
size_t *kernelstring_lens;
char **kernelstrings;
}ocl_program_type;
/**
* \brief OpenCL tools configuration parameters
*
* OpenCL configuration structure.
* This version supports single device, but multiple
* memory buffers, kernels and sources.
*/
typedef struct ocl_configuration_parameters{
cl_context oclcontext;
cl_device_id ocldevice;
cl_platform_id oclplatform;
cl_command_queue oclcmdqueue;
cl_mem *oclmemref;
//Active device and platform info
int devid;
int platfid;
ocl_plat_t platform_info;
ocl_dev_t device_info;
//If single .cl file:
cl_program oclprogram;
size_t *kernelstring_lens;
char **kernelstrings;
//if multiple .cl files:
ocl_program_type *prgs;
int nprgs;
char compiler_options[1000];
cl_kernel *oclkernels;
int fp64;
size_t work_dim[3];
size_t thread_dim[3];
cl_event t_s[20];
cl_int event_counter;
cl_ulong dev_mem;
cl_int Nbuffers;
cl_int Nkernels;
//Logging
logger_t *hLog;
int external_cLogger;
}ocl_config_type;
/* All production functions return 0 on success, -1 on OpenCL error and -2 on other errors.
when an error is encountered internally, it will print the message to stderr and fallback.
It is important for the user to decide how to continue.*/
/**
* \brief Initialises the internals of an ocl_config_type
*/
void ocl_tools_initialise(ocl_config_type *oclconfig);
void ocl_tools_initialise(ocl_config_type *oclconfig,logger_t *hLogIN);
logger_t *ocl_tools_initialise(ocl_config_type *oclconfig, FILE *stream, const char *fname,
int severity, enum_LOGTYPE type, enum_LOGDEPTH depth, int perf,
int timestamps);
/**
* \brief Deallocations inside ocl_config_type
*/
void ocl_tools_destroy(ocl_config_type *oclconfig);
/**
* \brief Simple check all platforms and devices and print information
*/
int ocl_check_platforms(logger_t *hLog);
/**
* \brief Simple check for a "device_type" device. Returns the first occurance only
*/
int ocl_find_devicetype(cl_device_type device_type, cl_platform_id &platform, cl_device_id &devid, logger_t *hLog);
/**
* \brief Simple check for a "device_type" device. Returns the first occurance that supports double precision only
*/
int ocl_find_devicetype_FP64(cl_device_type device_type, cl_platform_id& platform, cl_device_id& devid, logger_t *hLog);
/**
* \brief Probes OpenCL platforms & devices of a given cl_device_type. Keeps the selected device in oclconfig
*/
int ocl_probe(ocl_config_type *oclconfig,cl_device_type ocldevtype, int usefp64);
/**
* \brief Probes OpenCL device of a specific cl_device_type, platform and device number. Keeps the selected device in oclconfig
*/
int ocl_probe(ocl_config_type *oclconfig,cl_device_type ocldevtype,int preset_platform,int preset_device, int usefp64);
/**
* \brief Probes OpenCL device of a specific cl_platform_id and cl_device_id. Keeps the selected device in oclconfig
*/
int ocl_probe(ocl_config_type *oclconfig,cl_platform_id platform,cl_device_id device, int usefp64);
/**
* \brief Create an OpenCL context by device type
*/
/* Needs a string with the type of device: GPU,CPU,ACC,ALL,DEF. Runs ocl_probe and creates the context,
adding it on the appropriate ocl_config_type field*/
int ocl_init_context(ocl_config_type *oclconfig,const char *device_type, int usefp64);
/**
* \brief Create an OpenCL context by device type, platform and device number
*/
int ocl_init_context(ocl_config_type *oclconfig,const char *device_type,int preset_platform,int devid, int usefp64);
/**
* \brief Create an OpenCL context by cl_platform_id and cl_device_id
*/
int ocl_init_context(ocl_config_type *oclconfig,cl_platform_id platform,cl_device_id device, int usefp64);
/**
* \brief Destroy an OpenCL context
*/
int ocl_destroy_context(cl_context oclcontext, logger_t *hLog);
/**
* \brief Queries the fp64 capability of an OpenCL device that has been selected by ocl_probe
*/
/* Queries device capabilities to figure if it meets the minimum requirement for double precision*/
/* Returns 0 on successful FP64 evaluation and -1 if only FP32 */
int ocl_eval_FP64(ocl_config_type *oclconfig);
/**
* \brief Queries the fp64 capability of an OpenCL device directly via the cl_device_id
*/
/* Same as above but directly query a device (as not set in ocl_config_type)
* It is designed to be used while probing for devices so it does not print anything
* neither it sets the fp64 field */
int ocl_eval_FP64(cl_device_id devid, logger_t *hLog);
/**
* \brief Release N buffers referenced by oclconfig
*/
void ocl_relNbuffers_byref(ocl_config_type *oclconfig,int level);
/**
* \brief Release N kernels referenced by oclconfig
*/
void ocl_relNkernels_byref(ocl_config_type *oclconfig,int level);
/**
* \brief OpenCL sources compiler for a .cl file
*/
/* OpenCL Compiler for dynamic kernel creation. It will always report success or failure of the build.*/
int ocl_compiler(ocl_config_type *oclconfig,const char *kernelfilename,int BLOCK_SIZE,const char *optional=NULL);
/**
* \brief OpenCL sources compiler for multiple .cl files
*/
int ocl_compiler(ocl_config_type *oclconfig,const char **clList,int clNum,int BLOCK_SIZE,const char *optional);
/**
* \brief OpenCL sources compiler for cl string
*/
int ocl_compiler(ocl_config_type *oclconfig,unsigned char **clList,unsigned int *clLen,int clNum,int BLOCK_SIZE,const char *optional);
/**
* \brief Profiler function based on OpenCL events, for display
*/
/* A simple function to get OpenCL profiler information*/
float ocl_get_profT(cl_event *start, cl_event *stop, const char *message, logger_t *hLog);
/**
* \brief Profiler function based on OpenCL events, only return value
*/
float ocl_get_profT(cl_event *start, cl_event *stop, logger_t *hLog);
/**
* \brief Convert simple string to cl_device_type
*/
int ocl_string_to_cldevtype(const char *devicetype, cl_device_type &ocldevtype, logger_t *hLog);
/**
* \brief Initialise an ocl_plat_t struct
*/
void ocl_platform_info_init(ocl_plat_t &platinfo);
/**
* \brief Release the memory held by the strings inside an ocl_plat_t struct
*/
void ocl_platform_info_del(ocl_plat_t &platinfo);
/**
* \brief Initialise an ocl_dev_t struct
*/
void ocl_device_info_init(ocl_dev_t &devinfo);
/**
* \brief Release the memory held by the strings inside an ocl_dev_t struct
*/
void ocl_device_info_del(ocl_dev_t &devinfo);
/**
* \brief Populates an ocl_plat_t struct
*/
int ocl_current_platform_info(ocl_config_type *oclconfig);
/**
* \brief Populates an ocl_dev_t struct
*/
int ocl_current_device_info(ocl_config_type *oclconfig);
/**
* \brief Translate error code to error message
*/
/* This function get OpenCL error codes and returns the appropriate string with the error name. It is
REQUIRED by the error handling macros*/
/*inline*/ const char *ocl_perrc(cl_int err);
/**
* \brief OpenCL callback function
*/
/* Opencl error function. Some Opencl functions allow pfn_notify to report errors, by passing it as pointer.
Consult the OpenCL reference card for these functions. */
void __call_compat pfn_notify(const char *errinfo, const void *private_info, size_t cb, void *user_data);
/* Basic function to handle error messages */
void ocl_errmsg(const char *userstring, const char *file, const int line);
/* Vector creation functions. OpenCL allows reintepretation and conversion, but no direct clean vector initialisation like CUDA*/
#ifdef CL_HAS_NAMED_VECTOR_FIELDS
void make_int2(int x,int y, cl_int2 &conv);
cl_int2 make_int2(int x,int y);
void make_uint2(unsigned int x,unsigned int y, cl_uint2 &conv);
cl_uint2 make_uint2(unsigned int x,unsigned int y);
void make_float2(float x,float y, cl_float2 &conv);
cl_float2 make_float2(float x,float y);
void make_double2(double x,double y, cl_double2 &conv);
cl_double2 make_double2(double x,double y);
void make_uint4(unsigned int x,unsigned int y,unsigned int z,unsigned int w,cl_uint4 &conv);
cl_uint4 make_uint4(unsigned int x,unsigned int y,unsigned int z,unsigned int w);
#endif
#endif

View File

@ -20,7 +20,7 @@
* Grenoble, France
*
* Principal authors: D. Karkoulis (karkouli@esrf.fr)
* Last revision: 11/05/2012
* Last revision: 24/06/2012
*
* This program is free software: you can redistribute it and/or modify
* it under the terms of the GNU Lesser General Public License as published
@ -40,18 +40,10 @@
#ifndef OCL_XRPD1D_H
#define OCL_XRPD1D_H
#ifdef _WIN32
#ifndef _CRT_SECURE_NO_WARNINGS
#define _CRT_SECURE_NO_WARNINGS
#endif
#pragma warning(disable : 4996)
#endif
#include <iostream>
#include <CL/opencl.h>
#include "ocl_ckerr.h"
#include "ocl_tools.h"
#include "ocl_tools/ocl_tools.h"
#include "ocl_base.hpp"
/**
@ -70,7 +62,10 @@ public:
ocl_xrpd1D_fullsplit();
//Prints messages on file fname
explicit ocl_xrpd1D_fullsplit(const char* fname);
explicit ocl_xrpd1D_fullsplit(const char* fname, const char *identity=NULL);
//Complete logging functionality
explicit ocl_xrpd1D_fullsplit(FILE *stream, const char *fname, int safe, int depth, int perf_time, int timestamp, const char *identity=NULL);
~ocl_xrpd1D_fullsplit();

View File

@ -20,7 +20,7 @@
* Grenoble, France
*
* Principal authors: D. Karkoulis (karkouli@esrf.fr)
* Last revision: 11/05/2012
* Last revision: 24/06/2012
*
* This program is free software: you can redistribute it and/or modify
* it under the terms of the GNU Lesser General Public License as published
@ -44,18 +44,20 @@
#include "ocl_xrpd1d.hpp"
#define CE CL_CHECK_ERR_PR
#define C CL_CHECK_PR
#define CE(_expr) CL_CHECK_ERR_PR(_expr, &hLog)
#define C(_expr) CL_CHECK_PR(_expr, &hLog)
#define CEN CL_CHECK_ERR_PRN
#define CN CL_CHECK_PRN
#define CEN(_expr) CL_CHECK_ERR_PRN(_expr, &hLog)
#define CN(_expr) CL_CHECK_PRN(_expr, &hLog)
#define CER CL_CHECK_ERR_PR_RET
#define CR CL_CHECK_PR_RET
#define CER(_expr) CL_CHECK_ERR_PR_RET(_expr, &hLog)
#define CR(_expr) CL_CHECK_PR_RET(_expr, &hLog)
//#define silent
#ifdef _SILENT
#define fprintf(stream,...)
#ifdef _WIN32
#ifndef _CRT_SECURE_NO_WARNINGS
#define _CRT_SECURE_NO_WARNINGS
#endif
#pragma warning(disable : 4996)
#endif
#define BLOCK_SIZE 128
@ -117,14 +119,29 @@ ocl_xrpd1D_fullsplit::ocl_xrpd1D_fullsplit():ocl()
}
/**
* \brief Overloaded constructor for xrpd1d.
* \brief Overloaded constructor for xrpd1d with filename option.
*
* Output is set to filename "fname" and the docstring is set
*
* @param fname A const C-string with the name of the textfile to use as output
*
*/
ocl_xrpd1D_fullsplit::ocl_xrpd1D_fullsplit(const char* fname):ocl(fname)
ocl_xrpd1D_fullsplit::ocl_xrpd1D_fullsplit(const char* fname, const char* identity):ocl(fname, identity)
{
setDocstring("OpenCL 1d Azimuthal integrator. Check the readme file for more details\n","ocl_xrpd1d_fullsplit.readme");
}
/**
* \brief Overloaded constructor for xrpd1d with logging options.
*
* Output is set to filename "fname" and the docstring is set.
* cLogger is set according to arguments
*
* @param fname A const C-string with the name of the textfile to use as output
*
*/
ocl_xrpd1D_fullsplit::ocl_xrpd1D_fullsplit(FILE *stream, const char *fname, int safe, int depth, int perf_time, int timestamp, const char* identity):
ocl(stream, fname, safe, depth, perf_time, timestamp, identity)
{
setDocstring("OpenCL 1d Azimuthal integrator. Check the readme file for more details\n","ocl_xrpd1d_fullsplit.readme");
}
@ -159,11 +176,11 @@ int ocl_xrpd1D_fullsplit::getConfiguration(const int Nx,const int Nimage,const i
if(Nx < 1 || Nimage <1 || Nbins<1){
fprintf(stderr,"get_azim_args() parameters make no sense {%d %d %d}\n",Nx,Nimage,Nbins);
cLog_critical(&hLog,"get_azim_args() parameters make no sense {%d %d %d}\n",Nx,Nimage,Nbins);
return -2;
}
if(!(this->sgs)){
ocl_errmsg("Fatal error in get_azim_args(). Cannot allocate argument structure",__FILE__,__LINE__);
cLog_critical(&hLog,"Fatal error in get_azim_args(). Cannot allocate argument structure (%s:%d)\n",__FILE__,__LINE__);
return -1;
} else {
this->sgs->Nimage = Nimage;
@ -192,11 +209,11 @@ int ocl_xrpd1D_fullsplit::configure(const char* kernel_path)
{
if(!sgs->Nx || !sgs->Nimage || !sgs->Nbins){
fprintf(stderr,"You may not call config() at this point. Image and histogram parameters not set. (Hint: run get_azim_args())\n");
cLog_critical(&hLog,"You may not call config() at this point. Image and histogram parameters not set. (Hint: run get_azim_args())\n");
return -2;
}
if(!hasActiveContext){
fprintf(stderr,"You may not call config() at this point. There is no Active context. (Hint: run init())\n");
cLog_critical(&hLog,"You may not call config() at this point. There is no Active context. (Hint: run init())\n");
return -2;
}
@ -208,7 +225,7 @@ int ocl_xrpd1D_fullsplit::configure(const char* kernel_path)
//Next step after the creation of a context is to create a command queue. After this step we can enqueue command to the device
// such as memory copies, arguments, kernels etc.
oclconfig->oclcmdqueue = clCreateCommandQueue(oclconfig->oclcontext,oclconfig->ocldevice,CL_QUEUE_PROFILING_ENABLE,&err);
if(err){fprintf(stderr,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};
if(err){cLog_critical(&hLog,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};
hasQueue =1;
//Allocate device memory
@ -226,41 +243,41 @@ int ocl_xrpd1D_fullsplit::configure(const char* kernel_path)
sprintf(optional," -D BINS=%d -D NX=%u -D NN=%u -D ENABLE_FP64",sgs->Nbins,sgs->Nx,sgs->Nimage);
//The blocksize itself is set by the compiler function explicitly and then appends the string "optional"
fprintf(stream,"Will use kernel %s\n",kernel_path);
if(ocl_compiler(oclconfig,kernel_path,BLOCK_SIZE,optional,stream))return -1;
cLog_debug(&hLog,"Will use kernel %s\n",kernel_path);
if(ocl_compiler(oclconfig,kernel_path,BLOCK_SIZE,optional) )return -1;
hasProgram=1;
oclconfig->oclkernels = (cl_kernel*)malloc(8*sizeof(cl_kernel));
if(!oclconfig->oclkernels){
ocl_errmsg("Fatal error in ocl_config. Cannot allocate kernels",__FILE__,__LINE__);
cLog_critical(&hLog,"Fatal error in ocl_config. Cannot allocate kernels (%s:%d)",__FILE__,__LINE__);
return -2;
}
//Create the OpenCL kernels found in the compile OpenCL program
int i=0;
oclconfig->oclkernels[CLKERN_INTEGRATE] = clCreateKernel(oclconfig->oclprogram,"create_histo_binarray",&err);
if(err){fprintf(stderr,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
oclconfig->oclkernels[CLKERN_UIMEMSET2] = clCreateKernel(oclconfig->oclprogram,"uimemset2",&err);
if(err){fprintf(stderr,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
oclconfig->oclkernels[CLKERN_IMEMSET] = clCreateKernel(oclconfig->oclprogram,"imemset",&err);
if(err){fprintf(stderr,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
oclconfig->oclkernels[CLKERN_UI2F2] = clCreateKernel(oclconfig->oclprogram,"ui2f2",&err);
if(err){fprintf(stderr,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
oclconfig->oclkernels[CLKERN_GET_SPANS] = clCreateKernel(oclconfig->oclprogram,"get_spans",&err);
if(err){fprintf(stderr,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
oclconfig->oclkernels[CLKERN_GROUP_SPANS] = clCreateKernel(oclconfig->oclprogram,"group_spans",&err);
if(err){fprintf(stderr,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
oclconfig->oclkernels[CLKERN_SOLIDANGLE_CORRECTION] = clCreateKernel(oclconfig->oclprogram,"solidangle_correction",&err);
if(err){fprintf(stderr,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
oclconfig->oclkernels[CLKERN_DUMMYVAL_CORRECTION] = clCreateKernel(oclconfig->oclprogram,"dummyval_correction",&err);
if(err){fprintf(stderr,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
oclconfig->Nkernels=i;
hasKernels = 1;
@ -275,7 +292,7 @@ int ocl_xrpd1D_fullsplit::configure(const char* kernel_path)
CR(
clEnqueueNDRangeKernel(oclconfig->oclcmdqueue,oclconfig->oclkernels[CLKERN_IMEMSET],1,0,wdim,tdim,0,0,&oclconfig->t_s[0]) );
execTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[0], "Initialise Mask to 0", stream);
execTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[0], "Initialise Mask to 0", &hLog);
clReleaseEvent(oclconfig->t_s[0]);
return 0;
@ -297,18 +314,18 @@ return 0;
int ocl_xrpd1D_fullsplit::loadTth(float* tth, float* dtth, float tth_min, float tth_max)
{
fprintf(stream,"Loading Tth\n");
cLog_extended(&hLog,"Loading Tth\n");
float tthmm[2];
tthmm[0]=tth_min;
tthmm[1]=tth_max;
if(!hasActiveContext){
fprintf(stderr,"You may not call loadTth() at this point. There is no Active context. (Hint: run init())\n");
cLog_critical(&hLog,"You may not call loadTth() at this point. There is no Active context. (Hint: run init())\n");
return -2;
}
if(!oclconfig->Nbuffers || !isConfigured){
fprintf(stderr,"You may not call loadTth() at this point, OpenCL is not configured (Hint: run configure())\n");
cLog_critical(&hLog,"You may not call loadTth() at this point, OpenCL is not configured (Hint: run configure())\n");
return -2;
}
@ -321,7 +338,7 @@ int ocl_xrpd1D_fullsplit::loadTth(float* tth, float* dtth, float tth_min, float
CR(
clEnqueueWriteBuffer(oclconfig->oclcmdqueue,oclconfig->oclmemref[CLMEM_TTH_MIN_MAX],CL_TRUE,0,2*sizeof(cl_float),(void*)tthmm,0,0,&oclconfig->t_s[2]) );
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[2],"Load Tth",stream);
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[2],"Load Tth",oclconfig->hLog);
clReleaseEvent(oclconfig->t_s[0]);
clReleaseEvent(oclconfig->t_s[1]);
clReleaseEvent(oclconfig->t_s[2]);
@ -346,22 +363,22 @@ int ocl_xrpd1D_fullsplit::loadTth(float* tth, float* dtth, float tth_min, float
int ocl_xrpd1D_fullsplit::setSolidAngle(float *SolidAngle)
{
fprintf(stream,"Setting SolidAngle\n");
cLog_extended(&hLog,"Setting SolidAngle\n");
if(!oclconfig->Nbuffers || !isConfigured){
fprintf(stderr,"You may not call setSolidAngle() at this point, the required buffers are not allocated (Hint: run config())\n");
cLog_critical(&hLog,"You may not call setSolidAngle() at this point, the required buffers are not allocated (Hint: run config())\n");
return -2;
}
if(!hasActiveContext){
fprintf(stderr,"You may not call setSolidAngle() at this point. There is no Active context. (Hint: run init())\n");
cLog_critical(&hLog,"You may not call setSolidAngle() at this point. There is no Active context. (Hint: run init())\n");
return -2;
}
CR(
clEnqueueWriteBuffer(oclconfig->oclcmdqueue,oclconfig->oclmemref[CLMEM_SOLIDANGLE],CL_TRUE,0,sgs->Nimage*sizeof(cl_float),(void*)SolidAngle,0,0,&oclconfig->t_s[0]) );
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[0],"Load SolidAngle",stream);
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[0],"Load SolidAngle",oclconfig->hLog);
clReleaseEvent(oclconfig->t_s[0]);
useSolidAngle=1;
@ -377,7 +394,7 @@ int ocl_xrpd1D_fullsplit::setSolidAngle(float *SolidAngle)
*/
int ocl_xrpd1D_fullsplit::unsetSolidAngle()
{
fprintf(stream,"Unsetting SolidAngle\n");
cLog_extended(&hLog,"Unsetting SolidAngle\n");
if(useSolidAngle)
{
@ -399,22 +416,22 @@ int ocl_xrpd1D_fullsplit::unsetSolidAngle()
*/
int ocl_xrpd1D_fullsplit::setMask(int* Mask)
{
fprintf(stream,"Setting Mask\n");
cLog_extended(&hLog,"Setting Mask\n");
if(!oclconfig->Nbuffers || !isConfigured){
fprintf(stderr,"You may not call setMask() at this point, the required buffers are not allocated (Hint: run config())\n");
cLog_critical(&hLog,"You may not call setMask() at this point, the required buffers are not allocated (Hint: run config())\n");
return -2;
}
if(!hasActiveContext){
fprintf(stderr,"You may not call setMask() at this point. There is no Active context. (Hint: run init())\n");
cLog_critical(&hLog,"You may not call setMask() at this point. There is no Active context. (Hint: run init())\n");
return -2;
}
CR(
clEnqueueWriteBuffer(oclconfig->oclcmdqueue,oclconfig->oclmemref[CLMEM_MASK],CL_TRUE,0,sgs->Nimage*sizeof(cl_int),(void*)Mask,0,0,&oclconfig->t_s[0]) );
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[0],"Load Mask",stream);
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[0],"Load Mask",oclconfig->hLog);
clReleaseEvent(oclconfig->t_s[0]);
useMask=1;
@ -430,7 +447,7 @@ int ocl_xrpd1D_fullsplit::setMask(int* Mask)
*/
int ocl_xrpd1D_fullsplit::unsetMask()
{
fprintf(stream,"Unsetting Mask\n");
cLog_extended(&hLog,"Unsetting Mask\n");
if(useMask)
{
@ -439,7 +456,7 @@ int ocl_xrpd1D_fullsplit::unsetMask()
CR(
clEnqueueNDRangeKernel(oclconfig->oclcmdqueue,oclconfig->oclkernels[CLKERN_IMEMSET],1,0,wdim,tdim,0,0,&oclconfig->t_s[0]) );
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[0],"Reset Mask to 0");
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[0],"Reset Mask to 0", oclconfig->hLog);
clReleaseEvent(oclconfig->t_s[0]);
useMask=0;
@ -463,15 +480,15 @@ int ocl_xrpd1D_fullsplit::unsetMask()
*/
int ocl_xrpd1D_fullsplit::setDummyValue(float dummyVal, float deltaDummyVal)
{
fprintf(stream,"Setting Dummy Value\n");
cLog_extended(&hLog,"Setting Dummy Value\n");
if(!oclconfig->Nbuffers || !isConfigured){
fprintf(stderr,"You may not call setDummyValue() at this point, the required buffers are not allocated (Hint: run configure())\n");
cLog_critical(&hLog,"You may not call setDummyValue() at this point, the required buffers are not allocated (Hint: run configure())\n");
return -2;
}
if(!hasActiveContext){
fprintf(stderr,"You may not call setDummyValue() at this point. There is no Active context. (Hint: run init())\n");
cLog_critical(&hLog,"You may not call setDummyValue() at this point. There is no Active context. (Hint: run init())\n");
return -2;
}
@ -480,7 +497,7 @@ int ocl_xrpd1D_fullsplit::setDummyValue(float dummyVal, float deltaDummyVal)
CR(
clEnqueueWriteBuffer(oclconfig->oclcmdqueue,oclconfig->oclmemref[CLMEM_DUMMYVAL_DELTA],CL_TRUE,0,sizeof(cl_float),(void*)&deltaDummyVal,0,0,&oclconfig->t_s[1]) );
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[1],"Load Dummy Value",stream);
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[1],"Load Dummy Value", oclconfig->hLog);
clReleaseEvent(oclconfig->t_s[0]);
clReleaseEvent(oclconfig->t_s[1]);
@ -497,7 +514,7 @@ int ocl_xrpd1D_fullsplit::setDummyValue(float dummyVal, float deltaDummyVal)
*/
int ocl_xrpd1D_fullsplit::unsetDummyValue()
{
fprintf(stream,"Unsetting Dummy Value\n");
cLog_extended(&hLog,"Unsetting Dummy Value\n");
if(useDummyVal)
{
@ -525,18 +542,18 @@ int ocl_xrpd1D_fullsplit::unsetDummyValue()
*/
int ocl_xrpd1D_fullsplit::setRange(float lowerBound, float upperBound)
{
fprintf(stream,"Setting 2th Range\n");
cLog_extended(&hLog,"Setting 2th Range\n");
float tthrmm[2];
tthrmm[0]=lowerBound;
tthrmm[1]=upperBound;
if(!oclconfig->Nbuffers || !isConfigured){
fprintf(stderr,"You may not call setMask() at this point, the required buffers are not allocated (Hint: run config())\n");
cLog_critical(&hLog,"You may not call setMask() at this point, the required buffers are not allocated (Hint: run config())\n");
return -2;
}
if(!hasActiveContext){
fprintf(stderr,"You may not call setMask() at this point. There is no Active context. (Hint: run init())\n");
cLog_critical(&hLog,"You may not call setMask() at this point. There is no Active context. (Hint: run init())\n");
return -2;
}
@ -547,7 +564,7 @@ int ocl_xrpd1D_fullsplit::setRange(float lowerBound, float upperBound)
CR( clSetKernelArg(oclconfig->oclkernels[CLKERN_INTEGRATE],8,sizeof(cl_mem),&oclconfig->oclmemref[CLMEM_TTH_RANGE]) ); //TTH range user values
CR( clSetKernelArg(oclconfig->oclkernels[CLKERN_GET_SPANS],2,sizeof(cl_mem),&oclconfig->oclmemref[CLMEM_TTH_RANGE]) ); //TTH range user values
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[0],"Load 2th Range",stream);
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[0],"Load 2th Range", oclconfig->hLog);
clReleaseEvent(oclconfig->t_s[0]);
useTthRange=1;
@ -562,7 +579,7 @@ int ocl_xrpd1D_fullsplit::setRange(float lowerBound, float upperBound)
*/
int ocl_xrpd1D_fullsplit::unsetRange()
{
fprintf(stream,"Unsetting 2th Range\n");
cLog_extended(&hLog,"Unsetting 2th Range\n");
if(useTthRange)
{
@ -594,35 +611,35 @@ int ocl_xrpd1D_fullsplit::execute(float *im_inten,float *histogram,float *bins)
{
if(!isConfigured){
fprintf(stderr,"You may not call execute() at this point, kernels are not configured (Hint: run config())\n");
cLog_critical(&hLog,"You may not call execute() at this point, kernels are not configured (Hint: run config())\n");
return -2;
}
if(!hasActiveContext){
fprintf(stderr,"You may not call execute() at this point. There is no Active context. (Hint: run init())\n");
cLog_critical(&hLog,"You may not call execute() at this point. There is no Active context. (Hint: run init())\n");
return -2;
}
if(!hasTthLoaded){
fprintf(stderr,"You may not call execute() at this point. There is no 2th array loaded. (Hint: run loadTth())\n");
cLog_critical(&hLog,"You may not call execute() at this point. There is no 2th array loaded. (Hint: run loadTth())\n");
return -2;
}
//Setup the kernel execution parameters, grid,blocks and threads.
// Notice that in CUDA, a grid is measured in blocks, while in OpenCL is measured in threads.
fprintf(stream,"\n--Integration nr. %d\n",get_exec_count() + 1);
cLog_extended(&hLog,"\n--Integration nr. %d\n",get_exec_count() + 1);
size_t wdim_partialh[] = { (sgs->Nimage/BLOCK_SIZE) * BLOCK_SIZE + (sgs->Nimage%BLOCK_SIZE) * BLOCK_SIZE, 1, 1};
size_t tdim_partialh[] = {BLOCK_SIZE, 1, 1};
size_t wdim_reduceh[] = { (sgs->Nbins/BLOCK_SIZE) * BLOCK_SIZE + (sgs->Nbins%BLOCK_SIZE) * BLOCK_SIZE, 1, 1};
size_t tdim_reduceh[] = {BLOCK_SIZE, 1, 1};
fprintf(stream,"--Histo / Spans workdim %lu %lu %lu\n",(lui)wdim_partialh[0],(lui)wdim_partialh[1],(lui)wdim_partialh[2]);
fprintf(stream,"--Histo / Spans threadim %lu %lu %lu -- Blocks:%lu\n",(lui)tdim_partialh[0],(lui)tdim_partialh[1],(lui)tdim_partialh[2],\
cLog_debug(&hLog,"--Histo / Spans workdim %lu %lu %lu\n",(lui)wdim_partialh[0],(lui)wdim_partialh[1],(lui)wdim_partialh[2]);
cLog_debug(&hLog,"--Histo / Spans threadim %lu %lu %lu -- Blocks:%lu\n",(lui)tdim_partialh[0],(lui)tdim_partialh[1],(lui)tdim_partialh[2],\
(lui)wdim_partialh[0]/(lui)tdim_partialh[0]);
fprintf(stream,"--Memset / Convert workdim %lu %lu %lu\n",(lui)wdim_reduceh[0],(lui)wdim_reduceh[1],(lui)wdim_reduceh[2]);
fprintf(stream,"--Memset / Convert threadim %lu %lu %lu -- Blocks:%lu\n",(lui)tdim_reduceh[0],(lui)tdim_reduceh[1],(lui)tdim_reduceh[2],\
cLog_debug(&hLog,"--Memset / Convert workdim %lu %lu %lu\n",(lui)wdim_reduceh[0],(lui)wdim_reduceh[1],(lui)wdim_reduceh[2]);
cLog_debug(&hLog,"--Memset / Convert threadim %lu %lu %lu -- Blocks:%lu\n",(lui)tdim_reduceh[0],(lui)tdim_reduceh[1],(lui)tdim_reduceh[2],\
(lui)wdim_reduceh[0]/(lui)tdim_reduceh[0]);
@ -670,24 +687,24 @@ int ocl_xrpd1D_fullsplit::execute(float *im_inten,float *histogram,float *bins)
CR(
clEnqueueReadBuffer(oclconfig->oclcmdqueue,oclconfig->oclmemref[CLMEM_HISTOGRAM],CL_TRUE,0,sgs->Nbins*sizeof(cl_float),(void*)histogram,0,0,&oclconfig->t_s[7]) );
fprintf(stream,"--Waiting for the command queue to finish\n");
cLog_debug(&hLog,"--Waiting for the command queue to finish\n");
CR(clFinish(oclconfig->oclcmdqueue));
//Get execution time from first memory copy to last memory copy.
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[0],"copyIn ",stream);
execTime_ms += ocl_get_profT(&oclconfig->t_s[1], &oclconfig->t_s[1], "memset ",stream);
execTime_ms += ocl_get_profT(&oclconfig->t_s[2], &oclconfig->t_s[2], "getSpa ",stream);
execTime_ms += ocl_get_profT(&oclconfig->t_s[3], &oclconfig->t_s[3], "groupS ",stream);
execTime_ms += ocl_get_profT(&oclconfig->t_s[4], &oclconfig->t_s[4], "Azim GPU ",stream);
execTime_ms += ocl_get_profT(&oclconfig->t_s[5], &oclconfig->t_s[5], "convert ",stream);
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[0],"copyIn ",oclconfig->hLog);
execTime_ms += ocl_get_profT(&oclconfig->t_s[1], &oclconfig->t_s[1], "memset ",oclconfig->hLog);
execTime_ms += ocl_get_profT(&oclconfig->t_s[2], &oclconfig->t_s[2], "getSpa ",oclconfig->hLog);
execTime_ms += ocl_get_profT(&oclconfig->t_s[3], &oclconfig->t_s[3], "groupS ",oclconfig->hLog);
execTime_ms += ocl_get_profT(&oclconfig->t_s[4], &oclconfig->t_s[4], "Azim GPU ",oclconfig->hLog);
execTime_ms += ocl_get_profT(&oclconfig->t_s[5], &oclconfig->t_s[5], "convert ",oclconfig->hLog);
if(useSolidAngle)
execTime_ms += ocl_get_profT(&oclconfig->t_s[8], &oclconfig->t_s[8],"Solidan ",stream);
execTime_ms += ocl_get_profT(&oclconfig->t_s[8], &oclconfig->t_s[8],"Solidan ",oclconfig->hLog);
if(useDummyVal)
execTime_ms += ocl_get_profT(&oclconfig->t_s[9], &oclconfig->t_s[9],"dummyva ",stream);
execTime_ms += ocl_get_profT(&oclconfig->t_s[9], &oclconfig->t_s[9],"dummyva ",oclconfig->hLog);
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[6], &oclconfig->t_s[7],"copyOut ",stream);
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[6], &oclconfig->t_s[7],"copyOut ",oclconfig->hLog);
execCount++;
@ -721,12 +738,12 @@ int ocl_xrpd1D_fullsplit::allocate_CL_buffers()
cl_int err;
oclconfig->oclmemref = (cl_mem*)malloc(14*sizeof(cl_mem));
if(!oclconfig->oclmemref){
fprintf(stderr,"Fatal error in allocate_CL_buffers. Cannot allocate memrefs\n");
cLog_critical(&hLog,"Fatal error in allocate_CL_buffers. Cannot allocate memrefs\n");
return -2;
}
if(sgs->Nimage < BLOCK_SIZE){
fprintf(stderr,"Fatal error in allocate_CL_buffers. Nimage (%d) must be >= BLOCK_SIZE (%d)\n",sgs->Nimage,BLOCK_SIZE);
cLog_critical(&hLog,"Fatal error in allocate_CL_buffers. Nimage (%d) must be >= BLOCK_SIZE (%d)\n",sgs->Nimage,BLOCK_SIZE);
return -2;
}
@ -748,12 +765,12 @@ int ocl_xrpd1D_fullsplit::allocate_CL_buffers()
*/
if(ualloc >= oclconfig->dev_mem && oclconfig->dev_mem != 0){
fprintf(stderr,"Fatal error in allocate_CL_buffers. Not enough device memory for buffers (%lu requested, %lu available)\n",\
cLog_critical(&hLog,"Fatal error in allocate_CL_buffers. Not enough device memory for buffers (%lu requested, %lu available)\n",\
(lui)ualloc,(lui)oclconfig->dev_mem);
return -1;
} else {
if(oclconfig->dev_mem == 0){
fprintf(stream,"Caution: Device did not return the available memory size (%lu requested)\n",(lui)ualloc);
cLog_extended(&hLog,"Caution: Device did not return the available memory size (%lu requested)\n",(lui)ualloc);
}
}
@ -762,77 +779,77 @@ int ocl_xrpd1D_fullsplit::allocate_CL_buffers()
int i=0;
oclconfig->oclmemref[CLMEM_TTH]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_ONLY,(size_t)(sgs->Nimage*sizeof(cl_float)),0,&err);//tth array -0
if(err){fprintf(stderr,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);return -1;};i++;
oclconfig->oclmemref[CLMEM_IMAGE]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_ONLY,(size_t)(sgs->Nimage*sizeof(cl_float)),0,&err);//Image intensity -1
if(err){fprintf(stderr,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
oclconfig->oclmemref[CLMEM_SOLIDANGLE]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_ONLY,(size_t)(sgs->Nimage*sizeof(cl_float)),0,&err);//Solid Angle -2
if(err){fprintf(stderr,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
oclconfig->oclmemref[CLMEM_HISTOGRAM]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_WRITE,(size_t)(sgs->Nbins*sizeof(cl_float)),0,&err);//Histogram -3
if(err){fprintf(stderr,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
if(sgs->usefp64)
{
oclconfig->oclmemref[CLMEM_UHISTOGRAM]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_WRITE,(size_t)(sgs->Nbins*sizeof(cl_ulong)),0,&err);//ulHistogram -4
if(err){fprintf(stderr,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
}else
{
oclconfig->oclmemref[CLMEM_UHISTOGRAM]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_WRITE,(size_t)(sgs->Nbins*sizeof(cl_uint)),0,&err);//ulHistogram -4
if(err){fprintf(stderr,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
}
oclconfig->oclmemref[CLMEM_WEIGHTS]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_WRITE,(size_t)(sgs->Nbins*sizeof(cl_float)),0,&err);//Bin array -5
if(err){fprintf(stderr,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
if(sgs->usefp64)
{
oclconfig->oclmemref[CLMEM_UWEIGHTS]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_WRITE,(size_t)(sgs->Nbins*sizeof(cl_ulong)),0,&err);//uBinarray -6
if(err){fprintf(stderr,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
}else
{
oclconfig->oclmemref[CLMEM_UWEIGHTS]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_WRITE,(size_t)(sgs->Nbins*sizeof(cl_uint)),0,&err);//uBinarray -6
if(err){fprintf(stderr,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
}
oclconfig->oclmemref[CLMEM_SPAN_RANGES]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_WRITE,(size_t)((sgs->Nimage)*sizeof(cl_float)),0,&err);//span_ranges buffer -7
if(err){fprintf(stderr,"clCreateKernel error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateKernel error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
oclconfig->oclmemref[CLMEM_TTH_MIN_MAX]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_ONLY,(size_t)(2*sizeof(cl_float)),0,&err);//Min,Max values for tth -8
if(err){fprintf(stderr,"clCreateKernel error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateKernel error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
oclconfig->oclmemref[CLMEM_TTH_DELTA]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_ONLY,(size_t)(sgs->Nimage*sizeof(cl_float)),0,&err);//tth delta -9
if(err){fprintf(stderr,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);return -1;};i++;
oclconfig->oclmemref[CLMEM_MASK]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_ONLY,(size_t)(sgs->Nimage*sizeof(cl_int)),0,&err);//Mask -10
if(err){fprintf(stderr,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);return -1;};i++;
oclconfig->oclmemref[CLMEM_DUMMYVAL]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_ONLY,(size_t)(sizeof(cl_float)),0,&err);//Dummy Value -11
if(err){fprintf(stderr,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);return -1;};i++;
oclconfig->oclmemref[CLMEM_DUMMYVAL_DELTA]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_ONLY,(size_t)(sizeof(cl_float)),0,&err);//Dummy Value Delta -12
if(err){fprintf(stderr,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);return -1;};i++;
oclconfig->oclmemref[CLMEM_TTH_RANGE]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_ONLY,(size_t)(2*sizeof(cl_float)),0,&err);//TTH Range -13
if(err){fprintf(stderr,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);return -1;};i++;
fprintf(stream,"Allocated %d buffers (%.3f Mb) on device\n",i,(float)ualloc/1024./1024.);
cLog_extended(&hLog,"Allocated %d buffers (%.3f Mb) on device\n",i,(float)ualloc/1024./1024.);
oclconfig->Nbuffers = i;
return 0;
}
@ -926,7 +943,7 @@ int ocl_xrpd1D_fullsplit::clean(int preserve_context)
if(hasBuffers)
{
clean_clbuffers(oclconfig->Nbuffers);
fprintf(stream,"--released OpenCL buffers\n");
cLog_debug(&hLog,"--released OpenCL buffers\n");
hasBuffers = 0;
hasTthLoaded = 0;
useSolidAngle = 0;
@ -938,21 +955,21 @@ int ocl_xrpd1D_fullsplit::clean(int preserve_context)
if(hasKernels)
{
clean_clkernels(oclconfig->Nkernels);
fprintf(stream,"--released OpenCL kernels\n");
cLog_debug(&hLog,"--released OpenCL kernels\n");
hasKernels=0;
}
if(hasProgram)
{
CR(clReleaseProgram(oclconfig->oclprogram));
fprintf(stream,"--released OpenCL program\n");
cLog_debug(&hLog,"--released OpenCL program\n");
hasProgram=0;
}
if(hasQueue)
{
CR(clReleaseCommandQueue(oclconfig->oclcmdqueue));
fprintf(stream,"--released OpenCL queue\n");
cLog_debug(&hLog,"--released OpenCL queue\n");
hasQueue=0;
}
@ -966,18 +983,18 @@ int ocl_xrpd1D_fullsplit::clean(int preserve_context)
{
free(oclconfig->oclmemref);
oclconfig->oclmemref=NULL;
fprintf(stream,"--released OpenCL memory references\n");
cLog_debug(&hLog,"--released OpenCL memory references\n");
}
if(oclconfig->oclkernels)
{
free(oclconfig->oclkernels);
oclconfig->oclkernels=NULL;
fprintf(stream,"--released OpenCL kernel references\n");
cLog_debug(&hLog,"--released OpenCL kernel references\n");
}
if(hasActiveContext){
ocl_destroy_context(oclconfig->oclcontext);
ocl_destroy_context(oclconfig->oclcontext, &hLog);
hasActiveContext=0;
fprintf(stream,"--released OpenCL context\n");
cLog_debug(&hLog,"--released OpenCL context\n");
}
}
return 0;

View File

@ -20,7 +20,7 @@
* Grenoble, France
*
* Principal authors: D. Karkoulis (karkouli@esrf.fr)
* Last revision: 11/05/2012
* Last revision: 24/06/2012
*
* This program is free software: you can redistribute it and/or modify
* it under the terms of the GNU Lesser General Public License as published
@ -40,15 +40,10 @@
#ifndef OCL_XRPD2D_H
#define OCL_XRPD2D_H
#ifdef _WIN32
#define _CRT_SECURE_NO_WARNINGS 1
#endif
#include <iostream>
#include <CL/opencl.h>
#include "ocl_ckerr.h"
#include "ocl_tools.h"
#include "ocl_tools/ocl_tools.h"
#include "ocl_base.hpp"
typedef unsigned long lui;
@ -57,7 +52,8 @@ class ocl_xrpd2D_fullsplit: public ocl{
public:
ocl_xrpd2D_fullsplit();
explicit ocl_xrpd2D_fullsplit(const char* fname);
explicit ocl_xrpd2D_fullsplit(const char* fname, const char *identity=NULL);
explicit ocl_xrpd2D_fullsplit(FILE *stream, const char *fname, int safe, int depth, int perf_time, int timestamp, const char *identity=NULL);
~ocl_xrpd2D_fullsplit();
int getConfiguration(const int Nx,const int Nimage,const int NbinsTth, const int NbinsChi,const bool usefp64=false);

View File

@ -20,7 +20,7 @@
* Grenoble, France
*
* Principal authors: D. Karkoulis (karkouli@esrf.fr)
* Last revision: 11/05/2012
* Last revision: 24/06/2012
*
* This program is free software: you can redistribute it and/or modify
* it under the terms of the GNU Lesser General Public License as published
@ -43,22 +43,20 @@
#include "ocl_xrpd2d.hpp"
#define CE CL_CHECK_ERR_PR
#define C CL_CHECK_PR
#define CE(_expr) CL_CHECK_ERR_PR(_expr, &hLog)
#define C(_expr) CL_CHECK_PR(_expr, &hLog)
#define CEN CL_CHECK_ERR_PRN
#define CN CL_CHECK_PRN
#define CEN(_expr) CL_CHECK_ERR_PRN(_expr, &hLog)
#define CN(_expr) CL_CHECK_PRN(_expr, &hLog)
#define CER CL_CHECK_ERR_PR_RET
#define CR CL_CHECK_PR_RET
#define CER(_expr) CL_CHECK_ERR_PR_RET(_expr, &hLog)
#define CR(_expr) CL_CHECK_PR_RET(_expr, &hLog)
#ifdef _WIN32
#define _CRT_SECURE_NO_WARNINGS 1
#endif
//#define silent
#ifdef _SILENT
#define fprintf(stream,...)
#ifndef _CRT_SECURE_NO_WARNINGS
#define _CRT_SECURE_NO_WARNINGS
#endif
#pragma warning(disable : 4996)
#endif
#define BLOCK_SIZE 128
@ -102,7 +100,13 @@ ocl_xrpd2D_fullsplit::ocl_xrpd2D_fullsplit():ocl()
setDocstring("OpenCL 2d Azimuthal integrator. Check the readme file for more details\n","ocl_xrpd2d_fullsplit.readme");
}
ocl_xrpd2D_fullsplit::ocl_xrpd2D_fullsplit(const char* fname):ocl(fname)
ocl_xrpd2D_fullsplit::ocl_xrpd2D_fullsplit(const char* fname, const char *identity):ocl(fname, identity)
{
setDocstring("OpenCL 2d Azimuthal integrator. Check the readme file for more details\n","ocl_xrpd2d_fullsplit.readme");
}
ocl_xrpd2D_fullsplit::ocl_xrpd2D_fullsplit(FILE *stream, const char *fname, int safe, int depth, int perf_time, int timestamp, const char* identity):
ocl(stream, fname, safe, depth, perf_time, timestamp, identity)
{
setDocstring("OpenCL 2d Azimuthal integrator. Check the readme file for more details\n","ocl_xrpd2d_fullsplit.readme");
}
@ -119,11 +123,11 @@ int ocl_xrpd2D_fullsplit::getConfiguration(const int Nx,const int Nimage,const i
if(Nx < 1 || Nimage < 1 || NbinsTth < 1 || NbinsChi < 1){
fprintf(stderr,"get_azim_args() parameters make no sense {%d %d %d %d}\n",Nx,Nimage,NbinsTth,NbinsChi);
cLog_critical(&hLog,"get_azim_args() parameters make no sense {%d %d %d %d}\n",Nx,Nimage,NbinsTth,NbinsChi);
return -2;
}
if(!(this->sgs)){
ocl_errmsg("Fatal error in get_azim_args(). Cannot allocate argument structure",__FILE__,__LINE__);
cLog_critical(&hLog,"Fatal error in get_azim_args(). Cannot allocate argument structure (%s:%d)\n",__FILE__,__LINE__);
return -1;
} else {
this->sgs->Nimage = Nimage;
@ -143,11 +147,11 @@ int ocl_xrpd2D_fullsplit::configure()
//using namespace ocl_xrpd2D_fullsplit;
if(!sgs->Nx || !sgs->Nimage || !sgs->Nbinst || !sgs->Nbinsc){
fprintf(stderr,"You may not call config() at this point. Image and histogram parameters not set. (Hint: run get_azim_args())\n");
cLog_critical(&hLog,"You may not call config() at this point. Image and histogram parameters not set. (Hint: run get_azim_args())\n");
return -2;
}
if(!hasActiveContext){
fprintf(stderr,"You may not call config() at this point. There is no Active context. (Hint: run init())\n");
cLog_critical(&hLog,"You may not call config() at this point. There is no Active context. (Hint: run init())\n");
return -2;
}
@ -159,7 +163,7 @@ int ocl_xrpd2D_fullsplit::configure()
//Next step after the creation of a context is to create a command queue. After this step we can enqueue command to the device
// such as memory copies, arguments, kernels etc.
oclconfig->oclcmdqueue = clCreateCommandQueue(oclconfig->oclcontext,oclconfig->ocldevice,CL_QUEUE_PROFILING_ENABLE,&err);
if(err){fprintf(stderr,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};
if(err){cLog_critical(&hLog,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};
hasQueue =1;
//Allocate device memory
@ -179,40 +183,40 @@ int ocl_xrpd2D_fullsplit::configure()
//The blocksize itself is set by the compiler function explicitly and then appends the string "optional"
char kern_ver[100];
sprintf(kern_ver,"ocl_azim_kernel2d_%d.cl",2);
fprintf(stream,"Will use kernel %s\n",kern_ver);
if(ocl_compiler(oclconfig,kern_ver,BLOCK_SIZE,optional,stream))return -1;
cLog_debug(&hLog,"Will use kernel %s\n",kern_ver);
if(ocl_compiler(oclconfig,kern_ver,BLOCK_SIZE,optional))return -1;
hasProgram=1;
oclconfig->oclkernels = (cl_kernel*)malloc(8*sizeof(cl_kernel));
if(!oclconfig->oclkernels){
ocl_errmsg("Fatal error in ocl_config. Cannot allocate kernels",__FILE__,__LINE__);
cLog_critical(&hLog,"Fatal error in ocl_config. Cannot allocate kernels (%s:%d)\n",__FILE__,__LINE__);
return -2;
}
int i=0;
oclconfig->oclkernels[CLKERN_INTEGRATE] = clCreateKernel(oclconfig->oclprogram,"create_histo_binarray",&err);
if(err){fprintf(stderr,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
oclconfig->oclkernels[CLKERN_UIMEMSET2] = clCreateKernel(oclconfig->oclprogram,"uimemset2",&err);
if(err){fprintf(stderr,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
oclconfig->oclkernels[CLKERN_IMEMSET] = clCreateKernel(oclconfig->oclprogram,"imemset",&err);
if(err){fprintf(stderr,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
oclconfig->oclkernels[CLKERN_UI2F2] = clCreateKernel(oclconfig->oclprogram,"ui2f2",&err);
if(err){fprintf(stderr,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
oclconfig->oclkernels[CLKERN_GET_SPANS] = clCreateKernel(oclconfig->oclprogram,"get_spans",&err);
if(err){fprintf(stderr,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
oclconfig->oclkernels[CLKERN_GROUP_SPANS] = clCreateKernel(oclconfig->oclprogram,"group_spans",&err);
if(err){fprintf(stderr,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
oclconfig->oclkernels[CLKERN_SOLIDANGLE_CORRECTION] = clCreateKernel(oclconfig->oclprogram,"solidangle_correction",&err);
if(err){fprintf(stderr,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
oclconfig->oclkernels[CLKERN_DUMMYVAL_CORRECTION] = clCreateKernel(oclconfig->oclprogram,"dummyval_correction",&err);
if(err){fprintf(stderr,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateKernel error, %s\n",ocl_perrc(err));return -1;};i++;
oclconfig->Nkernels=i;
hasKernels = 1;
@ -228,7 +232,7 @@ int ocl_xrpd2D_fullsplit::configure()
CR(
clEnqueueNDRangeKernel(oclconfig->oclcmdqueue,oclconfig->oclkernels[CLKERN_IMEMSET],1,0,wdim,tdim,0,0,&oclconfig->t_s[0]) );
execTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[0], "Initialise Mask to 0");
execTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[0], "Initialise Mask to 0",&hLog);
clReleaseEvent(oclconfig->t_s[0]);
return 0;
@ -237,13 +241,13 @@ return 0;
int ocl_xrpd2D_fullsplit::loadTth(float* tth, float* dtth, float tth_min, float tth_max)
{
fprintf(stream,"Loading Tth\n");
cLog_extended(&hLog,"Loading Tth\n");
float tthmm[2];
tthmm[0]=tth_min;
tthmm[1]=tth_max;
if(!oclconfig->Nbuffers || !isConfigured){
fprintf(stderr,"You may not call loadTth() at this point, OpenCL is not configured (Hint: run config())\n");
cLog_critical(&hLog,"You may not call loadTth() at this point, OpenCL is not configured (Hint: run config())\n");
return -2;
}
@ -256,7 +260,7 @@ int ocl_xrpd2D_fullsplit::loadTth(float* tth, float* dtth, float tth_min, float
CR(
clEnqueueWriteBuffer(oclconfig->oclcmdqueue,oclconfig->oclmemref[CLMEM_TTH_MIN_MAX],CL_TRUE,0,2*sizeof(cl_float),(void*)tthmm,0,0,&oclconfig->t_s[2]) );
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[2],"Load Tth",stream);
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[2],"Load Tth",&hLog);
clReleaseEvent(oclconfig->t_s[0]);
clReleaseEvent(oclconfig->t_s[1]);
clReleaseEvent(oclconfig->t_s[2]);
@ -267,18 +271,18 @@ int ocl_xrpd2D_fullsplit::loadTth(float* tth, float* dtth, float tth_min, float
int ocl_xrpd2D_fullsplit::loadChi(float* chi, float* dchi, float chi_min, float chi_max)
{
fprintf(stream,"Loading Chi\n");
cLog_extended(&hLog,"Loading Chi\n");
float chimm[2];
chimm[0]=chi_min;
chimm[1]=chi_max;
if(!hasActiveContext){
fprintf(stderr,"You may not call loadChi() at this point. There is no Active context. (Hint: run init())\n");
cLog_critical(&hLog,"You may not call loadChi() at this point. There is no Active context. (Hint: run init())\n");
return -2;
}
if(!oclconfig->Nbuffers || !isConfigured){
fprintf(stderr,"You may not call loadChi() at this point, OpenCL is not configured (Hint: run configure())\n");
cLog_critical(&hLog,"You may not call loadChi() at this point, OpenCL is not configured (Hint: run configure())\n");
return -2;
}
@ -291,7 +295,7 @@ int ocl_xrpd2D_fullsplit::loadChi(float* chi, float* dchi, float chi_min, float
CR(
clEnqueueWriteBuffer(oclconfig->oclcmdqueue,oclconfig->oclmemref[CLMEM_CHI_MIN_MAX],CL_TRUE,0,2*sizeof(cl_float),(void*)chimm,0,0,&oclconfig->t_s[2]) );
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[2],"Load Chi",stream);
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[2],"Load Chi",&hLog);
clReleaseEvent(oclconfig->t_s[0]);
clReleaseEvent(oclconfig->t_s[1]);
clReleaseEvent(oclconfig->t_s[2]);
@ -302,22 +306,22 @@ int ocl_xrpd2D_fullsplit::loadChi(float* chi, float* dchi, float chi_min, float
int ocl_xrpd2D_fullsplit::setSolidAngle(float *SolidAngle)
{
fprintf(stream,"Setting SolidAngle\n");
cLog_extended(&hLog,"Setting SolidAngle\n");
if(!oclconfig->Nbuffers || !isConfigured){
fprintf(stderr,"You may not call setSolidAngle() at this point, the required buffers are not allocated (Hint: run config())\n");
cLog_critical(&hLog,"You may not call setSolidAngle() at this point, the required buffers are not allocated (Hint: run config())\n");
return -2;
}
if(!hasActiveContext){
fprintf(stderr,"You may not call setSolidAngle() at this point. There is no Active context. (Hint: run init())\n");
cLog_critical(&hLog,"You may not call setSolidAngle() at this point. There is no Active context. (Hint: run init())\n");
return -2;
}
CR(
clEnqueueWriteBuffer(oclconfig->oclcmdqueue,oclconfig->oclmemref[CLMEM_SOLIDANGLE],CL_TRUE,0,sgs->Nimage*sizeof(cl_float),(void*)SolidAngle,0,0,&oclconfig->t_s[0]) );
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[0],"Load SolidAngle");
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[0],"Load SolidAngle", &hLog);
clReleaseEvent(oclconfig->t_s[0]);
useSolidAngle=1;
@ -326,7 +330,7 @@ int ocl_xrpd2D_fullsplit::setSolidAngle(float *SolidAngle)
int ocl_xrpd2D_fullsplit::unsetSolidAngle()
{
fprintf(stream,"Unsetting SolidAngle\n");
cLog_extended(&hLog,"Unsetting SolidAngle\n");
if(useSolidAngle)
{
@ -338,22 +342,22 @@ int ocl_xrpd2D_fullsplit::unsetSolidAngle()
int ocl_xrpd2D_fullsplit::setMask(int* Mask)
{
fprintf(stream,"Setting Mask\n");
cLog_extended(&hLog,"Setting Mask\n");
if(!oclconfig->Nbuffers || !isConfigured){
fprintf(stderr,"You may not call setMask() at this point, the required buffers are not allocated (Hint: run config())\n");
cLog_critical(&hLog,"You may not call setMask() at this point, the required buffers are not allocated (Hint: run config())\n");
return -2;
}
if(!hasActiveContext){
fprintf(stderr,"You may not call setMask() at this point. There is no Active context. (Hint: run init())\n");
cLog_critical(&hLog,"You may not call setMask() at this point. There is no Active context. (Hint: run init())\n");
return -2;
}
CR(
clEnqueueWriteBuffer(oclconfig->oclcmdqueue,oclconfig->oclmemref[CLMEM_MASK],CL_TRUE,0,sgs->Nimage*sizeof(cl_int),(void*)Mask,0,0,&oclconfig->t_s[0]) );
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[0],"Load Mask");
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[0],"Load Mask", &hLog);
clReleaseEvent(oclconfig->t_s[0]);
useMask=1;
@ -363,7 +367,7 @@ int ocl_xrpd2D_fullsplit::setMask(int* Mask)
int ocl_xrpd2D_fullsplit::unsetMask()
{
fprintf(stream,"Unsetting Mask\n");
cLog_extended(&hLog,"Unsetting Mask\n");
if(useMask)
{
@ -372,7 +376,7 @@ int ocl_xrpd2D_fullsplit::unsetMask()
CR(
clEnqueueNDRangeKernel(oclconfig->oclcmdqueue,oclconfig->oclkernels[CLKERN_IMEMSET],1,0,wdim,tdim,0,0,&oclconfig->t_s[0]) );
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[0],"Reset Mask to 0");
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[0],"Reset Mask to 0", &hLog);
clReleaseEvent(oclconfig->t_s[0]);
useMask=0;
@ -385,22 +389,22 @@ int ocl_xrpd2D_fullsplit::unsetMask()
int ocl_xrpd2D_fullsplit::setDummyValue(float dummyVal)
{
fprintf(stream,"Setting Dummy Value\n");
cLog_extended(&hLog,"Setting Dummy Value\n");
if(!oclconfig->Nbuffers || !isConfigured){
fprintf(stderr,"You may not call setMask() at this point, the required buffers are not allocated (Hint: run config())\n");
cLog_critical(&hLog,"You may not call setMask() at this point, the required buffers are not allocated (Hint: run config())\n");
return -2;
}
if(!hasActiveContext){
fprintf(stderr,"You may not call setMask() at this point. There is no Active context. (Hint: run init())\n");
cLog_critical(&hLog,"You may not call setMask() at this point. There is no Active context. (Hint: run init())\n");
return -2;
}
CR(
clEnqueueWriteBuffer(oclconfig->oclcmdqueue,oclconfig->oclmemref[CLMEM_DUMMYVAL],CL_TRUE,0,sizeof(cl_float),(void*)&dummyVal,0,0,&oclconfig->t_s[0]) );
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[0],"Load Dummy Value");
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[0],"Load Dummy Value",&hLog);
clReleaseEvent(oclconfig->t_s[0]);
useDummyVal=1;
@ -410,7 +414,7 @@ int ocl_xrpd2D_fullsplit::setDummyValue(float dummyVal)
int ocl_xrpd2D_fullsplit::unsetDummyValue()
{
fprintf(stream,"Unsetting Dummy Value\n");
cLog_extended(&hLog,"Unsetting Dummy Value\n");
if(useDummyVal)
{
@ -437,18 +441,18 @@ int ocl_xrpd2D_fullsplit::unsetDummyValue()
*/
int ocl_xrpd2D_fullsplit::setTthRange(float lowerBound, float upperBound)
{
fprintf(stream,"Setting 2th Range\n");
cLog_extended(&hLog,"Setting 2th Range\n");
float tthrmm[2];
tthrmm[0]=lowerBound;
tthrmm[1]=upperBound;
if(!oclconfig->Nbuffers || !isConfigured){
fprintf(stderr,"You may not call setMask() at this point, the required buffers are not allocated (Hint: run config())\n");
cLog_critical(&hLog,"You may not call setMask() at this point, the required buffers are not allocated (Hint: run config())\n");
return -2;
}
if(!hasActiveContext){
fprintf(stderr,"You may not call setMask() at this point. There is no Active context. (Hint: run init())\n");
cLog_critical(&hLog,"You may not call setMask() at this point. There is no Active context. (Hint: run init())\n");
return -2;
}
@ -459,7 +463,7 @@ int ocl_xrpd2D_fullsplit::setTthRange(float lowerBound, float upperBound)
CR( clSetKernelArg(oclconfig->oclkernels[CLKERN_INTEGRATE],11,sizeof(cl_mem),&oclconfig->oclmemref[CLMEM_TTH_RANGE]) ); //TTH range user values
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[0],"Load 2th Range",stream);
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[0],"Load 2th Range", &hLog);
clReleaseEvent(oclconfig->t_s[0]);
useTthRange=1;
@ -474,7 +478,7 @@ int ocl_xrpd2D_fullsplit::setTthRange(float lowerBound, float upperBound)
*/
int ocl_xrpd2D_fullsplit::unsetTthRange()
{
fprintf(stream,"Unsetting 2th Range\n");
cLog_extended(&hLog,"Unsetting 2th Range\n");
if(useTthRange)
{
@ -503,18 +507,18 @@ int ocl_xrpd2D_fullsplit::unsetTthRange()
*/
int ocl_xrpd2D_fullsplit::setChiRange(float lowerBound, float upperBound)
{
fprintf(stream,"Setting 2th Range\n");
cLog_extended(&hLog,"Setting 2th Range\n");
float tthrmm[2];
tthrmm[0]=lowerBound;
tthrmm[1]=upperBound;
if(!oclconfig->Nbuffers || !isConfigured){
fprintf(stderr,"You may not call setMask() at this point, the required buffers are not allocated (Hint: run config())\n");
cLog_critical(&hLog,"You may not call setMask() at this point, the required buffers are not allocated (Hint: run config())\n");
return -2;
}
if(!hasActiveContext){
fprintf(stderr,"You may not call setMask() at this point. There is no Active context. (Hint: run init())\n");
cLog_critical(&hLog,"You may not call setMask() at this point. There is no Active context. (Hint: run init())\n");
return -2;
}
@ -524,7 +528,7 @@ int ocl_xrpd2D_fullsplit::setChiRange(float lowerBound, float upperBound)
//Set the tth_range argument of the kernels to point to tthRange instead of tth_min_max
CR( clSetKernelArg(oclconfig->oclkernels[CLKERN_INTEGRATE],12,sizeof(cl_mem),&oclconfig->oclmemref[CLMEM_CHI_RANGE]) ); //TTH range user values
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[0],"Load 2th Range",stream);
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[0],"Load 2th Range", &hLog);
clReleaseEvent(oclconfig->t_s[0]);
useChiRange=1;
@ -539,7 +543,7 @@ int ocl_xrpd2D_fullsplit::setChiRange(float lowerBound, float upperBound)
*/
int ocl_xrpd2D_fullsplit::unsetChiRange()
{
fprintf(stream,"Unsetting 2th Range\n");
cLog_extended(&hLog,"Unsetting 2th Range\n");
if(useChiRange)
{
@ -559,17 +563,17 @@ int ocl_xrpd2D_fullsplit::execute(float *im_inten,float *histogram,float *bins)
{
if(!isConfigured){
fprintf(stderr,"You may not call execute() at this point, kernels are not configured (Hint: run configure())\n");
cLog_critical(&hLog,"You may not call execute() at this point, kernels are not configured (Hint: run configure())\n");
return -2;
}
if(!hasActiveContext){
fprintf(stderr,"You may not call execute() at this point. There is no Active context. (Hint: run init())\n");
cLog_critical(&hLog,"You may not call execute() at this point. There is no Active context. (Hint: run init())\n");
return -2;
}
if(!hasTthLoaded || !hasChiLoaded){
fprintf(stderr,"You may not call execute() at this point. There is no 2th or chi array loaded. (Hint: run loadTth() or loadChi()))\n");
cLog_critical(&hLog,"You may not call execute() at this point. There is no 2th or chi array loaded. (Hint: run loadTth() or loadChi()))\n");
return -2;
}
@ -582,12 +586,12 @@ int ocl_xrpd2D_fullsplit::execute(float *im_inten,float *histogram,float *bins)
size_t wdim_reduceh[] = { (sgs->Nbins/BLOCK_SIZE) * BLOCK_SIZE + (sgs->Nbins%BLOCK_SIZE) * BLOCK_SIZE, 1, 1};
size_t tdim_reduceh[] = {BLOCK_SIZE, 1, 1};
fprintf(stream,"--Histo / Spans workdim %lu %lu %lu\n",(lui)wdim_partialh[0],(lui)wdim_partialh[1],(lui)wdim_partialh[2]);
fprintf(stream,"--Histo / Spans threadim %lu %lu %lu -- Blocks:%lu\n",(lui)tdim_partialh[0],(lui)tdim_partialh[1],(lui)tdim_partialh[2],\
cLog_debug(&hLog,"--Histo / Spans workdim %lu %lu %lu\n",(lui)wdim_partialh[0],(lui)wdim_partialh[1],(lui)wdim_partialh[2]);
cLog_debug(&hLog,"--Histo / Spans threadim %lu %lu %lu -- Blocks:%lu\n",(lui)tdim_partialh[0],(lui)tdim_partialh[1],(lui)tdim_partialh[2],\
(lui)wdim_partialh[0]/(lui)tdim_partialh[0]);
fprintf(stream,"--Memset / Convert workdim %lu %lu %lu\n",(lui)wdim_reduceh[0],(lui)wdim_reduceh[1],(lui)wdim_reduceh[2]);
fprintf(stream,"--Memset / Convert threadim %lu %lu %lu -- Blocks:%lu\n",(lui)tdim_reduceh[0],(lui)tdim_reduceh[1],(lui)tdim_reduceh[2],\
cLog_debug(&hLog,"--Memset / Convert workdim %lu %lu %lu\n",(lui)wdim_reduceh[0],(lui)wdim_reduceh[1],(lui)wdim_reduceh[2]);
cLog_debug(&hLog,"--Memset / Convert threadim %lu %lu %lu -- Blocks:%lu\n",(lui)tdim_reduceh[0],(lui)tdim_reduceh[1],(lui)tdim_reduceh[2],\
(lui)wdim_reduceh[0]/(lui)tdim_reduceh[0]);
@ -637,24 +641,24 @@ int ocl_xrpd2D_fullsplit::execute(float *im_inten,float *histogram,float *bins)
CR(
clEnqueueReadBuffer(oclconfig->oclcmdqueue,oclconfig->oclmemref[CLMEM_HISTOGRAM],CL_TRUE,0,sgs->Nbins*sizeof(cl_float),(void*)histogram,0,0,&oclconfig->t_s[7]) );
fprintf(stream,"--Waiting for the command queue to finish\n");
cLog_debug(&hLog,"--Waiting for the command queue to finish\n");
CR(clFinish(oclconfig->oclcmdqueue));
//Get execution time from first memory copy to last memory copy.
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[0],"copyIn ");
execTime_ms += ocl_get_profT(&oclconfig->t_s[1], &oclconfig->t_s[1], "memset ");
execTime_ms += ocl_get_profT(&oclconfig->t_s[2], &oclconfig->t_s[2], "getSpa ");
execTime_ms += ocl_get_profT(&oclconfig->t_s[3], &oclconfig->t_s[3], "groupS ");
execTime_ms += ocl_get_profT(&oclconfig->t_s[4], &oclconfig->t_s[4], "Azim GPU ");
execTime_ms += ocl_get_profT(&oclconfig->t_s[5], &oclconfig->t_s[5], "convert ");
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[0], &oclconfig->t_s[0],"copyIn ", &hLog);
execTime_ms += ocl_get_profT(&oclconfig->t_s[1], &oclconfig->t_s[1], "memset ", &hLog);
execTime_ms += ocl_get_profT(&oclconfig->t_s[2], &oclconfig->t_s[2], "getSpa ", &hLog);
execTime_ms += ocl_get_profT(&oclconfig->t_s[3], &oclconfig->t_s[3], "groupS ", &hLog);
execTime_ms += ocl_get_profT(&oclconfig->t_s[4], &oclconfig->t_s[4], "Azim GPU ", &hLog);
execTime_ms += ocl_get_profT(&oclconfig->t_s[5], &oclconfig->t_s[5], "convert ", &hLog);
if(useSolidAngle)
execTime_ms += ocl_get_profT(&oclconfig->t_s[8], &oclconfig->t_s[8],"Solidan ");
execTime_ms += ocl_get_profT(&oclconfig->t_s[8], &oclconfig->t_s[8],"Solidan ", &hLog);
if(useDummyVal)
execTime_ms += ocl_get_profT(&oclconfig->t_s[9], &oclconfig->t_s[9],"dummyva ");
execTime_ms += ocl_get_profT(&oclconfig->t_s[9], &oclconfig->t_s[9],"dummyva ", &hLog);
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[6], &oclconfig->t_s[7],"copyOut ");
memCpyTime_ms += ocl_get_profT(&oclconfig->t_s[6], &oclconfig->t_s[7],"copyOut ", &hLog);
for(int ievent=0;ievent<8 + useSolidAngle + useDummyVal;ievent++)clReleaseEvent(oclconfig->t_s[ievent]);
return 0;
@ -667,12 +671,12 @@ int ocl_xrpd2D_fullsplit::allocate_CL_buffers()
cl_int err;
oclconfig->oclmemref = (cl_mem*)malloc(16*sizeof(cl_mem));
if(!oclconfig->oclmemref){
fprintf(stderr,"Fatal error in ocl_azim_clbuffers. Cannot allocate memrefs\n");
cLog_critical(&hLog,"Fatal error in ocl_azim_clbuffers. Cannot allocate memrefs\n");
return -2;
}
if(sgs->Nimage < BLOCK_SIZE){
fprintf(stderr,"Fatal error in ocl_azim_clbuffers. Nimage (%d) must be >= BLOCK_SIZE (%d)\n",sgs->Nimage,BLOCK_SIZE);
cLog_critical(&hLog,"Fatal error in ocl_azim_clbuffers. Nimage (%d) must be >= BLOCK_SIZE (%d)\n",sgs->Nimage,BLOCK_SIZE);
return -2;
}
@ -687,13 +691,12 @@ int ocl_xrpd2D_fullsplit::allocate_CL_buffers()
ualloc += 2*sizeof(cl_float) * 4 + sizeof(cl_float);
if(ualloc >= oclconfig->dev_mem && oclconfig->dev_mem != 0){
fprintf(stderr,"Fatal error in ocl_azim_clbuffers. Not enough device memory for buffers (%lu requested, %lu available)\n",\
cLog_critical(&hLog,"Fatal error in ocl_azim_clbuffers. Not enough device memory for buffers (%lu requested, %lu available)\n",\
(lui)ualloc,(lui)oclconfig->dev_mem);
return -1;
} else {
if(oclconfig->dev_mem == 0){
fprintf(stream,"Caution: Device did not return the available memory size (%lu requested)\n",(lui)ualloc);
if(stream!=stdout) fprintf(stderr,"Caution: Device did not return the available memory size (%lu requested)\n",(lui)ualloc);
cLog_extended(&hLog,"Caution: Device did not return the available memory size (%lu requested)\n",(lui)ualloc);
}
}
@ -701,89 +704,89 @@ int ocl_xrpd2D_fullsplit::allocate_CL_buffers()
int i=0;
oclconfig->oclmemref[CLMEM_TTH]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_ONLY,(size_t)(sgs->Nimage*sizeof(cl_float)),0,&err);//tth array corners -0
if(err){fprintf(stderr,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);return -1;};i++;
oclconfig->oclmemref[CLMEM_CHI]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_ONLY,(size_t)(sgs->Nimage*sizeof(cl_float)),0,&err);//chi array corners -1
if(err){fprintf(stderr,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);return -1;};i++;
oclconfig->oclmemref[CLMEM_IMAGE]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_ONLY,(size_t)(sgs->Nimage*sizeof(cl_float)),0,&err);//Image intensity -2
if(err){fprintf(stderr,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
oclconfig->oclmemref[CLMEM_SOLIDANGLE]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_ONLY,(size_t)(sgs->Nimage*sizeof(cl_float)),0,&err);//Solid Angle -3
if(err){fprintf(stderr,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
oclconfig->oclmemref[CLMEM_HISTOGRAM]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_WRITE,(size_t)(sgs->Nbins*sizeof(cl_float)),0,&err);//Histogram -4
if(err){fprintf(stderr,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
if(sgs->usefp64)
{
oclconfig->oclmemref[CLMEM_UHISTOGRAM]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_WRITE,(size_t)(sgs->Nbins*sizeof(cl_ulong)),0,&err);//ulHistogram -5
if(err){fprintf(stderr,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
}else
{
oclconfig->oclmemref[CLMEM_UHISTOGRAM]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_WRITE,(size_t)(sgs->Nbins*sizeof(cl_uint)),0,&err);//ulHistogram -5
if(err){fprintf(stderr,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
}
oclconfig->oclmemref[CLMEM_WEIGHTS]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_WRITE,(size_t)(sgs->Nbins*sizeof(cl_float)),0,&err);//Bin array -6
if(err){fprintf(stderr,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
if(sgs->usefp64)
{
oclconfig->oclmemref[CLMEM_UWEIGHTS]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_WRITE,(size_t)(sgs->Nbins*sizeof(cl_ulong)),0,&err);//uBinarray -7
if(err){fprintf(stderr,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
}else
{
oclconfig->oclmemref[CLMEM_UWEIGHTS]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_WRITE,(size_t)(sgs->Nbins*sizeof(cl_uint)),0,&err);//uBinarray -7
if(err){fprintf(stderr,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
}
oclconfig->oclmemref[CLMEM_SPAN_RANGES]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_WRITE,(size_t)((sgs->Nimage * 2)*sizeof(cl_float)),0,&err);//span_ranges buffer -8
if(err){fprintf(stderr,"clCreateKernel error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateKernel error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
oclconfig->oclmemref[CLMEM_TTH_MIN_MAX]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_ONLY,(size_t)(2*sizeof(cl_float)),0,&err);//Min,Max values for tth -9
if(err){fprintf(stderr,"clCreateKernel error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateKernel error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
oclconfig->oclmemref[CLMEM_CHI_MIN_MAX]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_ONLY,(size_t)(2*sizeof(cl_float)),0,&err);//Min,Max values for chi -10
if(err){fprintf(stderr,"clCreateKernel error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateKernel error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
oclconfig->oclmemref[CLMEM_TTH_DELTA]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_WRITE,(size_t)(sgs->Nimage*sizeof(cl_float)),0,&err);//tth array min corners -12
if(err){fprintf(stderr,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);return -1;};i++;
oclconfig->oclmemref[CLMEM_CHI_DELTA]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_WRITE,(size_t)(sgs->Nimage*sizeof(cl_float)),0,&err);//tth array max corners -13
if(err){fprintf(stderr,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);return -1;};i++;
oclconfig->oclmemref[CLMEM_TTH_RANGE]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_ONLY,(size_t)(2*sizeof(cl_float)),0,&err);//Min,Max values for tth -9
if(err){fprintf(stderr,"clCreateKernel error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateKernel error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
oclconfig->oclmemref[CLMEM_CHI_RANGE]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_ONLY,(size_t)(2*sizeof(cl_float)),0,&err);//Min,Max values for chi -10
if(err){fprintf(stderr,"clCreateKernel error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateKernel error, %s (@%d)\n",ocl_perrc(err),i-1);clean_clbuffers(i-1);return -1;};i++;
oclconfig->oclmemref[CLMEM_MASK]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_WRITE,(size_t)(sgs->Nimage*sizeof(cl_int)),0,&err);//Mask -14
if(err){fprintf(stderr,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);return -1;};i++;
oclconfig->oclmemref[CLMEM_DUMMYVAL]=
clCreateBuffer(oclconfig->oclcontext,CL_MEM_READ_ONLY,(size_t)(sizeof(cl_float)),0,&err);//Dummy Value -15
if(err){fprintf(stderr,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);return -1;};i++;
if(err){cLog_critical(&hLog,"clCreateBuffer error, %s (@%d)\n",ocl_perrc(err),i-1);return -1;};i++;
fprintf(stream,"Allocated %d buffers (%.3f Mb) on device\n",i,(float)ualloc/1024./1024.);
cLog_extended(&hLog,"Allocated %d buffers (%.3f Mb) on device\n",i,(float)ualloc/1024./1024.);
oclconfig->Nbuffers = i;
return 0;
}
@ -861,7 +864,7 @@ int ocl_xrpd2D_fullsplit::clean(int preserve_context)
if(hasBuffers)
{
clean_clbuffers(oclconfig->Nbuffers);
fprintf(stream,"--released OpenCL buffers\n");
cLog_debug(&hLog,"--released OpenCL buffers\n");
hasBuffers = 0;
hasTthLoaded = 0;
hasChiLoaded = 0;
@ -875,21 +878,21 @@ int ocl_xrpd2D_fullsplit::clean(int preserve_context)
if(hasKernels)
{
clean_clkernels(oclconfig->Nkernels);
fprintf(stream,"--released OpenCL kernels\n");
cLog_debug(&hLog,"--released OpenCL kernels\n");
hasKernels=0;
}
if(hasProgram)
{
CR(clReleaseProgram(oclconfig->oclprogram));
fprintf(stream,"--released OpenCL program\n");
cLog_debug(&hLog,"--released OpenCL program\n");
hasProgram=0;
}
if(hasQueue)
{
CR(clReleaseCommandQueue(oclconfig->oclcmdqueue));
fprintf(stream,"--released OpenCL queue\n");
cLog_debug(&hLog,"--released OpenCL queue\n");
hasQueue=0;
}
@ -903,19 +906,19 @@ int ocl_xrpd2D_fullsplit::clean(int preserve_context)
{
free(oclconfig->oclmemref);
oclconfig->oclmemref=NULL;
fprintf(stream,"--released OpenCL memory references\n");
cLog_debug(&hLog,"--released OpenCL memory references\n");
}
if(oclconfig->oclkernels)
{
free(oclconfig->oclkernels);
oclconfig->oclkernels=NULL;
fprintf(stream,"--released OpenCL kernel references\n");
cLog_debug(&hLog,"--released OpenCL kernel references\n");
}
reset_time();
if(hasActiveContext){
ocl_destroy_context(oclconfig->oclcontext);
ocl_destroy_context(oclconfig->oclcontext, &hLog);
hasActiveContext=0;
fprintf(stream,"--released OpenCL context\n");
cLog_debug(&hLog,"--released OpenCL context\n");
return 0;
}
}