diff --git a/aten/src/ATen/cpu/vec/vec256/vec256_convert.h b/aten/src/ATen/cpu/vec/vec256/vec256_convert.h index 858d87a5d83f..d5910f957b6d 100644 --- a/aten/src/ATen/cpu/vec/vec256/vec256_convert.h +++ b/aten/src/ATen/cpu/vec/vec256/vec256_convert.h @@ -271,9 +271,9 @@ struct VecConvert< 1, int64_t, 2, - typename std::enable_if< + std::enable_if_t< std::is_same_v || - std::is_same_v>::type> { + std::is_same_v>> { static inline VectorizedN apply( const VectorizedN& src) { return VecConvert::apply( diff --git a/aten/src/ATen/cpu/vec/vec256/zarch/vec256_zarch.h b/aten/src/ATen/cpu/vec/vec256/zarch/vec256_zarch.h index 4ca57363ee4b..a85112f54a1e 100644 --- a/aten/src/ATen/cpu/vec/vec256/zarch/vec256_zarch.h +++ b/aten/src/ATen/cpu/vec/vec256/zarch/vec256_zarch.h @@ -2868,7 +2868,7 @@ std::pair, Vectorized> inline deinterleave2< } template -typename std::enable_if::value, at::vec::Vectorized>::type +std::enable_if_t, at::vec::Vectorized> inline convert_int8_to_float(const Vectorized &src) { // Note: this function only convert inputs number of elements equal to at::vec::Vectorized.size() // Only handle first 64 bits @@ -2878,7 +2878,7 @@ inline convert_int8_to_float(const Vectorized &src) { } template -typename std::enable_if::value, at::vec::Vectorized>::type +std::enable_if_t, at::vec::Vectorized> inline convert_float_to_int8(const Vectorized &src) { constexpr auto min_val = std::numeric_limits::min(); constexpr auto max_val = std::numeric_limits::max(); diff --git a/aten/src/ATen/cpu/vec/vec512/vec512_convert.h b/aten/src/ATen/cpu/vec/vec512/vec512_convert.h index cfb4ddb13732..af4801cccf48 100644 --- a/aten/src/ATen/cpu/vec/vec512/vec512_convert.h +++ b/aten/src/ATen/cpu/vec/vec512/vec512_convert.h @@ -281,9 +281,9 @@ struct VecConvert< 1, int64_t, 2, - typename std::enable_if< + std::enable_if_t< std::is_same_v || - std::is_same_v>::type> { + std::is_same_v>> { static inline VectorizedN apply( const VectorizedN& src) { return VecConvert::apply( diff --git a/aten/src/ATen/cpu/vec/vec512/vec512_mask.h b/aten/src/ATen/cpu/vec/vec512/vec512_mask.h index cdb433af2525..d32e1da1cf72 100644 --- a/aten/src/ATen/cpu/vec/vec512/vec512_mask.h +++ b/aten/src/ATen/cpu/vec/vec512/vec512_mask.h @@ -84,9 +84,9 @@ struct VecMaskLoad< dst_n, mask_t, dst_n, - typename std::enable_if< + std::enable_if_t< std::is_same_v || - std::is_same_v>::type> { + std::is_same_v>> { static inline VectorizedN apply( const data_t* ptr, const VecMask& vec_mask) { @@ -151,9 +151,9 @@ struct VecMaskLoad< 1, mask_t, 1, - typename std::enable_if< + std::enable_if_t< std::is_same_v || - std::is_same_v>::type> { + std::is_same_v>> { static inline VectorizedN apply( const data_t* ptr, const VecMask& vec_mask) { @@ -173,9 +173,9 @@ struct VecMaskLoad< 2, mask_t, 1, - typename std::enable_if< + std::enable_if_t< std::is_same_v || - std::is_same_v>::type> { + std::is_same_v>> { static inline VectorizedN apply( const data_t* ptr, const VecMask& vec_mask) { diff --git a/aten/src/ATen/native/EmbeddingBag.cpp b/aten/src/ATen/native/EmbeddingBag.cpp index fbff571fecec..58dc1b991d26 100644 --- a/aten/src/ATen/native/EmbeddingBag.cpp +++ b/aten/src/ATen/native/EmbeddingBag.cpp @@ -106,7 +106,7 @@ bool is_fast_path(const Tensor& src, const std::optional& scale, Tensor& // index_add (using add_indices as the index), without creating an intermediary // tensor to hold the selected embeddings template -static typename std::enable_if::value, void>::type +static std::enable_if_t, void> index_select_add( const Tensor& select_indices, const Tensor& add_indices, @@ -184,10 +184,9 @@ void fbgemm_spmdm_report_error_( } // namespace template -typename std::enable_if< - std::is_same::value || - std::is_same::value, - void>::type +std::enable_if_t< + std::is_same_v || std::is_same_v, + void> index_select_add( const Tensor& select_indices, const Tensor& add_indices, @@ -366,7 +365,7 @@ index_select_add( } } template -typename std::enable_if::value, void>::type +std::enable_if_t, void> index_select_add(const Tensor &select_indices, const Tensor &add_indices, const Tensor &src, @@ -493,7 +492,7 @@ index_select_add(const Tensor &select_indices, // mul (scaling by per_sample_weights) // index_add (using add_indices as the index) template -static typename std::enable_if::value, void>::type +static std::enable_if_t, void> index_select_scale_add( const Tensor& select_indices, const Tensor& add_indices, @@ -548,10 +547,9 @@ index_select_scale_add( } template -typename std::enable_if< - std::is_same::value || - std::is_same::value, - void>::type +std::enable_if_t< + std::is_same_v || std::is_same_v, + void> index_select_scale_add( const Tensor& select_indices, const Tensor& add_indices, @@ -741,7 +739,7 @@ index_select_scale_add( } } template -typename std::enable_if::value, void>::type +std::enable_if_t, void> index_select_scale_add(const Tensor &select_indices, const Tensor &add_indices, const Tensor &scale, diff --git a/aten/src/ATen/native/Pow.h b/aten/src/ATen/native/Pow.h index 76ddda846a59..f0e03b13f8f2 100644 --- a/aten/src/ATen/native/Pow.h +++ b/aten/src/ATen/native/Pow.h @@ -23,7 +23,7 @@ namespace native { // e.g. since 2**-1==0.5, the truncated integral result is zero. 1**negative_exponent is the // only non-zero result. template ::value, T>::type* = nullptr> + std::enable_if_t, T>* = nullptr> inline HOST_DEVICE __ubsan_ignore_signed_int_overflow__ T powi_impl(T a, T b) { T result = 1; while (b) { @@ -37,13 +37,13 @@ inline HOST_DEVICE __ubsan_ignore_signed_int_overflow__ T powi_impl(T a, T b) { } template ::value && !std::is_signed::value, T>::type* = nullptr> + std::enable_if_t && !std::is_signed_v, T>* = nullptr> inline HOST_DEVICE T powi(T a, T b) { return powi_impl(a, b); } template ::value && std::is_signed::value, T>::type* = nullptr> + std::enable_if_t && std::is_signed_v, T>* = nullptr> inline HOST_DEVICE T powi(T a, T b) { if ( b < 0 ) { if ( a == 1 ) { diff --git a/aten/src/ATen/native/ReduceOps.cpp b/aten/src/ATen/native/ReduceOps.cpp index e9201df2f336..766fa7f8d981 100644 --- a/aten/src/ATen/native/ReduceOps.cpp +++ b/aten/src/ATen/native/ReduceOps.cpp @@ -753,11 +753,11 @@ Tensor cumprod_backward(const Tensor& grad, const Tensor& input, int64_t dim, co namespace { #ifdef _MSC_VER template -inline typename std::enable_if::value, bool>::type isnan_(T x) { +inline std::enable_if_t, bool> isnan_(T x) { return false; } template -inline typename std::enable_if::value, bool>::type isnan_(T x) { +inline std::enable_if_t, bool> isnan_(T x) { return std::isnan(x); } #else diff --git a/aten/src/ATen/native/cpu/BlasKernel.cpp b/aten/src/ATen/native/cpu/BlasKernel.cpp index dd0375698219..d2475093bdd7 100644 --- a/aten/src/ATen/native/cpu/BlasKernel.cpp +++ b/aten/src/ATen/native/cpu/BlasKernel.cpp @@ -96,7 +96,7 @@ auto sum(int64_t N, Func f) { } template -typename std::enable_if::value, void>::type +std::enable_if_t, void> gemm_notrans_( int64_t m, int64_t n, @@ -132,7 +132,7 @@ gemm_notrans_( // std::is_same || std::is_same template -typename std::enable_if::value, void>::type +std::enable_if_t, void> gemm_notrans_( int64_t m, int64_t n, @@ -222,7 +222,7 @@ void gemm_transb_impl( } template -typename std::enable_if::value, void>::type +std::enable_if_t, void> gemm_transb_( TransposeType transb, int64_t m, @@ -244,7 +244,7 @@ gemm_transb_( // std::is_same || std::is_same template -typename std::enable_if::value, void>::type +std::enable_if_t, void> gemm_transb_( TransposeType transb, int64_t m, diff --git a/aten/src/ATen/native/cpu/FusedAdagradKernel.cpp b/aten/src/ATen/native/cpu/FusedAdagradKernel.cpp index e19915e0a4f2..24f04111f12c 100644 --- a/aten/src/ATen/native/cpu/FusedAdagradKernel.cpp +++ b/aten/src/ATen/native/cpu/FusedAdagradKernel.cpp @@ -12,10 +12,10 @@ namespace at::native { namespace{ template -typename std::enable_if< - std::is_same::value || std::is_same::value, - void>:: - type inline adagrad_math( +std::enable_if_t< + std::is_same_v || std::is_same_v, + void> + inline adagrad_math( scalar_t* param_ptr, scalar_t* grad_ptr, scalar_t* state_sum_ptr, @@ -81,10 +81,10 @@ typename std::enable_if< template -typename std::enable_if< - std::is_same::value || std::is_same::value, - void>:: - type inline adagrad_math( +std::enable_if_t< + std::is_same_v || std::is_same_v, + void> + inline adagrad_math( scalar_t* param_ptr, scalar_t* grad_ptr, scalar_t* state_sum_ptr, diff --git a/aten/src/ATen/native/cpu/FusedAdamKernel.cpp b/aten/src/ATen/native/cpu/FusedAdamKernel.cpp index 239cdc3b37ac..f583e089c6c0 100644 --- a/aten/src/ATen/native/cpu/FusedAdamKernel.cpp +++ b/aten/src/ATen/native/cpu/FusedAdamKernel.cpp @@ -12,10 +12,10 @@ namespace at::native { namespace{ template -typename std::enable_if< - std::is_same::value || std::is_same::value, - void>:: - type inline adam_math( +std::enable_if_t< + std::is_same_v || std::is_same_v, + void> + inline adam_math( scalar_t* param_ptr, scalar_t* exp_avg_ptr, scalar_t* exp_avg_sq_ptr, @@ -155,10 +155,10 @@ typename std::enable_if< template -typename std::enable_if< - std::is_same::value || std::is_same::value, - void>:: - type inline adam_math( +std::enable_if_t< + std::is_same_v || std::is_same_v, + void> + inline adam_math( scalar_t* param_ptr, scalar_t* exp_avg_ptr, scalar_t* exp_avg_sq_ptr, diff --git a/aten/src/ATen/native/cpu/FusedSGDKernel.cpp b/aten/src/ATen/native/cpu/FusedSGDKernel.cpp index 95e96ff5cf55..023feeb16fe0 100644 --- a/aten/src/ATen/native/cpu/FusedSGDKernel.cpp +++ b/aten/src/ATen/native/cpu/FusedSGDKernel.cpp @@ -12,10 +12,10 @@ namespace at::native { namespace{ template -typename std::enable_if< - std::is_same::value || std::is_same::value, - void>:: - type inline sgd_math( +std::enable_if_t< + std::is_same_v || std::is_same_v, + void> + inline sgd_math( scalar_t* param_ptr, scalar_t* grad_ptr, scalar_t* momentum_buf_ptr, @@ -104,10 +104,10 @@ typename std::enable_if< template -typename std::enable_if< - std::is_same::value || std::is_same::value, - void>:: - type inline sgd_math( +std::enable_if_t< + std::is_same_v || std::is_same_v, + void> + inline sgd_math( scalar_t* param_ptr, scalar_t* grad_ptr, scalar_t* momentum_buf_ptr, diff --git a/aten/src/ATen/native/cpu/IsContiguous.h b/aten/src/ATen/native/cpu/IsContiguous.h index ddbbb6fb8f5a..02d8f5dd78e4 100644 --- a/aten/src/ATen/native/cpu/IsContiguous.h +++ b/aten/src/ATen/native/cpu/IsContiguous.h @@ -31,14 +31,16 @@ struct IsContiguous<0, -1, traits, s> { }; // output and all inputs are contiguous -template ::value>::type* = nullptr> +template < + typename traits, + std::enable_if_t>* = + nullptr> static inline bool is_contiguous(const int64_t* strides) { return IsContiguous::eval(strides); } template ::value>::type* = nullptr> + std::enable_if_t>* = nullptr> static inline bool is_contiguous(const int64_t* strides) { return IsContiguous::eval(strides); } @@ -46,14 +48,14 @@ static inline bool is_contiguous(const int64_t* strides) { // input at `s` is scalar (stride 0); output and other inputs are contiguous // NB: output is typically at strides[0] so first input corresponds to s=1 template ::value>::type* = nullptr> + std::enable_if_t>* = nullptr> static inline bool is_contiguous_scalar(const int64_t* strides) { static_assert(s > 0 && s <= traits::arity, "scalar argument index out of bounds"); return IsContiguous::eval(strides); } template ::value>::type* = nullptr> + std::enable_if_t>* = nullptr> static inline bool is_contiguous_scalar(const int64_t* strides) { static_assert(s > 0 && s <= traits::arity, "scalar argument index out of bounds"); return IsContiguous::eval(strides); diff --git a/aten/src/ATen/native/cpu/MaxPoolKernel.cpp b/aten/src/ATen/native/cpu/MaxPoolKernel.cpp index c752106130fe..15b784f05521 100644 --- a/aten/src/ATen/native/cpu/MaxPoolKernel.cpp +++ b/aten/src/ATen/native/cpu/MaxPoolKernel.cpp @@ -64,7 +64,7 @@ vec::Vectorized is_nan_vec(vec::Vectorized vec) { template inline -typename std::enable_if::value, void>::type +std::enable_if_t, void> compute_internal( const scalar_t* input_data, scalar_t* out_data, @@ -139,7 +139,7 @@ compute_internal( // std::is_same || std::is_same template inline -typename std::enable_if::value, void>::type +std::enable_if_t, void> compute_internal( const scalar_t* input_data, scalar_t* out_data, diff --git a/aten/src/ATen/native/cpu/Reduce.h b/aten/src/ATen/native/cpu/Reduce.h index 4e512119d3fe..aeec6a3fda7e 100644 --- a/aten/src/ATen/native/cpu/Reduce.h +++ b/aten/src/ATen/native/cpu/Reduce.h @@ -129,13 +129,13 @@ static void set_results(const res_t result, const TensorIteratorBase &iter, cons } template -inline typename std::enable_if::type +inline std::enable_if_t for_each_in_tuple(const std::tuple& /*t*/, const TensorIteratorBase& /*iter*/, const int /*num_outputs*/) { return i; } template -inline typename std::enable_if::type +inline std::enable_if_t for_each_in_tuple(const std::tuple& t, const TensorIteratorBase &iter, const int num_outputs) { if (i < (size_t)num_outputs) { set_result(i, std::get(t), iter, num_outputs); diff --git a/aten/src/ATen/native/cpu/ReduceUtils.h b/aten/src/ATen/native/cpu/ReduceUtils.h index 8c6424f8b0ea..fd7c4a2750a6 100644 --- a/aten/src/ATen/native/cpu/ReduceUtils.h +++ b/aten/src/ATen/native/cpu/ReduceUtils.h @@ -106,7 +106,7 @@ inline void _init(scalar_t* self_ptr, at::opmath_type* buffer_ptr, int } template -inline typename std::enable_if::value, scalar_t>::type +inline std::enable_if_t, scalar_t> _max(const scalar_t& x, const scalar_t& y) { return at::_isnan(y) ? y : std::max(x, y); } @@ -118,14 +118,14 @@ inline Vectorized _max(const Vectorized& x, const Vectorized } template -inline typename std::enable_if::value, Vec2>::type +inline std::enable_if_t, Vec2> _max(const vec_t& x, const vec_t& y) { // vec::maximum propagates NaN return maximum(x, y); } template -inline typename std::enable_if::value, scalar_t>::type +inline std::enable_if_t, scalar_t> _min(const scalar_t& x, const scalar_t& y) { return at::_isnan(y) ? y : std::min(x, y); } @@ -137,7 +137,7 @@ inline Vectorized _min(const Vectorized& x, const Vectorized } template -inline typename std::enable_if::value, Vec2>::type +inline std::enable_if_t, Vec2> _min(const vec_t& x, const vec_t& y) { // vec::minimum propagates NaN return minimum(x, y); diff --git a/aten/src/ATen/native/cpu/group_norm_kernel.cpp b/aten/src/ATen/native/cpu/group_norm_kernel.cpp index f6b7f2a5d481..0aee364b49d8 100644 --- a/aten/src/ATen/native/cpu/group_norm_kernel.cpp +++ b/aten/src/ATen/native/cpu/group_norm_kernel.cpp @@ -85,8 +85,8 @@ void GroupNormKernelImplInternal( } template -typename std::enable_if>::value, - std::tuple>::type +std::enable_if_t>, + std::tuple> ColumnwiseMoments( const T* X_data, int64_t HxW, @@ -118,8 +118,8 @@ ColumnwiseMoments( // std::is_same || std::is_same template -typename std::enable_if>::value, - std::tuple, at::opmath_type>>::type +std::enable_if_t>, + std::tuple, at::opmath_type>> ColumnwiseMoments( const T* X_data, int64_t HxW, @@ -160,7 +160,7 @@ ColumnwiseMoments( } template -inline typename std::enable_if::value, void>::type +inline std::enable_if_t, void> CalcMeanVar( const T* X_ptr, opmath_t* mean_ptr, @@ -183,7 +183,7 @@ CalcMeanVar( // std::is_same || std::is_same template -inline typename std::enable_if::value, void>::type +inline std::enable_if_t, void> CalcMeanVar( const T* X_ptr, opmath_t* mean_ptr, @@ -227,7 +227,7 @@ CalcMeanVar( } template -inline typename std::enable_if::value, void>::type +inline std::enable_if_t, void> ApplyScaleBias( T* Y_ptr, const T* X_ptr, @@ -246,7 +246,7 @@ ApplyScaleBias( // std::is_same || std::is_same template -inline typename std::enable_if::value, void>::type +inline std::enable_if_t, void> ApplyScaleBias( T* Y_ptr, const T* X_ptr, @@ -529,7 +529,7 @@ void GroupNormKernelImpl( template -typename std::enable_if::value, void>::type +std::enable_if_t, void> ComputeInternalGradients( int64_t N, int64_t C, @@ -556,7 +556,7 @@ ComputeInternalGradients( } template -typename std::enable_if::value, void>::type +std::enable_if_t, void> ComputeInternalGradients( int64_t N, int64_t C, @@ -603,7 +603,7 @@ ComputeInternalGradients( } template -inline typename std::enable_if::value, void>::type +inline std::enable_if_t, void> CalcDsDb( const opmath_t* ds_ptr, const opmath_t* db_ptr, @@ -626,7 +626,7 @@ CalcDsDb( } template -inline typename std::enable_if::value, void>::type +inline std::enable_if_t, void> CalcDsDb( const opmath_t* ds_ptr, const opmath_t* db_ptr, @@ -708,7 +708,7 @@ void GroupNormInputBackward( } template -typename std::enable_if::value, void>::type +std::enable_if_t, void> GammaBackward( int64_t N, int64_t C, @@ -755,7 +755,7 @@ GammaBackward( } template -typename std::enable_if::value, void>::type +std::enable_if_t, void> GammaBackward( int64_t N, int64_t C, @@ -817,7 +817,7 @@ GammaBackward( } template -typename std::enable_if::value, void>::type +std::enable_if_t, void> BetaBackward(int64_t N, int64_t C, const opmath_t* db, PT* dbeta) { using Vec = at::vec::Vectorized; constexpr int64_t K = Vec::size(); @@ -841,7 +841,7 @@ BetaBackward(int64_t N, int64_t C, const opmath_t* db, PT* dbeta) { } template -typename std::enable_if::value, void>::type +std::enable_if_t, void> BetaBackward(int64_t N, int64_t C, const opmath_t* db, PT* dbeta) { using Vec = at::vec::Vectorized; using fVec = at::vec::Vectorized; @@ -937,7 +937,7 @@ void GroupNormBackwardKernelImplInternal( } template -inline typename std::enable_if::value, void>::type +inline std::enable_if_t, void> DsDbRowwiseMomentsChannelsLast( const T* dY_ptr, const T* X_ptr, @@ -972,7 +972,7 @@ DsDbRowwiseMomentsChannelsLast( } template -inline typename std::enable_if::value, void>::type +inline std::enable_if_t, void> DsDbRowwiseMomentsChannelsLast( const T* dY_ptr, const T* X_ptr, @@ -1024,10 +1024,10 @@ DsDbRowwiseMomentsChannelsLast( } template -inline typename std::enable_if>::value, +inline std::enable_if_t>, std::tuple< vec::Vectorized, - vec::Vectorized>>::type + vec::Vectorized>> load_util(const T* data_ptr, int64_t n) { using Vec = vec::Vectorized; auto vec0 = Vec::loadu(data_ptr, n > Vec::size() ? Vec::size() : n); @@ -1037,11 +1037,11 @@ load_util(const T* data_ptr, int64_t n) { } template -inline typename std::enable_if>::value, +inline std::enable_if_t>, std::tuple< vec::Vectorized>, vec::Vectorized>> - >::type + > load_util(const T* data_ptr, int64_t n) { using Vec = vec::Vectorized; auto vec = Vec::loadu(data_ptr, n); @@ -1049,7 +1049,7 @@ load_util(const T* data_ptr, int64_t n) { } template -inline typename std::enable_if::value, void>::type +inline std::enable_if_t, void> ApplyInputGradientsChannelsLastColMov( const T* dY_data, const T* X_data, @@ -1097,7 +1097,7 @@ ApplyInputGradientsChannelsLastColMov( } template -inline typename std::enable_if::value, void>::type +inline std::enable_if_t, void> ApplyInputGradientsChannelsLastColMov( const T* dY_data, const T* X_data, @@ -1154,7 +1154,7 @@ ApplyInputGradientsChannelsLastColMov( } template -inline typename std::enable_if::value, void>::type +inline std::enable_if_t, void> ApplyInputGradientsChannelsLastRowMov( const T* dY_data, const T* X_data, @@ -1190,7 +1190,7 @@ ApplyInputGradientsChannelsLastRowMov( } template -inline typename std::enable_if::value, void>::type +inline std::enable_if_t, void> ApplyInputGradientsChannelsLastRowMov( const T* dY_data, const T* X_data, diff --git a/torch/csrc/cuda/GdsFile.cpp b/torch/csrc/cuda/GdsFile.cpp index b95b86b3374f..945da3be6510 100644 --- a/torch/csrc/cuda/GdsFile.cpp +++ b/torch/csrc/cuda/GdsFile.cpp @@ -12,8 +12,7 @@ namespace { // filesystem error and a negative CUfileOpError enum value otherwise). template < class T, - typename std::enable_if::value, std::nullptr_t>::type = - nullptr> + std::enable_if_t, std::nullptr_t> = nullptr> std::string cuGDSFileGetErrorString(T status) { status = std::abs(status); return IS_CUFILE_ERR(status) ? std::string(CUFILE_ERRSTR(status)) @@ -24,8 +23,7 @@ std::string cuGDSFileGetErrorString(T status) { // CUfileError_t template < class T, - typename std::enable_if::value, std::nullptr_t>::type = - nullptr> + std::enable_if_t, std::nullptr_t> = nullptr> std::string cuGDSFileGetErrorString(T status) { std::string errStr = cuGDSFileGetErrorString(static_cast(status.err)); if (IS_CUDA_ERR(status)) diff --git a/torch/csrc/distributed/c10d/CUDASymmetricMemory-inl.h b/torch/csrc/distributed/c10d/CUDASymmetricMemory-inl.h index 36b3591bbdba..0515dc9fca75 100644 --- a/torch/csrc/distributed/c10d/CUDASymmetricMemory-inl.h +++ b/torch/csrc/distributed/c10d/CUDASymmetricMemory-inl.h @@ -418,9 +418,8 @@ __device__ __inline__ Vec add_vec( // With world_size specialization: perform balanced load from all peers before // performing reduction. template -__device__ inline - typename std::enable_if<(k_world_size > 0), Vec>::type - load_and_reduce(T** ptrs, size_t rank, size_t world_size, size_t offset) { +__device__ inline std::enable_if_t<(k_world_size > 0), Vec> +load_and_reduce(T** ptrs, size_t rank, size_t world_size, size_t offset) { Vec vecs[k_world_size]; #pragma unroll k_world_size for (size_t step = 0; step < k_world_size; ++step) { @@ -438,9 +437,8 @@ __device__ inline // Without world_size specialization: perform ordered (unbalanced) load and // accumulate on each load. template -__device__ inline - typename std::enable_if<(k_world_size <= 0), Vec>::type - load_and_reduce(T** ptrs, size_t rank, size_t world_size, size_t offset) { +__device__ inline std::enable_if_t<(k_world_size <= 0), Vec> +load_and_reduce(T** ptrs, size_t rank, size_t world_size, size_t offset) { Vec acc{}; for (size_t step = 0; step < world_size; ++step) { auto vec = ld_vec(ptrs[step] + offset); diff --git a/torch/csrc/jit/tensorexpr/cpp_intrinsics.h b/torch/csrc/jit/tensorexpr/cpp_intrinsics.h index caeeed693ff3..3149335ea30f 100644 --- a/torch/csrc/jit/tensorexpr/cpp_intrinsics.h +++ b/torch/csrc/jit/tensorexpr/cpp_intrinsics.h @@ -8,13 +8,13 @@ constexpr auto cpp_intrinsics_definition = R"( namespace std { template ::value, int>::type = 0> + std::enable_if_t, int> = 0> T rsqrt(T v) { return 1.0f / std::sqrt(v); } template ::value, int>::type = 0> + std::enable_if_t, int> = 0> T frac(T v) { T intpart; return std::modf(v, &intpart); diff --git a/torch/csrc/jit/tensorexpr/llvm_codegen.cpp b/torch/csrc/jit/tensorexpr/llvm_codegen.cpp index a9f24139e029..5fe52a0ff9e0 100644 --- a/torch/csrc/jit/tensorexpr/llvm_codegen.cpp +++ b/torch/csrc/jit/tensorexpr/llvm_codegen.cpp @@ -1075,14 +1075,16 @@ void LLVMCodeGenImpl::visit(const CompareSelectPtr& v) { } template -typename std::enable_if::value, llvm::Value*>::type -getFromType(llvm::Type* type, T value) { +std::enable_if_t, llvm::Value*> getFromType( + llvm::Type* type, + T value) { return llvm::ConstantInt::get(type, value, std::is_signed::value); } template -typename std::enable_if::value, llvm::Value*>::type -getFromType(llvm::Type* type, T value) { +std::enable_if_t, llvm::Value*> getFromType( + llvm::Type* type, + T value) { return llvm::ConstantFP::get(type, value); }