Skip to content
体验新版
项目
组织
正在加载...
登录
切换导航
打开侧边栏
正统之独孤求败
mindspore
提交
e1b78c1f
M
mindspore
项目概览
正统之独孤求败
/
mindspore
与 Fork 源项目一致
Fork自
MindSpore / mindspore
通知
1
Star
0
Fork
0
代码
文件
提交
分支
Tags
贡献者
分支图
Diff
Issue
0
列表
看板
标记
里程碑
合并请求
0
Wiki
0
Wiki
分析
仓库
DevOps
项目成员
Pages
M
mindspore
项目概览
项目概览
详情
发布
仓库
仓库
文件
提交
分支
标签
贡献者
分支图
比较
Issue
0
Issue
0
列表
看板
标记
里程碑
合并请求
0
合并请求
0
Pages
分析
分析
仓库分析
DevOps
Wiki
0
Wiki
成员
成员
收起侧边栏
关闭侧边栏
动态
分支图
创建新Issue
提交
Issue看板
体验新版 GitCode,发现更多精彩内容 >>
提交
e1b78c1f
编写于
8月 12, 2020
作者:
C
chenzupeng
浏览文件
操作
浏览文件
下载
电子邮件补丁
差异文件
add opencl reshape op
上级
874972ca
变更
7
显示空白变更内容
内联
并排
Showing
7 changed file
with
190 addition
and
1 deletion
+190
-1
mindspore/lite/src/runtime/kernel/opencl/CMakeLists.txt
mindspore/lite/src/runtime/kernel/opencl/CMakeLists.txt
+1
-0
mindspore/lite/src/runtime/kernel/opencl/cl/fp16/reshape.cl
mindspore/lite/src/runtime/kernel/opencl/cl/fp16/reshape.cl
+14
-0
mindspore/lite/src/runtime/kernel/opencl/cl/fp32/reshape.cl
mindspore/lite/src/runtime/kernel/opencl/cl/fp32/reshape.cl
+14
-0
mindspore/lite/src/runtime/kernel/opencl/kernel/convolution.cc
...pore/lite/src/runtime/kernel/opencl/kernel/convolution.cc
+2
-1
mindspore/lite/src/runtime/kernel/opencl/kernel/reshape.cc
mindspore/lite/src/runtime/kernel/opencl/kernel/reshape.cc
+114
-0
mindspore/lite/src/runtime/kernel/opencl/kernel/reshape.h
mindspore/lite/src/runtime/kernel/opencl/kernel/reshape.h
+44
-0
mindspore/lite/test/CMakeLists.txt
mindspore/lite/test/CMakeLists.txt
+1
-0
未找到文件。
mindspore/lite/src/runtime/kernel/opencl/CMakeLists.txt
浏览文件 @
e1b78c1f
...
...
@@ -10,4 +10,5 @@ set(OPENCL_KERNEL_SRC
${
CMAKE_CURRENT_SOURCE_DIR
}
/kernel/concat.cc
${
CMAKE_CURRENT_SOURCE_DIR
}
/kernel/conv2d_transpose.cc
${
CMAKE_CURRENT_SOURCE_DIR
}
/kernel/transpose.cc
${
CMAKE_CURRENT_SOURCE_DIR
}
/kernel/reshape.cc
)
mindspore/lite/src/runtime/kernel/opencl/cl/fp16/reshape.cl
0 → 100644
浏览文件 @
e1b78c1f
#
define
FLT
half
#
define
FLT4
half4
#
define
READ_IMAGE
read_imageh
#
define
WRITE_IMAGE
write_imageh
__constant
sampler_t
smp_zero
=
CLK_NORMALIZED_COORDS_FALSE
| CLK_ADDRESS_CLAMP |
CLK_FILTER_NEAREST
;
__kernel
void
reshape
(
__read_only
image2d_t
src_data,
__write_only
image2d_t
dst_data,
int4
size
)
{
int
X
=
get_global_id
(
0
)
;
int
Y
=
get_global_id
(
1
)
;
int
Z
=
get_global_id
(
2
)
;
if
(
X
>=
size.x
|
| Y >= size.y |
|
Z
>=
size.z
)
{
return
;
}
WRITE_IMAGE
(
dst_data,
(
int2
)(
Y
*
size.z
+
Z,
X
)
,
READ_IMAGE
(
src_data,
smp_zero,
(
int2
)(
Y
*
size.z
+
Z,
X
)))
;
}
mindspore/lite/src/runtime/kernel/opencl/cl/fp32/reshape.cl
0 → 100644
浏览文件 @
e1b78c1f
#
define
FLT
float
#
define
FLT4
float4
#
define
READ_IMAGE
read_imagef
#
define
WRITE_IMAGE
write_imagef
__constant
sampler_t
smp_zero
=
CLK_NORMALIZED_COORDS_FALSE
| CLK_ADDRESS_CLAMP |
CLK_FILTER_NEAREST
;
__kernel
void
reshape
(
__read_only
image2d_t
src_data,
__write_only
image2d_t
dst_data,
int4
size
)
{
int
X
=
get_global_id
(
0
)
;
int
Y
=
get_global_id
(
1
)
;
int
Z
=
get_global_id
(
2
)
;
if
(
X
>=
size.x
|
| Y >= size.y |
|
Z
>=
size.z
)
{
return
;
}
WRITE_IMAGE
(
dst_data,
(
int2
)(
Y
*
size.z
+
Z,
X
)
,
READ_IMAGE
(
src_data,
smp_zero,
(
int2
)(
Y
*
size.z
+
Z,
X
)))
;
}
mindspore/lite/src/runtime/kernel/opencl/kernel/convolution.cc
浏览文件 @
e1b78c1f
...
...
@@ -30,7 +30,6 @@ namespace mindspore::kernel {
int
ConvolutionOpenCLKernel
::
Init
()
{
static
int
count
=
0
;
std
::
cout
<<
"ConvolutionOpenCLKernel::Init()
\n
"
;
std
::
set
<
std
::
string
>
build_options
;
std
::
string
source
=
CodeGen
();
std
::
string
program_name
=
"convolution"
+
std
::
to_string
(
count
);
...
...
@@ -41,6 +40,8 @@ int ConvolutionOpenCLKernel::Init() {
ocl_runtime
->
LoadSource
(
program_name
,
source
);
ocl_runtime
->
BuildKernel
(
kernel_
,
program_name
,
kernel_name
,
build_options
);
this
->
InitBuffer
();
out_tensors_
[
0
]
->
SetFormat
(
schema
::
Format_NHWC4
);
MS_LOG
(
DEBUG
)
<<
kernel_name
<<
" Init Done!"
;
return
RET_OK
;
}
...
...
mindspore/lite/src/runtime/kernel/opencl/kernel/reshape.cc
0 → 100644
浏览文件 @
e1b78c1f
/**
* Copyright 2019 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#include <set>
#include <string>
#include "include/errorcode.h"
#include "src/kernel_registry.h"
#include "src/runtime/opencl/opencl_runtime.h"
#include "src/runtime/kernel/opencl/kernel/reshape.h"
#include "src/runtime/kernel/opencl/cl/fp16/reshape.cl.inc"
#include "src/runtime/kernel/opencl/cl/fp32/reshape.cl.inc"
using
mindspore
::
kernel
::
KERNEL_ARCH
::
kGPU
;
using
mindspore
::
lite
::
KernelRegistrar
;
using
mindspore
::
lite
::
RET_ERROR
;
using
mindspore
::
lite
::
RET_OK
;
using
mindspore
::
schema
::
PrimitiveType_Reshape
;
namespace
mindspore
::
kernel
{
int
ReshapeOpenCLKernel
::
Init
()
{
std
::
string
kernel_name
=
"reshape"
;
auto
ocl_runtime
=
lite
::
opencl
::
OpenCLRuntime
::
GetInstance
();
#ifdef PROGRAM_WITH_IL
ocl_runtime
->
CreateKernelFromIL
(
kernel_
(),
kernel_name
);
#else
std
::
set
<
std
::
string
>
build_options
;
#ifdef ENABLE_FP16
std
::
string
source
=
reshape_source_fp16
;
#else
std
::
string
source
=
reshape_source_fp32
;
#endif
std
::
string
program_name
=
"reshape"
;
ocl_runtime
->
LoadSource
(
program_name
,
source
);
ocl_runtime
->
BuildKernel
(
kernel_
,
program_name
,
kernel_name
,
build_options
);
#endif
out_tensors_
[
0
]
->
SetFormat
(
schema
::
Format_NHWC
);
MS_LOG
(
DEBUG
)
<<
kernel_name
<<
" Init Done!"
;
return
RET_OK
;
}
int
ReshapeOpenCLKernel
::
ReSize
()
{
return
0
;
}
int
ReshapeOpenCLKernel
::
GetImageSize
(
size_t
idx
,
std
::
vector
<
size_t
>
*
img_size
)
{
size_t
im_dst_x
,
im_dst_y
;
std
::
vector
<
int
>
shapex
=
in_tensors_
[
0
]
->
shape
();
int
h
=
shapex
[
1
];
int
w
=
shapex
[
2
];
int
c
=
shapex
[
3
];
im_dst_x
=
UP_DIV
(
w
*
c
,
C4NUM
);
im_dst_y
=
h
;
#ifdef ENABLE_FP16
size_t
img_dtype
=
CL_HALF_FLOAT
;
#else
size_t
img_dtype
=
CL_FLOAT
;
#endif
img_size
->
clear
();
std
::
vector
<
size_t
>
vec
{
im_dst_x
,
im_dst_y
,
img_dtype
};
*
img_size
=
vec
;
return
RET_OK
;
}
int
ReshapeOpenCLKernel
::
Run
()
{
MS_LOG
(
DEBUG
)
<<
this
->
name
()
<<
" Running!"
;
std
::
vector
<
int
>
shapex
=
in_tensors_
[
0
]
->
shape
();
int
h
=
shapex
[
1
];
int
w
=
shapex
[
2
];
int
c
=
shapex
[
3
];
int
c4
=
UP_DIV
(
c
,
C4NUM
);
auto
ocl_runtime
=
lite
::
opencl
::
OpenCLRuntime
::
GetInstance
();
// local size should less than MAX_GROUP_SIZE
std
::
vector
<
size_t
>
local
=
{};
std
::
vector
<
size_t
>
global
=
{(
size_t
)
h
,
(
size_t
)
w
,
(
size_t
)
c4
};
cl_int4
size
=
{
h
,
w
,
c4
,
1
};
ocl_runtime
->
SetKernelArg
(
kernel_
,
0
,
in_tensors_
[
0
]
->
Data
());
ocl_runtime
->
SetKernelArg
(
kernel_
,
1
,
out_tensors_
[
0
]
->
Data
());
ocl_runtime
->
SetKernelArg
(
kernel_
,
2
,
size
);
ocl_runtime
->
RunKernel
(
kernel_
,
global
,
local
,
nullptr
);
return
RET_OK
;
}
kernel
::
LiteKernel
*
OpenCLReshapeKernelCreator
(
const
std
::
vector
<
lite
::
tensor
::
Tensor
*>
&
inputs
,
const
std
::
vector
<
lite
::
tensor
::
Tensor
*>
&
outputs
,
OpParameter
*
opParameter
,
const
lite
::
Context
*
ctx
,
const
kernel
::
KernelKey
&
desc
,
const
lite
::
Primitive
*
primitive
)
{
auto
*
kernel
=
new
(
std
::
nothrow
)
ReshapeOpenCLKernel
(
reinterpret_cast
<
OpParameter
*>
(
opParameter
),
inputs
,
outputs
);
if
(
kernel
==
nullptr
)
{
MS_LOG
(
ERROR
)
<<
"kernel "
<<
opParameter
->
name_
<<
" create failed."
;
return
nullptr
;
}
auto
ret
=
kernel
->
Init
();
if
(
ret
!=
RET_OK
)
{
delete
kernel
;
return
nullptr
;
}
return
kernel
;
}
REG_KERNEL
(
kGPU
,
kNumberTypeFloat32
,
PrimitiveType_Reshape
,
OpenCLReshapeKernelCreator
)
}
// namespace mindspore::kernel
mindspore/lite/src/runtime/kernel/opencl/kernel/reshape.h
0 → 100644
浏览文件 @
e1b78c1f
/**
* Copyright 2020 Huawei Technologies Co., Ltd
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef MINDSPORE_LITE_SRC_BACKEND_OPENCL_RESHAPE_H_
#define MINDSPORE_LITE_SRC_BACKEND_OPENCL_RESHAPE_H_
#include <vector>
#include "src/lite_kernel.h"
#include "src/runtime/opencl/opencl_runtime.h"
#include "src/runtime/kernel/opencl/opencl_kernel.h"
namespace
mindspore
::
kernel
{
class
ReshapeOpenCLKernel
:
public
OpenCLKernel
{
public:
explicit
ReshapeOpenCLKernel
(
OpParameter
*
parameter
,
const
std
::
vector
<
lite
::
tensor
::
Tensor
*>
&
inputs
,
const
std
::
vector
<
lite
::
tensor
::
Tensor
*>
&
outputs
)
:
OpenCLKernel
(
parameter
,
inputs
,
outputs
)
{}
~
ReshapeOpenCLKernel
()
override
{};
int
Init
()
override
;
int
ReSize
()
override
;
int
Run
()
override
;
int
GetImageSize
(
size_t
idx
,
std
::
vector
<
size_t
>
*
img_size
)
override
;
private:
cl
::
Kernel
kernel_
;
};
}
// namespace mindspore::kernel
#endif // MINDSPORE_LITE_SRC_BACKEND_OPENCL_RESHAPE_H_
mindspore/lite/test/CMakeLists.txt
浏览文件 @
e1b78c1f
...
...
@@ -145,6 +145,7 @@ if (SUPPORT_GPU)
# ${LITE_DIR}/src/runtime/kernel/opencl/kernel/leaky_relu.cc
${
LITE_DIR
}
/src/runtime/kernel/opencl/kernel/conv2d_transpose.cc
${
LITE_DIR
}
/src/runtime/kernel/opencl/kernel/transpose.cc
${
LITE_DIR
}
/src/runtime/kernel/opencl/kernel/reshape.cc
)
endif
()
### minddata lite
...
...
编辑
预览
Markdown
is supported
0%
请重试
或
添加新附件
.
添加附件
取消
You are about to add
0
people
to the discussion. Proceed with caution.
先完成此消息的编辑!
取消
想要评论请
注册
或
登录