`ops.reshape` errors with `!fusion->hasDynamicTransform()`

[kevinstephano]

Error:

Traceback (most recent call last):
  File "/workspace/simple_dl_models/repro.py", line 22, in <module>
    out = fd.execute(inputs)
  File "/opt/pytorch/nvfuser/nvfuser/__init__.py", line 76, in execute
    result = self._execute(inputs, override_user_schedule)
RuntimeError: !fusion->hasDynamicTransform() INTERNAL ASSERT FAILED at "/opt/pytorch/nvfuser/csrc/kernel_cache.cpp":619, please report a bug to PyTorch. Fusion must be concretized before constructing FusionKernelRuntime

This is a pattern we will find in the GroupNorm operation where the input tensors is reshaped. I am not quite sure what the error indicates since this pattern will have a dynamic transform.

This repro is looking into a future that uses a branch: add_new_reshape. This branch has some new operations not found in TOT in the Python API.

Repro:

import torch
from nvfuser import FusionDefinition, DataType

inputs = [
    torch.randn(256, 128, 28, 28, device='cuda'),
    32 
]

def func(fd: FusionDefinition) :
    T0 = fd.from_pytorch(inputs[0])
    S0 = fd.define_scalar(dtype=DataType.Int)
    V0 = T0.shape()
    T1 = fd.ops.reshape(T0, [V0[0], V0[1] / S0, S0, V0[2], V0[3]])
    var, mean = fd.ops.var_mean(T1, axes=[2, 3, 4], correction=0, keepdim=True)
    fd.add_output(var)
    fd.add_output(mean)

with FusionDefinition() as fd:
    func(fd)

for _ in range(5):
    out = fd.execute(inputs)

[jacobhinkle]

Thanks for reporting. I am looking into this. The issue has to do with how we traverse during concretization to concretize tensors downstream of a dynamic op like the reshape here. We are properly concretizing the variance and mean in the Welford op, but not the (unused) count output. I am experimenting with fixes now.

[kevinstephano]

This is the full example I am interested in:

import torch
from nvfuser import FusionDefinition, DataType

inputs = [
    torch.randn(256, 128, 28, 28, device='cuda'),
    torch.randn(1, 128, 1, 1, device='cuda'),
    torch.randn(1, 128, 1, 1, device='cuda'),
    32 
]

def func(fd: FusionDefinition) :
    T0 = fd.from_pytorch(inputs[0])
    weight = fd.from_pytorch(inputs[1])
    bias = fd.from_pytorch(inputs[2])
    S0 = fd.define_scalar(dtype=DataType.Int)
    S1 = fd.define_scalar(1.0e-5)
    V0 = T0.shape()
    T1 = fd.ops.reshape(T0, [V0[0], V0[1] / S0, S0, V0[2], V0[3]])
    var, mean = fd.ops.var_mean(T1, axes=[2, 3, 4], correction=0, keepdim=True)
    T2 = (T1 - mean) / fd.ops.sqrt(var + S1)
    T3 = fd.ops.reshape(T2, shape=V0)
    T4 = T3 * weight + bias 
    fd.add_output(T4)

with FusionDefinition() as fd:
    func(fd)

for _ in range(5):
    out = fd.execute(inputs)

[kevinstephano]

For completeness, here is the NHWC version.

import torch
from nvfuser import FusionDefinition, DataType

inputs = [
    torch.randn(256, 28, 28, 128, device='cuda'),
    torch.randn(1, 1, 1, 128, device='cuda'),
    torch.randn(1, 1, 1, 128, device='cuda'),
    32 
]

def func(fd: FusionDefinition) :
    T0 = fd.from_pytorch(inputs[0])
    weight = fd.from_pytorch(inputs[1])
    bias = fd.from_pytorch(inputs[2])
    S0 = fd.define_scalar(dtype=DataType.Int)
    S1 = fd.define_scalar(1.0e-5)
    V0 = T0.shape()
    T1 = fd.ops.reshape(T0, [V0[0], V0[1], V0[2], V0[3] / S0, S0])
    var, mean = fd.ops.var_mean(T1, axes=[1, 2, 4], correction=0, keepdim=True)
    T2 = (T1 - mean) / fd.ops.sqrt(var + S1)
    T3 = fd.ops.reshape(T2, shape=V0)
    T4 = T3 * weight + bias
    fd.add_output(T4)

