diff --git a/mlir/lib/Dialect/MIOpen/Generator/Conv2dGenerator.cpp b/mlir/lib/Dialect/MIOpen/Generator/Conv2dGenerator.cpp index 56817cae8a78..41e5e070733c 100644 --- a/mlir/lib/Dialect/MIOpen/Generator/Conv2dGenerator.cpp +++ b/mlir/lib/Dialect/MIOpen/Generator/Conv2dGenerator.cpp @@ -120,10 +120,10 @@ LogicalResult Conv2dGenerator::parseConvConfig(const char *arguments) { auto isValid = [&argMap]() { // only require tensor configs static const std::vector validKeys = { - "batchsize", "groupsize", "in_layout", "in_type", - "in_channels", "in_h", "in_w", "out_layout", - "out_type", "out_channels", "out_h", "out_w", - "fil_layout", "fil_type", "fil_w", "fil_h"}; + "batchsize", "groupsize", "in_layout", "in_type", "in_channels", + "in_h", "in_w", "out_layout", "out_type", "out_channels", + "out_h", "out_w", "fil_layout", "fil_type", "fil_w", + "fil_h", "kernel_id"}; return std::all_of( validKeys.cbegin(), validKeys.cend(), [&argMap](const std::string &key) { return argMap.count(key) > 0; }); diff --git a/mlir/test/mlir-miopen-lib/populate_bwd.mlir b/mlir/test/mlir-miopen-lib/populate_bwd.mlir index 22c1814f71af..4a41817d33a2 100644 --- a/mlir/test/mlir-miopen-lib/populate_bwd.mlir +++ b/mlir/test/mlir-miopen-lib/populate_bwd.mlir @@ -1,9 +1,9 @@ -// RUN: mlir-miopen-lib-test --args " --operation conv2d_bwd_data --arch gfx906 --num_cu 64 --in_type fp32 --fil_type fp32 --out_type fp32 --fil_layout GNCHW --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name conv2d_bwd --groupsize 1" --option cflags | FileCheck %s --check-prefix=CFLAGS -// RUN: mlir-miopen-lib-test --args " --operation conv2d_bwd_data --arch gfx906 --num_cu 64 --in_type fp32 --fil_type fp32 --out_type fp32 --fil_layout GNCHW --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name foo --groupsize 1" --option source | FileCheck %s --check-prefix=SOURCE -// RUN: mlir-miopen-lib-test --args " --operation conv2d_bwd_data --arch gfx906 --num_cu 64 --in_type fp32 --fil_type fp32 --out_type fp32 --fil_layout GNCHW --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name bar --groupsize 1" --option header | FileCheck %s --check-prefix=HEADER -// RUN: mlir-miopen-lib-test --args " --operation conv2d_bwd_data --arch gfx906 --num_cu 64 --in_type fp32 --fil_type fp32 --out_type fp32 --fil_layout GNCHW --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name conv2d_nchw_nchw_nchw --groupsize 1" --option bin | FileCheck %s --check-prefix=BIN -// RUN: mlir-miopen-lib-test --args " --operation conv2d_bwd_data --arch gfx906 --num_cu 64 --in_type fp32 --fil_type fp32 --out_type fp32 --fil_layout GNCHW --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name bar --groupsize 1" --option tuningparams | FileCheck %s --check-prefix=TUNING -// RUN: mlir-miopen-driver --conv-config "--operation conv2d_bwd_data --arch gfx906 --num_cu 64 --in_type fp32 --fil_type fp32 --out_type fp32 --fil_layout GNCHW --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name bar --groupsize 1 " | FileCheck %s --check-prefix=DRIVER +// RUN: mlir-miopen-lib-test --args " --operation conv2d_bwd_data --arch gfx906 --num_cu 64 --in_type fp32 --fil_type fp32 --out_type fp32 --fil_layout GNCHW --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name conv2d_bwd --groupsize 1 --kernel_id 0" --option cflags | FileCheck %s --check-prefix=CFLAGS +// RUN: mlir-miopen-lib-test --args " --operation conv2d_bwd_data --arch gfx906 --num_cu 64 --in_type fp32 --fil_type fp32 --out_type fp32 --fil_layout GNCHW --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name foo --groupsize 1 --kernel_id 0" --option source | FileCheck %s --check-prefix=SOURCE +// RUN: mlir-miopen-lib-test --args " --operation conv2d_bwd_data --arch gfx906 --num_cu 64 --in_type fp32 --fil_type fp32 --out_type fp32 --fil_layout GNCHW --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name bar --groupsize 1 --kernel_id 0" --option header | FileCheck %s --check-prefix=HEADER +// RUN: mlir-miopen-lib-test --args " --operation conv2d_bwd_data --arch gfx906 --num_cu 64 --in_type fp32 --fil_type fp32 --out_type fp32 --fil_layout GNCHW --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name conv2d_nchw_nchw_nchw --groupsize 1 --kernel_id 0" --option bin | FileCheck %s --check-prefix=BIN +// RUN: mlir-miopen-lib-test --args " --operation conv2d_bwd_data --arch gfx906 --num_cu 64 --in_type fp32 --fil_type fp32 --out_type fp32 --fil_layout GNCHW --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name bar --groupsize 1 --kernel_id 0" --option tuningparams | FileCheck %s --check-prefix=TUNING +// RUN: mlir-miopen-driver --conv-config "--operation conv2d_bwd_data --arch gfx906 --num_cu 64 --in_type fp32 --fil_type fp32 --out_type fp32 --fil_layout GNCHW --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name bar --groupsize 1 --kernel_id 0" | FileCheck %s --check-prefix=DRIVER // CFLAGS: conv2d_bwd // SOURCE: void mlir_gen_igemm_conv2d_cpp_v4r1_bwd diff --git a/mlir/test/mlir-miopen-lib/populate_bww.mlir b/mlir/test/mlir-miopen-lib/populate_bww.mlir index 0ef0810c891b..7e9a86768d5e 100644 --- a/mlir/test/mlir-miopen-lib/populate_bww.mlir +++ b/mlir/test/mlir-miopen-lib/populate_bww.mlir @@ -1,9 +1,9 @@ -// RUN: mlir-miopen-lib-test --args " --operation conv2d_bwd_weight --arch gfx906 --num_cu 64 --in_type fp32 --fil_type fp32 --out_type fp32 --fil_layout GNCHW --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name conv2d_bww --groupsize 1" --option cflags | FileCheck %s --check-prefix=CFLAGS -// RUN: mlir-miopen-lib-test --args " --operation conv2d_bwd_weight --arch gfx906 --num_cu 64 --in_type fp32 --fil_type fp32 --out_type fp32 --fil_layout GNCHW --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name foo --groupsize 1" --option source | FileCheck %s --check-prefix=SOURCE -// RUN: mlir-miopen-lib-test --args " --operation conv2d_bwd_weight --arch gfx906 --num_cu 64 --in_type fp32 --fil_type fp32 --out_type fp32 --fil_layout GNCHW --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name bar --groupsize 1" --option header | FileCheck %s --check-prefix=HEADER -// RUN: mlir-miopen-lib-test --args " --operation conv2d_bwd_weight --arch gfx906 --num_cu 64 --in_type fp32 --fil_type fp32 --out_type fp32 --fil_layout GNCHW --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name conv2d --groupsize 1" --option bin | FileCheck %s --check-prefix=BIN -// RUN: mlir-miopen-lib-test --args " --operation conv2d_bwd_weight --arch gfx906 --num_cu 64 --in_type fp32 --fil_type fp32 --out_type fp32 --fil_layout GNCHW --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name conv2d --groupsize 1" --option tuningparams | FileCheck %s --check-prefix=TUNING -// RUN: mlir-miopen-driver --conv-config " --operation conv2d_bwd_weight --arch gfx906 --num_cu 64 --in_type fp32 --fil_type fp32 --out_type fp32 --fil_layout GNCHW --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name conv2d --groupsize 1 " | FileCheck %s --check-prefix=DRIVER +// RUN: mlir-miopen-lib-test --args " --operation conv2d_bwd_weight --arch gfx906 --num_cu 64 --in_type fp32 --fil_type fp32 --out_type fp32 --fil_layout GNCHW --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name conv2d_bww --groupsize 1 --kernel_id 0" --option cflags | FileCheck %s --check-prefix=CFLAGS +// RUN: mlir-miopen-lib-test --args " --operation conv2d_bwd_weight --arch gfx906 --num_cu 64 --in_type fp32 --fil_type fp32 --out_type fp32 --fil_layout GNCHW --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name foo --groupsize 1 --kernel_id 0" --option source | FileCheck %s --check-prefix=SOURCE +// RUN: mlir-miopen-lib-test --args " --operation conv2d_bwd_weight --arch gfx906 --num_cu 64 --in_type fp32 --fil_type fp32 --out_type fp32 --fil_layout GNCHW --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name bar --groupsize 1 --kernel_id 0" --option header | FileCheck %s --check-prefix=HEADER +// RUN: mlir-miopen-lib-test --args " --operation conv2d_bwd_weight --arch gfx906 --num_cu 64 --in_type fp32 --fil_type fp32 --out_type fp32 --fil_layout GNCHW --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name conv2d --groupsize 1 --kernel_id 0" --option bin | FileCheck %s --check-prefix=BIN +// RUN: mlir-miopen-lib-test --args " --operation conv2d_bwd_weight --arch gfx906 --num_cu 64 --in_type fp32 --fil_type fp32 --out_type fp32 --fil_layout GNCHW --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name conv2d --groupsize 1 --kernel_id 0" --option tuningparams | FileCheck %s --check-prefix=TUNING +// RUN: mlir-miopen-driver --conv-config " --operation conv2d_bwd_weight --arch gfx906 --num_cu 64 --in_type fp32 --fil_type fp32 --out_type fp32 --fil_layout GNCHW --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name conv2d --groupsize 1 --kernel_id 0" | FileCheck %s --check-prefix=DRIVER // CFLAGS: conv2d_bww diff --git a/mlir/test/mlir-miopen-lib/populate_fw.mlir b/mlir/test/mlir-miopen-lib/populate_fw.mlir index d9ab15da6c15..fa173996f924 100644 --- a/mlir/test/mlir-miopen-lib/populate_fw.mlir +++ b/mlir/test/mlir-miopen-lib/populate_fw.mlir @@ -1,9 +1,9 @@ -// RUN: mlir-miopen-lib-test --args " --operation conv2d --arch gfx906 --num_cu 64 --fil_layout GNCHW --in_type fp32 --fil_type fp32 --out_type fp32 --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name conv2d_fwd --groupsize 1" --option cflags | FileCheck %s --check-prefix=CFLAGS -// RUN: mlir-miopen-lib-test --args " --operation conv2d --arch gfx906 --num_cu 64 --in_type fp32 --fil_type fp32 --out_type fp32 --fil_layout GNCHW --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name foo --groupsize 1" --option source | FileCheck %s --check-prefix=SOURCE -// RUN: mlir-miopen-lib-test --args " --operation conv2d --arch gfx906 --num_cu 64 --in_type fp32 --fil_type fp32 --out_type fp32 --fil_layout GNCHW --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name bar --groupsize 1" --option header | FileCheck %s --check-prefix=HEADER -// RUN: mlir-miopen-lib-test --args " --operation conv2d --arch gfx906 --num_cu 64 --in_type fp32 --fil_type fp32 --out_type fp32 --fil_layout GNCHW --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name conv2d_nchw_kcyx_nkhw --groupsize 1" --option bin | FileCheck %s --check-prefix=BIN -// RUN: mlir-miopen-lib-test --args " --operation conv2d --arch gfx906 --num_cu 64 --fil_layout GNCHW --in_type fp32 --fil_type fp32 --out_type fp32 --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name conv2d_fwd --groupsize 1" --option tuningparams | FileCheck %s --check-prefix=TUNING -// RUN: mlir-miopen-driver --conv-config "--operation conv2d --arch gfx906 --num_cu 64 --fil_layout GNCHW --in_type fp32 --fil_type fp32 --out_type fp32 --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name conv2d_fwd --groupsize 1" | FileCheck %s --check-prefix=DRIVER +// RUN: mlir-miopen-lib-test --args " --operation conv2d --arch gfx906 --num_cu 64 --fil_layout GNCHW --in_type fp32 --fil_type fp32 --out_type fp32 --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name conv2d_fwd --groupsize 1 --kernel_id 0" --option cflags | FileCheck %s --check-prefix=CFLAGS +// RUN: mlir-miopen-lib-test --args " --operation conv2d --arch gfx906 --num_cu 64 --in_type fp32 --fil_type fp32 --out_type fp32 --fil_layout GNCHW --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name foo --groupsize 1 --kernel_id 0" --option source | FileCheck %s --check-prefix=SOURCE +// RUN: mlir-miopen-lib-test --args " --operation conv2d --arch gfx906 --num_cu 64 --in_type fp32 --fil_type fp32 --out_type fp32 --fil_layout GNCHW --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name bar --groupsize 1 --kernel_id 0" --option header | FileCheck %s --check-prefix=HEADER +// RUN: mlir-miopen-lib-test --args " --operation conv2d --arch gfx906 --num_cu 64 --in_type fp32 --fil_type fp32 --out_type fp32 --fil_layout GNCHW --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name conv2d_nchw_kcyx_nkhw --groupsize 1 --kernel_id 0" --option bin | FileCheck %s --check-prefix=BIN +// RUN: mlir-miopen-lib-test --args " --operation conv2d --arch gfx906 --num_cu 64 --fil_layout GNCHW --in_type fp32 --fil_type fp32 --out_type fp32 --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name conv2d_fwd --groupsize 1 --kernel_id 0" --option tuningparams | FileCheck %s --check-prefix=TUNING +// RUN: mlir-miopen-driver --conv-config "--operation conv2d --arch gfx906 --num_cu 64 --fil_layout GNCHW --in_type fp32 --fil_type fp32 --out_type fp32 --in_layout NGCHW --out_layout NGCHW --batchsize 64 --in_channels 1024 --out_channels 1024 --in_h 14 --in_w 14 --out_h 14 --out_w 14 --fil_h 1 --fil_w 1 --dilation_h 1 --dilation_w 1 --conv_stride_h 1 --conv_stride_w 1 --padding_h 0 --padding_w 0 --kernel_name conv2d_fwd --groupsize 1 --kernel_id 0" | FileCheck %s --check-prefix=DRIVER // CFLAGS: conv2d_fwd // SOURCE: void mlir_gen_igemm_conv2d_cpp_v4r4_fwd