diff --git a/paddle/fluid/operators/elementwise/elementwise_add_op.h b/paddle/fluid/operators/elementwise/elementwise_add_op.h index 731cef3d3662f866e04321bfbc4857bb8993768d..41e97a3946695ec9384d10293b1e54f7b2a5cfad 100644 --- a/paddle/fluid/operators/elementwise/elementwise_add_op.h +++ b/paddle/fluid/operators/elementwise/elementwise_add_op.h @@ -179,6 +179,7 @@ __global__ void MatrixColReduce(const T *__restrict__ in, T *__restrict__ out, } } +#if CUDA_VERSION >= 10000 template __global__ void VecFP16MatrixColReduce(const __half2 *__restrict__ in, __half2 *__restrict__ out, size_t width, @@ -199,6 +200,7 @@ __global__ void VecFP16MatrixColReduce(const __half2 *__restrict__ in, } #endif } +#endif template __global__ void MatrixReduceLongWidth(const T *__restrict__ in, T *out, @@ -365,6 +367,7 @@ class ElementwiseAddGradKernel : public ElemwiseGradKernel { int max_blocks = std::max(max_physical_threads / (block_x * block_y), 1); int theory_block = (width + blocks.x - 1) / blocks.x; dim3 grids(std::min(theory_block, max_blocks)); +#if CUDA_VERSION >= 10000 if (std::is_same::value && width < 2048 && width % 2 == 0 && height % 64 == 0) { auto &dev_ctx = @@ -382,6 +385,7 @@ class ElementwiseAddGradKernel : public ElemwiseGradKernel { width, height); return; } +#endif if (width / height < 32) { MatrixColReduce<<>>(