double VolumeTetrapore::calculateNumberInside( const Vector& cpos, Vector& derivatives, Tensor& vir, std::vector<Vector>& rderiv ) const {
// Setup the histogram bead
HistogramBead bead; bead.isNotPeriodic(); bead.setKernelType( getKernelType() );
// Calculate distance of atom from origin of new coordinate frame
Vector datom=pbcDistance( origin, cpos );
double ucontr, uder, vcontr, vder, wcontr, wder;
// Calculate contribution from integral along bi
bead.set( 0, len_bi, sigma );
double upos=dotProduct( datom, bi );
ucontr=bead.calculate( upos, uder );
double udlen=bead.uboundDerivative( upos );
double uder2 = bead.lboundDerivative( upos ) - udlen;
// Calculate contribution from integral along cross
bead.set( 0, len_cross, sigma );
double vpos=dotProduct( datom, cross );
vcontr=bead.calculate( vpos, vder );
double vdlen=bead.uboundDerivative( vpos );
double vder2 = bead.lboundDerivative( vpos ) - vdlen;
// Calculate contribution from integral along perp
bead.set( 0, len_perp, sigma );
double wpos=dotProduct( datom, perp );
wcontr=bead.calculate( wpos, wder );
double wdlen=bead.uboundDerivative( wpos );
double wder2 = bead.lboundDerivative( wpos ) - wdlen;
Vector dfd; dfd[0]=uder*vcontr*wcontr; dfd[1]=ucontr*vder*wcontr; dfd[2]=ucontr*vcontr*wder;
derivatives[0] = (dfd[0]*bi[0]+dfd[1]*cross[0]+dfd[2]*perp[0]);
derivatives[1] = (dfd[0]*bi[1]+dfd[1]*cross[1]+dfd[2]*perp[1]);
derivatives[2] = (dfd[0]*bi[2]+dfd[1]*cross[2]+dfd[2]*perp[2]);
double tot = ucontr*vcontr*wcontr*jacob_det;
// Add reference atom derivatives
dfd[0]=uder2*vcontr*wcontr; dfd[1]=ucontr*vder2*wcontr; dfd[2]=ucontr*vcontr*wder2;
Vector dfld; dfld[0]=udlen*vcontr*wcontr; dfld[1]=ucontr*vdlen*wcontr; dfld[2]=ucontr*vcontr*wdlen;
rderiv[0] = dfd[0]*matmul(datom,dbi[0]) + dfd[1]*matmul(datom,dcross[0]) + dfd[2]*matmul(datom,dperp[0]) +
dfld[0]*dlbi[0] + dfld[1]*dlcross[0] + dfld[2]*dlperp[0] - derivatives;
rderiv[1] = dfd[0]*matmul(datom,dbi[1]) + dfd[1]*matmul(datom,dcross[1]) + dfd[2]*matmul(datom,dperp[1]) +
dfld[0]*dlbi[1] + dfld[1]*dlcross[1] + dfld[2]*dlperp[1];
rderiv[2] = dfd[0]*matmul(datom,dbi[2]) + dfd[1]*matmul(datom,dcross[2]) + dfd[2]*matmul(datom,dperp[2]) +
dfld[0]*dlbi[2] + dfld[1]*dlcross[2] + dfld[2]*dlperp[2];
rderiv[3] = dfld[0]*dlbi[3] + dfld[1]*dlcross[3] + dfld[2]*dlperp[3];
vir.zero(); vir-=Tensor( cpos,derivatives );
for(unsigned i=0;i<4;++i){
vir -= Tensor( getPosition(i), rderiv[i] );
}
return tot;
}
bool OpenCLParser::convert(std::string fileNameIN, std::string fileNameOUT) {
if ( access( fileNameIN.c_str(), F_OK ) == -1 ) {
LOG(ERROR) << "kernel source file = '" << fileNameIN.c_str() << "' doesn't exist";
return false;
}
if ( access( fileNameIN.c_str(), R_OK ) == -1 ) {
LOG(ERROR) << "kernel source file = '" << fileNameIN.c_str() << "' isn't readable";
return false;
}
if ( access( fileNameOUT.c_str(), F_OK ) == 0 ) {
struct stat statIN;
if (stat(fileNameIN.c_str(), &statIN) == -1) {
perror(fileNameIN.c_str());
return false;
}
struct stat statOUT;
if (stat(fileNameOUT.c_str(), &statOUT) == -1) {
perror(fileNameOUT.c_str());
return false;
}
if ( statOUT.st_mtime > statIN.st_mtime ) {
DLOG(INFO) << "kernel source file = '" << fileNameOUT.c_str() << "' up-to-date";
return true;
}
}
std::ifstream file;
file.open(fileNameIN.c_str(), std::ifstream::in );
if ( ! file.is_open() ) {
LOG(ERROR) << "failed to open file = '" << fileNameIN.c_str() << "' for reading";
return false;
}
std::string line;
std::string kernel_buffer;
std::string kernel_name;
std::string kernel_type;
std::string kernel_name_typed;
std::string kernel_line_typed;
std::string kernel_modified;
std::string type_replace;
std::string stdOpenCL;
stdOpenCL += "// This file was auto-generated from file '" + fileNameIN + "' to conform to standard OpenCL\n";
bool recording = false;
while (std::getline(file, line)) {
if ( isAttributeLine(line) ) {
if ( recording ) {
recording = false;
}
kernel_name_typed = getTypedKernelName(line);
kernel_line_typed = "__kernel void " + kernel_name_typed + getTypedKernelLine(line) + " {";
if ( isFloatType(kernel_name_typed) ) {
type_replace = "float";
}
if ( isDoubleType(kernel_name_typed) ) {
type_replace = "double";
}
kernel_modified = kernel_line_typed + "\n" + kernel_buffer;
boost::regex re;
re = boost::regex("\\sT\\s", boost::regex::perl);
kernel_modified = boost::regex_replace(kernel_modified, re, " "+type_replace+" ");
re = boost::regex("\\sT\\*\\s", boost::regex::perl);
kernel_modified = boost::regex_replace(kernel_modified, re, " "+type_replace+"* ");
stdOpenCL += kernel_modified;
continue;
}
if ( isTemplateKernelLine(line) ) {
kernel_name = getKernelName(line);
kernel_type = getKernelType(line);
DLOG(INFO)<<"found template kernel '"<<kernel_name<<"' with type '"<<kernel_type<<"'";
if ( recording == false ) {
recording = true;
} else {
LOG(ERROR) << "error parsing kernel source file = '" << fileNameIN.c_str() << "'";
return false;
}
continue;
}
if ( recording ) {
kernel_buffer += line + "\n";
} else {
kernel_buffer = "";
stdOpenCL += line + "\n";
}
}
std::ofstream out(fileNameOUT.c_str());
out << stdOpenCL;
out.close();
DLOG(INFO) << "convert AMD OpenCL '"<<fileNameIN.c_str()<<"' to standard OpenCL '"<<fileNameOUT.c_str()<<"'";
return true;
}
