ConvBinWinogradRxSf2x3 bug #411

alexandraBara · 2020-08-31T20:54:25Z

Found Winograd failing in Generic search during tuning:

 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvBinWinogradRxSf2x3: /root/dMIOpen/src/include/miopen/generic_search.hpp:408: GenericSearch: top_ocl_ptr == nullptr || bot_ocl_ptr == nullptr || wei_ocl_ptr == nullptr

Here is a sample:

INFO - Executing: sudo docker run --device='/dev/kfd' --device='/dev/dri' -w /home/miopenpdb -v /home/miopenpdb:/home/miopenpdb --user=root --group-add video --privileged=true --rm miopentuna bash  -c "export MIOPEN_LOG_LEVEL=7 && export MIOPEN_FIND_ENFORCE=4 && export MIOPEN_DISABLE_CACHE=1 && export HIP_VISIBLE_DEVICES=1 && MIOpenDriver conv -V 0 -i 1 --forw 1 --pad_h 0 --out_channels 1024 --fil_w 1 --dilation_w 1 --fil_h 1 --in_h 28 --conv_stride_w 1 --group_count 1 --in_channels 256 --in_w 28 --dilation_h 1 --conv_stride_h 1 --pad_w 0 --batchsize 64 --pad_mode default --mode conv --fil_d 1 --in_d 1 --spatial_dim 2 --conv_stride_d 1 --dilation_d 1 --pad_d 0 --trans_output_pad_d 0 2>&1 " on machine: 10.216.64.100 for GPU ID: 1
 INFO - Setting job id 11065938 state to running
 INFO - MIOpenDriver conv -V 0 -i 1 --forw 1 --pad_h 0 --out_channels 1024 --fil_w 1 --dilation_w 1 --fil_h 1 --in_h 28 --conv_stride_w 1 --group_count 1 --in_channels 256 --in_w 28 --dilation_h 1 --conv_stride_h 1 --pad_w 0 --batchsize 64 --pad_mode default --mode conv --fil_d 1 --in_d 1 --spatial_dim 2 --conv_stride_d 1 --dilation_d 1 --pad_d 0 --trans_output_pad_d 0
 INFO - MIOpen(HIP): Info [Handle] stream: 0x1c0ba60, device_id: 0
 INFO - MIOpen(HIP): Info [ForwardGetWorkSpaceSize]
 INFO - MIOpen(HIP): Info2 [HipCompilerVersionImpl] Read version information from HIP package...
 INFO - MIOpen(HIP): Info [HipCompilerVersionImpl] 3.7.20315
 INFO - MIOpen(HIP): Info [AmdRocmMetadataVersionDetect] ROCm MD version AMDHSA_COv3, MIOpen version 2.7.0.8175-13470ceb
 INFO - MIOpen(HIP): Info2 [ValidateGcnAssemblerImpl] Running: '/opt/rocm/llvm/bin/clang --version'
 INFO - MIOpen(HIP): Info2 [ValidateGcnAssemblerImpl] clang version 11.0.0 (/src/external/llvm-project/clang ee4e4ebbadcc8ea14ce99e34ed31ab31e94827ac)
 INFO - MIOpen(HIP): Info2 [ValidateGcnAssemblerImpl] Target: x86_64-unknown-linux-gnu
 INFO - MIOpen(HIP): Info2 [ValidateGcnAssemblerImpl] Thread model: posix
 INFO - MIOpen(HIP): Info2 [ValidateGcnAssemblerImpl] InstalledDir: /opt/rocm/llvm/bin
 INFO - MIOpen(HIP): Info2 [ValidateGcnAssemblerImpl]
 INFO - MIOpen(HIP): Info2 [SQLiteBase] Initializing system database file /opt/rocm/miopen/share/miopen/db/miopen.db
 INFO - MIOpen(HIP): Trace [Exec] 140632355535360:PRAGMA table_info(config);
 INFO - MIOpen(HIP): Trace [Exec] 140632355535360:PRAGMA table_info(perf_db);
 INFO - MIOpen(HIP): Info2 [SQLiteBase] Initializing user database file /home/miopenpdb/.config/miopen/miopen_1.0.0.udb
 INFO - MIOpen(HIP): Trace [Exec] 140632355535360:SELECT name FROM sqlite_master WHERE type = 'table' AND (name = 'config');
 INFO - MIOpen(HIP): Trace [Exec] 140632355535360:SELECT name FROM sqlite_master WHERE type = 'table' AND (name = 'perf_db');
 INFO - MIOpen(HIP): Trace [SQLitePerfDb] Database created successfully
 INFO - MIOpen(HIP): Trace [Exec] 140632355535360:PRAGMA table_info(config);
 INFO - MIOpen(HIP): Trace [Exec] 140632355535360:PRAGMA table_info(perf_db);
 INFO - MIOpen(HIP): Info2 [SearchForAllSolutions] ConvBinWinograd3x3U: Not applicable
 INFO - MIOpen(HIP): Info [FindSolutionImpl] ConvBinWinogradRxSf3x2 (not searchable)
 INFO - MIOpen(HIP): Info2 [SearchForAllSolutions] ConvBinWinogradRxSf3x2: Success.
 INFO - MIOpen(HIP): Info [FindSolutionImpl] ConvBinWinogradRxSf2x3
 INFO - MIOpen(HIP): Warning [FindSolutionImpl] Perf Db: load skipped: ConvBinWinogradRxSf2x3, enforce: SEARCH_DB_UPDATE(4), ALL(1)
 INFO - MIOpen(HIP): Info [FindSolutionImpl] Starting search: ConvBinWinogradRxSf2x3, enforce: SEARCH_DB_UPDATE(4), ALL(1)
 INFO - MIOpen(HIP): Info [GetPerformanceConfig] 56
 **INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvBinWinogradRxSf2x3: /root/dMIOpen/src/include/miopen/generic_search.hpp:408: GenericSearch: top_ocl_ptr == nullptr || bot_ocl_ptr == nullptr || wei_ocl_ptr == nullptr**

@atamazov can you please add IIya and Kamil to this issue, i dont know their github usernames

The text was updated successfully, but these errors were encountered:

daniellowell · 2020-08-31T21:26:24Z

I've seen a similar issue (probably side effect from this issue) here:

MIOpen(HIP): Info [CompileWrwSolution] solver_id = ConvBinWinogradRxSf2x3
MIOpen(HIP): Info [FindConvBwdWeightsAlgorithm] miopenConvolutionBwdWeightsAlgoWinograd	0.05104	0
MIOpen(HIP): Info [FindConvBwdWeightsAlgorithm] BWrW Chosen Algorithm: ConvBinWinogradRxSf2x3 , 0, 0.05104
MIOpen(HIP): Info [ConvolutionBackwardWeights] algo = 3, workspace = 0
1x1x1x2x2x1x1x1x1xNCHWxFP32x0x0x1x1x1x1x1x0
/var/jenkins/workspace/MLLibs_MIOpen_zyin-layouts/build/bin/test_conv2d --float --cmode conv --pmode same --group-count 1 --input 1, 1, 1, 1 --weights 1, 1, 2, 2 --pads_strides_dilations 2 2 1 1 1 1 --trans_output_pads 0 0 
FAILED: /var/jenkins/workspace/MLLibs_MIOpen_zyin-layouts/src/ocl/convolutionocl.cpp:3934: Error running Winograd WrW. Was Find() run previously?
Backward weights convolution: ConvBinWinogradRxSf2x3

