Skip to content
体验新版
项目
组织
正在加载...
登录
切换导航
打开侧边栏
BaiXuePrincess
Paddle
提交
01c37c8e
P
Paddle
项目概览
BaiXuePrincess
/
Paddle
与 Fork 源项目一致
Fork自
PaddlePaddle / Paddle
通知
1
Star
1
Fork
0
代码
文件
提交
分支
Tags
贡献者
分支图
Diff
Issue
0
列表
看板
标记
里程碑
合并请求
0
Wiki
0
Wiki
分析
仓库
DevOps
项目成员
Pages
P
Paddle
项目概览
项目概览
详情
发布
仓库
仓库
文件
提交
分支
标签
贡献者
分支图
比较
Issue
0
Issue
0
列表
看板
标记
里程碑
合并请求
0
合并请求
0
Pages
分析
分析
仓库分析
DevOps
Wiki
0
Wiki
成员
成员
收起侧边栏
关闭侧边栏
动态
分支图
创建新Issue
提交
Issue看板
未验证
提交
01c37c8e
编写于
12月 22, 2020
作者:
W
wangchaochaohu
提交者:
GitHub
12月 22, 2020
浏览文件
操作
浏览文件
下载
电子邮件补丁
差异文件
refine the compiler error for half2 operation (#29816)
上级
82630408
变更
1
隐藏空白更改
内联
并排
Showing
1 changed file
with
2 addition
and
2 deletion
+2
-2
paddle/fluid/operators/elementwise/elementwise_add_op.h
paddle/fluid/operators/elementwise/elementwise_add_op.h
+2
-2
未找到文件。
paddle/fluid/operators/elementwise/elementwise_add_op.h
浏览文件 @
01c37c8e
...
@@ -183,6 +183,7 @@ template <int SIZE>
...
@@ -183,6 +183,7 @@ template <int SIZE>
__global__
void
VecFP16MatrixColReduce
(
const
__half2
*
__restrict__
in
,
__global__
void
VecFP16MatrixColReduce
(
const
__half2
*
__restrict__
in
,
__half2
*
__restrict__
out
,
size_t
width
,
__half2
*
__restrict__
out
,
size_t
width
,
size_t
height
)
{
size_t
height
)
{
#if CUDA_ARCH_FP16_SUPPORTED(__CUDA_ARCH__)
int
idx
=
threadIdx
.
x
+
blockIdx
.
x
*
blockDim
.
x
;
int
idx
=
threadIdx
.
x
+
blockIdx
.
x
*
blockDim
.
x
;
int
by
=
blockIdx
.
y
;
int
by
=
blockIdx
.
y
;
__half2
zero
=
__half2half2
(
static_cast
<
__half
>
(
0
));
__half2
zero
=
__half2half2
(
static_cast
<
__half
>
(
0
));
...
@@ -196,6 +197,7 @@ __global__ void VecFP16MatrixColReduce(const __half2 *__restrict__ in,
...
@@ -196,6 +197,7 @@ __global__ void VecFP16MatrixColReduce(const __half2 *__restrict__ in,
atomicAdd
(
&
(
out
[
idx
]),
sum
);
atomicAdd
(
&
(
out
[
idx
]),
sum
);
}
}
#endif
}
}
template
<
typename
T
>
template
<
typename
T
>
...
@@ -363,7 +365,6 @@ class ElementwiseAddGradKernel : public ElemwiseGradKernel<T> {
...
@@ -363,7 +365,6 @@ class ElementwiseAddGradKernel : public ElemwiseGradKernel<T> {
int
max_blocks
=
std
::
max
(
max_physical_threads
/
(
block_x
*
block_y
),
1
);
int
max_blocks
=
std
::
max
(
max_physical_threads
/
(
block_x
*
block_y
),
1
);
int
theory_block
=
(
width
+
blocks
.
x
-
1
)
/
blocks
.
x
;
int
theory_block
=
(
width
+
blocks
.
x
-
1
)
/
blocks
.
x
;
dim3
grids
(
std
::
min
(
theory_block
,
max_blocks
));
dim3
grids
(
std
::
min
(
theory_block
,
max_blocks
));
#if CUDA_ARCH_FP16_SUPPORTED(__CUDA_ARCH__)
if
(
std
::
is_same
<
T
,
paddle
::
platform
::
float16
>::
value
&&
width
<
2048
&&
if
(
std
::
is_same
<
T
,
paddle
::
platform
::
float16
>::
value
&&
width
<
2048
&&
width
%
2
==
0
&&
height
%
64
==
0
)
{
width
%
2
==
0
&&
height
%
64
==
0
)
{
auto
&
dev_ctx
=
auto
&
dev_ctx
=
...
@@ -381,7 +382,6 @@ class ElementwiseAddGradKernel : public ElemwiseGradKernel<T> {
...
@@ -381,7 +382,6 @@ class ElementwiseAddGradKernel : public ElemwiseGradKernel<T> {
width
,
height
);
width
,
height
);
return
;
return
;
}
}
#endif
if
(
width
/
height
<
32
)
{
if
(
width
/
height
<
32
)
{
MatrixColReduce
<
T
,
block_x
,
block_y
><<<
grids
,
blocks
,
0
,
stream
>>>
(
MatrixColReduce
<
T
,
block_x
,
block_y
><<<
grids
,
blocks
,
0
,
stream
>>>
(
...
...
编辑
预览
Markdown
is supported
0%
请重试
或
添加新附件
.
添加附件
取消
You are about to add
0
people
to the discussion. Proceed with caution.
先完成此消息的编辑!
取消
想要评论请
注册
或
登录