Replace THError() check in THCTensorMathReduce.cu with C10_CUDA_KERNE…

…L_LAUNCH_CHECK() (#49424)

Summary:
Pull Request resolved: #49424

As per conversation in this [comment](https://www.internalfb.com/intern/diff/D25541113 (https://github.com/pytorch/pytorch/commit/e2510a0b60232aba5160ceb18b6ece8c59a9b79d)/?dest_fbid=393026838623691&transaction_id=3818008671564312) on D25541113 (e2510a0), although THError does more than just log any errors associated cuda kernel launches, we're going to go ahead and replace it with C10_CUDA_KERNEL_LAUNCH_CHECK, so as to be consistent throughout the code base.
Standardization FTW.

This commit is purposefully sent in as a single file change so it can be easily reverted if it introduces a regression.

Test Plan:
Checked that the code still builds with
```
buck build //caffe2/aten:ATen-cu
```
Also ran basic aten tests
```
buck test //caffe2/aten:atest
```

Reviewed By: r-barnes

Differential Revision: D25567863

fbshipit-source-id: 1093bfe2b6ca6b9a3bfb79dcdc5d713f6025eb77

aten/src/THC/generic/THCTensorMathReduce.cu

@@ -2,6 +2,8 @@
 #define THC_GENERIC_FILE "THC/generic/THCTensorMathReduce.cu"
 #else
 
+#include <c10/cuda/CUDAException.h>
+
 #if !defined(THC_REAL_IS_BOOL)
 
 void THCTensor_(prod)(THCState* state, THCTensor *self, THCTensor *src, int dimension, int keepdim) {
@@ -43,12 +45,7 @@ void THCTensor_(renorm)(THCState *state, THCTensor* self, THCTensor* src, scalar
     THCTensor_kernel_renorm<scalar_t, accreal>
       <<<grid, threads, 0, c10::cuda::getCurrentCUDAStream()>>>(THCTensor_(data)(state, data),
         scalar_cast<accreal>(value), size, scalar_cast<accreal>(maxnorm));
-
-    // Do not replace with C10_CUDA_KERNEL_LAUNCH_CHECK() yet as it exhibits different behaviour from THError().
-    // THError() calls the an error handler, or throws std::runtime_error if a custom handler hasn't been registered.
-    cudaError_t errcode = cudaGetLastError();
-    if(errcode != cudaSuccess)
-      THError(cudaGetErrorString(errcode));
+    C10_CUDA_KERNEL_LAUNCH_CHECK();
   }
 
   THCTensor_(free)(state, src_);