Update README and fix bug

README.md

@@ -10,11 +10,7 @@ arXiv: <https://arxiv.org/abs/2112.05131>
 
 [Featured at Two Minute Papers YouTube](https://youtu.be/yptwRRpPEBM) 2022-01-11
 
-**Note:** This is a preliminary release. We have not carefully tested everything,
-but feel that it would be better to first put the code out there.
-
-Also, despite the name, it's not strictly intended to be a successor of svox
-
+Despite the name, it's not strictly intended to be a successor of svox
 
 Citation:
 ```
@@ -34,7 +30,10 @@ A JAX implementation is also available at <https://github.com/sarafridov/plenoxe
 
 ![Overview](https://raw.githubusercontent.com/sxyu/svox2/master/github_img/pipeline.png)
 
-Please also check out PeRFCeption [Jeong, Shin, Lee, et al], which uses Plenoxels with tuned parameters:
+### Examples
+
+Check out PeRFCeption [Jeong, Shin, Lee, et al], which uses Plenoxels with tuned parameters to generate a large
+dataset of radiance fields:
 https://github.com/POSTECH-CVLab/PeRFception
 
 Artistic Radiance Fields by Kai Zhang et al
@@ -76,6 +75,9 @@ Note this data should be identical to that in NeRF++
 Finally, the real Lego capture can be downloaded from:
 https://drive.google.com/file/d/1PG-KllCv4vSRPO7n5lpBjyTjlUyT8Nag/view?usp=sharing
 
+**Note: we currently do not support the instant-ngp format data (since the project was released before NGP). Using it will trigger the nerf-synthetic (Blender) data loader
+due to similarity, but will not train properly. For real data we use the NSVF format.**
+
 ## Voxel Optimization (aka Training)
 
 For training a single scene, see `opt/opt.py`. The launch script makes this easier.

svox2/csrc/render_lerp_kernel_cuvol.cu

@@ -249,19 +249,20 @@ __device__ __inline__ void trace_ray_cuvol_backward(
                       fmaf(color_cache[1], grad_output[1],
                            color_cache[2] * grad_output[2]));
 
+    if (ray.tmin > ray.tmax) {
+        if (accum_out != nullptr) { *accum_out = accum; }
+        if (log_transmit_out != nullptr) { *log_transmit_out = 0.f; }
+        // printf("accum_end_fg_fast=%f\n", accum);
+        return;
+    }
+
     if (beta_loss > 0.f) {
         const float transmit_in = _EXP(log_transmit_in);
         beta_loss *= (1 - transmit_in / (1 - transmit_in + 1e-3)); // d beta_loss / d log_transmit_in
         accum += beta_loss;
         // Interesting how this loss turns out, kinda nice?
     }
 
-    if (ray.tmin > ray.tmax) {
-        if (accum_out != nullptr) { *accum_out = accum; }
-        if (log_transmit_out != nullptr) { *log_transmit_out = 0.f; }
-        // printf("accum_end_fg_fast=%f\n", accum);
-        return;
-    }
     float t = ray.tmin;
 
     const float gout = grad_output[lane_colorgrp];