Skip to content
体验新版
项目
组织
正在加载...
登录
切换导航
打开侧边栏
wzpei9
Paddle
提交
90d73c79
P
Paddle
项目概览
wzpei9
/
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看板
体验新版 GitCode,发现更多精彩内容 >>
提交
90d73c79
编写于
5月 02, 2018
作者:
C
chengduoZH
浏览文件
操作
浏览文件
下载
电子邮件补丁
差异文件
fix shfl_sync for CUDA8.0
上级
46c90ea6
变更
4
隐藏空白更改
内联
并排
Showing
4 changed file
with
28 addition
and
10 deletion
+28
-10
paddle/cuda/include/hl_base.h
paddle/cuda/include/hl_base.h
+15
-0
paddle/cuda/src/hl_cuda_lstm.cu
paddle/cuda/src/hl_cuda_lstm.cu
+9
-5
paddle/cuda/src/hl_top_k.cu
paddle/cuda/src/hl_top_k.cu
+4
-1
paddle/fluid/platform/cuda_primitives.h
paddle/fluid/platform/cuda_primitives.h
+0
-4
未找到文件。
paddle/cuda/include/hl_base.h
浏览文件 @
90d73c79
...
...
@@ -228,6 +228,21 @@ extern __thread cudaStream_t default_stream;
<< "CUDA error: " << hl_get_device_error_string((size_t)err); \
}
// __shfl has been deprecated as of CUDA 9.0.
#if CUDA_VERSION < 9000
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
#endif
/* __NVCC__ */
#endif
/* HL_BASE_H_ */
paddle/cuda/src/hl_cuda_lstm.cu
浏览文件 @
90d73c79
...
...
@@ -341,12 +341,15 @@ void hl_lstm_parallel_forward(real *gateValue,
}
__device__
__forceinline__
void
transpose_32x32
(
real
a
[],
const
int
idx
)
{
int
addr
=
idx
%
32
;
const
int
warp_size
=
32
;
int
addr
=
idx
%
warp_size
;
unsigned
mask
=
0u
;
CREATE_SHFL_MASK
(
mask
,
addr
<
warp_size
);
#pragma unroll
for
(
int
k
=
1
;
k
<
32
;
k
++
)
{
// rSrc[k] = __shfl_sync(rSrc[k], (threadIdx.x + k) % 32, 32);
addr
=
__shfl_sync
(
addr
,
(
idx
+
1
)
%
32
,
32
);
a
[
k
]
=
__shfl_sync
(
a
[
k
],
addr
,
32
);
addr
=
__shfl_sync
(
mask
,
addr
,
(
idx
+
1
)
%
32
,
32
);
a
[
k
]
=
__shfl_sync
(
mask
,
a
[
k
],
addr
,
32
);
}
#pragma unroll
...
...
@@ -360,10 +363,11 @@ __device__ __forceinline__ void transpose_32x32(real a[], const int idx) {
}
addr
=
(
32
-
idx
)
%
32
;
CREATE_SHFL_MASK
(
mask
,
idx
%
32
<
warp_size
);
#pragma unroll
for
(
int
k
=
0
;
k
<
32
;
k
++
)
{
a
[
k
]
=
__shfl_sync
(
a
[
k
],
addr
,
32
);
addr
=
__shfl_sync
(
addr
,
(
idx
+
31
)
%
32
,
32
);
a
[
k
]
=
__shfl_sync
(
mask
,
a
[
k
],
addr
,
32
);
addr
=
__shfl_sync
(
mask
,
addr
,
(
idx
+
31
)
%
32
,
32
);
}
}
...
...
paddle/cuda/src/hl_top_k.cu
浏览文件 @
90d73c79
...
...
@@ -244,13 +244,16 @@ __device__ __forceinline__ void blockReduce(Pair* shTopK,
if
(
--
beamSize
==
0
)
break
;
__syncthreads
();
unsigned
mask
=
0u
;
// CREATE_SHFL_MASK(mask, tid < len);
if
(
tid
==
maxId
[
0
])
{
if
(
beam
<
maxLength
)
{
shTopK
[
tid
]
=
topK
[
beam
];
}
}
if
(
maxId
[
0
]
/
32
==
warp
)
{
if
(
__shfl_sync
(
beam
,
(
maxId
[
0
])
%
32
,
32
)
==
maxLength
)
break
;
if
(
__shfl_sync
(
mask
,
beam
,
(
maxId
[
0
])
%
32
,
32
)
==
maxLength
)
break
;
}
}
}
...
...
paddle/fluid/platform/cuda_primitives.h
浏览文件 @
90d73c79
...
...
@@ -74,10 +74,6 @@ __forceinline__ __device__ T __shfl_down_sync(unsigned, T val, int delta) {
}
#define CREATE_SHFL_MASK(mask, predicate) mask = 0u;
#else
template
<
typename
T
>
__forceinline__
__device__
T
__shfl_down_sync
(
unsigned
mask
,
T
val
,
int
delta
)
{
return
__shfl_down
(
mask
,
val
,
delta
);
}
#define FULL_WARP_MASK 0xFFFFFFFF
#define CREATE_SHFL_MASK(mask, predicate) \
mask = __ballot_sync(FULL_WARP_MASK, (predicate))
...
...
编辑
预览
Markdown
is supported
0%
请重试
或
添加新附件
.
添加附件
取消
You are about to add
0
people
to the discussion. Proceed with caution.
先完成此消息的编辑!
取消
想要评论请
注册
或
登录