From 16f4b0ed6b7f94f8afb89a75bc5c1a3c8c26202d Mon Sep 17 00:00:00 2001 From: Amogh Akshintala Date: Tue, 15 Dec 2020 13:57:54 -0800 Subject: [PATCH] Replace THError() check in THCTensorMathReduce.cu with C10_CUDA_KERNEL_LAUNCH_CHECK() (#49424) Summary: Pull Request resolved: https://github.com/pytorch/pytorch/pull/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 (https://github.com/pytorch/pytorch/commit/e2510a0b60232aba5160ceb18b6ece8c59a9b79d), 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 | 9 +++------ 1 file changed, 3 insertions(+), 6 deletions(-) diff --git a/aten/src/THC/generic/THCTensorMathReduce.cu b/aten/src/THC/generic/THCTensorMathReduce.cu index ce2f124215ca..13108491b0ec 100644 --- a/aten/src/THC/generic/THCTensorMathReduce.cu +++ b/aten/src/THC/generic/THCTensorMathReduce.cu @@ -2,6 +2,8 @@ #define THC_GENERIC_FILE "THC/generic/THCTensorMathReduce.cu" #else +#include + #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 <<>>(THCTensor_(data)(state, data), scalar_cast(value), size, scalar_cast(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_);