Hi everyone, I’m modifying a library to make it possible to use the GPU. I modified the code to ensure that various functions can be executed directly within the GPU when called from my kernel.
- the kernel :
__global__ void SetTDState_rhoe_Kernel(CCoolProp * d_CoolProp, Helmholtz *d_HELMHOLTZ, su2double *d_DENSITY, su2double *d_ENERGY,su2double *d_PRESSURE, su2double *d_TEMPERATURE,
su2double *d_SOUNDSPEED2, su2double *d_DPDRHOE, su2double *d_DPDERHO, su2double *d_DTDERHO, su2double *d_DTDRHOE, su2double *d_ENTROPY,
unsigned long nPoint){
int id= blockIdx.x * blockDim.x + threadIdx.x;
// VERSIONE BASE
if(id<nPoint){
su2double Density = d_DENSITY[id];
su2double StaticEnergy = d_ENERGY[id];
d_CoolProp->update(d_HELMHOLTZ[id], DmassUmass_INPUTS, Density,
StaticEnergy);
__syncthreads();
d_PRESSURE[id] = d_HELMHOLTZ[id].p();
d_TEMPERATURE[id] = d_HELMHOLTZ[id].T();
d_ENTROPY[id] = d_HELMHOLTZ[id].smass();
d_DPDRHOE[id] = d_CoolProp[id].calc_first_partial_deriv(d_HELMHOLTZ[id],
iP, iDmass, iUmass);
d_DPDERHO[id] = d_CoolProp[id].calc_first_partial_deriv(d_HELMHOLTZ[id],
iP, iUmass, iDmass);
d_DTDRHOE[id] = d_CoolProp[id].calc_first_partial_deriv(d_HELMHOLTZ[id],
iT, iDmass, iUmass);
d_DTDERHO[id] = d_CoolProp[id].calc_first_partial_deriv(d_HELMHOLTZ[id],
iT, iUmass, iDmass);
if (d_HELMHOLTZ[id]._phase == iphase_twophase) {
// impose gas phase
d_TEMPERATURE[id] = d_TEMPERATURE[id] + 0.1;
d_CoolProp[id].CheckPressure1(d_PRESSURE[id]);
d_CoolProp[id].CheckTemperature1(d_TEMPERATURE[id]);
d_CoolProp[id].update(d_HELMHOLTZ[id], PT_INPUTS, d_PRESSURE[id],
d_TEMPERATURE[id]);
d_SOUNDSPEED2[id] = pow(
d_CoolProp[id].calc_speed_sound(d_HELMHOLTZ[id]), 2);
} else {
d_SOUNDSPEED2[id] = pow(
d_CoolProp[id].calc_speed_sound(d_HELMHOLTZ[id]), 2);
}
}
}
- The update function:
__device__ void CCoolProp::update(Helmholtz& HEOS, input_pairs pairs, su2double value1, su2double value2){
su2double ld_value1 = value1, ld_value2 = value2;
pre_update(HEOS, pairs, ld_value1, ld_value2);
value1 = ld_value1;
value2 = ld_value2;
switch (pairs) {
case PT_INPUTS:
HEOS._p = value1;
HEOS._T = value2;
PT_flash(HEOS);
break;
case DmolarT_INPUTS:
HEOS._rhomolar = value1;
HEOS._T = value2;
DHSU_T_flash(HEOS, iDmolar);
break;
case SmolarT_INPUTS:
HEOS._smolar = value1;
HEOS._T = value2;
DHSU_T_flash(HEOS, iSmolar);
break;
case DmolarUmolar_INPUTS:
HEOS._rhomolar = value1;
HEOS._umolar = value2;
HSU_D_flash(HEOS, iUmolar);
break;
case QT_INPUTS:
HEOS._Q = value1;
HEOS._T = value2;
if ((HEOS._Q < 0) || (HEOS._Q > 1)) printf("Input vapor quality [Q] must be between 0 and 1 \n");
QT_flash(HEOS);
break;
}
}
This is only a small part of the code, but if I run the program with the compute-sanitizer debugger I get:
Exit_CompSan.txt (106.2 KB)
For example if you catch the first mistake reported in the log file:
========= COMPUTE-SANITIZER
========= Stack overflow
========= at 0x4a9580 in /home/marco/CFD/SU2_CUDA_HELMHOLTZ/SU2_CFD/src/fluid/CCoolProp.cu:1384:SetTDState_rhoe_Kernel(CCoolProp *, Helmholtz *, double *, double *, double *, double *, double *, double *, double *, double *, double *, double *, unsigned long)
========= by thread (0,0,0) in block (12,0,0)
Where the function at line 1384 of CCoolProp.cu is:
__device__ void CCoolProp::QT_flash(Helmholtz& HEOS) {
// The maximum possible saturation temperature
// Critical point for pure fluids, slightly different for pseudo-pure, very different for mixtures
CoolPropDbl Tmax_sat = HEOS.calc_Tmax_sat() + 1e-13;
// Check what the minimum limits for the equation of state are
CoolPropDbl Tmin_sat;
// calc_Tmin_sat(HEOS, Tmin_satL, Tmin_satV);
Tmin_sat = max(HEOS.sat_min_liquid.T, HEOS.sat_min_vapor.T) - 1e-13;
// // Get a reference to keep the code a bit cleaner
// const CriticalRegionSplines& splines = HEOS.components[0].EOS().critical_region_splines;
// Set some input options
saturation_T_pure_Akasaka_options options(false);
// Actually call the solver
saturation_T_pure_Maxwell(HEOS, HEOS._T, options);
HEOS._p = 0.5 * HEOS.SatV->p() + 0.5 * HEOS.SatL->p();
HEOS._rhomolar = 1 / (HEOS._Q / HEOS.SatV->rhomolar() + (1 - HEOS._Q) / HEOS.SatL->rhomolar());
// Load the outputs
HEOS._phase = iphase_twophase;
}
My question is : what kind of error is this :
========= COMPUTE-SANITIZER
========= Stack overflow
========= at 0x4a9580 in /home/marco/CFD/SU2_CUDA_HELMHOLTZ/SU2_CFD/src/fluid/CCoolProp.cu:1384:SetTDState_rhoe_Kernel(CCoolProp *, Helmholtz *, double *, double *, double *, double *, double *, double *, double *, double *, double *, double *, unsigned long)
========= by thread (0,0,0) in block (12,0,0)