Skip to content
体验新版
项目
组织
正在加载...
登录
切换导航
打开侧边栏
Crayon鑫
Paddle
提交
4558c0ec
P
Paddle
项目概览
Crayon鑫
/
Paddle
与 Fork 源项目一致
Fork自
PaddlePaddle / Paddle
通知
1
Star
1
Fork
0
代码
文件
提交
分支
Tags
贡献者
分支图
Diff
Issue
1
列表
看板
标记
里程碑
合并请求
0
Wiki
0
Wiki
分析
仓库
DevOps
项目成员
Pages
P
Paddle
项目概览
项目概览
详情
发布
仓库
仓库
文件
提交
分支
标签
贡献者
分支图
比较
Issue
1
Issue
1
列表
看板
标记
里程碑
合并请求
0
合并请求
0
Pages
分析
分析
仓库分析
DevOps
Wiki
0
Wiki
成员
成员
收起侧边栏
关闭侧边栏
动态
分支图
创建新Issue
提交
Issue看板
未验证
提交
4558c0ec
编写于
5月 04, 2018
作者:
C
chengduo
提交者:
GitHub
5月 04, 2018
浏览文件
操作
浏览文件
下载
差异文件
Merge pull request #10414 from chengduoZH/wrap_shfl_x_sync
Wrap shfl_x_sync
上级
2d98a418
d36af62c
变更
3
显示空白变更内容
内联
并排
Showing
3 changed file
with
21 addition
and
21 deletion
+21
-21
paddle/fluid/operators/row_conv_op.cu
paddle/fluid/operators/row_conv_op.cu
+2
-2
paddle/fluid/operators/top_k_op.cu
paddle/fluid/operators/top_k_op.cu
+2
-1
paddle/fluid/platform/cuda_device_function.h
paddle/fluid/platform/cuda_device_function.h
+17
-18
未找到文件。
paddle/fluid/operators/row_conv_op.cu
浏览文件 @
4558c0ec
...
...
@@ -224,7 +224,7 @@ __global__ void RowConvGradFilterImproved(const T *in, const T *dout,
for
(
int
offset
=
16
;
offset
>
0
;
offset
=
offset
/
2
)
{
// blockDim.x is 32.
val
+=
platform
::
__shfl_down_s
ync
(
mask
,
val
,
offset
);
val
+=
platform
::
CudaShuffleDownS
ync
(
mask
,
val
,
offset
);
}
__syncthreads
();
...
...
@@ -284,7 +284,7 @@ __global__ void RowConvGradFilter(const T *in, const T *dout, int num_sequence,
for
(
int
offset
=
16
;
offset
>
0
;
offset
=
offset
/
2
)
{
// blockDim.x is 32.
val
+=
platform
::
__shfl_down_s
ync
(
mask
,
val
,
offset
);
val
+=
platform
::
CudaShuffleDownS
ync
(
mask
,
val
,
offset
);
}
__syncthreads
();
...
...
paddle/fluid/operators/top_k_op.cu
浏览文件 @
4558c0ec
...
...
@@ -241,7 +241,8 @@ __device__ __forceinline__ void BlockReduce(Pair<T>* sh_topk, int* maxid,
CREATE_SHFL_MASK
(
mask
,
true
);
if
(
maxid
[
0
]
/
32
==
warp
)
{
if
(
platform
::
__shfl_sync
(
mask
,
*
beam
,
(
maxid
[
0
])
%
32
,
32
)
==
MaxLength
)
if
(
platform
::
CudaShuffleSync
(
mask
,
*
beam
,
(
maxid
[
0
])
%
32
,
32
)
==
MaxLength
)
break
;
}
}
...
...
paddle/fluid/platform/cuda_device_function.h
浏览文件 @
4558c0ec
...
...
@@ -18,34 +18,33 @@ limitations under the License. */
namespace
paddle
{
namespace
platform
{
// __shfl_down and __shfl have been deprecated as of CUDA 9.0.
#if CUDA_VERSION < 9000
template
<
typename
T
>
__forceinline__
__device__
T
__shfl_down_sync
(
unsigned
,
T
val
,
int
delta
)
{
return
__shfl_down
(
val
,
delta
);
}
template
<
typename
T
>
__forceinline__
__device__
T
__shfl_sync
(
unsigned
,
T
val
,
int
src_line
,
int
width
)
{
return
__shfl
(
val
,
src_line
,
width
);
}
#define CREATE_SHFL_MASK(mask, predicate) mask = 0u;
#else
#define FULL_WARP_MASK 0xFFFFFFFF
#define CREATE_SHFL_MASK(mask, predicate) \
mask = __ballot_sync(FULL_WARP_MASK, (predicate))
#endif
template
<
typename
T
>
__forceinline__
__device__
T
__shfl_down_sync
(
unsigned
mask
,
T
val
,
int
delta
)
{
return
__shfl_down_sync
(
mask
,
val
,
delta
);
__forceinline__
__device__
T
CudaShuffleDownSync
(
unsigned
mask
,
T
val
,
int
delta
,
int
width
=
32
)
{
#if CUDA_VERSION < 9000
return
__shfl_down
(
val
,
delta
,
width
);
#else
return
__shfl_down_sync
(
mask
,
val
,
delta
,
width
);
#endif
}
template
<
typename
T
>
__forceinline__
__device__
T
__shfl_sync
(
unsigned
mask
,
T
val
,
int
src_line
,
int
width
)
{
__forceinline__
__device__
T
CudaShuffleSync
(
unsigned
mask
,
T
val
,
int
src_line
,
int
width
=
32
)
{
#if CUDA_VERSION < 9000
return
__shfl
(
val
,
src_line
,
width
);
#else
return
__shfl_sync
(
mask
,
val
,
src_line
,
width
);
}
#endif
}
template
<
typename
T
>
__device__
T
reduceSum
(
T
val
,
int
tid
,
int
len
)
{
...
...
@@ -61,7 +60,7 @@ __device__ T reduceSum(T val, int tid, int len) {
CREATE_SHFL_MASK
(
mask
,
tid
<
len
);
for
(
int
offset
=
warpSize
/
2
;
offset
>
0
;
offset
/=
2
)
val
+=
platform
::
__shfl_down_s
ync
(
mask
,
val
,
offset
);
val
+=
platform
::
CudaShuffleDownS
ync
(
mask
,
val
,
offset
);
if
(
tid
<
warpSize
)
shm
[
tid
]
=
0
;
...
...
@@ -75,7 +74,7 @@ __device__ T reduceSum(T val, int tid, int len) {
if
(
tid
<
warpSize
)
{
val
=
shm
[
tid
];
for
(
int
offset
=
warpSize
/
2
;
offset
>
0
;
offset
/=
2
)
val
+=
platform
::
__shfl_down_s
ync
(
mask
,
val
,
offset
);
val
+=
platform
::
CudaShuffleDownS
ync
(
mask
,
val
,
offset
);
}
return
val
;
}
...
...
编辑
预览
Markdown
is supported
0%
请重试
或
添加新附件
.
添加附件
取消
You are about to add
0
people
to the discussion. Proceed with caution.
先完成此消息的编辑!
取消
想要评论请
注册
或
登录