workSpaceSize check assertion fails in ConvolutionBackwardWeights() - DEBUG builds only

[syoyo]

System

• Ubuntu 18.04
• ROCm 2.0
• MIOpen master branch + Debug build
• VEGA56(gfx900)

I have encountered assertion failure when trying to run keithito's tacotron.

https://github.com/keithito/tacotron

How to reproduce

Download LJ-Speech data and setup as described in Tacotron's README.

Set MIOPEN_ENABLE_LOGGING=1 and MIOPEN_LOG_LEVEL=6

Then, run train.py

$ gdb python
$ set args train.py
...
python: /home/syoyo/work/MIOpen/src/ocl/convolutionocl.cpp:4111: auto miopen::ConvolutionDescriptor::ConvolutionBackwardWeights(miopen::Handle &, const void *, const miopen::TensorDescriptor &, ConstData_t, const miopen::TensorDescriptor &, ConstData_t, miopenConvBwdWeightsAlgorithm_t, const void *, const miopen::TensorDescriptor &, Data_t, Data_t, size_t)::(anonymous class)::operator()(type-parameter-0-0) const: Assertion `(workSpace != nullptr) && (workSpaceSize >= workSpaceSizeDirect)' failed.
Thread 47 "python" received signal SIGABRT, Aborted.
...
$1 = (Data_t &) @0x7fff067f8f98: 0x60f235400
(gdb) p workSpaceSize
$2 = (size_t &) @0x7fff067f8fa0: 1048576
(gdb) p workSpaceSizeDirect
$3 = 4194304

Full log: issue_72_full_log.txt

workSpaceSizeDirect is 4x larger(due to sizeof(float)?) than workSpaceSize.

I have modified MIOpen source code to print the value of BackwardWeightsGetWorkSpaceSizeDirect

[atamazov]

Explanation:

• More than one Solution is suitable for Direct WrW algo
• One (ConvOclBwdWrW2) wants 4MB workspace, another one (ConvOclBwdWrW53) needs only 1MB
• The second Solution (ConvOclBwdWrW53) is faster and then it is selected
• The application provides 1MB workspace, just as needed for ConvOclBwdWrW53.
• The assertion checks for max of all available Direct Solutions, which is 4MB, and assertion fails.

driver command to reproduce:

./bin/MIOpenDriver conv -n 32 -c 128 -H 1 -W 111 -k 128 -c 128 -y 1 -x 2 -p 0 -q 0 -V 0 -t 1 -w 1

I'll provide a fix soon.

[atamazov]

@syoyo The patch which shall resolve the issue resides here. Please close the issue if it fixes the problem. Thanks.

[syoyo]

@atamazov At least the patch now avoids assertion failure, but an application now goes into infinite loop of compiling MIOpen conv kernels(by glimpsing MIOpen logs, it looks it infinitely try to search conv kernel)

[atamazov]

@syoyo Theoretically, this could be not an infinite loop but auto-tuning (exhaustive search) process. Please check if autotuning is enabled in the environment.

Anyway, I am unable to reproduce the "loop" problem with this config (./bin/MIOpenDriver conv -n 32 -c 128 -H 1 -W 111 -k 128 -c 128 -y 1 -x 2 -p 0 -q 0 -V 0 -t 1 -w 1). Please attach log. Thanks.

[syoyo]

I am unable to reproduce the "loop" problem with this config (./bin/MIOpenDriver conv -n 32 -c 128 -H 1 -W 111 -k 128 -c 128 -y 1 -x 2 -p 0 -q 0 -V 0 -t 1 -w 1).

Ah, I meant "loop" occurs when I run tacotron app, not MIOpenDriver solely.
I'll also look into autotuning document.

Currently I'm running a learning task for #70 regression. I'll attach log of #72 loop issue after that(it may take half or one day)

[atamazov]

@syoyo I highly recommend opening another issue for the loop problem. Let's keep this issue dedicated to workspace size check assertion. Thanks.

[syoyo]

@atamazov I understand. I'll open another issue once ready to submit log file

I highly recommend opening another issue for the loop problem

So this issue is ready to close once PR #73 has been merged(Inherently its @atamazov's patch)

[atamazov]

I am not sure that #73 will be merged. AFAIK we are going to release 1.7.1 instead, which shall contain a set of fixes & improvements, including fix for #70 and #72.

[atamazov]

Note that the issue only occurs with Debug builds of MIOpen (i.e. those which have assertions enabled).

[syoyo]

Fix has been included in 1.7.1: a478ac8