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

[Metal] Fix arg max #8384

Merged
merged 3 commits into from
Feb 6, 2022
Merged

[Metal] Fix arg max #8384

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
ed7dc35
fix
xiaoxiaohehe001 Jan 29, 2022
fa87573
fix_argmax
xiaoxiaohehe001 Jan 29, 2022
3c36414
fix_argmax
xiaoxiaohehe001 Jan 29, 2022
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
70 changes: 68 additions & 2 deletions lite/backends/metal/metal_kernel/texture/MaxKernel.metal
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24,9 +24,9 @@ struct ArgParam {
inline int max_index(texture2d_array<ftype, access::read> inTexture[[texture(0)]], uint2 gid) {
int index = 0;
#if LITE_WITH_METAL_FULL
float omax = FLT_MIN;
float omax = -FLT_MAX;
#else
half omax = HALF_MIN;
half omax = -HALF_MAX;
#endif
uint iAL = inTexture.get_array_size();
for (uint i = 0; i < iAL; i++) {
Expand Down Expand Up @@ -77,3 +77,69 @@ kernel void arg_max_c(texture2d_array<ftype, access::read> inTexture[[texture(0)
outTexture.write(ftype4(index_r, index_g, index_b, index_a), gid.xy, gid.z);
}
}

kernel void arg_max_h(texture2d_array<ftype, access::read> inTexture[[texture(0)]],
texture2d_array<ftype, access::write> outTexture[[texture(1)]],
constant ArgParam& param[[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;

#if LITE_WITH_METAL_FULL
float omax = -FLT_MAX;
#else
half omax = -HALF_MAX;
#endif

if (param.orank == 4) {
uint iAL = inTexture.get_height();
ftype4 cur_data = omax;
ftype4 max_data = omax;
uint4 cur_idx = 0;
uint4 max_idx = 0;
bool4 flag_v = false;

for (uint i = 0; i < iAL; i++) {
cur_data = inTexture.read(uint2(gid.x, i), gid.z);
cur_idx = uint4(i);
flag_v = bool4(max_data >= cur_data);
max_data = select(cur_data, max_data, flag_v);
max_idx = select(cur_idx, max_idx, flag_v);
}
outTexture.write(ftype4(max_idx), gid.xy, gid.z);
}
}

kernel void arg_max_w(texture2d_array<ftype, access::read> inTexture[[texture(0)]],
texture2d_array<ftype, access::write> outTexture[[texture(1)]],
constant ArgParam& param[[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;

#if LITE_WITH_METAL_FULL
float omax = -FLT_MAX;
#else
half omax = -HALF_MAX;
#endif

if (param.orank == 4) {
uint iAL = inTexture.get_width();
ftype4 cur_data = omax;
ftype4 max_data = omax;
uint4 cur_idx = 0;
uint4 max_idx = 0;
bool4 flag_v = false;

for (uint i = 0; i < iAL; i++) {
cur_data = inTexture.read(uint2(i, gid.y), gid.z);
cur_idx = uint4(i);
flag_v = bool4(max_data >= cur_data);
max_data = select(cur_data, max_data, flag_v);
max_idx = select(cur_idx, max_idx, flag_v);
}
outTexture.write(ftype4(max_idx), gid.xy, gid.z);
}
}
16 changes: 10 additions & 6 deletions lite/kernels/metal/image_op/argmax_image_compute.mm
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -44,10 +44,9 @@
bool should_use_mps = false;
if (@available(iOS 12.0, *)) {
if (metal_context_->use_mps()) {
should_use_mps = true;
if (param.Axis == 1) should_use_mps = true;
}
}

use_mps_ = should_use_mps;
if (use_mps_) {
setup_with_mps();
Expand Down Expand Up @@ -85,17 +84,22 @@
const auto& param = this->Param<param_t>();
auto irank = input_buffer_->tensor_dim_.size();
auto orank = output_buffer_->tensor_dim_.size();

// axis
if (param.Axis == 1 && irank == 4) {
if (irank == 4 && param.Axis == 1) {
function_name_ = "arg_max_c";
} else if (irank == 4 && param.Axis == 2) {
function_name_ = "arg_max_h";
} else if (irank == 4 && param.Axis == 3) {
function_name_ = "arg_max_w";
} else {
LOG(FATAL) << "arg_max: max only support by channel";
LOG(FATAL) << "This input is not supported by arg_max";
}

ArgMetalParam metal_params{(int)orank};
params_buffer_ =
std::make_shared<MetalBuffer>(metal_context_, sizeof(metal_params), &metal_params);

function_name_ = "arg_max_c";
// pipline
auto backend = (__bridge MetalContextImp*)metal_context_->backend();
pipline_ = [backend pipline:function_name_];
Expand Down Expand Up @@ -147,7 +151,7 @@
mps_op_ = (__bridge_retained void*)[[MPSNNReduceFeatureChannelsArgumentMax alloc]
initWithDevice:backend.device];
// MPS input and output
auto input_c = static_cast<int>(input_buffer_->tensor_dim_[1]);
auto input_c = fmax(4, static_cast<int>(input_buffer_->tensor_dim_[1]));
mps_input_image_ =
(__bridge_retained void*)[[MPSImage alloc] initWithTexture:input_buffer_->image()
featureChannels:input_c];
Expand Down