zeroRK constant pressure GPU reactor runs thrust::transform on the host over device memory #2007

zeroRK (CHIP-SPV/zero-rk branch hip-chipStar at a0783ad1) segfaults on the GPU for every constant pressure case (constant_volume 0 in the tester input), on both Level Zero and OpenCL. Constant volume with the same settings passes. Seen on an Intel Arc A770 while adding the zeroRK GPU checks in #1945.

applications/cfd_plugin/reactor_constant_pressure_gpu.cpp lines 74 to 82 pass raw device pointers to thrust with no execution policy:

double *tmp3_ptr = N_VGetDeviceArrayPointer_Hip(tmp3_);
double *scaled_temp_ptr = &y_ptr_dev[num_species_*num_reactors_];
thrust::copy(initial_temperatures_dev_.begin(), initial_temperatures_dev_.end(), tmp3_ptr);
thrust::transform(tmp3_ptr, tmp3_ptr + num_reactors_, scaled_temp_ptr, thrust::placeholders::_1*inv_reference_temperature);

A raw double* selects thrust's host (cpp) system, so the transform runs on the host over device memory. The backtrace ends in thrust::system::detail::sequential::for_each at thrust/functional.h:373, called from ReactorConstantPressureGPU::InitializeState, from ZeroRKReactorManager::SolveReactors at zerork_reactor_manager.cpp:1028, with a host load from 0xffffd556a9800000. The constant volume reactor (reactor_constant_volume_gpu.cpp:78) uses device_vector iterators and dispatches to the device.

The fix is in zero-rk: wrap both destinations in thrust::device_pointer_cast (the thrust::copy destination has the same problem as the transform). CHIP-SPV/zero-rk has issues disabled, so this is tracked here. No upstream report exists in llnl/zero-rk, whose CUDA version uses device iterators here.

Reproduce: build zero-rk hip-chipStar with chipStar, run zerork_cfd_plugin_tester_gpu.x on the hydrogen example with constant_volume 0 in the tester input and 2048 reactors.

This issue was generated using AI