[MV-1195] Implement Pdist Backward #62
Conversation
| { | ||
| pdist_backward_contiguous<INPUT_TYPE>( | ||
| input, output, grad, input_grad, p, n2, n2_squared_minus_1, N, NO, M); | ||
| } |
There was a problem hiding this comment.
add new line at the end
There was a problem hiding this comment.
[Partial Review] I found no kernel rightly implemented for non-contiguous case, please add and re-benchmark the performance. Additionally, why you dont implement Forward case, I found pdist kernel (forward case) in CL code and before all of the modification you will make, please change base branch to develop-moreh for the most update code. I will continue reviewing after you add non-contiguous and forward case
| for(size_t i = 0; i < input_numel; i++) | ||
| { | ||
| dinputHost[i] = 0; | ||
| } | ||
|
|
||
| for(int i = 0; i < N; ++i) | ||
| { | ||
| for(int j = i + 1; j < N; ++j) |
There was a problem hiding this comment.
pard_for for better perf
| double grad_k = static_cast<double>(doutput[k]); | ||
| double output_k = static_cast<double>(output[k]); | ||
|
|
||
| for(int m = 0; m < M; ++m) |
There was a problem hiding this comment.
this may lead to CI/CD failed because M is size_t while m is int (comparation between different type)
| double input_first = static_cast<double>(input[i * M + m]); | ||
| double input_second = static_cast<double>(input[j * M + m]); | ||
| double diff = input_first - input_second; | ||
|
|
||
| Tcheck res = | ||
| static_cast<Tcheck>(miopen::solver::pdist::backward(diff, grad_k, output_k, p)); | ||
|
|
||
| dinputHost[i * M + m] += res; | ||
| dinputHost[j * M + m] -= res; |
There was a problem hiding this comment.
use tensor_view instead: tv.get_tensor_view_idx({i, m})
| { | ||
| for(int j = i + 1; j < N; ++j) | ||
| { | ||
| long k = j + N * i - i * (i + 1) / 2 - i - 1; |
There was a problem hiding this comment.
can k be negative? if not, use uint64_t
There was a problem hiding this comment.
if can be, check for k before use it as index
| #include <cstdint> | ||
| #include <cstdlib> | ||
| #include <iostream> | ||
| #include <memory> | ||
| #include <vector> |
There was a problem hiding this comment.
are those lib necessary ?
There was a problem hiding this comment.
follow this order (AMD's convention):
`#include "InputFlags.hpp"
#include "driver.hpp"
#include "mloMarginRakningLossHost.hpp"
#include "random.hpp"
#include "tensor_driver.hpp"
#include "timer.hpp"
#include <../test/tensor_holder.hpp>
#include <../test/verify.hpp>
#include <miopen/env.hpp>
#include <miopen/handle.hpp>
#include <miopen/miopen.h>
#include <miopen/tensor.hpp>
#include `
| #include "miopen/buffer_info.hpp" | ||
| #include "miopen/errors.hpp" | ||
| #include "miopen/execution_context.hpp" | ||
| #include "miopen/invoke_params.hpp" | ||
| #include "miopen/tensor.hpp" | ||
| #include "miopen/tensor_view_utils.hpp" |
| #include <miopen/pdist.hpp> | ||
| #include <miopen/pdist/solvers.hpp> | ||
| #include <miopen/pdist/invoke_params.hpp> | ||
| #include "miopen/pdist/problem_description.hpp" |
| if(!problem.IsAllContiguous()) | ||
| { | ||
| return false; | ||
| } |
There was a problem hiding this comment.
I found no kernel implemented for non-contiguous case, please add and re-benchmark the performance. btw, I see you define problem.IsAllContiguou() as not supported, this is not a clear meaning. If you did benchmark the right implementation of non-contiguous kernel, this IsAllContiguou() should not THROW, just return false and should be in IsOverRocm function
| auto input_dtype = miopen::GetDataType(problem.GetdInputDesc().GetType()); | ||
| auto output_dtype = miopen::GetDataType(problem.GetdOutputDesc().GetType()); | ||
| auto dinput_dtype = miopen::GetDataType(problem.GetdInputDesc().GetType()); |
There was a problem hiding this comment.
you have IsSameType so this is redundant
| {"INPUT_TYPE", input_dtype == "bfloat16" ? "ushort" : input_dtype}, | ||
| {"OUTPUT_TYPE", input_dtype == "bfloat16" ? "ushort" : output_dtype}, |
Summary of Changes
PdistBackwardoperation and kernel.PdistBackward.Additional Notes
pdistdoesn't support fordtype=[fp16, bfp16]yet. This PR already includes support for those two types.input_dtype=fp16, the result may experience underflow or overflow in the following cases: (1) wheninput.dims[0]is large, indicating a high number of points for distance calculation, or (2) when the value ofpis large.Benchmark Results
Average improvement over ROCm
Detail Benchmark
float32