MIOpen is generating Vega ISA for gfx803 with fp16/MP enabled

[Delaunay]

I had the error below when running resnet18 on pytorch.
This only happens for a specific batch size.

To reproduce the error one can execute the scripts provided
fp16util.py.txt
micro_benchmarking_pytorch.py.txt
with the following arguments
python micro_benchmarking_pytorch.py --network resnet18 --batch-size 64 --iterations 20 --fp16 1

<stdin>:216:3: error: instruction not supported on this GPU
  v_add_u32 v3, s8, v0    
  ^
<stdin>:230:5: error: instruction not supported on this GPU
    v_add_u32 v6, 0+MIO_BN_CHW, v3  
    ^
<stdin>:236:3: error: instruction not supported on this GPU
  v_add_co_u32 v8, s[4:5], s32, v8   
  ^
<stdin>:238:3: error: instruction not supported on this GPU
  v_addc_co_u32 v9, s[4:5], v4, v9, s[4:5]   
  ^
<stdin>:242:3: error: instruction not supported on this GPU
  v_add_co_u32 v6, s[4:5], s32, v6   
  ^
<stdin>:243:3: error: instruction not supported on this GPU
  v_addc_co_u32 v7, s[4:5], v10, v7, s[4:5]   
  ^
<stdin>:245:3: error: instruction not supported on this GPU
  global_load_ushort v4, v[6:7], off     
  ^
<stdin>:246:3: error: instruction not supported on this GPU
  global_load_ushort v6, v[8:9], off   
  ^
<stdin>:249:3: error: instruction not supported on this GPU
  v_add_u32 v11, 4, v5         
  ^
<stdin>:251:3: error: instruction not supported on this GPU
  v_add_u32 v5, 4, v5         
  ^
<stdin>:256:9: error: instruction not supported on this GPU
        v_add_u32 v3, 0+MIO_BN_CHW*2, v3       
        ^
<stdin>:357:3: error: instruction not supported on this GPU
  v_add_u32 v3, 4, v3   
  ^
<stdin>:402:3: error: instruction not supported on this GPU
  v_add_u32 v7, s8, v0   
  ^
<stdin>:404:3: error: instruction not supported on this GPU
  v_add_u32 v0, 4, v4   
  ^
<stdin>:413:3: error: instruction not supported on this GPU
  v_add_u32 v4, 2, v4    
  ^
<stdin>:420:3: error: instruction not supported on this GPU
  v_add_co_u32 v8, s[4:5], s24, v8    
  ^
<stdin>:422:3: error: instruction not supported on this GPU
  v_add_u32 v7, 0+MIO_BN_CHW, v7  
  ^
<stdin>:426:3: error: instruction not supported on this GPU
  v_addc_co_u32 v9, s[4:5], v10, v9, s[4:5]   
  ^
<stdin>:446:3: error: instruction not supported on this GPU
  global_store_short v[8:9], v0, off   
  ^
<stdin>:475:6: error: instruction not supported on this GPU
     global_store_dword v[4:5], v1, off  
     ^
<stdin>:478:6: error: instruction not supported on this GPU
     global_store_dword v[6:7], v3, off  
     ^
<stdin>:490:6: error: instruction not supported on this GPU
     global_load_dword v8, v[3:4], off    
     ^
<stdin>:492:6: error: instruction not supported on this GPU
     global_load_dword v9, v[5:6], off   
     ^
<stdin>:522:6: error: instruction not supported on this GPU
     global_store_dword v[3:4], v8, off      
     ^
<stdin>:524:6: error: instruction not supported on this GPU
     global_store_dword v[5:6], v2, off     
     ^
