Fixed uninitialized entries in the gradient of `bary_img` for unused pixels

Summary:
For unused warps we always write zeros to `bary_img_grad`.
However it is possible that the warp is used, but only portion of the threads are used. In this case, for unused threads we do not write zeros to `bary_img_grad`.

For efficiency, `bary_img_grad` is created with `at::empty`, thus  those before mentioned entries, will still have uninitialized values.

This is not an issue, because the `render` function will never read those entries, however it is possible that the uninitialized values will coincide with `nan` and it will trigger a false positive detection in auto grad anomaly detection.  Please see more details in D60904848 about the issue.

Differential Revision: D60912474

fbshipit-source-id: 6eda5a07789db456c17eb60de222dd4c7b1c53d2

src/interpolate/interpolate_kernel.cu

@@ -234,11 +234,16 @@ __global__ void interpolate_backward_kernel(
                 vert_2_grad_ptr, i * vert_attributes_grad_sC, memory_span, grad_v_2, true);
         }
       }
-
-      if (thread_is_used && bary_img_requires_grad) {
-        bary_grad_ptr[0 * bary_img_grad_sB] = bary_0_grad;
-        bary_grad_ptr[1 * bary_img_grad_sB] = bary_1_grad;
-        bary_grad_ptr[2 * bary_img_grad_sB] = bary_2_grad;
+      if (bary_img_requires_grad) {
+        if (thread_is_used) {
+          bary_grad_ptr[0 * bary_img_grad_sB] = bary_0_grad;
+          bary_grad_ptr[1 * bary_img_grad_sB] = bary_1_grad;
+          bary_grad_ptr[2 * bary_img_grad_sB] = bary_2_grad;
+        } else {
+          bary_grad_ptr[0 * bary_img_grad_sB] = scalar_t(0.);
+          bary_grad_ptr[1 * bary_img_grad_sB] = scalar_t(0.);
+          bary_grad_ptr[2 * bary_img_grad_sB] = scalar_t(0.);
+        }
       }
     } else if ((index < nthreads) && bary_img_requires_grad) {
       bary_grad_ptr[0 * bary_img_grad_sB] = scalar_t(0.);