Replace THCudaCheck with C10_CUDA_CHECK

apex/contrib/csrc/multihead_attn/additive_masked_softmax_dropout_cuda.cu

@@ -5,14 +5,9 @@
 #include <cuda.h>
 #include <cuda_runtime.h>
 #include <cuda_fp16.h>
-<<<<<<< HEAD
 //#include <cuda_profiler_api.h>
-#include "THC/THC.h"
-=======
-#include <cuda_profiler_api.h>
 
 #include <ATen/ATen.h>
->>>>>>> 0c7d8e3 (remove THC headers/functions (#1192))
 #include <ATen/cuda/CUDAContext.h>
 #include <torch/extension.h>
 

apex/contrib/csrc/multihead_attn/dropout.h

@@ -220,7 +220,7 @@ void apex_fused_dropout_cuda(scalar_t const *inputs,
   }
 
   apex_fused_dropout_kernel<scalar_t, accscalar_t, IndexType><<<grid, dim_block, 0, at::cuda::getCurrentCUDAStream()>>>(inputs, outputs, mask, totalElements, p, rng_engine_inputs);
-  THCudaCheck(cudaGetLastError());
+  C10_CUDA_CHECK(cudaGetLastError());
 }
 
 template <
@@ -258,7 +258,7 @@ void apex_dropout_add_cuda(scalar_t const *inputs,
   }
 
   apex_dropout_add_kernel<scalar_t, accscalar_t, IndexType><<<grid, dim_block, 0, at::cuda::getCurrentCUDAStream()>>>(inputs, add_inputs, outputs, mask, totalElements, p, rng_engine_inputs);
-  THCudaCheck(cudaGetLastError());
+  C10_CUDA_CHECK(cudaGetLastError());
 }
 
 template <
@@ -279,7 +279,7 @@ void apex_add_cuda(scalar_t const *inputs,
   grid.x = std::min((unsigned int)at::cuda::getCurrentDeviceProperties()->multiProcessorCount * blocks_per_sm, grid.x);
 
   apex_add_kernel<scalar_t, accscalar_t, IndexType><<<grid, dim_block, 0, at::cuda::getCurrentCUDAStream()>>>(inputs, add_inputs, outputs, totalElements);
-  THCudaCheck(cudaGetLastError());
+  C10_CUDA_CHECK(cudaGetLastError());
 }
 
 template<typename scalar_t, 
@@ -300,7 +300,7 @@ void apex_masked_scale_cuda(scalar_t const *inputs,
   grid.x = std::min((unsigned int)at::cuda::getCurrentDeviceProperties()->multiProcessorCount * blocks_per_sm, grid.x);
 
   apex_masked_scale_kernel<scalar_t, accscalar_t, IndexType><<<grid, dim_block, 0, at::cuda::getCurrentCUDAStream()>>>(inputs, outputs, mask, totalElements, scale);
-  THCudaCheck(cudaGetLastError());
+  C10_CUDA_CHECK(cudaGetLastError());
 }
 
 

apex/contrib/csrc/multihead_attn/self_multihead_attn_cuda.cu

@@ -502,7 +502,7 @@ std::vector<torch::Tensor> bwd_cuda(
                              algo,
                              solution_index,
                              flags));
-  TORCH_CUDABLAS_CHECK(cublasSetMathMode(handle, CUBLAS_DEFAULT_MATH));
+  //TORCH_CUDABLAS_CHECK(cublasSetMathMode(handle, CUBLAS_DEFAULT_MATH));
   
   return { 
            input_grads, 

apex/contrib/csrc/optimizers/fused_adam_cuda_kernel.cu

@@ -276,7 +276,7 @@ void fused_adam_cuda(
                         decay);
             );
       }
-      THCudaCheck(cudaGetLastError());
+      C10_CUDA_CHECK(cudaGetLastError());
 
 }
 
@@ -383,7 +383,7 @@ void fused_adam_cuda_mt(
             );
         }
     }
