fix asserts in cuda code

pytorch · May 27, 2020 · 79ee1f5 · 79ee1f5
1 parent c6e9e93
commit 79ee1f5
Show file tree

Hide file tree

Showing 6 changed files with 11 additions and 11 deletions.
diff --git a/aten/src/ATen/native/cuda/EmbeddingBag.cu b/aten/src/ATen/native/cuda/EmbeddingBag.cu
@@ -54,7 +54,7 @@ __global__ void EmbeddingBag_updateOutputKernel(
       scalar_t *weightFeat = weight + featureDim * weight_stride1;
       int64_t begin = bag == 0 ? 0 : offsets[bag]; // forces first offset to be 0 instead of asserting on it
       int64_t end = (bag < numBags - 1) ? (offsets[bag + 1]) : numIndices;
-      assert(end >= begin);
+      CUDA_KERNEL_ASSERT(end >= begin);
 
       accscalar_t weightFeatSum = 0;
       scalar_t weightFeatMax;

diff --git a/aten/src/ATen/native/cuda/SortingRadixSelect.cuh b/aten/src/ATen/native/cuda/SortingRadixSelect.cuh
@@ -65,7 +65,7 @@ struct TopKTypeConfig<int16_t> {
   typedef uint32_t RadixType;
 
   static inline __device__ RadixType convert(int16_t v) {
-    assert(sizeof(short) == 2);
+    static_assert(sizeof(short) == 2);
     return 32768u + v;
   }
 
@@ -79,7 +79,7 @@ struct TopKTypeConfig<int32_t> {
   typedef uint32_t RadixType;
 
   static inline __device__ RadixType convert(int32_t v) {
-    assert(sizeof(int) == 4);
+    static_assert(sizeof(int) == 4);
     return 2147483648u + v;
   }
 
@@ -93,7 +93,7 @@ struct TopKTypeConfig<int64_t> {
   typedef uint64_t RadixType;
 
   static inline __device__ RadixType convert(int64_t v) {
-    assert(sizeof(int64_t) == 8);
+    static_assert(sizeof(int64_t) == 8);
     return 9223372036854775808ull + v;
   }
 

diff --git a/aten/src/THC/THCTensorInfo.cuh b/aten/src/THC/THCTensorInfo.cuh
@@ -73,15 +73,15 @@ TensorInfo<T, IndexType>::TensorInfo(T* p,
 template <typename T, typename IndexType>
 void
 TensorInfo<T, IndexType>::reduceDim(int dim) {
-  assert(dim < dims && dim >= 0);
+  TORCH_INTERNAL_ASSERT(dim < dims && dim >= 0);
   sizes[dim] = 1;
 }
 
 template <typename T, typename IndexType>
 int
 TensorInfo<T, IndexType>::collapseDims(const int excludeDim) {
 
-  assert(excludeDim >= -1 && excludeDim < dims);
+  TORCH_INTERNAL_ASSERT(excludeDim >= -1 && excludeDim < dims);
 
   int stopDim = (excludeDim == -1) ? dims : excludeDim;
   int newIndex = -1;

diff --git a/aten/src/THC/generic/THCTensorMode.cu b/aten/src/THC/generic/THCTensorMode.cu
@@ -269,7 +269,7 @@ void THCTensor_(mode)(THCState *state,
         break;
       case 1:
       default:
-        assert(false);
+        TORCH_INTERNAL_ASSERT(false);
     }
     THCudaCheck(cudaGetLastError());
 

diff --git a/aten/src/THC/generic/THCTensorSort.cu b/aten/src/THC/generic/THCTensorSort.cu
@@ -101,7 +101,7 @@ void THCTensor_(sortKeyValueInplace)(THCState* state,
       /* Nothing to do, data already sorted */          \
       break;                                            \
       default:                                          \
-      assert(false);                                    \
+      TORCH_INTERNAL_ASSERT(false);                                    \
     }                                                   \
   }
 

diff --git a/caffe2/operators/top_k_radix_selection.cuh b/caffe2/operators/top_k_radix_selection.cuh
@@ -76,7 +76,7 @@ struct TopKTypeConfig<short> {
   typedef unsigned int RadixType;
 
   static inline __device__ RadixType convert(short v) {
-    CUDA_KERNEL_ASSERT(sizeof(short) == 2);
+    static_assert(sizeof(short) == 2);
     return 32768u + v;
   }
 
@@ -90,7 +90,7 @@ struct TopKTypeConfig<int> {
   typedef unsigned int RadixType;
 
   static inline __device__ RadixType convert(int v) {
-    CUDA_KERNEL_ASSERT(sizeof(int) == 4);
+    static_assert(sizeof(int) == 4);
     return 2147483648u + v;
   }
 
@@ -104,7 +104,7 @@ struct TopKTypeConfig<long> {
   typedef unsigned long long int RadixType;
 
   static inline __device__ RadixType convert(long v) {
-    CUDA_KERNEL_ASSERT(sizeof(long) == 8);
+    static_assert(sizeof(long) == 8);
     return 9223372036854775808ull + v;
   }