atamazov · 2020-08-31T22:27:33Z

@DrizztDoUrden Please have a glance at these error messages. Anything that looks familiar?

DrizztDoUrden · 2020-08-31T22:33:29Z

Not familiar to me. I may have messed up with something in the generic search, though. But I won't bet on that without further investigation if it is only Winograd related. And the message looks like it is from develop. I'll check tomorrow.

daniellowell · 2020-09-01T00:25:31Z

AFAWK this is the only solver that is affected: ConvBinWinogradRxSf2x3

DrizztDoUrden · 2020-09-01T08:25:41Z

Does it happen on develop or on branch based from invokers auto-tune?

alexandraBara · 2020-09-01T12:18:35Z

Does it happen on develop or on branch based from invokers auto-tune?

It's in develop

DrizztDoUrden · 2020-09-01T14:50:04Z

Then I am not sure what causes it, but it may be fixed in my branch due to the removal of almost all buffer mangling in the generic search.
I pass them uniformly everywhere in the autotune branch. But it's interesting where do they disappear in this specific case.

alexandraBara · 2020-09-01T15:45:55Z

On gfx908 i see this error on more than just the Winograd solver:

gfx908 LOG

 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvBinWinogradRxSf2x3: /root/dMIOpen/src/include/miopen/generic_search.hpp:408: GenericSearch: top_ocl_ptr == nullptr || bot_ocl_ptr == nullptr || wei_ocl_ptr == nullptr
 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvBinWinogradRxSf2x3: /root/dMIOpen/src/include/miopen/generic_search.hpp:408: GenericSearch: top_ocl_ptr == nullptr || bot_ocl_ptr == nullptr || wei_ocl_ptr == nullptr
 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvHipImplicitGemmBwdDataV1R1Xdlops: /root/dMIOpen/src/include/miopen/generic_search.hpp:575: Search failed
- MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvHipImplicitGemmBwdDataV1R1Xdlops: /root/dMIOpen/src/include/miopen/generic_search.hpp:575: Search failed
- INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvBinWinogradRxSf2x3: /root/dMIOpen/src/include/miopen/generic_search.hpp:408: GenericSearch: top_ocl_ptr == nullptr || bot_ocl_ptr == nullptr || wei_ocl_ptr == nullptr
 - INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvBinWinogradRxSf2x3: /root/dMIOpen/src/include/miopen/generic_search.hpp:408: GenericSearch: top_ocl_ptr == nullptr || bot_ocl_ptr == nullptr || wei_ocl_ptr == nullptr
 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvBinWinogradRxSf2x3: /root/dMIOpen/src/include/miopen/generic_search.hpp:408: GenericSearch: top_ocl_ptr == nullptr || bot_ocl_ptr == nullptr || wei_ocl_ptr == nullptr
 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvHipImplicitGemmBwdDataV1R1Xdlops: /root/dMIOpen/src/include/miopen/generic_search.hpp:575: Search failed
 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvHipImplicitGemmBwdDataV1R1Xdlops: /root/dMIOpen/src/include/miopen/generic_search.hpp:575: Search failed
 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvHipImplicitGemmBwdDataV1R1Xdlops: /root/dMIOpen/src/include/miopen/generic_search.hpp:575: Search failed
INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvHipImplicitGemmBwdDataV1R1Xdlops: /root/dMIOpen/src/include/miopen/generic_search.hpp:575: Search failed
 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvBinWinogradRxSf2x3: /root/dMIOpen/src/include/miopen/generic_search.hpp:408: GenericSearch: top_ocl_ptr == nullptr || bot_ocl_ptr == nullptr || wei_ocl_ptr == nullptr
 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvBinWinogradRxSf2x3: /root/dMIOpen/src/include/miopen/generic_search.hpp:408: GenericSearch: top_ocl_ptr == nullptr || bot_ocl_ptr == nullptr || wei_ocl_ptr == nullptr
 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvBinWinogradRxSf2x3: /root/dMIOpen/src/include/miopen/generic_search.hpp:408: GenericSearch: top_ocl_ptr == nullptr || bot_ocl_ptr == nullptr || wei_ocl_ptr == nullptr
 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvHipImplicitGemmBwdDataV1R1Xdlops: /root/dMIOpen/src/include/miopen/generic_search.hpp:575: Search failed
 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvHipImplicitGemmBwdDataV1R1Xdlops: /root/dMIOpen/src/include/miopen/generic_search.hpp:575: Search failed
 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvHipImplicitGemmBwdDataV1R1Xdlops: /root/dMIOpen/src/include/miopen/generic_search.hpp:575: Search failed
 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvHipImplicitGemmBwdDataV1R1Xdlops: /root/dMIOpen/src/include/miopen/generic_search.hpp:575: Search failed
 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvHipImplicitGemmBwdDataV1R1Xdlops: /root/dMIOpen/src/include/miopen/generic_search.hpp:575: Search failed
 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvBinWinogradRxSf2x3: /root/dMIOpen/src/include/miopen/generic_search.hpp:408: GenericSearch: top_ocl_ptr == nullptr || bot_ocl_ptr == nullptr || wei_ocl_ptr == nullptr
 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvBinWinogradRxSf2x3: /root/dMIOpen/src/include/miopen/generic_search.hpp:408: GenericSearch: top_ocl_ptr == nullptr || bot_ocl_ptr == nullptr || wei_ocl_ptr == nullptr
 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvHipImplicitGemmBwdDataV1R1Xdlops: /root/dMIOpen/src/include/miopen/generic_search.hpp:575: Search failed
 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvHipImplicitGemmBwdDataV1R1Xdlops: /root/dMIOpen/src/include/miopen/generic_search.hpp:575: Search failed
 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvHipImplicitGemmBwdDataV1R1Xdlops: /root/dMIOpen/src/include/miopen/generic_search.hpp:575: Search failed
 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvBinWinogradRxSf2x3: /root/dMIOpen/src/include/miopen/generic_search.hpp:408: GenericSearch: top_ocl_ptr == nullptr || bot_ocl_ptr == nullptr || wei_ocl_ptr == nullptr
INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvBinWinogradRxSf2x3: /root/dMIOpen/src/include/miopen/generic_search.hpp:408: GenericSearch: top_ocl_ptr == nullptr || bot_ocl_ptr == nullptr || wei_ocl_ptr == nullptr
 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvBinWinogradRxSf2x3: /root/dMIOpen/src/include/miopen/generic_search.hpp:408: GenericSearch: top_ocl_ptr == nullptr || bot_ocl_ptr == nullptr || wei_ocl_ptr == nullptr
 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvHipImplicitGemmBwdDataV1R1Xdlops: /root/dMIOpen/src/include/miopen/generic_search.hpp:575: Search failed
 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvBinWinogradRxSf2x3: /root/dMIOpen/src/include/miopen/generic_search.hpp:408: GenericSearch: top_ocl_ptr == nullptr || bot_ocl_ptr == nullptr || wei_ocl_ptr == nullptr
 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvBinWinogradRxSf2x3: /root/dMIOpen/src/include/miopen/generic_search.hpp:408: GenericSearch: top_ocl_ptr == nullptr || bot_ocl_ptr == nullptr || wei_ocl_ptr == nullptr
 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvBinWinogradRxSf2x3: /root/dMIOpen/src/include/miopen/generic_search.hpp:408: GenericSearch: top_ocl_ptr == nullptr || bot_ocl_ptr == nullptr || wei_ocl_ptr == nullptr
 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvBinWinogradRxSf2x3: /root/dMIOpen/src/include/miopen/generic_search.hpp:408: GenericSearch: top_ocl_ptr == nullptr || bot_ocl_ptr == nullptr || wei_ocl_ptr == nullptr
 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvHipImplicitGemmBwdDataV1R1Xdlops: /root/dMIOpen/src/include/miopen/generic_search.hpp:575: Search failed
 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvHipImplicitGemmBwdDataV1R1Xdlops: /root/dMIOpen/src/include/miopen/generic_search.hpp:575: Search failed
 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvHipImplicitGemmBwdDataV1R1Xdlops: /root/dMIOpen/src/include/miopen/generic_search.hpp:575: Search failed
 INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvHipImplicitGemmBwdDataV1R1Xdlops: /root/dMIOpen/src/include/miopen/generic_search.hpp:575: Search failed

daniellowell · 2020-09-02T14:57:54Z

@DrizztDoUrden Any progress here? @shurale-nkn Can you take a look as well?

shurale-nkn · 2020-09-02T15:59:16Z

@daniellowell
Yes, I'm already working on it.

I think this is a driver error, not solver.

In my tests this error is printed, but execution continues and the winograd 'ConvBinWinogradRxSf2x3' was tuned later.

log