-    THCudaCheck(cudaGetLastError());
+    C10_CUDA_CHECK(cudaGetLastError());
 }
 
 template <typename FROM_T, typename TO_T> 
@@ -808,7 +808,7 @@ void fused_strided_check_finite(
                     stride,
                     clear_overflow_first);
                 );
-	THCudaCheck(cudaGetLastError());
+	C10_CUDA_CHECK(cudaGetLastError());
 }
 
 void fused_reversible_adam_cuda(
@@ -909,7 +909,7 @@ void fused_reversible_adam_cuda(
                       decay);
                   );
       }
-      THCudaCheck(cudaGetLastError());
+      C10_CUDA_CHECK(cudaGetLastError());
 }
 
 void maybe_cast_cuda(
@@ -933,7 +933,7 @@ void maybe_cast_cuda(
                       p_in.DATA_PTR<scalar_t_0>(),
                       p_out.DATA_PTR<scalar_t_1>(),
                       tsize); ))
-      THCudaCheck(cudaGetLastError());
+      C10_CUDA_CHECK(cudaGetLastError());
 }
 
 void maybe_cast_cuda_mt(
@@ -955,7 +955,7 @@ void maybe_cast_cuda_mt(
                     overflow_flag,
                     tensor_lists,
                     MaybeCastFunctor<2, scalar_t_0, scalar_t_1>()); ))
-    THCudaCheck(cudaGetLastError());
+    C10_CUDA_CHECK(cudaGetLastError());
 }
 
 void fused_maybe_adam_undo_cuda(
@@ -1033,5 +1033,5 @@ void fused_maybe_adam_undo_cuda(
                     decay);
                 );
     }
-    THCudaCheck(cudaGetLastError());
+    C10_CUDA_CHECK(cudaGetLastError());
 }

apex/contrib/csrc/optimizers/multi_tensor_distopt_adam_kernel.cu

@@ -224,5 +224,5 @@ void multi_tensor_fused_adam_cuda(
         (adamMode_t) mode);
     );
   }
-  THCudaCheck(cudaGetLastError());
+  C10_CUDA_CHECK(cudaGetLastError());
 }

apex/contrib/csrc/transducer/transducer_joint_kernel.cu

@@ -823,7 +823,7 @@ std::vector<torch::Tensor> transducer_joint_cuda_forward(
         }));  
     }
  
-    THCudaCheck(cudaGetLastError());
+    C10_CUDA_CHECK(cudaGetLastError());
     if (masked) 
         return {sum, mask};
     else

apex/contrib/csrc/transducer/transducer_loss_kernel.cu

@@ -640,7 +640,7 @@ std::vector<torch::Tensor> transducer_loss_cuda_forward(
                 loss.data_ptr<scalar_t>());  
 
     }));
-    THCudaCheck(cudaGetLastError());
+    C10_CUDA_CHECK(cudaGetLastError());
 
     return {alpha, beta, loss};
 }
@@ -761,7 +761,7 @@ torch::Tensor transducer_loss_cuda_backward(
                 xGrad.data_ptr<scalar_t>());
         }));
     }
-    THCudaCheck(cudaGetLastError());
+    C10_CUDA_CHECK(cudaGetLastError());
     
     return xGrad;
 }

apex/contrib/csrc/xentropy/xentropy_kernel.cu

@@ -634,7 +634,7 @@ std::vector<Tensor> host_softmax_xentropy(
     }
   );
 
-  THCudaCheck(cudaGetLastError());
+  C10_CUDA_CHECK(cudaGetLastError());
 
   std::vector<at::Tensor> ret = {losses, max_log_sum_exp};
   return ret;
@@ -704,7 +704,7 @@ Tensor host_softmax_xentropy_backward(
     }
   );
 
-  THCudaCheck(cudaGetLastError());
+  C10_CUDA_CHECK(cudaGetLastError());
   return gI;
 }