with FusionDefinition() as fd:
    func(fd)

print(fd)

for _ in range(5):
    out = fd.execute(inputs)

[kevinstephano]

The real examples are not working. I might guess it is something segmentation related?

import torch                                                                                                                                                                                                                                                                               
from nvfuser import FusionDefinition, DataType                                                                                                                                                                                                                                             
                                                                                                                                                                                                                                                                                           
def nvfuser_fusion_id0(fd : FusionDefinition) -> None :                                                                                                                                                                                                                                    
    T0 = fd.define_tensor(shape=[-1, -1, -1, -1], contiguity=[True, True, True, True], dtype=DataType.Float, is_cpu=False)                                                                                                                                                                 
    T1 = fd.define_tensor(shape=[1, -1, 1, 1], contiguity=[None, True, None, None], dtype=DataType.Float, is_cpu=False)                                                                                                                                                                    
    T2 = fd.define_tensor(shape=[1, -1, 1, 1], contiguity=[None, True, None, None], dtype=DataType.Float, is_cpu=False)                                                                                                                                                                    
    S3 = fd.define_scalar(None, dtype=DataType.Int)                                                                                                                                                                                                                                        
    S4 = fd.define_scalar(1.00000e-05, dtype=DataType.Double)                                                                                                                                                                                                                              
    V5 = fd.ops.shape(T0)                                                                                                                                                                                                                                                                  
    S6 = fd.ops.at(V5, index=0)                                                                                                                                                                                                                                                            
    S7 = fd.ops.at(V5, index=1)                                                                                                                                                                                                                                                            
    S8 = fd.ops.div(S7, S3)                                          
    S9 = fd.ops.at(V5, index=2)                                                                                                              
    S10 = fd.ops.at(V5, index=3)                                                                                                             
    V11 = fd.define_vector([S6, S8, S3, S9, S10], dtype=DataType.Int)                                                                        
    T12 = fd.ops.reshape(T0, shape=V11)                                                                                                      
    T13, T14 = fd.ops.var_mean(T12, axes=[2, 3, 4], correction=0, keepdim=True)                
    T15 = fd.ops.sub(T12, T14)                                                                                                               
    T16 = fd.ops.add(T13, S4)                                                                                                                
    T17 = fd.ops.sqrt(T16)                                                                                                                   
    T18 = fd.ops.div(T15, T17)                                                                                                               
    T19 = fd.ops.reshape(T18, shape=V5)                                                                                                      
    T20 = fd.ops.mul(T19, T1)                                       
    T21 = fd.ops.add(T20, T2)                                                                                                                
    fd.add_output(T21)                                                                                                                       
                                                                                                                                             
with FusionDefinition() as fd:                                                                                                               
    nvfuser_fusion_id0(fd)                                                                                                                   
                                                                      
inputs = [                                                                                                                                   
    torch.randn((256, 128, 28, 28), dtype=torch.float32, device='cuda:0').as_strided((256, 128, 28, 28), (100352, 784, 28, 1)),
    torch.randn((1, 128, 1, 1), dtype=torch.float32, device='cuda:0').as_strided((1, 128, 1, 1), (128, 1, 1, 1)),
    torch.randn((1, 128, 1, 1), dtype=torch.float32, device='cuda:0').as_strided((1, 128, 1, 1), (128, 1, 1, 1)),
    32,                                                                                                                                      
]                                                                                                                                            
fd.execute(inputs)                                           
                                                                                                                                       
Traceback (most recent call last):                                                                                                           
  File "/opt/pytorch/nvfuser/nvfuser/__init__.py", line 109, in execute                    
    result = self._execute(inputs, override_user_schedule, device=device)                                                                                                                                                                                                                  
