Backward data

[ltqin]

This PR is backward data muti-kernels version, if merge with oob_write(#227). stride=2 case can pass

[whchung]

@asleepzzz / @ltqin is logic in getGemmA correct about OOB dims?

[whchung]

@ltqin please fix merge conflicts.

and now since #227 has been merged, it's better to do a rebase.

[whchung]

jenkins: retest this please.

[jerryyin]

Two major issues needs attention/resolution:

• gemmId should be guaranteed to be consecutive.
• You need to have end-to-end interface unit test to test the client perspective, since there are fairly fundamental changes. You need to add test such that it take advantage of mlir-miopen-lib-test.cpp to compile at least one case of bwd that can populate multiple kernels.

[ltqin]

• You need to have end-to-end interface unit test to test the client perspective, since there are fairly fundamental changes. You need to add test such that it take advantage of mlir-miopen-lib-test.cpp to compile at least one case of bwd that can populate multiple kernels.

@jerryyin please check unit test: mlir/test/mlir-miopen-driver/populate_bwd_multi_kernels.mlir

[jerryyin]

@ltqin That file is not enough because it is not end-to-end.

You'd need to add a for loop here such that it can cover multi-kernel case:

https://github.com/ROCmSoftwarePlatform/llvm-project-mlir/blob/miopen-dialect/mlir/tools/mlir-miopen-lib/mlir-miopen-lib-test.cpp#L62-L87

The add unit test like here:

https://github.com/ROCmSoftwarePlatform/llvm-project-mlir/blob/miopen-dialect/mlir/test/mlir-miopen-lib/populate_fw.mlir#L1

[whchung]

jenkins: retest this please.

[ltqin]

@asroy According information(ROCm/MIOpen#917), I test the cases that fail on MIOpen . but all case pass on mlir-miopen-driver

[ltqin]

@ltqin That file is not enough because it is not end-to-end.

You'd need to add a for loop here such that it can cover multi-kernel case:

@jerryyin please check the update, there are some different from what we talked

[jerryyin]

@ltqin We synced offline, I'd prefer to reconstruct the arguments instead of forcing the client to take the kernel_id in all scenarios.

[asroy]

LGTM.

@ltqin I'm not sure about the logic for vector load, but since you have test some configs and they passed. I'm going to approve it.

Please paste the configs you have tested for future records.

Please also watch this task, I will write done the derivation for bwd-data-v4r1 when I got time, and then you can check with current implementation.
https://github.com/ROCmSoftwarePlatform/miopen-kernel-project-private/issues/45

[ltqin]

@asroy I have tested configs for padding issue:
./bin/mlir-miopen-driver -p=false -fil_layout=gkcyx -in_layout=ngchw -out_layout=ngkhw -batchsize=32 -in_channels=32 -out_channels=32 -in_h=14 -in_w=14 -fil_h=1 -fil_w=1 --dilation_h=1 --dilation_w=1 --padding_h=1 --padding_w=1 --conv_stride_h=1 --conv_stride_w=1 --groupsize=1 -pv --operation=conv2d_bwd_data -c -rand 1| ./bin/mlir-rocm-runner --shared-libs=./lib/librocm-runtime-wrappers.so,./lib/libmlir_runner_utils.so --entry-point-result=void

./bin/mlir-miopen-driver -p=false -fil_layout=gkcyx -in_layout=ngchw -out_layout=ngkhw -batchsize=32 -in_channels=32 -out_channels=32 -in_h=14 -in_w=14 -fil_h=3 -fil_w=3 --dilation_h=1 --dilation_w=1 --padding_h=1 --padding_w=1 --conv_stride_h=1 --conv_stride_w=1 --groupsize=1 -pv --operation=conv2d_bwd_data -c -rand 1| ./bin/mlir-rocm-runner --shared-libs=./lib/librocm-runtime-wrappers.so,./lib/libmlir_runner_utils.so --entry-point-result=void

./bin/mlir-miopen-driver -p=false -fil_layout=gkcyx -in_layout=ngchw -out_layout=ngkhw -batchsize=32 -in_channels=32 -out_channels=32 -in_h=15 -in_w=15 -fil_h=1 -fil_w=1 --dilation_h=1 --dilation_w=1 --padding_h=1 --padding_w=1 --conv_stride_h=1 --conv_stride_w=1 --groupsize=1 -pv --operation=conv2d_bwd_data -c -rand 1| ./bin/mlir-rocm-runner --shared-libs=./lib/librocm-runtime-wrappers.so,./lib/libmlir_runner_utils.so --entry-point-result=void

./bin/mlir-miopen-driver -p=false -fil_layout=gkcyx -in_layout=ngchw -out_layout=ngkhw -batchsize=32 -in_channels=32 -out_channels=32 -in_h=16 -in_w=16 -fil_h=1 -fil_w=1 --dilation_h=1 --dilation_w=1 --padding_h=1 --padding_w=1 --conv_stride_h=1 --conv_stride_w=1 --groupsize=1 -pv --operation=conv2d_bwd_data -c | ./bin/mlir-rocm-runner --shared-libs=./lib/librocm-runtime-wrappers.so,./lib/libmlir_runner_utils.so --entry-point-result=void

./bin/mlir-miopen-driver -p=false -fil_layout=gkcyx -in_layout=ngchw -out_layout=ngkhw -batchsize=32 -in_channels=32 -out_channels=32 -in_h=18 -in_w=18 -fil_h=1 -fil_w=1 --dilation_h=1 --dilation_w=1 --padding_h=1 --padding_w=1 --conv_stride_h=1 --conv_stride_w=1 --groupsize=1 -pv --operation=conv2d_bwd_data -c -rand 1| ./bin/mlir-rocm-runner --shared-libs=./lib/librocm-runtime-wrappers.so,./lib/libmlir_runner_utils.so --entry-point-result=void

./bin/mlir-miopen-driver -p=false -fil_layout=gkcyx -in_layout=ngchw -out_layout=ngkhw -batchsize=32 -in_channels=32 -out_channels=32 -in_h=20 -in_w=20 -fil_h=1 -fil_w=1 --dilation_h=1 --dilation_w=1 --padding_h=1 --padding_w=1 --conv_stride_h=1 --conv_stride_w=1 --groupsize=1 -pv --operation=conv2d_bwd_data -c -rand 1| ./bin/mlir-rocm-runner --shared-libs=./lib/librocm-runtime-wrappers.so,./lib/libmlir_runner_utils.so --entry-point-result=void