GitLab

Merge branch 'develop' of github.com:ROCmSoftwarePlatform/AMDMIGraphX into test_runner_match_input_output

* Check jit kernels files with clang-tidy

This provides a helper script to run rocTX markers with migraphx-driver and reduces the number of steps a user would go through running rocTX knob.
Run:
python roctx.py --run --onnx_file <ONNX_FILE> --migraphx_args "--onnx --gpu --fp16 --batch 16" --out outputfolder
Runs and parses the run output (JSON file). An example output is given below:

                                                     SUM  MIN  MAX
Marker start: gpu::convolution                      5272   10  563
Marker start: gpu::add_relu                          605   12   18
Marker start: gpu::gather                            299  145  154
Marker start: gpu::mul_add                           227   14   57
Marker start: gpu::sub                               177   13   42
Marker start: gpu::concat                            169   22   31
Marker start: gpu::triadd_relu                       163   15   18
Marker start: load                                   141    0    3
Marker start: hip::hip_copy_literal                  111    0    3
Marker start: gpu::add                                58   13   17
Marker start: broadcast                               52    0    3
Marker start: gpu::convert                            31   15   16
Marker start: slice                                   11    0    1
Marker start: gpu::pooling                             9    9    9
Marker start: step                                     2    2    2
Marker start: @param                                   2    0    1
Marker start: reshape                                  1    0    1
Marker start: hip::hip_allocate_memory                 1    1    1
Marker start: check_context::migraphx::version_...     0  ERR  ERR

TOTAL TIME: 7331 us

JSON FILE PATH: [...]/rpl_data_211019_195229_9369/input_results_211019_195229/trace.json
Parse:
python roctx.py --parse --json_path <JSON PATH FROM RUN>
Note: The parse knob is made available if the user wants to parse an already existing JSON output.

Allows --fp16 to be used in the driver to compare the target fp16 result and the ref fp32 result.

Do compilation in parallel

Currently, eliminate_contiguous will never remove contiguous for operators that use module inputs due to the fact that it doesn't pass the module inputs to compute_shape.

- Update to pass the module inputs correctly to compute_shape
- Fix the overloads of compute_shape so that when passed an empty vector of module inputs it will call the overload without module inputs
- Add tests with contiguous and pointwise module function.
- Move add_pointwise function to a seperate header to reuse across different tests

Merge branch 'develop' of github.com:ROCmSoftwarePlatform/AMDMIGraphX into test_runner_match_input_output

Currently we have the option of passing in --batch to the driver to change the batch size when the model has a dynamic dim value. We can use this flag to adjust the perf report's rate.

This enables the pointwise fusions using the MIGRAPHX_ENABLE_POINTWISE_FUSION env variable. Its disabled by default since MIOpen fusions need to be refactored.

This also adds a compile_ops pass to compile the pointwise modules. All tests except test_gpu_fast_math passes with MIGRAPHX_ENABLE_POINTWISE_FUSION=1 set.

This PR is to turn on a few gemm unit test with int8 input datatype. Before rocm4.4, int8 input data type requires matrix size to be no less than 4 in rocblas implementation. Because of this limitation, we turned off a few gemm unit tests with int8 input data type.

This limitation is removed in rocm4.4, so after we upgrade to rocm4.5, we can turn on these unit tests. Also we change to unit test conv_bn_add to adding instructions to module instead of program.
Co-authored-by: kahmed10 <15948690+kahmed10@users.noreply.github.com>

* Add workaround for devices that do not support miopen conv fusions

* Install pcre from github since the ftp.pcre.org site is no longer available

Moving our Docker file from ROCm 4.3 to 4.5 
Add Navi base GPUs in to the CI infrastructure 

In migraphx, DepthToSpace (d2s) is implemented as reshape --> transpose --> contiguous --> reshape.

If there is trailing binary pointwise operator after depthToSpace then, migraphx can move binary operator before contiguous and reshape of the depthtospce.

So, it becomes reshape-->transpose-->binary_op-->contiguous-->reshape.

Explicit contiguous wouldn't be required since binary_op outputs standard shape. So, it becomes reshape-->transpose-->binary-->reshape.

simplify_reshapes already has matcher that can do this transformation. This PR adds test for cases like depthtospace +binary op.

solves #905

This PR is the ref implementation of the nonmaxsuppression operator. It always returns the max possible output shape, which is the problem tracked in issue #948.

In migraphx, DepthToSpace (d2s) is implemented as reshape --> transpose --> contiguous --> reshape.

This PR adds matcher to find d2s + unary pointwise ops.

Application of the matcher moves the pointwise unary operation before the contiguous and reshape of the d2s.
So it becomes
reshape --> transpose --> unary --> contiguous --> reshape.

Motivation is that, later pointwise module would be created out of unary --> contiguous --> reshape. Codegen for this pointwise module can write out buffer such that explicit contiguous and reshape wouldn't be required.

This transformation is not always guaranteed to improve performance, since unary op will operate on non-standard shape. So, we would need some tuning mechanism to make decision.

#905 pending PR for binary operations.

GPU implementation of the roialign operator, using the jit approach to reduce the lib size.

Updates the theme of our documentation so that it matches the rest of the ROCm libraries.

Implementation of the roialign operator. For now, we have only the ref implementation. When we run a model on the GPU, we fall back the execution to use the ref implementation.

pthread linking errors on SLES. 

Adds a pass to fuse pointwise operators into one "pointwsie" op that has a submodule which does the calculation.

Designed to allow a user to format the values needed for the json_string: migraphx::operation("reduce_mean", "{axes : [%i, %i, %i, %i]}", axes[0], axes[1], axes[2], axes[3]) instead of needing to use string concat or stringstream

Enable a cppcheck rule to catch these redundant casts in the future

Added features:
This enables wrapping each migraphx operator with rocTX markers.
It adds new knob trace to migraphx-driver binary.

Limitation:

rocTX standalone does not output a file, it needs to be used with rocprof. Example command line:

/opt/rocm/bin/rocprof -i ./in.txt --hip-trace --roctx-trace --flush-rate 10ms --timestamp on -d cagri_out --obj-tracking on /opt/rocm/bin/migraphx-driver trace ./resnet50-v2-7.onnx --onnx --gpu
Co-authored-by: Shucai Xiao <shucai@gmail.com>

Inverse of DepthToSpace op
Co-authored-by: Shucai Xiao <shucai@gmail.com>

 when running a model on GPU, migraphx tries to print out content from gpu memory, which causes a segfault. The solution is to copy the gpu memory content back to CPU before the print.