Passthru activation different in cuDNN and MIOpen

[mattsinc]

I have been working on fixing a bug with a benchmark (bwd_bypass from DNNMark) that is bizarrely using 0 vector registers. After doing some digging, I found the problem is that HCC/HIP is generating no code for the kernel. This comes down to the fact that the behavior of cuDNN and MIOpen for the passthru activation is different -- in cuDNN an explicit copy from the input array to output array will happen, but in MIOpen the copy does not happen. My conversation with the DNNMark developers that documents all of this in more detail is here: shidong-ai/DNNMark#20.

I was wondering if this is a known issue with MIOpen? If so, is there a different setting that should be used to achieve the same effect? Or is this a problem with the activation (passthru) code that needs to be fixed?

Thanks,
Matt

[pfultz2]

It looks like in the test we need to zero the input data, which will catch this issue.

[daniellowell]

This is addressed in the next release.

[mattsinc]

Thanks for getting back to me. @daniellowell, thanks. Is there something I need to do in the meantime as a stopgap though? @pfultz2, it sounds like you are saying that the DNNMark test is not zeroing the input data, and that's why this is happening.

Matt

[pfultz2]

it sounds like you are saying that the DNNMark test is not zeroing the input data, and that's why this is happening.

No, I am saying that miopen tests are not zeroing the data, which is why we didnt catch the error.

[mattsinc]

I see, thanks.

Matt

[daniellowell]

Fixed in 1.4.0