remove dynamic initialization warning (#13913) (#13967)

Summary:
removed assignment in default constructor.
removed static shared memory and used dynamic shared memory.
Pull Request resolved: https://github.com/pytorch/pytorch/pull/13967

Differential Revision: D13089996

Pulled By: soumith

fbshipit-source-id: 2a218b909c849bed39636b45a02d10ebc279a0b0

diff --git a/aten/src/THC/THCTensorMathReduce.cuh b/aten/src/THC/THCTensorMathReduce.cuh
index 358f926..ce736a4 100644 (file)
--- a/aten/src/THC/THCTensorMathReduce.cuh
+++ b/aten/src/THC/THCTensorMathReduce.cuh
@@ -27,6 +27,12 @@ struct WelfordData {
   int count_; // do we need int64_t?
 
   __host__ __device__ WelfordData() {
+  }
+
+  // stripping initialization from default constructor to avoid dynamic
+  // initialization warning thrown from using this data structure in CUDA kernel
+  // as static shared memory.
+  __host__ __device__ void reset() {
     mean_ = T(0);
     m_2_n_ = T(0);
     count_ = 0;

diff --git a/aten/src/THC/generic/THCTensorMathReduce.cu b/aten/src/THC/generic/THCTensorMathReduce.cu
index c54b38a..9341b53 100644 (file)
--- a/aten/src/THC/generic/THCTensorMathReduce.cu
+++ b/aten/src/THC/generic/THCTensorMathReduce.cu
@@ -88,11 +88,13 @@ void THCTensor_(std)(THCState *state, THCTensor *self_, THCTensor *src, int dime
 {
   THCAssertSameGPU(THCTensor_(checkGPU)(state, 2, self_, src));
 
+  WelfordData<accreal, scalar_t> init;
+  init.reset();
   if (!THC_reduceDim<scalar_t>(state, self_, src,
                            ModifyWelford<WelfordData<accreal, scalar_t>>{},
                            ReduceWelford<accreal, scalar_t>{},
                            VarianceWelford<accreal, scalar_t>{biased, true},
-                           WelfordData<accreal, scalar_t>{},
+                           init,
                            dimension,
                            keepdim)) {
     THArgCheck(false, 2, CUTORCH_DIM_WARNING);
@@ -105,11 +107,13 @@ void THCTensor_(var)(THCState *state, THCTensor *self_, THCTensor *src, int dime
 {
   THCAssertSameGPU(THCTensor_(checkGPU)(state, 2, self_, src));
 
+  WelfordData<accreal, scalar_t> init;
+  init.reset();
   if (!THC_reduceDim<scalar_t>(state, self_, src,
                            ModifyWelford<WelfordData<accreal, scalar_t>>{},
                            ReduceWelford<accreal, scalar_t>{},
                            VarianceWelford<accreal, scalar_t>{biased, false},
-                           WelfordData<accreal, scalar_t>{},
+                           init,
                            dimension,
                            keepdim)) {
     THArgCheck(false, 2, CUTORCH_DIM_WARNING);