RuntimeError: inferred_val.has_value() INTERNAL ASSERT FAILED at "/opt/pytorch/nvfuser/csrc/executor.cpp":475, please report a bug to PyTorch. Could not launch kernel as program could not infer i1(i1) for the buffer T16_g[ iS133{i0}, iS82{i1}, iS134{i2}, iS135{i3} ]                 
Traceback (most recent call last):               
  File "/workspace/test.py", line 29, in <module>                                                                                                                                                                                                                                          
    out = fd.execute(inputs)                                                                                                                 
  File "/opt/pytorch/nvfuser/nvfuser/__init__.py", line 109, in execute                                                                                                                                                                                                                    
    result = self._execute(inputs, override_user_schedule, device=device)                                                                                                                                                                                                                  
RuntimeError: inferred_val.has_value() INTERNAL ASSERT FAILED at "/opt/pytorch/nvfuser/csrc/executor.cpp":475, please report a bug to PyTorch. Could not launch kernel as program could not infer i1(i1) for the buffer T16_g[ iS133{i0}, iS82{i1}, iS134{i2}, iS135{i3} ]                 
[E thread_pool.cpp:110] Exception in thread pool task: val.has_value() INTERNAL ASSERT FAILED at "/opt/pytorch/nvfuser/csrc/executor.cpp":1041, please report a bug to PyTorch. Tried to evaluate the extent, ( fmax(( fmax(( T11.size[0] * i1 ), ( T11.size[0] * 4 )) ), T1.size[1]) ) for
 the ptype: blockIdx.y to set launch bounds but could not.                                                                                   
Exception raised from computeLaunchParams at /opt/pytorch/nvfuser/csrc/executor.cpp:1041 (most recent call first):                                                                                                                                                                         
frame #0: c10::Error::Error(c10::SourceLocation, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >) + 0xae (0x7fac943b011e in /opt/pytorch/pytorch/torch/lib/libc10.so)                                                                                      
frame #1: c10::detail::torchCheckFail(char const*, char const*, unsigned int, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&) + 0xf3 (0x7fac94365bab in /opt/pytorch/pytorch/torch/lib/libc10.so)                                                  
frame #2: c10::detail::torchInternalAssertFail(char const*, char const*, unsigned int, char const*, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&) + 0x53 (0x7fac943ada23 in /opt/pytorch/pytorch/torch/lib/libc10.so)                            
frame #3: nvfuser::FusionExecutor::computeLaunchParams(nvfuser::LaunchParams const&, nvfuser::ExpressionEvaluator&, long) + 0x9b3 (0x7fac137de6e3 in /opt/pytorch/pytorch/torch/lib/libnvfuser_codegen.so)                                                                                 
frame #4: nvfuser::FusionExecutor::compileFusion(nvfuser::Fusion*, nvfuser::KernelArgumentHolder const&, nvfuser::LaunchParams const&, nvfuser::CompileParams) + 0x1087 (0x7fac137ebbc7 in /opt/pytorch/pytorch/torch/lib/libnvfuser_codegen.so)                                           
frame #5: nvfuser::FusionKernelRuntime::compileKernel(nvfuser::KernelArgumentHolder const&, nvfuser::SegmentedGroup*) + 0x199 (0x7fac13933049 in /opt/pytorch/pytorch/torch/lib/libnvfuser_codegen.so)                                                                                     
frame #6: <unknown function> + 0x42a206 (0x7fac13933206 in /opt/pytorch/pytorch/torch/lib/libnvfuser_codegen.so)        
frame #7: c10::ThreadPool::main_loop(unsigned long) + 0x2bd (0x7fac943a088d in /opt/pytorch/pytorch/torch/lib/libc10.so)
frame #8: <unknown function> + 0xdc2b3 (0x7fac64ab22b3 in /usr/lib/x86_64-linux-gnu/libstdc++.so.6)
frame #9: <unknown function> + 0x94b43 (0x7facab510b43 in /usr/lib/x86_64-linux-gnu/libc.so.6)                                               
frame #10: <unknown function> + 0x126a00 (0x7facab5a2a00 in /usr/lib/x86_64-linux-gnu/libc.so.6) 

[kevinstephano]

If this helps, here is a container with everything merged together:

gitlab-master.nvidia.com/kstephano/pytorch:test_fix418

[jacobhinkle]

If this helps, here is a container with everything merged together:

gitlab-master.nvidia.com/kstephano/pytorch:test_fix418

Thanks! I will look at it. I suspect this might be fixed by adding back the change I had that added output extent scalars to the list of ViewOp inputs. I will post here the results.