NVFUSER_DUMP=fusion_ir dumps Kernel IR

[naoyam]

NVFUSER_DUMP=fusion_ir seems currently broken as it dumps a Kernel IR. For example, NVFUSER_DUMP=fusion_ir ./bin/nvfuser_tests --gtest_filter='*FusionBNRepro_CUDA' results in like:

Running main() from /home/nmaruyama/nvfuser/debug5/third_party/googletest/googletest/src/gtest_main.cc
Note: Google Test filter = *FusionBNRepro_CUDA
[==========] Running 1 test from 1 test suite.
[----------] Global test environment set-up.
[----------] 1 test from NVFuserTest
[ RUN      ] NVFuserTest.FusionBNRepro_CUDA

%Kernel { (T0_g_float[iS314{( (( (( getMetaData(T0) )).logical_size ))[1] )}, iS153{( ceilDiv(( ceilDiv(( (( (( getMetaData(T0) )).logical_size ))[0] ), 4) ), 4) )}, iS160{( ceilDiv(( ( (( (( getMetaData(T0) )).logical_size ))[2] ) * ( (( (( getMetaData(T0) )).logical_size ))[3] ) ), 2) )}, iS150{4}, iS155{1}, iS157{1}, iS152{4}, iS159{2}], T1_g_float[iS364{( (( (( getMetaData(T0) )).logical_size ))[1] )}], T2_g_float[iS371{( (( (( getMetaData(T0) )).logical_size ))[1] )}], T3_g_float[iS311{( (( (( getMetaData(T0) )).logical_size ))[1] )}], T4_g_float[iS331{( (( (( getMetaData(T0) )).logical_size ))[1] )}]) -> (T10_g_float[iblockIdx.x330{( (( (( getMetaData(T0) )).logical_size ))[1] )}] ca_pos( 1 ) produce_pos( 1 ), T14_g_float[iblockIdx.x351{( (( (( getMetaData(T0) )).logical_size ))[1] )}] ca_pos( 1 ) produce_pos( 1 ), T25_g_float[iblockIdx.x378{( (( (( getMetaData(T0) )).logical_size ))[1] )}, ithreadIdx.z252{( ceilDiv(( ceilDiv(( (( (( getMetaData(T0) )).logical_size ))[0] ), 4) ), 4) )}, ithreadIdx.x259{( ceilDiv(( ( (( (( getMetaData(T0) )).logical_size ))[2] ) * ( (( (( getMetaData(T0) )).logical_size ))[3] ) ), 2) )}_p, iS249{4}, iS254{1}, iUS256{1}, iUR251{4}, iUR258{2}] ca_pos( 6 ) produce_pos( 6 ), T37_g_float[iblockIdx.x383{( (( (( getMetaData(T0) )).logical_size ))[1] )}] ca_pos( 1 ) produce_pos( 1 ), T39_g_float[iblockIdx.x385{( (( (( getMetaData(T0) )).logical_size ))[1] )}] ca_pos( 1 ) produce_pos( 1 )) :
NVFUSER_DEFINE_MAGIC_ZERO;
d29 = ALLOCATE(buffer=d29, mem_type=register, size=1, zero_init=false, resets_to_zero=false)
d29 = double(1) - double(0.10000000149011612);
s162 = ALLOCATE(buffer=s162, mem_type=register, size=1, zero_init=false, resets_to_zero=false)
s162 = getMetaData(T0_g_float[iS314{( (( (( getMetaData(T0) )).logical_size ))[1] )}, iS153{( ceilDiv(( ceilDiv(( (( (( getMetaData(T0) )).logical_size ))[0] ), 4) ), 4) )}, iS160{( ceilDiv(( ( (( (( getMetaData(T0) )).logical_size ))[2] ) * ( (( (( getMetaData(T0) )).logical_size ))[3] ) ), 2) )}, iS150{4}, iS155{1}, iS157{1}, iS152{4}, iS159{2}])
i3036 = ALLOCATE(buffer=i3036, mem_type=register, size=1, zero_init=false, resets_to_zero=false)
i3036 = i2992 + i3035;
d40 = ALLOCATE(buffer=d40, mem_type=register, size=1, zero_init=false, resets_to_zero=false)
d40 = (double)(1);
...

I suspect this is another side effect of #3468.