From f340f37f021274598c2561e2346987ec50463541 Mon Sep 17 00:00:00 2001 From: liaogang Date: Wed, 30 Nov 2016 20:00:51 +0800 Subject: [PATCH] Change atomicAdd to paddleAtomicAdd --- paddle/cuda/src/hl_cuda_cnn.cu | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/paddle/cuda/src/hl_cuda_cnn.cu b/paddle/cuda/src/hl_cuda_cnn.cu index ae387a8bc0e..7f2f6897b49 100644 --- a/paddle/cuda/src/hl_cuda_cnn.cu +++ b/paddle/cuda/src/hl_cuda_cnn.cu @@ -16,6 +16,7 @@ limitations under the License. */ #include #include "hl_base.h" #include "hl_cnn.h" +#include "hl_device_functions.cuh" __global__ void KeFeature2col(size_t n, size_t height, const real* data_im, size_t blockH, size_t blockW, size_t width, @@ -641,10 +642,10 @@ __global__ void KeBilinearInterpBw(real* in, real* inPos = &in[outIdH * inputW + channelId * inImgSize + inImgIdy * inImgW + inImgIdx]; const real* outPos = &out[outIdH * outputW + outIdW]; - atomicAdd(&inPos[0], h2lambda * w2lambda * outPos[0]); - atomicAdd(&inPos[wId], h2lambda * w1lambda * outPos[0]); - atomicAdd(&inPos[hId * inImgW], h1lambda * w2lambda * outPos[0]); - atomicAdd(&inPos[hId * inImgW + wId], h1lambda * w1lambda * outPos[0]); + paddle::paddleAtomicAdd(&inPos[0], h2lambda * w2lambda * outPos[0]); + paddle::paddleAtomicAdd(&inPos[wId], h2lambda * w1lambda * outPos[0]); + paddle::paddleAtomicAdd(&inPos[hId * inImgW], h1lambda * w2lambda * outPos[0]); + paddle::paddleAtomicAdd(&inPos[hId * inImgW + wId], h1lambda * w1lambda * outPos[0]); } } -- GitLab