MIOpenDriver conv -V 0 -i 1 --forw 1 --pad_h 0 --out_channels 1024 --fil_w 1 --dilation_w 1 --fil_h 1 --in_h 28 --conv_stride_w 1 --group_count 1 --in_channels 256 --in_w 28 --dilation_h 1 --conv_stride_h 1 --pad_w 0 --batchsize 64 --pad_mode default --mode conv --fil_d 1 --in_d 1 --spatial_dim 2 --conv_stride_d 1 --dilation_d 1 --pad_d 0 --trans_output_pad_d 0 -F 1
MIOpen(HIP): Info [Handle] stream: 0x2bef800, device_id: 0
MIOpen(HIP): Info [ForwardGetWorkSpaceSize] 
MIOpen(HIP): Info2 [HipCompilerVersionImpl] Read version information from HIP package...
MIOpen(HIP): Info [HipCompilerVersionImpl] 3.5.20231
MIOpen(HIP): Info [AmdRocmMetadataVersionDetect] ROCm MD version AMDHSA_COv3, MIOpen version 2.7.0.8219-1dd7d90c
MIOpen(HIP): Info2 [ValidateGcnAssemblerImpl] Running: '/opt/rocm/llvm/bin/clang --version'
MIOpen(HIP): Info2 [ValidateGcnAssemblerImpl] clang version 11.0.0 (/src/external/llvm-project/clang 0383ad1cfb0a8e05b0a020e8632400194628b243)
MIOpen(HIP): Info2 [ValidateGcnAssemblerImpl] Target: x86_64-unknown-linux-gnu
MIOpen(HIP): Info2 [ValidateGcnAssemblerImpl] Thread model: posix
MIOpen(HIP): Info2 [ValidateGcnAssemblerImpl] InstalledDir: /opt/rocm/llvm/bin
MIOpen(HIP): Info2 [ValidateGcnAssemblerImpl] 
MIOpen(HIP): Info2 [SQLiteBase] Initializing system database file /home/kamil/vf/MLOpen-Devel/MIOpen/src/kernels/miopen.db
MIOpen(HIP): Info2 [SQLiteBase] Initializing user database file /home/kamil/vf/MLOpen-Devel/MIOpen/src/kernels/miopen_1.0.0.udb
MIOpen(HIP): Info2 [SearchForAllSolutions] ConvBinWinograd3x3U: Not applicable
MIOpen(HIP): Info [FindSolutionImpl] ConvBinWinogradRxSf3x2 (not searchable)
MIOpen(HIP): Info2 [SearchForAllSolutions] ConvBinWinogradRxSf3x2: Success.
MIOpen(HIP): Info [FindSolutionImpl] ConvBinWinogradRxSf2x3
MIOpen(HIP): Warning [FindSolutionImpl] Perf Db: load skipped: ConvBinWinogradRxSf2x3, enforce: SEARCH_DB_UPDATE(4), ALL(1)
MIOpen(HIP): Info [FindSolutionImpl] Starting search: ConvBinWinogradRxSf2x3, enforce: SEARCH_DB_UPDATE(4), ALL(1)
MIOpen(HIP): Info [Search] SEARCH PROBLEMS START
MIOpen(HIP): Info [GetPerformanceConfig] 60
MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvBinWinogradRxSf2x3: /home/kamil/vf/MLOpen-Devel/MIOpen/src/include/miopen/generic_search.hpp:408: GenericSearch: top_ocl_ptr == nullptr || bot_ocl_ptr == nullptr || wei_ocl_ptr == nullptr
MIOpen(HIP): Info [GetPerformanceConfig] 60
MIOpen(HIP): Info2 [SearchForAllSolutions] ConvBinWinogradRxSf2x3: Success.
MIOpen(HIP): Info [FindSolutionImpl] ConvBinWinogradRxS (not searchable)
MIOpen(HIP): Info2 [SearchForAllSolutions] ConvBinWinogradRxS: Success.
MIOpen(HIP): Info2 [IsApplicable] Workspace required: 845414400, limit: 2000000000
MIOpen(HIP): Info2 [SearchForAllSolutions] ConvMPBidirectWinograd<3-3>: Not applicable
MIOpen(HIP): Info2 [IsApplicable] Workspace required: 615776256, limit: 2000000000
MIOpen(HIP): Info2 [SearchForAllSolutions] ConvMPBidirectWinograd<4-3>: Not applicable
MIOpen(HIP): Info2 [IsApplicable] Workspace required: 629407744, limit: 2000000000
MIOpen(HIP): Info2 [SearchForAllSolutions] ConvMPBidirectWinograd<5-3>: Not applicable
MIOpen(HIP): Info2 [IsApplicable] Workspace required: 591396864, limit: 2000000000
MIOpen(HIP): Info2 [SearchForAllSolutions] ConvMPBidirectWinograd<6-3>: Not applicable
MIOpen(HIP): Info [GetFindModeValueImpl] MIOPEN_FIND_MODE = NORMAL(1)
MIOpen(HIP): Info2 [GetWorkspaceSize] ConvAsm3x3U: Not applicable
MIOpen(HIP): Info2 [GetWorkspaceSize] ConvAsm1x1UV2: Not applicable
MIOpen(HIP): Info2 [GetWorkspaceSize] ConvAsm5x10u2v2f1: Not applicable
MIOpen(HIP): Info2 [GetWorkspaceSize] ConvAsm7x7c3h224w224k64u2v2p3q3f1: Not applicable
MIOpen(HIP): Info2 [GetWorkspaceSize] ConvAsm5x10u2v2b1: Not applicable
MIOpen(HIP): Info2 [GetWorkspaceSize] ConvOclDirectFwd11x11: Not applicable
MIOpen(HIP): Info2 [GetWorkspaceSize] ConvOclDirectFwdGen: Not applicable
MIOpen(HIP): Info2 [GetWorkspaceSize] ConvOclDirectFwd3x3: Not applicable
MIOpen(HIP): Info2 [GetWorkspaceSize] ConvOclDirectFwd: Not applicable
MIOpen(HIP): Info2 [SearchForAllSolutions] ConvHipImplicitGemmForwardV4R4Xdlops: Not applicable
MIOpen(HIP): Info2 [SearchForAllSolutions] ConvHipImplicitGemmV4R4GenXdlopsFwdFp32: Not applicable
MIOpen(HIP): Info2 [SearchForAllSolutions] ConvHipImplicitGemmV4R4GenFwdXdlops: Not applicable
MIOpen(HIP): Info2 [SearchForAllSolutions] ConvHipImplicitGemmBwdDataV1R1Xdlops: Not applicable
MIOpen(HIP): Info2 [SearchForAllSolutions] ConvHipImplicitGemmBwdDataV4R1Xdlops: Not applicable
MIOpen(HIP): Info [FindSolutionImpl] ConvHipImplicitGemmV4R1Fwd (db access disabled)
MIOpen(HIP): Info [EuristicInit] 16,128,8,2,4,4,4,4,4,4,8,2,16,1,2,128
MIOpen(HIP): Info [GetPerformanceConfigBase] 16,128,8,2,4,4,4,4,4,4,8,2,16,1,2,128
MIOpen(HIP): Info2 [SearchForAllSolutions] ConvHipImplicitGemmV4R1Fwd: Success.
MIOpen(HIP): Info2 [ForwardGetWorkSpaceSize] 0
MIOpen(HIP): Info [FindConvFwdAlgorithm] requestAlgoCount = 2, workspace = 0
MIOpen(HIP): Info2 [FindRecordUnsafe] Looking for key 256-28-28-1x1-1024-28-28-64-0x0-1x1-1x1-0-NCHW-FP32-F in file /home/kamil/vf/MLOpen-Devel/MIOpen/src/kernels/gfx906_60.HIP.2_7_0.ufdb.txt
MIOpen(HIP): Info2 [FindRecordUnsafe] Key match: 256-28-28-1x1-1024-28-28-64-0x0-1x1-1x1-0-NCHW-FP32-F
MIOpen(HIP): Info2 [FindRecordUnsafe] Contents found: miopenConvolutionFwdAlgoGEMM:gemm,3.3136,0,rocBlas,<unused>
MIOpen(HIP): Info2 [Measure] Db::FindRecord time: 0.257243 ms
MIOpen(HIP): Info [TryLoad] Find-db regenerating.
MIOpen(HIP): Info2 [CallGemmStridedBatched] gemm_desc: {isColMajor 0, transA 0, transB 0, m 1024, n 784, k 256, lda 256, ldb 784, ldc 784, batch_count 64, strideA 0, strideB 200704, strideC 802816, alpha 1, beta 0, dataType 1} 
MIOpen(HIP): Info2 [CallGemmStridedBatched] gemm_desc: {isColMajor 0, transA 0, transB 0, m 1024, n 784, k 256, lda 256, ldb 784, ldc 784, batch_count 64, strideA 0, strideB 200704, strideC 802816, alpha 1, beta 0, dataType 1} 
MIOpen(HIP): Info2 [dummy_memset] dummy gpu memset
MIOpen(HIP): Info [SetValues] 256-28-28-1x1-1024-28-28-64-0x0-1x1-1x1-0-NCHW-FP32-F, content inserted: miopenConvolutionFwdAlgoGEMM:gemm,3.28863,0,rocBlas,<unused>
MIOpen(HIP): Info2 [SearchForAllSolutions] ConvBinWinograd3x3U: Not applicable
MIOpen(HIP): Info [FindSolutionImpl] ConvBinWinogradRxSf3x2 (not searchable)
MIOpen(HIP): Info2 [SearchForAllSolutions] ConvBinWinogradRxSf3x2: Success.
MIOpen(HIP): Info [FindSolutionImpl] ConvBinWinogradRxSf2x3
MIOpen(HIP): Warning [FindSolutionImpl] Perf Db: load skipped: ConvBinWinogradRxSf2x3, enforce: SEARCH_DB_UPDATE(4), ALL(1)
MIOpen(HIP): Info [FindSolutionImpl] Starting search: ConvBinWinogradRxSf2x3, enforce: SEARCH_DB_UPDATE(4), ALL(1)
MIOpen(HIP): Info [Search] SEARCH PROBLEMS START
MIOpen(HIP): Info [GetPerformanceConfig] 60
MIOpen(HIP): Warning [GenericSearch] ConvBinWinogradRxSf2x3: Searching the best solution among 60...
MIOpen(HIP): Info2 [GenericSearch] #0/0/60 1
MIOpen(HIP): Info2 [AmdgcnAssembleQuiet] /opt/rocm/llvm/bin/clang  -x assembler -target amdgcn--amdhsa -mcpu=gfx900 /tmp/8bc1-cf47-137d-b821 -o /dev/null 2>&1
MIOpen(HIP): Info2 [AmdgcnAssemble] ' -x assembler -target amdgcn--amdhsa  -Wa,-defsym,ROCM_METADATA_VERSION=5 -mcpu=gfx906 - -o /tmp/miopen-tmp-8682-021d-d137-64d8/amdgcn-asm-out-XXXXXX'
MIOpen(HIP): Info2 [Log] Kernel Conv_Winograd_v21_1_0_gfx9_fp32_stride1.s Compile Time, ms: 123.258
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 544.505 / 3.40282e+38 = 1.60016e-36
MIOpen(HIP): Info [GenericSearch] #0/0/60 543.554 < 3.40282e+38 1
MIOpen(HIP): Info2 [GenericSearch] #1/0/60 2
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 281.042 / 543.554 = 0.517046
MIOpen(HIP): Info [GenericSearch] #1/0/60 275.456 < 543.554 2
MIOpen(HIP): Warning [Monitor] 1/0/60 275.456, best within recent 2: 275.456 #1 2, ETA:249.106 sec.
MIOpen(HIP): Info2 [GenericSearch] #2/0/60 3
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 181.174 / 275.456 = 0.657725
MIOpen(HIP): Info [GenericSearch] #2/0/60 181.17 < 275.456 3
MIOpen(HIP): Info2 [GenericSearch] #3/0/60 4
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 135.852 / 181.17 = 0.74986
MIOpen(HIP): Info [GenericSearch] #3/0/60 135.849 < 181.17 4
MIOpen(HIP): Info2 [GenericSearch] #4/0/60 5
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 108.693 / 135.849 = 0.800104
MIOpen(HIP): Info [GenericSearch] #4/0/60 108.701 < 135.849 5
MIOpen(HIP): Info2 [GenericSearch] #5/0/60 6
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 90.586 / 108.701 = 0.833353
MIOpen(HIP): Info [GenericSearch] #5/0/60 90.5871 < 108.701 6
MIOpen(HIP): Info2 [GenericSearch] #6/0/60 7
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 77.833 / 90.5871 = 0.859206
MIOpen(HIP): Info [GenericSearch] #6/0/60 77.8362 < 90.5871 7
MIOpen(HIP): Info2 [GenericSearch] #7/0/60 8
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 67.9314 / 77.8362 = 0.872748
MIOpen(HIP): Info [GenericSearch] #7/0/60 67.9299 < 77.8362 8
MIOpen(HIP): Warning [Monitor] 7/0/60 67.9299, best within recent 6: 67.9299 #7 8, ETA:57.0982 sec.
MIOpen(HIP): Info2 [GenericSearch] #8/0/60 9
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 60.3993 / 67.9299 = 0.889141
MIOpen(HIP): Info [GenericSearch] #8/0/60 60.3973 < 67.9299 9
MIOpen(HIP): Info2 [GenericSearch] #9/0/60 10
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 54.3742 / 60.3973 = 0.900275
MIOpen(HIP): Info [GenericSearch] #9/0/60 54.3755 < 60.3973 10
MIOpen(HIP): Info2 [GenericSearch] #10/0/60 11
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 49.4367 / 54.3755 = 0.909174
MIOpen(HIP): Info [GenericSearch] #10/0/60 49.438 < 54.3755 11
MIOpen(HIP): Info2 [GenericSearch] #11/0/60 12
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 45.3275 / 49.438 = 0.916854
MIOpen(HIP): Info [GenericSearch] #11/0/60 45.3257 < 49.438 12
MIOpen(HIP): Info2 [GenericSearch] #12/0/60 13
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 41.8289 / 45.3257 = 0.922852
MIOpen(HIP): Info [GenericSearch] #12/0/60 43.5882 < 45.3257 13
MIOpen(HIP): Info2 [GenericSearch] #13/0/60 14
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 39.6808 / 43.5882 = 0.910355
MIOpen(HIP): Info [GenericSearch] #13/0/60 39.1021 < 43.5882 14
MIOpen(HIP): Info2 [GenericSearch] #14/0/60 15
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 36.3273 / 39.1021 = 0.929038
MIOpen(HIP): Info [GenericSearch] #14/0/60 36.3267 < 39.1021 15
MIOpen(HIP): Info2 [GenericSearch] #15/0/60 16
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 34.0648 / 36.3267 = 0.937732
MIOpen(HIP): Info [GenericSearch] #15/0/60 34.064 < 36.3267 16
MIOpen(HIP): Info2 [GenericSearch] #16/0/60 17
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 32.1393 / 34.064 = 0.943499
MIOpen(HIP): Info [GenericSearch] #16/0/60 32.1322 < 34.064 17
MIOpen(HIP): Info2 [GenericSearch] #17/0/60 18
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 30.4116 / 32.1322 = 0.946452
MIOpen(HIP): Info [GenericSearch] #17/0/60 30.4334 < 32.1322 18
MIOpen(HIP): Info2 [GenericSearch] #18/0/60 19
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 28.9644 / 30.4334 = 0.951732
MIOpen(HIP): Info [GenericSearch] #18/0/60 28.9576 < 30.4334 19
MIOpen(HIP): Info2 [GenericSearch] #19/0/60 20
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 27.5923 / 28.9576 = 0.952851
MIOpen(HIP): Info [GenericSearch] #19/0/60 27.5988 < 28.9576 20
MIOpen(HIP): Info2 [GenericSearch] #20/0/60 21
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 26.4016 / 27.5988 = 0.95662
MIOpen(HIP): Info [GenericSearch] #20/0/60 26.401 < 27.5988 21
MIOpen(HIP): Info2 [GenericSearch] #21/0/60 22
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 25.2052 / 26.401 = 0.95471
MIOpen(HIP): Info [GenericSearch] #21/0/60 25.209 < 26.401 22
MIOpen(HIP): Info2 [GenericSearch] #22/0/60 23
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 24.1158 / 25.209 = 0.956635
MIOpen(HIP): Info [GenericSearch] #22/0/60 24.1212 < 25.209 23
MIOpen(HIP): Info2 [GenericSearch] #23/0/60 24
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 23.147 / 24.1212 = 0.959612
MIOpen(HIP): Info [GenericSearch] #23/0/60 23.15 < 24.1212 24
MIOpen(HIP): Info2 [GenericSearch] #24/0/60 25
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 22.2464 / 23.15 = 0.960968
MIOpen(HIP): Info [GenericSearch] #24/0/60 22.2402 < 23.15 25
MIOpen(HIP): Warning [Monitor] 24/0/60 22.2402, best within recent 17: 22.2402 #24 25, ETA:15.8673 sec.
MIOpen(HIP): Info2 [GenericSearch] #25/0/60 26
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 21.4081 / 22.2402 = 0.962586
MIOpen(HIP): Info [GenericSearch] #25/0/60 21.5991 < 22.2402 26
MIOpen(HIP): Info2 [GenericSearch] #26/0/60 27
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 20.6437 / 21.5991 = 0.955763
MIOpen(HIP): Info [GenericSearch] #26/0/60 20.657 < 21.5991 27
MIOpen(HIP): Info2 [GenericSearch] #27/0/60 28
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 19.983 / 20.657 = 0.967374
MIOpen(HIP): Info [GenericSearch] #27/0/60 19.9398 < 20.657 28
MIOpen(HIP): Info2 [GenericSearch] #28/0/60 29
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 19.2093 / 19.9398 = 0.963365
MIOpen(HIP): Info [GenericSearch] #28/0/60 19.2076 < 19.9398 29
MIOpen(HIP): Info2 [GenericSearch] #29/0/60 30
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 18.5921 / 19.2076 = 0.967959
MIOpen(HIP): Info [GenericSearch] #29/0/60 18.5927 < 19.2076 30
MIOpen(HIP): Info2 [GenericSearch] #30/0/60 31
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 17.9773 / 18.5927 = 0.966896
MIOpen(HIP): Info [GenericSearch] #30/0/60 17.9785 < 18.5927 31
MIOpen(HIP): Info2 [GenericSearch] #31/0/60 32
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 17.4014 / 17.9785 = 0.967901
MIOpen(HIP): Info [GenericSearch] #31/0/60 17.4042 < 17.9785 32
MIOpen(HIP): Info2 [GenericSearch] #32/0/60 33
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 16.9254 / 17.4042 = 0.97249
MIOpen(HIP): Info [GenericSearch] #32/0/60 16.9251 < 17.4042 33
MIOpen(HIP): Info2 [GenericSearch] #33/0/60 34
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 16.4478 / 16.9251 = 0.971802
MIOpen(HIP): Info [GenericSearch] #33/0/60 16.4845 < 16.9251 34
MIOpen(HIP): Info2 [GenericSearch] #34/0/60 35
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 16.1621 / 16.4845 = 0.98044
MIOpen(HIP): Info [GenericSearch] #34/0/60 16.1712 < 16.4845 35
MIOpen(HIP): Info2 [GenericSearch] #35/0/60 36
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 15.7688 / 16.1712 = 0.975116
MIOpen(HIP): Info [GenericSearch] #35/0/60 15.7903 < 16.1712 36
MIOpen(HIP): Info2 [GenericSearch] #36/0/60 37
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 15.4627 / 15.7903 = 0.979254
MIOpen(HIP): Info [GenericSearch] #36/0/60 15.4689 < 15.7903 37
MIOpen(HIP): Info2 [GenericSearch] #37/0/60 38
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 15.1158 / 15.4689 = 0.977172
MIOpen(HIP): Info [GenericSearch] #37/0/60 15.1378 < 15.4689 38
MIOpen(HIP): Info2 [GenericSearch] #38/0/60 39
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 14.8059 / 15.1378 = 0.978075
MIOpen(HIP): Info [GenericSearch] #38/0/60 14.8265 < 15.1378 39
MIOpen(HIP): Info2 [GenericSearch] #39/0/60 40
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 14.5294 / 14.8265 = 0.979962
MIOpen(HIP): Info [GenericSearch] #39/0/60 14.5329 < 14.8265 40
MIOpen(HIP): Info2 [GenericSearch] #40/0/60 41
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 14.2569 / 14.5329 = 0.981013
MIOpen(HIP): Info [GenericSearch] #40/0/60 14.247 < 14.5329 41
MIOpen(HIP): Info2 [GenericSearch] #41/0/60 42
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 14.0045 / 14.247 = 0.982977
MIOpen(HIP): Info [GenericSearch] #41/0/60 14.0123 < 14.247 42
MIOpen(HIP): Info2 [GenericSearch] #42/0/60 43
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 13.74 / 14.0123 = 0.980566
MIOpen(HIP): Info [GenericSearch] #42/0/60 13.7395 < 14.0123 43
MIOpen(HIP): Info2 [GenericSearch] #43/0/60 44
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 13.5334 / 13.7395 = 0.984999
MIOpen(HIP): Info [GenericSearch] #43/0/60 13.5253 < 13.7395 44
MIOpen(HIP): Info2 [GenericSearch] #44/0/60 45
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 13.2425 / 13.5253 = 0.979095
MIOpen(HIP): Info [GenericSearch] #44/0/60 13.2438 < 13.5253 45
MIOpen(HIP): Info2 [GenericSearch] #45/0/60 46
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 13.0621 / 13.2438 = 0.986276
MIOpen(HIP): Info [GenericSearch] #45/0/60 13.0524 < 13.2438 46
MIOpen(HIP): Info2 [GenericSearch] #46/0/60 47
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 12.9377 / 13.0524 = 0.991213
MIOpen(HIP): Info [GenericSearch] #46/0/60 12.896 < 13.0524 47
MIOpen(HIP): Info2 [GenericSearch] #47/0/60 48
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 12.7009 / 12.896 = 0.984878
MIOpen(HIP): Info [GenericSearch] #47/0/60 12.7146 < 12.896 48
MIOpen(HIP): Info2 [GenericSearch] #48/0/60 49
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 12.5104 / 12.7146 = 0.98394
MIOpen(HIP): Info [GenericSearch] #48/0/60 12.5165 < 12.7146 49
MIOpen(HIP): Info2 [GenericSearch] #49/0/60 50
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 12.3614 / 12.5165 = 0.987611
MIOpen(HIP): Info [GenericSearch] #49/0/60 12.3206 < 12.5165 50
MIOpen(HIP): Info2 [GenericSearch] #50/0/60 51
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 12.0849 / 12.3206 = 0.980871
MIOpen(HIP): Info [GenericSearch] #50/0/60 12.1473 < 12.3206 51
MIOpen(HIP): Info2 [GenericSearch] #51/0/60 52
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 11.9825 / 12.1473 = 0.986438
MIOpen(HIP): Info [GenericSearch] #51/0/60 11.9964 < 12.1473 52
MIOpen(HIP): Info2 [GenericSearch] #52/0/60 53
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 11.831 / 11.9964 = 0.986217
MIOpen(HIP): Info [GenericSearch] #52/0/60 11.8212 < 11.9964 53
MIOpen(HIP): Info2 [GenericSearch] #53/0/60 54
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 11.7355 / 11.8212 = 0.992751
MIOpen(HIP): Info [GenericSearch] #53/0/60 11.7022 < 11.8212 54
MIOpen(HIP): Info2 [GenericSearch] #54/0/60 55
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 11.5669 / 11.7022 = 0.988438
MIOpen(HIP): Info [GenericSearch] #54/0/60 11.5571 < 11.7022 55
MIOpen(HIP): Info2 [GenericSearch] #55/0/60 56
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 11.3552 / 11.5571 = 0.982526
MIOpen(HIP): Info [GenericSearch] #55/0/60 11.3876 < 11.5571 56
MIOpen(HIP): Info2 [GenericSearch] #56/0/60 57
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 11.2781 / 11.3876 = 0.990378
MIOpen(HIP): Info [GenericSearch] #56/0/60 11.2709 < 11.3876 57
MIOpen(HIP): Info2 [GenericSearch] #57/0/60 58
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 11.0971 / 11.2709 = 0.984581
MIOpen(HIP): Info [GenericSearch] #57/0/60 11.1089 < 11.2709 58
MIOpen(HIP): Info2 [GenericSearch] #58/0/60 59
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 10.9547 / 11.1089 = 0.986124
MIOpen(HIP): Info [GenericSearch] #58/0/60 10.9499 < 11.1089 59
MIOpen(HIP): Info2 [GenericSearch] #59/0/60 60
MIOpen(HIP): Info2 [GenericSearch] Finding average for: 10.7792 / 10.9499 = 0.984409
MIOpen(HIP): Info [GenericSearch] #59/0/60 10.8274 < 10.9499 60
MIOpen(HIP): Warning [GenericSearch] Done: 60/0/60, best #59 10.8274 60
MIOpen(HIP): Warning [GenericSearch] ...Score: 1.00168 (default time 10.8456)
MIOpen(HIP): Info2 [Prepare] INSERT OR IGNORE INTO config( layout,data_type,direction,spatial_dim,in_channels,in_h,in_w,in_d,fil_h,fil_w,fil_d,out_channels,batchsize,pad_h,pad_w,pad_d,conv_stride_h,conv_stride_w,conv_stride_d,dilation_h,dilation_w,dilation_d,bias,group_count ) VALUES( ?,?,?,?,?,?,?,?,?,?,?,?,?,?,?,?,?,?,?,?,?,?,?,?);
MIOpen(HIP): Info2 [impl] [NCHW,FP32,F,2,256,28,28,1,1,1,1,1024,64,0,0,0,1,1,0,1,1,0,0,1]
MIOpen(HIP): Info2 [UpdateUnsafe] 0 rows updated
MIOpen(HIP): Info2 [Prepare] INSERT OR REPLACE INTO perf_db(config, solver, params, arch, num_cu) VALUES((SELECT id FROM config WHERE ( (layout = ? ) AND (data_type = ? ) AND (direction = ? ) AND (spatial_dim = ? ) AND (in_channels = ? ) AND (in_h = ? ) AND (in_w = ? ) AND (in_d = ? ) AND (fil_h = ? ) AND (fil_w = ? ) AND (fil_d = ? ) AND (out_channels = ? ) AND (batchsize = ? ) AND (pad_h = ? ) AND (pad_w = ? ) AND (pad_d = ? ) AND (conv_stride_h = ? ) AND (conv_stride_w = ? ) AND (conv_stride_d = ? ) AND (dilation_h = ? ) AND (dilation_w = ? ) AND (dilation_d = ? ) AND (bias = ? ) AND (group_count = ? ) ) ) , ? , ? , ? , ?);
MIOpen(HIP): Info2 [impl] [NCHW,FP32,F,2,256,28,28,1,1,1,1,1024,64,0,0,0,1,1,0,1,1,0,0,1,ConvBinWinogradRxSf2x3,60,gfx906,60]
MIOpen(HIP): Info [SetValues] , content inserted: ConvBinWinogradRxSf2x3:60
MIOpen(HIP): Info2 [Measure] Db::Update time: 13.5204 ms
MIOpen(HIP): Info2 [SearchForAllSolutions] ConvBinWinogradRxSf2x3: Success.
MIOpen(HIP): Info [FindSolutionImpl] ConvBinWinogradRxS (not searchable)
MIOpen(HIP): Info2 [SearchForAllSolutions] ConvBinWinogradRxS: Success.
MIOpen(HIP): Info2 [IsApplicable] Workspace required: 845414400, limit: 2000000000
MIOpen(HIP): Info2 [SearchForAllSolutions] ConvMPBidirectWinograd<3-3>: Not applicable
MIOpen(HIP): Info2 [IsApplicable] Workspace required: 615776256, limit: 2000000000
MIOpen(HIP): Info2 [SearchForAllSolutions] ConvMPBidirectWinograd<4-3>: Not applicable
.....