MIOpen(HIP): Warning [AmdgcnAssemble]  -x assembler -target amdgcn--amdhsa -mno-code-object-v3  -Wa,-defsym,MIOPEN_USE_FP16=0 -Wa,-defsym,MIOPEN_USE_FP32=0 -Wa,-defsym,MIOPEN_USE_FPMIX=1 -Wa,-defsym,MIO_SAVE_MEAN_VARIANCE=1 -Wa,-defsym,MIO_RUNNING_RESULT=1 -Wa,-defsym,MIO_BN_N=64 -Wa,-defsym,MIO_BN_C=128 -Wa,-defsym,MIO_BN_HW=784 -Wa,-defsym,MIO_BN_NHW=50176 -Wa,-defsym,MIO_BN_NHW_DIV=1065353383 -Wa,-defsym,MIO_BN_CHW=100352 -Wa,-defsym,MIO_BN_NCHW=6422528 -Wa,-defsym,MIO_BN_LDS_SIZE=832 -Wa,-defsym,MIO_BN_LDSGCN_SIZE=13 -Wa,-defsym,MIO_BN_VARIANT=3 -Wa,-defsym,MIO_BN_GRP0=832 -Wa,-defsym,MIO_BN_GRP1=1 -Wa,-defsym,MIO_BN_GRP2=1 -mcpu=gfx803 - -o /tmp/miopen-tmp-90cd-a69e-6095-dc78/amdgcn-asm-out-XXXXXX
MIOpen Error: /root/driver/MLOpen/src/ocl/gcn_asm_utils.cpp:228: Assembly error(1)
Traceback (most recent call last):
  File "/home/setepenre/dev/bin/convnets", line 11, in <module>
    load_entry_point('MixedPrecision==0.0.0', 'console_scripts', 'convnets')()
  File "/home/setepenre/dev/lib/python3.6/site-packages/MixedPrecision-0.0.0-py3.6.egg/MixedPrecision/pytorch/convnet.py", line 246, in main
  File "/home/setepenre/dev/lib/python3.6/site-packages/MixedPrecision-0.0.0-py3.6.egg/MixedPrecision/pytorch/convnet.py", line 237, in generic_main
  File "/home/setepenre/dev/lib/python3.6/site-packages/MixedPrecision-0.0.0-py3.6.egg/MixedPrecision/pytorch/convnet.py", line 108, in train
  File "/home/setepenre/rocm_pytorch/torch/nn/modules/module.py", line 494, in __call__
    result = self.forward(*input, **kwargs)
  File "/home/setepenre/rocm_pytorch/torch/nn/modules/container.py", line 97, in forward
    input = module(input)
  File "/home/setepenre/rocm_pytorch/torch/nn/modules/module.py", line 494, in __call__
    result = self.forward(*input, **kwargs)
  File "/home/setepenre/dev/lib/python3.6/site-packages/torchvision/models/resnet.py", line 156, in forward
    x = self.layer2(x)
  File "/home/setepenre/rocm_pytorch/torch/nn/modules/module.py", line 494, in __call__
    result = self.forward(*input, **kwargs)
  File "/home/setepenre/rocm_pytorch/torch/nn/modules/container.py", line 97, in forward
    input = module(input)
  File "/home/setepenre/rocm_pytorch/torch/nn/modules/module.py", line 494, in __call__
    result = self.forward(*input, **kwargs)
  File "/home/setepenre/dev/lib/python3.6/site-packages/torchvision/models/resnet.py", line 46, in forward
    out = self.bn1(out)
  File "/home/setepenre/rocm_pytorch/torch/nn/modules/module.py", line 494, in __call__
    result = self.forward(*input, **kwargs)
  File "/home/setepenre/rocm_pytorch/torch/nn/modules/batchnorm.py", line 83, in forward
    exponential_average_factor, self.eps)
  File "/home/setepenre/rocm_pytorch/torch/nn/functional.py", line 1691, in batch_norm
    training, momentum, eps, torch.backends.cudnn.enabled
RuntimeError: miopenStatusUnknownError

With a batch size of 256 this does not happen.

python micro_benchmarking_pytorch.py --network resnet18 --batch-size 256 --iterations 20 --fp16 1
INFO: running forward and backward for warmup.
INFO: running the benchmark..
OK: finished running benchmark..
--------------------SUMMARY--------------------------
Microbenchmark for network : resnet18
Mini batch size [img] : 256
Time per mini-batch : 1.5049312353134154
Throughput [img/sec] : 170.10744012279451

[atamazov]

@Delaunay Please tell us how important is this for you. This question is do you need a hotfix patch ASAP ot not.

[Delaunay]

Not important

[atamazov]

@Delaunay Thanks. If so, then we'll will provide the fix with the next release.

/cc @daniellowell @muralinr

[daniellowell]

Fixed in MIOpen 2.0.0