423A35C7/pytorch3d
1
0
Fork 0
mirror of https://github.com/facebookresearch/pytorch3d.git synced 2025-08-02 20:02:49 +08:00
Code Issues Packages Projects Releases Wiki Activity

simple warning for bin overflow

Browse Source 
Summary: Since coarse rasterization on cuda can overflow bins, we detect when this happens for memory safety. See https://github.com/facebookresearch/pytorch3d/issues/348 . Also try to print a warning.

Reviewed By: patricklabatut

Differential Revision: D33065604

fbshipit-source-id: 99b3c576d01b78e6d77776cf1a3e95984506c93a
This commit is contained in:
Jeremy Reizenstein 2022-01-06 02:29:50 -08:00 committed by Facebook GitHub Bot
parent d6a12afbe7
commit 6726500ad3
1 changed files with 16 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

16
pytorch3d/csrc/rasterize_coarse/rasterize_coarse.cu
View File

@ -183,6 +183,22 @@ __global__ void RasterizeCoarseCudaKernel(
// this effectively allocates space in the bin_faces array for the // this effectively allocates space in the bin_faces array for the
// elems in the current chunk that fall into this bin. // elems in the current chunk that fall into this bin.
const int start = atomicAdd(elems_per_bin + elems_per_bin_idx, count); const int start = atomicAdd(elems_per_bin + elems_per_bin_idx, count);
if (start + count > M) {
// The number of elems in this bin is so big that they won't fit.
// We print a warning using CUDA's printf. This may be invisible
// to notebook users, but apparent to others. It would be nice to
// also have a Python-friendly warning, but it is not obvious
// how to do this without slowing down the normal case.
const char* warning =
"Bin size was too small in the coarse rasterization phase. "
"This caused an overflow, meaning output may be incomplete. "
"To solve, "
"try increasing max_faces_per_bin / max_points_per_bin, "
"decreasing bin_size, "
"or setting bin_size to -1 to use the naive rasterization.";
printf(warning);
continue;
}
// Now loop over the binmask and write the active bits for this bin // Now loop over the binmask and write the active bits for this bin
// out to bin_faces. // out to bin_faces.