atamazov · 2020-09-02T21:39:29Z

[off-topic] The following construct can be used to collapse long logs on github pages:

<details><summary><b><i>DESCRIPTION</i></b></summary><hr>

MANY MANY LINES
***
<hr></details>

atamazov · 2020-09-02T23:23:52Z

Together with Kamil we have managed to find the error. Please expect the fix tomorrow morning.

daniellowell · 2020-09-10T15:43:45Z

Looks like this issue is not fixed.

daniellowell · 2020-09-10T15:44:39Z

@alexandraBara Please add additional details.
@shurale-nkn Please look into this as soon as possible.

alexandraBara · 2020-09-10T15:53:06Z

I am still seeing this issue in develop today, here is a log: looks like the same issue

LOG

Executing: sudo docker run --device='/dev/kfd' --device='/dev/dri' -w /home/miopenpdb -v /home/miopenpdb:/home/miopenpdb --user=root --group-add video --privileged=true --rm miopentuna bash -c "export MIOPEN_LOG_LEVEL=7 && export MIOPEN_FIND_ENFORCE=3 && export HIP_VISIBLE_DEVICES=1 && MIOpenDriver conv -V 0 -i 1 --forw 1 --pad_h 0 --out_channels 96 --fil_w 1 --dilation_w 1 --fil_h 1 --in_h 14 --conv_stride_w 1 --group_count 1 --in_channels 480 --in_w 14 --dilation_h 1 --conv_stride_h 1 --pad_w 0 --batchsize 128 --pad_mode default --mode conv --fil_d 1 --in_d 1 --spatial_dim 2 --conv_stride_d 1 --dilation_d 1 --pad_d 0 --trans_output_pad_d 0 2>&1 "
INFO - Setting job id 11098755 state to running
INFO - MIOpenDriver conv -V 0 -i 1 --forw 1 --pad_h 0 --out_channels 96 --fil_w 1 --dilation_w 1 --fil_h 1 --in_h 14 --conv_stride_w 1 --group_count 1 --in_channels 480 --in_w 14 --dilation_h 1 --conv_stride_h 1 --pad_w 0 --batchsize 128 --pad_mode default --mode conv --fil_d 1 --in_d 1 --spatial_dim 2 --conv_stride_d 1 --dilation_d 1 --pad_d 0 --trans_output_pad_d 0
INFO - MIOpen(HIP): Info [Handle] stream: 0x1b77550, device_id: 0
INFO - MIOpen(HIP): Info [ForwardGetWorkSpaceSize]
INFO - MIOpen(HIP): Info2 [HipCompilerVersionImpl] Read version information from HIP package...
INFO - MIOpen(HIP): Info [HipCompilerVersionImpl] 3.7.20315
INFO - MIOpen(HIP): Info [AmdRocmMetadataVersionDetect] ROCm MD version AMDHSA_COv3, MIOpen version 2.7.0.8175-13470ceb
INFO - MIOpen(HIP): Info2 [ValidateGcnAssemblerImpl] Running: '/opt/rocm/llvm/bin/clang --version'
INFO - MIOpen(HIP): Info2 [ValidateGcnAssemblerImpl] clang version 11.0.0 (/src/external/llvm-project/clang ee4e4ebbadcc8ea14ce99e34ed31ab31e94827ac)
INFO - MIOpen(HIP): Info2 [ValidateGcnAssemblerImpl] Target: x86_64-unknown-linux-gnu
INFO - MIOpen(HIP): Info2 [ValidateGcnAssemblerImpl] Thread model: posix
INFO - MIOpen(HIP): Info2 [ValidateGcnAssemblerImpl] InstalledDir: /opt/rocm/llvm/bin
INFO - MIOpen(HIP): Info2 [ValidateGcnAssemblerImpl]
INFO - MIOpen(HIP): Info2 [SQLiteBase] Initializing system database file /opt/rocm/miopen/share/miopen/db/miopen.db
INFO - MIOpen(HIP): Trace [Exec] 139695531477504:PRAGMA table_info(config);
INFO - MIOpen(HIP): Trace [Exec] 139695531477504:PRAGMA table_info(perf_db);
INFO - MIOpen(HIP): Info2 [SQLiteBase] Initializing user database file /home/miopenpdb/.config/miopen/miopen_1.0.0.udb
INFO - MIOpen(HIP): Trace [Exec] 139695531477504:SELECT name FROM sqlite_master WHERE type = 'table' AND (name = 'config');
INFO - MIOpen(HIP): Trace [Exec] 139695531477504:SELECT name FROM sqlite_master WHERE type = 'table' AND (name = 'perf_db');
INFO - MIOpen(HIP): Trace [SQLitePerfDb] Database created successfully
INFO - MIOpen(HIP): Trace [Exec] 139695531477504:PRAGMA table_info(config);
INFO - MIOpen(HIP): Trace [Exec] 139695531477504:PRAGMA table_info(perf_db);
INFO - MIOpen(HIP): Info2 [SearchForAllSolutions] ConvBinWinograd3x3U: Not applicable
INFO - MIOpen(HIP): Info [FindSolutionImpl] ConvBinWinogradRxSf3x2 (not searchable)
INFO - MIOpen(HIP): Info2 [SearchForAllSolutions] ConvBinWinogradRxSf3x2: Success.
INFO - MIOpen(HIP): Info [FindSolutionImpl] ConvBinWinogradRxSf2x3

