1x1 kernel hangs indefinitely in miopenFindConvolutionBackwardWeightsAlgorithm

[patflick]

The following configuration with a 1x1 kernel fails the forward verification and then seems to hang indefinitely inside the miopenFindConvolutionBackwardWeightsAlgorithm function.

$ ./MIOpenDriver conv -H 14 -W 14 -P 1 -k 512 -c 256 -n 128 -p 0 -q 0 -u 2 -v 2 -x 1 -y 1 -t 1
MIOpenDriver: conv -H 14 -W 14 -P 1 -k 512 -c 256 -n 128 -p 0 -q 0 -u 2 -v 2 -x 1 -y 1 -t 1
MIOpen Forward Conv. Algorithm: 1
GPU Kernel Time Forward Conv. Elapsed: 3.433880 ms
Forward Convolution Verifies on CPU and GPU
MIOpen Backward Data Conv. Algorithm: 0
GPU Kernel Time Backward Data Conv. Elapsed: 104.385124 ms
^C

(waited a good 10 minutes). Interrupting in gdb shows that this hangs inside the miopenFindConvolutionBackwardWeightsAlgorithm function.

[patflick]

Sorry, it fails the forward verification only without -t 1:

$ ./MIOpenDriver conv -H 14 -W 14 -P 1 -k 512 -c 256 -n 128 -p 0 -q 0 -u 2 -v 2 -x 1 -y 1     
MIOpenDriver: conv -H 14 -W 14 -P 1 -k 512 -c 256 -n 128 -p 0 -q 0 -u 2 -v 2 -x 1 -y 1
Forward Convolution Failed: 0.206861
^C

[dagamayank]

Closing this for now as the common problem is missing synchronization without -t 1 flag. This is the same issue as #12

[patflick]

Even with the -t 1 flag, it still hangs indefinitely. The failed verification was just an additional thing.

[dagamayank]

@patflick I narrowed it down to a bug in the Direct algorithm. If you disable the Direct algorithm with this env. variable, the test passes.
MIOPEN_DEBUG_CONV_DIRECT=0

MIOPEN_DEBUG_CONV_DIRECT=0 ./bin/MIOpenDriver conv -H 14 -W 14 -P 1 -k 512 -c 256 -n 128 -p 0 -q 0 -u 2 -v 2 -x 1 -y 1 
MIOpenDriver: conv -H 14 -W 14 -P 1 -k 512 -c 256 -n 128 -p 0 -q 0 -u 2 -v 2 -x 1 -y 1
Forward Convolution Verifies on CPU and GPU
Backward Convolution Data Verifies on CPU and GPU
Backward Convolution Weights Verifies on CPU and GPU

But the downside is that the Direct algorithm is disabled for fwd as well resulting in slower fwd time.

Let us look into fixing it. Thanks for reporting.
cc\ @alyashev

[dagamayank]

Disabling Direct algorithm for all stride=2 cases fixes this issue for now 756f73f. We will update and optimize the Direct algorithm in subsequent releases.