alus-3 Copernicus DEM 30m COG support #19
Conversation
kautlenbachs
commented
Feb 10, 2023
kautlenbachs
commented
- Support for calibration and coherence routines
- Somewhat generalized DEM handling introduced - should be easier for future support of different DEMs
| void Srtm3ElevationModel::DeviceToHost() { CHECK_CUDA_ERR(cudaErrorNotYetImplemented); } | ||
|
|
||
| void Srtm3ElevationModel::DeviceFree() { | ||
| void Srtm3ElevationModel::ReleaseFromDevice() { | ||
| if (device_formated_srtm_buffers_info_ != nullptr) { | ||
| LOGI << "Unloading SRTM3 tiles from GPU"; | ||
| REPORT_WHEN_CUDA_ERR(cudaFree(this->device_formated_srtm_buffers_info_)); |
There was a problem hiding this comment.
Due to this exception can be thrown from a destructor, e.g terminate()
| dem_assistant->GetElevationManager()->TransferToDevice(); | ||
| backgeocoding.SetElevationData( | ||
| dem_assistant->GetEgm96Manager()->GetDeviceValues(), | ||
| {const_cast<alus::PointerHolder*>(dem_assistant->GetElevationManager()->GetBuffers()), |
There was a problem hiding this comment.
Could it be reasonable to get rid of const_casts now that major changes to dem handling have been made?
| return nullptr; | ||
| } | ||
|
|
||
| inline __device__ double CopDemCog30mGetElevation(double geo_pos_lat, double geo_pos_lon, PointerArray* p_array, |
There was a problem hiding this comment.
For future reference:
I think there is a potential optimization here with regards to PointerArray* and Property* arrays in regards to limiting global memory accesses:
They basic idea would be to have a reasonable upper bound number for the number of tiles, what ever it is - let's say 24. Now if we changed the actual Kernel that calls this function as:
struct DemArgs
{
int count;
PointerArray[24] tiles;
dem::Property[24] dem_prop;
};
If the original kernel had arguments like this, then all the for loops each kernels do on these would be constant memory accesses instead of global memory reads. Another way to achieve a similar result would be to use these directly via __ constant __ memory usage. I think pretty much only the actual kernel calls would need to change, but most underlying utility functions would stay the same interface side.
I think this could improve the memory performance on these kernels - however no point doing this yet, as likely these kernels are limited by double FLOPS on consumer GPUs - however the same idea may became relevant in a similar use case.
util/dem/include/dem_property.h
Outdated
| struct Property { | ||
| double tile_pixel_count_inverted_x; // NUM_PIXELS_PER_TILE_INVERTED | ||
| double tile_pixel_count_inverted_y; // NUM_PIXELS_PER_TILE_INVERTED | ||
| size_t tile_pixel_count_x; // NUM_PIXELS_PER_TILE |
There was a problem hiding this comment.
By default most integers used by Cuda kernels should default to int or unsigned int, not size_t as there is are no native 64 integers on GPUs.