INFO - MIOpen(HIP): Info2 [impl] [NCHW,FP32,F,2,480,14,14,1,1,1,1,96,128,0,0,0,1,1,0,1,1,0,0,1]
INFO - MIOpen(HIP): Info [SetValues] , content inserted: ConvActivAsm1x1U:2,16,2,64,2,1
INFO - MIOpen(HIP): Info [SetValues] , content inserted: ConvOclDirectFwd1x1:1,64,1,1,1,4,16,2048,0
INFO - MIOpen(HIP): Info [SetValues] , content inserted: ConvAsm1x1U:3,8,3,16,2,2,1,4
INFO - MIOpen(HIP): Info [SetValues] , content inserted: ConvBiasActivAsm1x1U:2,16,2,64,2,2,1,1
INFO - MIOpen(HIP): Info [GetValues] =ConvBinWinogradRxSf2x3:
INFO - MIOpen(HIP): Info2 [Measure] Db::Load time: 127.36 ms
INFO - MIOpen(HIP): Info [FindSolutionImpl] Perf Db: record not found for: ConvBinWinogradRxSf2x3
INFO - MIOpen(HIP): Info [FindSolutionImpl] Starting search: ConvBinWinogradRxSf2x3, enforce: SEARCH(3), ALL(1)

INFO - MIOpen(HIP): Error [FindSolutionImpl] Search failed for: ConvBinWinogradRxSf2x3: /root/dMIOpen/src/include/miopen/generic_search.hpp:408: GenericSearch: top_ocl_ptr == nullptr || bot_ocl_ptr == nullptr || wei_ocl_ptr == nullptr

