Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

[metal] fix_elementwise #7467

Merged
merged 2 commits into from
Nov 3, 2021
Merged

[metal] fix_elementwise #7467

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
8f33074
fix_elementwise
xiaoxiaohehe001 Oct 26, 2021
87b1aad
fix_elementwise
xiaoxiaohehe001 Nov 1, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 4 additions & 0 deletions lite/backends/metal/metal_kernel/texture/Common.metal
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -79,6 +79,10 @@ struct ElementwiseAddParam {
int32_t xtrans[4]; // input_x transpose dim
int32_t ydim[4]; // input_y dim on gpu
int32_t ytrans[4]; // input_x transpose dim
int32_t ByNum; // only one number
int32_t ByHW; // only HW
int32_t ByW; // only W
int32_t arithmetic_type;
};

struct ElementwiseParam {
Expand Down
87 changes: 25 additions & 62 deletions lite/backends/metal/metal_kernel/texture/Elementwise.metal
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,7 +18,7 @@

using namespace metal;

kernel void elementwise_add(texture2d_array<ftype, access::read> inputX[[texture(0)]],
kernel void elementwise(texture2d_array<ftype, access::read> inputX[[texture(0)]],
texture2d_array<ftype, access::read> inputY[[texture(1)]],
texture2d_array<ftype, access::write> outTexture[[texture(2)]],
constant ElementwiseAddParam& pm[[buffer(0)]],
Expand All @@ -28,13 +28,25 @@ kernel void elementwise_add(texture2d_array<ftype, access::read> inputX[[texture
return;
ftype4 rx, ry;
rx = inputX.read(gid.xy, gid.z);
// elementwise add
// elementwise
if (pm.fast == 1) {
ry = inputY.read(gid.xy, gid.z);
}
// add at one number
else if (pm.ByNum == 1) {
ry = inputY.read(uint2(0, 0), gid.z);
}
// add at C channel
else if (pm.addByChannel == 1) {
ry = inputY.read(uint2(0, 0), gid.z);
}
// add at HW
else if (pm.ByHW == 1) {
ry = inputY.read(gid.xy, 0);
}
// add at W
else if (pm.ByW == 1) {
ry = inputY.read(uint2(gid.x, 0), 0);
} else {
// X coordinate GPU NHNC
int32_t x_xyzn[4] = {int32_t(gid.x), int32_t(gid.y), int32_t(gid.z), 0};
Expand Down Expand Up @@ -73,68 +85,19 @@ kernel void elementwise_add(texture2d_array<ftype, access::read> inputX[[texture
ry[n] = inputY.read(uint2(y_xyzn[0], y_xyzn[1]), y_xyzn[2])[y_xyzn[3]];
}
}
ftype4 r = rx + ry;
outTexture.write(r, gid.xy, gid.z);
}

kernel void elementwise_sub(texture2d_array<ftype, access::read> inputX[[texture(0)]],
texture2d_array<ftype, access::read> inputY[[texture(1)]],
texture2d_array<ftype, access::write> outTexture[[texture(2)]],
constant ElementwiseParam& pm[[buffer(0)]],
uint3 gid[[thread_position_in_grid]]) {
if (gid.x >= outTexture.get_width() || gid.y >= outTexture.get_height() ||
gid.z >= outTexture.get_array_size())
return;
ftype4 rx, ry;
rx = inputX.read(gid.xy, gid.z);
if (pm.byChannel == 1) {
ry = inputY.read(uint2(0, 0), gid.z);
} else {
ry = inputY.read(gid.xy, gid.z);
}
ftype4 r = rx - ry;
outTexture.write(r, gid.xy, gid.z);
}

kernel void elementwise_mul(texture2d_array<ftype, access::read> inputX[[texture(0)]],
texture2d_array<ftype, access::read> inputY[[texture(1)]],
texture2d_array<ftype, access::write> outTexture[[texture(2)]],
constant ElementwiseParam& pm[[buffer(0)]],
uint3 gid[[thread_position_in_grid]]) {
if (gid.x >= outTexture.get_width() || gid.y >= outTexture.get_height() ||
gid.z >= outTexture.get_array_size())
return;
ftype4 rx, ry;
rx = inputX.read(gid.xy, gid.z);
if (pm.byChannel == 1) {
ry = inputY.read(uint2(0, 0), gid.z);
} else {
ry = inputY.read(gid.xy, gid.z);
}
ftype4 r = rx * ry;
ftype4 r;
if (pm.arithmetic_type == 0)
r = rx + ry;
else if (pm.arithmetic_type == 1)
r = rx - ry;
else if (pm.arithmetic_type == 2)
r = rx * ry;
else if (pm.arithmetic_type == 3)
r = rx / ry;
outTexture.write(r, gid.xy, gid.z);
}

kernel void elementwise_div(texture2d_array<ftype, access::read> inputX[[texture(0)]],
texture2d_array<ftype, access::read> inputY[[texture(1)]],
texture2d_array<ftype, access::write> outTexture[[texture(2)]],
constant ElementwiseParam& pm[[buffer(0)]],
uint3 gid[[thread_position_in_grid]]) {
if (gid.x >= outTexture.get_width() || gid.y >= outTexture.get_height() ||
gid.z >= outTexture.get_array_size())
return;
ftype4 rx, ry;
rx = inputX.read(gid.xy, gid.z);
if (pm.byChannel == 1) {
ry = inputY.read(uint2(0, 0), gid.z);
} else {
ry = inputY.read(gid.xy, gid.z);
}
ftype4 r = rx / ry;
outTexture.write(r, gid.xy, gid.z);
}

kernel void elementwise_add_relu(texture2d_array<ftype, access::read> inputX[[texture(0)]],
kernel void elementwise_relu(texture2d_array<ftype, access::read> inputX[[texture(0)]],
texture2d_array<ftype, access::read> inputY[[texture(1)]],
texture2d_array<ftype, access::write> outTexture[[texture(2)]],
constant ElementwiseAddParam& pm[[buffer(0)]],
Expand Down Expand Up @@ -173,4 +136,4 @@ kernel void elementwise_add_relu(texture2d_array<ftype, access::read> inputX[[te
output = fmax(output, 0.0);

outTexture.write(output, gid.xy, gid.z);
}
}
4 changes: 1 addition & 3 deletions lite/kernels/metal/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -27,9 +27,7 @@ add_kernel(conv2d_metal_image METAL basic SRCS image_op/conv2d_image_compute.mm)
add_kernel(conv2d_transpose_metal_image METAL basic SRCS image_op/conv2d_transpose_image_compute.mm)
add_kernel(depthwise_conv2d_metal_image METAL basic SRCS image_op/depthwise_conv2d_image_compute.mm)
add_kernel(dropout_metal_image METAL basic SRCS image_op/dropout_image_compute.mm)
add_kernel(elementwise_add_metal_image METAL basic SRCS image_op/elementwise_add_image_compute.mm)
add_kernel(elementwise_mul_metal_image METAL basic SRCS image_op/elementwise_mul_image_compute.mm)
add_kernel(elementwise_sub_metal_image METAL basic SRCS image_op/elementwise_sub_image_compute.mm)
add_kernel(elementwise_metal_image METAL basic SRCS image_op/elementwise_image_compute.mm)
add_kernel(exp_metal_image METAL basic SRCS image_op/exp_image_compute.mm)
add_kernel(fc_metal_image METAL basic SRCS image_op/fc_image_compute.mm)
add_kernel(feed_metal_image METAL basic SRCS image_op/feed_image_compute.mm)
Expand Down
64 changes: 0 additions & 64 deletions lite/kernels/metal/image_op/elementwise_div_image_compute.h
View file
Open in desktop

This file was deleted.

161 changes: 0 additions & 161 deletions lite/kernels/metal/image_op/elementwise_div_image_compute.mm
View file
Open in desktop

This file was deleted.

Loading