// -----------------------------------------------------------------------------
//
// -----------------------------------------------------------------------------
void FindKernelAvgMisorientations::setupFilterParameters()
{
FilterParameterVector parameters;
parameters.push_back(IntVec3FilterParameter::New("Kernel Radius", "KernelSize", getKernelSize(), FilterParameter::Parameter));
parameters.push_back(SeparatorFilterParameter::New("Cell Data", FilterParameter::RequiredArray));
{
DataArraySelectionFilterParameter::RequirementType req = DataArraySelectionFilterParameter::CreateRequirement(DREAM3D::TypeNames::Int32, 1, DREAM3D::AttributeMatrixType::Cell, DREAM3D::GeometryType::ImageGeometry);
parameters.push_back(DataArraySelectionFilterParameter::New("Feature Ids", "FeatureIdsArrayPath", getFeatureIdsArrayPath(), FilterParameter::RequiredArray, req));
}
{
DataArraySelectionFilterParameter::RequirementType req = DataArraySelectionFilterParameter::CreateRequirement(DREAM3D::TypeNames::Int32, 1, DREAM3D::AttributeMatrixType::Cell, DREAM3D::GeometryType::ImageGeometry);
parameters.push_back(DataArraySelectionFilterParameter::New("Phases", "CellPhasesArrayPath", getCellPhasesArrayPath(), FilterParameter::RequiredArray, req));
}
{
DataArraySelectionFilterParameter::RequirementType req = DataArraySelectionFilterParameter::CreateRequirement(DREAM3D::TypeNames::Float, 4, DREAM3D::AttributeMatrixType::Cell, DREAM3D::GeometryType::ImageGeometry);
parameters.push_back(DataArraySelectionFilterParameter::New("Quaternions", "QuatsArrayPath", getQuatsArrayPath(), FilterParameter::RequiredArray, req));
}
parameters.push_back(SeparatorFilterParameter::New("Cell Ensemble Data", FilterParameter::RequiredArray));
{
DataArraySelectionFilterParameter::RequirementType req = DataArraySelectionFilterParameter::CreateRequirement(DREAM3D::TypeNames::UInt32, 1, DREAM3D::AttributeMatrixType::CellEnsemble, DREAM3D::GeometryType::ImageGeometry);
parameters.push_back(DataArraySelectionFilterParameter::New("Crystal Structures", "CrystalStructuresArrayPath", getCrystalStructuresArrayPath(), FilterParameter::RequiredArray, req));
}
parameters.push_back(SeparatorFilterParameter::New("Cell Data", FilterParameter::CreatedArray));
parameters.push_back(StringFilterParameter::New("Kernel Average Misorientations", "KernelAverageMisorientationsArrayName", getKernelAverageMisorientationsArrayName(), FilterParameter::CreatedArray));
setFilterParameters(parameters);
}
// -----------------------------------------------------------------------------
void FindKernelAvgMisorientations::readFilterParameters(AbstractFilterParametersReader* reader, int index)
{
reader->openFilterGroup(this, index);
setKernelAverageMisorientationsArrayName(reader->readString("KernelAverageMisorientationsArrayName", getKernelAverageMisorientationsArrayName() ) );
setQuatsArrayPath(reader->readDataArrayPath("QuatsArrayPath", getQuatsArrayPath() ) );
setCrystalStructuresArrayPath(reader->readDataArrayPath("CrystalStructuresArrayPath", getCrystalStructuresArrayPath() ) );
setCellPhasesArrayPath(reader->readDataArrayPath("CellPhasesArrayPath", getCellPhasesArrayPath() ) );
setFeatureIdsArrayPath(reader->readDataArrayPath("FeatureIdsArrayPath", getFeatureIdsArrayPath() ) );
setKernelSize( reader->readIntVec3("KernelSize", getKernelSize() ) );
reader->closeFilterGroup();
}
// -----------------------------------------------------------------------------
//
// -----------------------------------------------------------------------------
void FindKernelAvgMisorientations::setupFilterParameters()
{
FilterParameterVector parameters;
parameters.push_back(IntVec3FilterParameter::New("Kernel Radius", "KernelSize", getKernelSize(), FilterParameter::Parameter));
parameters.push_back(SeparatorFilterParameter::New("Cell Data", FilterParameter::RequiredArray));
parameters.push_back(DataArraySelectionFilterParameter::New("Feature Ids", "FeatureIdsArrayPath", getFeatureIdsArrayPath(), FilterParameter::RequiredArray));
parameters.push_back(DataArraySelectionFilterParameter::New("Phases", "CellPhasesArrayPath", getCellPhasesArrayPath(), FilterParameter::RequiredArray));
parameters.push_back(DataArraySelectionFilterParameter::New("Quaternions", "QuatsArrayPath", getQuatsArrayPath(), FilterParameter::RequiredArray));
parameters.push_back(SeparatorFilterParameter::New("Cell Ensemble Data", FilterParameter::RequiredArray));
parameters.push_back(DataArraySelectionFilterParameter::New("Crystal Structures", "CrystalStructuresArrayPath", getCrystalStructuresArrayPath(), FilterParameter::RequiredArray));
parameters.push_back(SeparatorFilterParameter::New("Cell Data", FilterParameter::CreatedArray));
parameters.push_back(StringFilterParameter::New("Kernel Average Misorientations", "KernelAverageMisorientationsArrayName", getKernelAverageMisorientationsArrayName(), FilterParameter::CreatedArray));
setFilterParameters(parameters);
}
string PoolingLayer::convet2CaffeFormat()
{
string layerStrStart = "layer\n{\n";
string layerStrEnd = "}";
string nameStrStart = "\tname: \"";
string nameStrEnd = "\"\n";
string typeStrStart = "\ttype: \"";
string typeStrEnd = "\"\n";
string topStrStart = "\ttop: \"";
string topStrEnd = "\"\n";
string bottomStrStart = "\tbottom: \"";
string bottomStrEnd = "\"\n";
string poolingParamStrStart = "\tpooling_param\n\t{\n";
string poolingParamStrEnd = "\t}\n";
string kernelSizeStrStart = "\t\tkernel_size: ";
string kernelSizeStrEnd = "\n";
/*string kernelHStrStart = "\t\tkernel_h: ";
string kernelHStrEnd = "\n";
string kernelWStrStart = "\t\tkernel_w: ";
string kernelWStrEnd = "\n";*/
string strideStrStart = "\t\tstride: ";
string strideStrEnd = "\n";
/*string strideHStrStart = "\t\tstride_h: ";
string strideHStrEnd = "\n";
string strideWStrStart = "\t\tstride_w: ";
string strideWStrEnd = "\n";*/
string padStrStart = "\t\tpad: ";
string padStrEnd = "\n";
/*string padHStrStart = "\t\tpad_h: ";
string padHStrEnd = "\n";
string padWStrStart = "\t\tpad_w: ";
string padWStrEnd = "\n";*/
string poolStrStart = "\t\tpool: ";
string poolStrEnd = "\n";
string engineStrStart = "\t\tengine: ";
string engineStrEnd = "\n";
string globalPoolingStrStart = "\t\tglobal_pooling: ";
string globalPoolingStrEnd = "\n";
string phaseStrStart = "\tinclude:\n\t{\n";
string phaseStrEnd = "\t}\n";
string phaseStateStrStart = "\t\tphase: ";
string phaseStateStrEnd = "\n";
string outStr = layerStrStart+
nameStrStart + mName + nameStrEnd +
typeStrStart + getLayerType() + typeStrEnd;
for(size_t i = 0; i < mTops->size(); i++)
{
outStr = outStr + topStrStart + (*mTops)[i] + topStrEnd;
}
for(size_t i = 0; i < mBottoms->size(); i++)
{
outStr = outStr + bottomStrStart + (*mBottoms)[i] + bottomStrEnd;
}
outStr += poolingParamStrStart;
if (((PoolingParam*)mParam)->mPool != MMALab::POOLMETHOD_MAX)
{
switch (((PoolingParam*)mParam)->mPool)
{
case MMALab::POOLMETHOD_AVE:
outStr += poolStrStart + "AVE" + poolStrEnd;
break;
case MMALab::POOLMETHOD_STOCHASTIC:
outStr += poolStrStart + "STOCHASTIC" + poolStrEnd;
break;
}
}
if (((PoolingParam*)mParam)->mPad != 0)
{
outStr += padStrStart + to_string(getPad()) + padStrEnd;
}
/*if (((PoolingParam*)mParam)->mPadH != 0)
{
outStr += padHStrStart + to_string(getPadH()) + padHStrEnd;
}
if (((PoolingParam*)mParam)->mPadW != 0)
{
outStr += padWStrStart + to_string(getPadW()) + padWStrEnd;
}*/
outStr += kernelSizeStrStart + to_string(getKernelSize()) + kernelSizeStrEnd;
/*outStr += kernelHStrStart + to_string(getKernelH()) + kernelHStrEnd;
outStr += kernelWStrStart + to_string(getKernelW()) + kernelWStrEnd;*/
if (((PoolingParam*)mParam)->mStride != 1)
{
outStr += strideStrStart + to_string(getStride()) + strideStrEnd;
}
/*if (((PoolingParam*)mParam)->mStrideH != 1)
{
outStr += strideHStrStart + to_string(getStrideH()) + strideHStrEnd;
}
if (((PoolingParam*)mParam)->mStrideW != 1)
{
outStr += strideWStrStart + to_string(getStrideW()) + strideWStrEnd;
}*/
if (((PoolingParam*)mParam)->mEngine != MMALab::DEFAULT)
{
switch (((PoolingParam*)mParam)->mEngine)
{
case MMALab::CAFFE:
outStr += engineStrStart + "CAFFE" + engineStrEnd;
break;
case MMALab::CUDNN:
outStr += engineStrStart + "CUDNN" + engineStrEnd;
break;
}
}
if (((PoolingParam*)mParam)->mGlobalPooling != false)
{
outStr += globalPoolingStrStart + "true" + globalPoolingStrEnd;
}
outStr += poolingParamStrEnd;
if (mPhase != Phase::BOTH)
{
outStr += phaseStrStart + phaseStateStrStart;
if (mPhase == Phase::TRAIN)
{
outStr += "TRAIN";
}
else if (mPhase == Phase::TEST)
{
outStr += "TEST";
}
outStr += phaseStateStrEnd + phaseStrEnd;
}
outStr += layerStrEnd;
return outStr;
}
int main(int argc, char** argv)
{
ocd_init(&argc, &argv, NULL);
ocd_initCL();
cl_int err;
size_t global_size;
size_t local_size;
cl_program program;
cl_kernel kernel_compute_flux;
cl_kernel kernel_compute_flux_contributions;
cl_kernel kernel_compute_step_factor;
cl_kernel kernel_time_step;
cl_kernel kernel_initialize_variables;
cl_mem ff_variable;
cl_mem ff_fc_momentum_x;
cl_mem ff_fc_momentum_y;
cl_mem ff_fc_momentum_z;
cl_mem ff_fc_density_energy;
if (argc < 2)
{
printf("Usage ./cfd <data input file>\n");
return 0;
}
const char* data_file_name = argv[1];
// set far field conditions and load them into constant memory on the gpu
{
float h_ff_variable[NVAR];
const float angle_of_attack = (float)(3.1415926535897931 / 180.0) * (float)(deg_angle_of_attack);
h_ff_variable[VAR_DENSITY] = (float)(1.4);
float ff_pressure = (float)(1.0);
float ff_speed_of_sound = sqrt(GAMMA*ff_pressure / h_ff_variable[VAR_DENSITY]);
float ff_speed = (float)(ff_mach)*ff_speed_of_sound;
float3 ff_velocity;
ff_velocity.x = ff_speed*(float)(cos((float)angle_of_attack));
ff_velocity.y = ff_speed*(float)(sin((float)angle_of_attack));
ff_velocity.z = 0.0;
h_ff_variable[VAR_MOMENTUM+0] = h_ff_variable[VAR_DENSITY] * ff_velocity.x;
h_ff_variable[VAR_MOMENTUM+1] = h_ff_variable[VAR_DENSITY] * ff_velocity.y;
h_ff_variable[VAR_MOMENTUM+2] = h_ff_variable[VAR_DENSITY] * ff_velocity.z;
h_ff_variable[VAR_DENSITY_ENERGY] = h_ff_variable[VAR_DENSITY]*((float)(0.5)*(ff_speed*ff_speed)) + (ff_pressure / (float)(GAMMA-1.0));
float3 h_ff_momentum;
h_ff_momentum.x = *(h_ff_variable+VAR_MOMENTUM+0);
h_ff_momentum.y = *(h_ff_variable+VAR_MOMENTUM+1);
h_ff_momentum.z = *(h_ff_variable+VAR_MOMENTUM+2);
float3 h_ff_fc_momentum_x;
float3 h_ff_fc_momentum_y;
float3 h_ff_fc_momentum_z;
float3 h_ff_fc_density_energy;
compute_flux_contribution(&h_ff_variable[VAR_DENSITY], &h_ff_momentum,
&h_ff_variable[VAR_DENSITY_ENERGY], ff_pressure, &ff_velocity,
&h_ff_fc_momentum_x, &h_ff_fc_momentum_y, &h_ff_fc_momentum_z,
&h_ff_fc_density_energy);
// copy far field conditions to the gpu
ff_variable = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(float) * NVAR, h_ff_variable, &err);
CHKERR(err, "Unable to allocate ff data");
ff_fc_momentum_x = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(float3), &h_ff_fc_momentum_x, &err);
CHKERR(err, "Unable to allocate ff data");
ff_fc_momentum_y = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(float3), &h_ff_fc_momentum_y, &err);
CHKERR(err, "Unable to allocate ff data");
ff_fc_momentum_z = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(float3), &h_ff_fc_momentum_z, &err);
CHKERR(err, "Unable to allocate ff data");
ff_fc_density_energy = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(float3), &h_ff_fc_density_energy, &err);
CHKERR(err, "Unable to allocate ff data");
}
int nel;
int nelr;
// read in domain geometry
cl_mem areas;
cl_mem elements_surrounding_elements;
cl_mem normals;
{
std::ifstream file(data_file_name);
file >> nel;
nelr = block_length*((nel / block_length )+ std::min(1, nel % block_length));
float* h_areas = new float[nelr];
int* h_elements_surrounding_elements = new int[nelr*NNB];
float* h_normals = new float[nelr*NDIM*NNB];
// read in data
for(int i = 0; i < nel; i++)
{
file >> h_areas[i];
for(int j = 0; j < NNB; j++)
{
file >> h_elements_surrounding_elements[i + j*nelr];
if(h_elements_surrounding_elements[i+j*nelr] < 0) h_elements_surrounding_elements[i+j*nelr] = -1;
h_elements_surrounding_elements[i + j*nelr]--; //it's coming in with Fortran numbering
for(int k = 0; k < NDIM; k++)
{
file >> h_normals[i + (j + k*NNB)*nelr];
h_normals[i + (j + k*NNB)*nelr] = -h_normals[i + (j + k*NNB)*nelr];
}
}
}
// fill in remaining data
int last = nel-1;
for(int i = nel; i < nelr; i++)
{
h_areas[i] = h_areas[last];
for(int j = 0; j < NNB; j++)
{
// duplicate the last element
h_elements_surrounding_elements[i + j*nelr] = h_elements_surrounding_elements[last + j*nelr];
for(int k = 0; k < NDIM; k++) h_normals[last + (j + k*NNB)*nelr] = h_normals[last + (j + k*NNB)*nelr];
}
}
areas = alloc<float>(context, nelr);
upload<float>(commands, areas, h_areas, nelr);
elements_surrounding_elements = alloc<int>(context, nelr*NNB);
upload<int>(commands, elements_surrounding_elements, h_elements_surrounding_elements, nelr*NNB);
normals = alloc<float>(context, nelr*NDIM*NNB);
upload<float>(commands, normals, h_normals, nelr*NDIM*NNB);
delete[] h_areas;
delete[] h_elements_surrounding_elements;
delete[] h_normals;
}
// Get program source.
long kernelSize = getKernelSize();
char* kernelSource = new char[kernelSize];
getKernelSource(kernelSource, kernelSize);
// Create the compute program from the source buffer
program = clCreateProgramWithSource(context, 1, (const char **) &kernelSource, NULL, &err);
CHKERR(err, "Failed to create a compute program!");
// Build the program executable
err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
if (err == CL_BUILD_PROGRAM_FAILURE)
{
char *log;
size_t logLen;
err = clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, 0, NULL, &logLen);
log = (char *) malloc(sizeof(char)*logLen);
err = clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, logLen, (void *) log, NULL);
fprintf(stderr, "CL Error %d: Failed to build program! Log:\n%s", err, log);
free(log);
exit(1);
}
CHKERR(err, "Failed to build program!");
delete[] kernelSource;
// Create the compute kernel in the program we wish to run
kernel_compute_flux = clCreateKernel(program, "compute_flux", &err);
CHKERR(err, "Failed to create a compute kernel!");
// Create the reduce kernel in the program we wish to run
kernel_compute_flux_contributions = clCreateKernel(program, "compute_flux_contributions", &err);
CHKERR(err, "Failed to create a compute_flux_contributions kernel!");
// Create the reduce kernel in the program we wish to run
kernel_compute_step_factor = clCreateKernel(program, "compute_step_factor", &err);
CHKERR(err, "Failed to create a compute_step_factor kernel!");
// Create the reduce kernel in the program we wish to run
kernel_time_step = clCreateKernel(program, "time_step", &err);
CHKERR(err, "Failed to create a time_step kernel!");
// Create the reduce kernel in the program we wish to run
kernel_initialize_variables = clCreateKernel(program, "initialize_variables", &err);
CHKERR(err, "Failed to create a initialize_variables kernel!");
// Create arrays and set initial conditions
cl_mem variables = alloc<cl_float>(context, nelr*NVAR);
err = 0;
err = clSetKernelArg(kernel_initialize_variables, 0, sizeof(int), &nelr);
err |= clSetKernelArg(kernel_initialize_variables, 1, sizeof(cl_mem),&variables);
err |= clSetKernelArg(kernel_initialize_variables, 2, sizeof(cl_mem),&ff_variable);
CHKERR(err, "Failed to set kernel arguments!");
// Get the maximum work group size for executing the kernel on the device
//err = clGetKernelWorkGroupInfo(kernel_initialize_variables, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), (void *) &local_size, NULL);
CHKERR(err, "Failed to retrieve kernel_initialize_variables work group info!");
local_size = 1;//std::min(local_size, (size_t)nelr);
global_size = nelr;
err = clEnqueueNDRangeKernel(commands, kernel_initialize_variables, 1, NULL, &global_size, NULL, 0, NULL, &ocdTempEvent);
err = clFinish(commands);
START_TIMER(ocdTempEvent, OCD_TIMER_KERNEL, "CFD Init Kernels", ocdTempTimer)
END_TIMER(ocdTempTimer)
CHKERR(err, "Failed to execute kernel [kernel_initialize_variables]! 0");
cl_mem old_variables = alloc<float>(context, nelr*NVAR);
cl_mem fluxes = alloc<float>(context, nelr*NVAR);
cl_mem step_factors = alloc<float>(context, nelr);
clFinish(commands);
cl_mem fc_momentum_x = alloc<float>(context, nelr*NDIM);
cl_mem fc_momentum_y = alloc<float>(context, nelr*NDIM);
cl_mem fc_momentum_z = alloc<float>(context, nelr*NDIM);
cl_mem fc_density_energy = alloc<float>(context, nelr*NDIM);
clFinish(commands);
// make sure all memory is floatly allocated before we start timing
err = 0;
err = clSetKernelArg(kernel_initialize_variables, 0, sizeof(int), &nelr);
err |= clSetKernelArg(kernel_initialize_variables, 1, sizeof(cl_mem),&old_variables);
err |= clSetKernelArg(kernel_initialize_variables, 2, sizeof(cl_mem),&ff_variable);
CHKERR(err, "Failed to set kernel arguments!");
// Get the maximum work group size for executing the kernel on the device
err = clGetKernelWorkGroupInfo(kernel_initialize_variables, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), (void *) &local_size, NULL);
CHKERR(err, "Failed to retrieve kernel_initialize_variables work group info!");
err = clEnqueueNDRangeKernel(commands, kernel_initialize_variables, 1, NULL, &global_size, NULL, 0, NULL, &ocdTempEvent);
clFinish(commands);
START_TIMER(ocdTempEvent, OCD_TIMER_KERNEL, "CFD Init Kernels", ocdTempTimer)
END_TIMER(ocdTempTimer)
CHKERR(err, "Failed to execute kernel [kernel_initialize_variables]! 1");
err = 0;
err = clSetKernelArg(kernel_initialize_variables, 0, sizeof(int), &nelr);
err |= clSetKernelArg(kernel_initialize_variables, 1, sizeof(cl_mem),&fluxes);
err |= clSetKernelArg(kernel_initialize_variables, 2, sizeof(cl_mem),&ff_variable);
CHKERR(err, "Failed to set kernel arguments!");
// Get the maximum work group size for executing the kernel on the device
err = clGetKernelWorkGroupInfo(kernel_compute_step_factor, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), (void *) &local_size, NULL);
CHKERR(err, "Failed to retrieve kernel_compute_step_factor work group info!");
err = clEnqueueNDRangeKernel(commands, kernel_initialize_variables, 1, NULL, &global_size, NULL, 0, NULL, &ocdTempEvent);
clFinish(commands);
START_TIMER(ocdTempEvent, OCD_TIMER_KERNEL, "CFD Init Kernels", ocdTempTimer)
END_TIMER(ocdTempTimer)
CHKERR(err, "Failed to execute kernel [kernel_initialize_variables]! 2");
std::cout << "About to memcopy" << std::endl;
err = clReleaseMemObject(step_factors);
float temp[nelr];
for(int i = 0; i < nelr; i++)
temp[i] = 0;
step_factors = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(float) * nelr, temp, &err);
CHKERR(err, "Unable to memset step_factors");
// make sure CUDA isn't still doing something before we start timing
clFinish(commands);
// these need to be computed the first time in order to compute time step
std::cout << "Starting..." << std::endl;
// Begin iterations
for(int i = 0; i < iterations; i++)
{
copy<float>(commands, old_variables, variables, nelr*NVAR);
// for the first iteration we compute the time step
err = 0;
err = clSetKernelArg(kernel_compute_step_factor, 0, sizeof(int), &nelr);
err |= clSetKernelArg(kernel_compute_step_factor, 1, sizeof(cl_mem),&variables);
err |= clSetKernelArg(kernel_compute_step_factor, 2, sizeof(cl_mem), &areas);
err |= clSetKernelArg(kernel_compute_step_factor, 3, sizeof(cl_mem), &step_factors);
CHKERR(err, "Failed to set kernel arguments!");
// Get the maximum work group size for executing the kernel on the device
err = clGetKernelWorkGroupInfo(kernel_compute_step_factor, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), (void *) &local_size, NULL);
CHKERR(err, "Failed to retrieve kernel_compute_step_factor work group info!");
err = clEnqueueNDRangeKernel(commands, kernel_compute_step_factor, 1, NULL, &global_size, NULL, 0, NULL, &ocdTempEvent);
clFinish(commands);
START_TIMER(ocdTempEvent, OCD_TIMER_KERNEL, "CFD Step Factor Kernel", ocdTempTimer)
END_TIMER(ocdTempTimer)
CHKERR(err, "Failed to execute kernel[kernel_compute_step_factor]!");
for(int j = 0; j < RK; j++)
{
err = 0;
err = clSetKernelArg(kernel_compute_flux_contributions, 0, sizeof(int), &nelr);
err |= clSetKernelArg(kernel_compute_flux_contributions, 1, sizeof(cl_mem),&variables);
err |= clSetKernelArg(kernel_compute_flux_contributions, 2, sizeof(cl_mem), &fc_momentum_x);
err |= clSetKernelArg(kernel_compute_flux_contributions, 3, sizeof(cl_mem), &fc_momentum_y);
err |= clSetKernelArg(kernel_compute_flux_contributions, 4, sizeof(cl_mem), &fc_momentum_z);
err |= clSetKernelArg(kernel_compute_flux_contributions, 5, sizeof(cl_mem), &fc_density_energy);
CHKERR(err, "Failed to set kernel arguments!");
// Get the maximum work group size for executing the kernel on the device
err = clGetKernelWorkGroupInfo(kernel_compute_flux_contributions, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), (void *) &local_size, NULL);
CHKERR(err, "Failed to retrieve kernel_compute_flux_contributions work group info!");
err = clEnqueueNDRangeKernel(commands, kernel_compute_flux_contributions, 1, NULL, &global_size, NULL, 0, NULL, &ocdTempEvent);
clFinish(commands);
START_TIMER(ocdTempEvent, OCD_TIMER_KERNEL, "CFD Flux Contribution Kernel", ocdTempTimer)
//compute_flux_contributions(nelr, variables, fc_momentum_x, fc_momentum_y, fc_momentum_z, fc_density_energy);
END_TIMER(ocdTempTimer)
CHKERR(err, "Failed to execute kernel [kernel_compute_flux_contributions]!");
err = 0;
err = clSetKernelArg(kernel_compute_flux, 0, sizeof(int), &nelr);
err |= clSetKernelArg(kernel_compute_flux, 1, sizeof(cl_mem), &elements_surrounding_elements);
err |= clSetKernelArg(kernel_compute_flux, 2, sizeof(cl_mem), &normals);
err |= clSetKernelArg(kernel_compute_flux, 3, sizeof(cl_mem), &variables);
err |= clSetKernelArg(kernel_compute_flux, 4, sizeof(cl_mem), &fc_momentum_x);
err |= clSetKernelArg(kernel_compute_flux, 5, sizeof(cl_mem), &fc_momentum_y);
err |= clSetKernelArg(kernel_compute_flux, 6, sizeof(cl_mem), &fc_momentum_z);
err |= clSetKernelArg(kernel_compute_flux, 7, sizeof(cl_mem), &fc_density_energy);
err |= clSetKernelArg(kernel_compute_flux, 8, sizeof(cl_mem), &fluxes);
err |= clSetKernelArg(kernel_compute_flux, 9, sizeof(cl_mem), &ff_variable);
err |= clSetKernelArg(kernel_compute_flux, 10, sizeof(cl_mem), &ff_fc_momentum_x);
err |= clSetKernelArg(kernel_compute_flux, 11, sizeof(cl_mem), &ff_fc_momentum_y);
err |= clSetKernelArg(kernel_compute_flux, 12, sizeof(cl_mem), &ff_fc_momentum_z);
err |= clSetKernelArg(kernel_compute_flux, 13, sizeof(cl_mem), &ff_fc_density_energy);
CHKERR(err, "Failed to set kernel arguments!");
// Get the maximum work group size for executing the kernel on the device
err = clGetKernelWorkGroupInfo(kernel_compute_flux, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), (void *) &local_size, NULL);
CHKERR(err, "Failed to retrieve kernel_compute_flux work group info!");
err = clEnqueueNDRangeKernel(commands, kernel_compute_flux, 1, NULL, &global_size, NULL, 0, NULL, &ocdTempEvent);
clFinish(commands);
START_TIMER(ocdTempEvent, OCD_TIMER_KERNEL, "CFD Flux Kernel", ocdTempTimer)
END_TIMER(ocdTempTimer)
CHKERR(err, "Failed to execute kernel [kernel_compute_flux]!");
err = 0;
err = clSetKernelArg(kernel_time_step, 0, sizeof(int), &j);
err |= clSetKernelArg(kernel_time_step, 1, sizeof(int), &nelr);
err |= clSetKernelArg(kernel_time_step, 2, sizeof(cl_mem), &old_variables);
err |= clSetKernelArg(kernel_time_step, 3, sizeof(cl_mem), &variables);
err |= clSetKernelArg(kernel_time_step, 4, sizeof(cl_mem), &step_factors);
err |= clSetKernelArg(kernel_time_step, 5, sizeof(cl_mem), &fluxes);
CHKERR(err, "Failed to set kernel arguments!");
// Get the maximum work group size for executing the kernel on the device
err = clGetKernelWorkGroupInfo(kernel_time_step, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), (void *) &local_size, NULL);
CHKERR(err, "Failed to retrieve kernel_time_step work group info!");
err = clEnqueueNDRangeKernel(commands, kernel_time_step, 1, NULL, &global_size, NULL, 0, NULL, &ocdTempEvent);
clFinish(commands);
START_TIMER(ocdTempEvent, OCD_TIMER_KERNEL, "CFD Time Step Kernel", ocdTempTimer)
END_TIMER(ocdTempTimer)
CHKERR(err, "Failed to execute kernel [kernel_time_step]!");
}
}
clFinish(commands);
std::cout << "Finished" << std::endl;
std::cout << "Saving solution..." << std::endl;
dump(commands, variables, nel, nelr);
std::cout << "Saved solution..." << std::endl;
std::cout << "Cleaning up..." << std::endl;
clReleaseProgram(program);
clReleaseKernel(kernel_compute_flux);
clReleaseKernel(kernel_compute_flux_contributions);
clReleaseKernel(kernel_compute_step_factor);
clReleaseKernel(kernel_time_step);
clReleaseKernel(kernel_initialize_variables);
clReleaseCommandQueue(commands);
clReleaseContext(context);
dealloc<float>(areas);
dealloc<int>(elements_surrounding_elements);
dealloc<float>(normals);
dealloc<float>(variables);
dealloc<float>(old_variables);
dealloc<float>(fluxes);
dealloc<float>(step_factors);
dealloc<float>(fc_momentum_x);
dealloc<float>(fc_momentum_y);
dealloc<float>(fc_momentum_z);
dealloc<float>(fc_density_energy);
std::cout << "Done..." << std::endl;
ocd_finalize();
return 0;
}