alexandraBara assigned atamazov Aug 31, 2020

alexandraBara added the bug label Aug 31, 2020

daniellowell added the value_high label Aug 31, 2020

daniellowell assigned Kirpich30000 and shurale-nkn Aug 31, 2020

daniellowell added urgency_blocker and removed value_high labels Aug 31, 2020

atamazov assigned DrizztDoUrden Aug 31, 2020

shurale-nkn mentioned this issue Sep 3, 2020

[fix] Context configuration in GetWorkSpaceSize #415

Merged

shurale-nkn linked a pull request Sep 3, 2020 that will close this issue

[fix] Context configuration in GetWorkSpaceSize #415

Merged

daniellowell closed this as completed in #415 Sep 3, 2020

alexandraBara mentioned this issue Sep 3, 2020

Tuning set for 2.7 ROCm/MITuna#223

Closed

daniellowell reopened this Sep 10, 2020

alexandraBara closed this as completed Sep 11, 2020

This was referenced Sep 11, 2020

Solver generic_search fail: ConvHipImplicitGemmBwdDataV1R1Xdlops and ConvHipImplicitGemmForwardV4R4Xdlops #427

Closed

Re-ran find_db for 2.7 tuning #438

Merged

junliume mentioned this issue Nov 18, 2021

Investigating CI Instability #1147

Open

11 tasks

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

ConvBinWinogradRxSf2x3 bug #411

ConvBinWinogradRxSf2x3 bug #411

alexandraBara commented Aug 31, 2020 •

edited by atamazov

Loading

daniellowell commented Aug 31, 2020

atamazov commented Aug 31, 2020

DrizztDoUrden commented Aug 31, 2020

daniellowell commented Sep 1, 2020

DrizztDoUrden commented Sep 1, 2020

alexandraBara commented Sep 1, 2020

DrizztDoUrden commented Sep 1, 2020 •

edited

Loading

alexandraBara commented Sep 1, 2020 •

edited by atamazov

Loading

daniellowell commented Sep 2, 2020

shurale-nkn commented Sep 2, 2020 •

edited

Loading

atamazov commented Sep 2, 2020

atamazov commented Sep 2, 2020

daniellowell commented Sep 10, 2020

daniellowell commented Sep 10, 2020

alexandraBara commented Sep 10, 2020 •

edited by shurale-nkn

Loading

ConvBinWinogradRxSf2x3 bug #411

ConvBinWinogradRxSf2x3 bug #411

Comments

alexandraBara commented Aug 31, 2020 • edited by atamazov Loading

daniellowell commented Aug 31, 2020

atamazov commented Aug 31, 2020

DrizztDoUrden commented Aug 31, 2020

daniellowell commented Sep 1, 2020

DrizztDoUrden commented Sep 1, 2020

alexandraBara commented Sep 1, 2020

DrizztDoUrden commented Sep 1, 2020 • edited Loading

alexandraBara commented Sep 1, 2020 • edited by atamazov Loading

daniellowell commented Sep 2, 2020

shurale-nkn commented Sep 2, 2020 • edited Loading

atamazov commented Sep 2, 2020

atamazov commented Sep 2, 2020

daniellowell commented Sep 10, 2020

daniellowell commented Sep 10, 2020

alexandraBara commented Sep 10, 2020 • edited by shurale-nkn Loading

alexandraBara commented Aug 31, 2020 •

edited by atamazov

Loading

DrizztDoUrden commented Sep 1, 2020 •

edited

Loading

alexandraBara commented Sep 1, 2020 •

edited by atamazov

Loading

shurale-nkn commented Sep 2, 2020 •

edited

Loading

alexandraBara commented Sep 10, 2020 •

edited by shurale-nkn

Loading