`FusionMatmulSchedulerEpilogueBias_CUDA` is failing with misaligned address on RTX4090

[zasdfgbnm]

❯ ./build/nvfuser_tests --gtest_filter=*./build/nvfuser_tests --gtest_filter=*NVFuserTest.FusionMatmulSchedulerEpilogueBias_CUDA*
Note: Google Test filter = *NVFuserTest.FusionMatmulSchedulerEpilogueBias_CUDA*
[==========] Running 1 test from 1 test suite.
[----------] Global test environment set-up.
[----------] 1 test from NVFuserTest
[ RUN      ] NVFuserTest.FusionMatmulSchedulerEpilogueBias_CUDA
unknown file: Failure
C++ exception with description "CUDA error: misaligned address
CUDA kernel errors might be asynchronously reported at some other API call, so the stacktrace below might be incorrect.
For debugging consider passing CUDA_LAUNCH_BLOCKING=1.
Compile with `TORCH_USE_CUDA_DSA` to enable device-side assertions.

Exception raised from c10_cuda_check_implementation at ../c10/cuda/CUDAException.cpp:44 (most recent call first):
frame #0: c10::Error::Error(c10::SourceLocation, std::string) + 0x57 (0x7f2f6f3a8477 in /home/gaoxiang/.local/lib/python3.11/site-packages/torch/lib/libc10.so)
frame #1: c10::detail::torchCheckFail(char const*, char const*, unsigned int, std::string const&) + 0x64 (0x7f2f6f36489b in /home/gaoxiang/.local/lib/python3.11/site-packages/torch/lib/libc10.so)
frame #2: c10::cuda::c10_cuda_check_implementation(int, char const*, char const*, int, bool) + 0x118 (0x7f2f6d538298 in /home/gaoxiang/.local/lib/python3.11/site-packages/torch/lib/libc10_cuda.so)
frame #3: void at::native::gpu_kernel_impl<at::native::BinaryFunctor<float, float, bool, at::native::(anonymous namespace)::CompareEqFunctor<float> > >(at::TensorIteratorBase&, at::native::BinaryFunctor<float, float, bool, at::native::(anonymous namespace)::CompareEqFunctor<float> > const&) + 0xc9f (0x7f2f0bc3187f in /home/gaoxiang/.local/lib/python3.11/site-packages/torch/lib/libtorch_cuda.so)
frame #4: void at::native::gpu_kernel<at::native::BinaryFunctor<float, float, bool, at::native::(anonymous namespace)::CompareEqFunctor<float> > >(at::TensorIteratorBase&, at::native::BinaryFunctor<float, float, bool, at::native::(anonymous namespace)::CompareEqFunctor<float> > const&) + 0x32b (0x7f2f0bc3209b in /home/gaoxiang/.local/lib/python3.11/site-packages/torch/lib/libtorch_cuda.so)
frame #5: void at::native::opmath_symmetric_gpu_kernel_with_scalars<float, bool, at::native::(anonymous namespace)::CompareEqFunctor<float> >(at::TensorIteratorBase&, at::native::(anonymous namespace)::CompareEqFunctor<float> const&) + 0x105 (0x7f2f0bc479b5 in /home/gaoxiang/.local/lib/python3.11/site-packages/torch/lib/libtorch_cuda.so)
frame #6: at::native::compare_eq_ne_kernel(at::TensorIteratorBase&, at::native::(anonymous namespace)::EqOpType) + 0x1a9 (0x7f2f0bc24979 in /home/gaoxiang/.local/lib/python3.11/site-packages/torch/lib/libtorch_cuda.so)
frame #7: <unknown function> + 0x303ee53 (0x7f2f0d43ee53 in /home/gaoxiang/.local/lib/python3.11/site-packages/torch/lib/libtorch_cuda.so)
frame #8: <unknown function> + 0x303eef0 (0x7f2f0d43eef0 in /home/gaoxiang/.local/lib/python3.11/site-packages/torch/lib/libtorch_cuda.so)
frame #9: at::_ops::eq_Tensor::call(at::Tensor const&, at::Tensor const&) + 0x161 (0x7f2f51bd5951 in /home/gaoxiang/.local/lib/python3.11/site-packages/torch/lib/libtorch_cpu.so)
frame #10: at::native::isclose(at::Tensor const&, at::Tensor const&, double, double, bool) + 0xa5 (0x7f2f516a5da5 in /home/gaoxiang/.local/lib/python3.11/site-packages/torch/lib/libtorch_cpu.so)
frame #11: <unknown function> + 0x2fe38c4 (0x7f2f525e38c4 in /home/gaoxiang/.local/lib/python3.11/site-packages/torch/lib/libtorch_cpu.so)
frame #12: at::_ops::isclose::call(at::Tensor const&, at::Tensor const&, double, double, bool) + 0x18b (0x7f2f521b47fb in /home/gaoxiang/.local/lib/python3.11/site-packages/torch/lib/libtorch_cpu.so)
frame #13: at::native::allclose(at::Tensor const&, at::Tensor const&, double, double, bool) + 0x21 (0x7f2f516a3f21 in /home/gaoxiang/.local/lib/python3.11/site-packages/torch/lib/libtorch_cpu.so)
frame #14: at::_ops::allclose::redispatch(c10::DispatchKeySet, at::Tensor const&, at::Tensor const&, double, double, bool) + 0x91 (0x7f2f51b8ac11 in /home/gaoxiang/.local/lib/python3.11/site-packages/torch/lib/libtorch_cpu.so)
frame #15: <unknown function> + 0x41e0136 (0x7f2f537e0136 in /home/gaoxiang/.local/lib/python3.11/site-packages/torch/lib/libtorch_cpu.so)
frame #16: at::_ops::allclose::call(at::Tensor const&, at::Tensor const&, double, double, bool) + 0x17f (0x7f2f51bce45f in /home/gaoxiang/.local/lib/python3.11/site-packages/torch/lib/libtorch_cpu.so)
frame #17: <unknown function> + 0x58e13e (0x55fdc218413e in ./build/nvfuser_tests)
frame #18: <unknown function> + 0x687f17 (0x55fdc227df17 in ./build/nvfuser_tests)
frame #19: <unknown function> + 0x67a7cd (0x55fdc22707cd in ./build/nvfuser_tests)
frame #20: <unknown function> + 0x67a9a5 (0x55fdc22709a5 in ./build/nvfuser_tests)
frame #21: <unknown function> + 0x67ab80 (0x55fdc2270b80 in ./build/nvfuser_tests)
frame #22: <unknown function> + 0x67e35d (0x55fdc227435d in ./build/nvfuser_tests)
frame #23: <unknown function> + 0x6884a7 (0x55fdc227e4a7 in ./build/nvfuser_tests)
frame #24: <unknown function> + 0x67adb2 (0x55fdc2270db2 in ./build/nvfuser_tests)
frame #25: <unknown function> + 0x18ff89 (0x55fdc1d85f89 in ./build/nvfuser_tests)
frame #26: <unknown function> + 0x23850 (0x7f2f09e39850 in /usr/lib/libc.so.6)
frame #27: __libc_start_main + 0x8a (0x7f2f09e3990a in /usr/lib/libc.so.6)
frame #28: <unknown function> + 0x1c7ce5 (0x55fdc1dbdce5 in ./build/nvfuser_tests)
" thrown in the test body.
[  FAILED  ] NVFuserTest.FusionMatmulSchedulerEpilogueBias_CUDA (993 ms)

[drzejan2]

I reproduced the issue with other ad102 device with 530 driver installed. Below is an example from compute-sanitizer from that device:

Note: Google Test filter = *FusionMatmulSchedulerEpilogueBias_CUDA
[==========] Running 1 test from 1 test suite.
[----------] Global test environment set-up.
[----------] 1 test from NVFuserTest
[ RUN      ] NVFuserTest.FusionMatmulSchedulerEpilogueBias_CUDA
========= Invalid __global__ read of size 16 bytes
=========     at 0x2090 in CudaCodeGen::kernel1(CudaCodeGen::Tensor<CudaCodeGen::__half, (int)2, (int)2>, CudaCodeGen::Tensor<CudaCodeGen::__half, (int)2, (int)2>, CudaCodeGen::Tensor<float, (int)1, (int)1>, CudaCodeGen::Tensor<float, (int)2, (int)2>)
=========     by thread (8,0,0) in block (0,0,0)
=========     Address 0x14ed8e540004 is misaligned
=========     and is inside the nearest allocation at 0x14ed8e400000 of size 2,097,152 bytes
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame: [0x30b442]
=========                in /lib/x86_64-linux-gnu/libcuda.so.1
=========     Host Frame:/tmp/Fuser/csrc/executor.cpp:1723:nvfuser::FusionExecutor::runFusion(nvfuser::KernelArgumentHolder&, nvfuser::LaunchParams const&, nvfuser::CompileParams, std::vector<at::Tensor, std::allocator<at::Tensor> >) [0xdd7400]
=========                in /tmp/libnvfuser_codegen.so
=========     Host Frame:/tmp/Fuser/csrc/kernel_cache.cpp:753:nvfuser::FusionKernelRuntime::runKernelWithInput(nvfuser::KernelArgumentHolder&, nvfuser::SegmentedGroup*) [0xfe80e4]
=========                in /tmp/libnvfuser_codegen.so
=========     Host Frame:/tmp/Fuser/csrc/kernel_cache.cpp:1063:nvfuser::FusionKernelRuntime::runSegmentsWithInputs(nvfuser::KernelArgumentHolder&) [0xfea701]
=========                in /tmp/libnvfuser_codegen.so
=========     Host Frame:/tmp/Fuser/csrc/kernel_cache.cpp:959:nvfuser::FusionKernelRuntime::runWithInputs(nvfuser::KernelArgumentHolder&) [0xfe9caf]
=========                in /tmp/libnvfuser_codegen.so
=========     Host Frame:/tmp/Fuser/csrc/kernel_cache.cpp:430:nvfuser::FusionExecutorCache::runFusionWithInputs(c10::ArrayRef<c10::IValue> const&, std::optional<nvfuser::PrimDataType>, std::optional<signed char>) [0xfe5c9e]
=========                in /tmp/libnvfuser_codegen.so
=========     Host Frame:/tmp/Fuser/test/test_gpu_tensorcore.cpp:4286:nvfuser::NVFuserTest_FusionMatmulSchedulerEpilogueBias_CUDA_Test::TestBody() [0xade47e]
=========                in /tmp/nvfuser_tests
=========     Host Frame:/tmp/pytorch/third_party/googletest/googletest/src/gtest.cc:2607:void testing::internal::HandleSehExceptionsInMethodIfSupported<testing::Test, void>(testing::Test*, void (testing::Test::*)(), char const*) [0xcf7f18]
=========                in /tmp/nvfuser_tests
=========     Host Frame:/tmp/pytorch/third_party/googletest/googletest/src/gtest.cc:2643:void testing::internal::HandleExceptionsInMethodIfSupported<testing::Test, void>(testing::Test*, void (testing::Test::*)(), char const*) [0xcf1097]
=========                in /tmp/nvfuser_tests
=========     Host Frame:/tmp/pytorch/third_party/googletest/googletest/src/gtest.cc:2689:testing::Test::Run() [0xcc6b96]
=========                in /tmp/nvfuser_tests
=========     Host Frame:/tmp/pytorch/third_party/googletest/googletest/src/gtest.cc:2864:testing::TestInfo::Run() [0xcc75d8]
=========                in /tmp/nvfuser_tests
=========     Host Frame:/tmp/pytorch/third_party/googletest/googletest/src/gtest.cc:3016:testing::TestSuite::Run() [0xcc7ed5]
=========                in /tmp/nvfuser_tests
=========     Host Frame:/tmp/pytorch/third_party/googletest/googletest/src/gtest.cc:5856:testing::internal::UnitTestImpl::RunAllTests() [0xcd7873]
=========                in /tmp/nvfuser_tests
=========     Host Frame:/tmp/pytorch/third_party/googletest/googletest/src/gtest.cc:2609:bool testing::internal::HandleSehExceptionsInMethodIfSupported<testing::internal::UnitTestImpl, bool>(testing::internal::UnitTestImpl*, bool (testing::internal::UnitTestImpl::*)(), char const*) [0xcf8e2f]
=========                in /tmp/nvfuser_tests
=========     Host Frame:/tmp/pytorch/third_party/googletest/googletest/src/gtest.cc:2643:bool testing::internal::HandleExceptionsInMethodIfSupported<testing::internal::UnitTestImpl, bool>(testing::internal::UnitTestImpl*, bool (testing::internal::UnitTestImpl::*)(), char const*) [0xcf204b]
=========                in /tmp/nvfuser_tests
=========     Host Frame:/tmp/pytorch/third_party/googletest/googletest/src/gtest.cc:5441:testing::UnitTest::Run() [0xcd5fd8]
=========                in /tmp/nvfuser_tests
=========     Host Frame:/tmp/pytorch/third_party/googletest/googletest/include/gtest/gtest.h:2491:RUN_ALL_TESTS() [0x2ebf5b]
=========                in /tmp/nvfuser_tests
=========     Host Frame:/tmp/Fuser/test/main.cpp:32:main [0x2eb4d0]
=========                in /tmp/nvfuser_tests
=========     Host Frame:../sysdeps/nptl/libc_start_call_main.h:58:__libc_start_call_main [0x29d90]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame:../csu/libc-start.c:379:__libc_start_main [0x29e40]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame:_start [0x2954fe]
=========                in /tmp/nvfuser_tests
...
...
...
...
========= Invalid __global__ read of size 16 bytes
=========     at 0x2090 in CudaCodeGen::kernel1(CudaCodeGen::Tensor<CudaCodeGen::__half, (int)2, (int)2>, CudaCodeGen::Tensor<CudaCodeGen::__half, (int)2, (int)2>, CudaCodeGen::Tensor<float, (int)1, (int)1>, CudaCodeGen::Tensor<float, (int)2, (int)2>)
=========     by thread (11,0,0) in block (1,0,0)
=========     Address 0x14ed8e540404 is misaligned
=========     and is inside the nearest allocation at 0x14ed8e400000 of size 2,097,152 bytes
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame: [0x30b442]
=========                in /lib/x86_64-linux-gnu/libcuda.so.1
=========     Host Frame:/tmp/Fuser/csrc/executor.cpp:1723:nvfuser::FusionExecutor::runFusion(nvfuser::KernelArgumentHolder&, nvfuser::LaunchParams const&, nvfuser::CompileParams, std::vector<at::Tensor, std::allocator<at::Tensor> >) [0xdd7400]
=========                in /tmp/libnvfuser_codegen.so
=========     Host Frame:/tmp/Fuser/csrc/kernel_cache.cpp:753:nvfuser::FusionKernelRuntime::runKernelWithInput(nvfuser::KernelArgumentHolder&, nvfuser::SegmentedGroup*) [0xfe80e4]
=========                in /tmp/libnvfuser_codegen.so
=========     Host Frame:/tmp/Fuser/csrc/kernel_cache.cpp:1063:nvfuser::FusionKernelRuntime::runSegmentsWithInputs(nvfuser::KernelArgumentHolder&) [0xfea701]
=========                in /tmp/libnvfuser_codegen.so
=========     Host Frame:/tmp/Fuser/csrc/kernel_cache.cpp:959:nvfuser::FusionKernelRuntime::runWithInputs(nvfuser::KernelArgumentHolder&) [0xfe9caf]
=========                in /tmp/libnvfuser_codegen.so
=========     Host Frame:/tmp/Fuser/csrc/kernel_cache.cpp:430:nvfuser::FusionExecutorCache::runFusionWithInputs(c10::ArrayRef<c10::IValue> const&, std::optional<nvfuser::PrimDataType>, std::optional<signed char>) [0xfe5c9e]
=========                in /tmp/libnvfuser_codegen.so
=========     Host Frame:/tmp/Fuser/test/test_gpu_tensorcore.cpp:4286:nvfuser::NVFuserTest_FusionMatmulSchedulerEpilogueBias_CUDA_Test::TestBody() [0xade47e]
=========                in /tmp/nvfuser_tests
=========     Host Frame:/tmp/pytorch/third_party/googletest/googletest/src/gtest.cc:2607:void testing::internal::HandleSehExceptionsInMethodIfSupported<testing::Test, void>(testing::Test*, void (testing::Test::*)(), char const*) [0xcf7f18]
=========                in /tmp/nvfuser_tests
=========     Host Frame:/tmp/pytorch/third_party/googletest/googletest/src/gtest.cc:2643:void testing::internal::HandleExceptionsInMethodIfSupported<testing::Test, void>(testing::Test*, void (testing::Test::*)(), char const*) [0xcf1097]
=========                in /tmp/nvfuser_tests
=========     Host Frame:/tmp/pytorch/third_party/googletest/googletest/src/gtest.cc:2689:testing::Test::Run() [0xcc6b96]
=========                in /tmp/nvfuser_tests
=========     Host Frame:/tmp/pytorch/third_party/googletest/googletest/src/gtest.cc:2864:testing::TestInfo::Run() [0xcc75d8]
=========                in /tmp/nvfuser_tests
=========     Host Frame:/tmp/pytorch/third_party/googletest/googletest/src/gtest.cc:3016:testing::TestSuite::Run() [0xcc7ed5]
=========                in /tmp/nvfuser_tests
=========     Host Frame:/tmp/pytorch/third_party/googletest/googletest/src/gtest.cc:5856:testing::internal::UnitTestImpl::RunAllTests() [0xcd7873]
=========                in /tmp/nvfuser_tests
=========     Host Frame:/tmp/pytorch/third_party/googletest/googletest/src/gtest.cc:2609:bool testing::internal::HandleSehExceptionsInMethodIfSupported<testing::internal::UnitTestImpl, bool>(testing::internal::UnitTestImpl*, bool (testing::internal::UnitTestImpl::*)(), char const*) [0xcf8e2f]
=========                in /tmp/nvfuser_tests
=========     Host Frame:/tmp/pytorch/third_party/googletest/googletest/src/gtest.cc:2643:bool testing::internal::HandleExceptionsInMethodIfSupported<testing::internal::UnitTestImpl, bool>(testing::internal::UnitTestImpl*, bool (testing::internal::UnitTestImpl::*)(), char const*) [0xcf204b]
=========                in /tmp/nvfuser_tests
=========     Host Frame:/tmp/pytorch/third_party/googletest/googletest/src/gtest.cc:5441:testing::UnitTest::Run() [0xcd5fd8]
=========                in /tmp/nvfuser_tests
=========     Host Frame:/tmp/pytorch/third_party/googletest/googletest/include/gtest/gtest.h:2491:RUN_ALL_TESTS() [0x2ebf5b]
=========                in /tmp/nvfuser_tests
=========     Host Frame:/tmp/Fuser/test/main.cpp:32:main [0x2eb4d0]
=========                in /tmp/nvfuser_tests
=========     Host Frame:../sysdeps/nptl/libc_start_call_main.h:58:__libc_start_call_main [0x29d90]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame:../csu/libc-start.c:379:__libc_start_main [0x29e40]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame:_start [0x2954fe]
=========                in /tmp/nvfuser_tests

I also checked the same binary with the same tools on a a100 GPU and there are no unaligned reads.

I will proceed with further investigation.

[drzejan2]

For the time being the failing tests are disabled in #683

[drzejan2]

A WAR solution (disabled propagation of vectorization for epilogue inputs) has added with #707

[naoyam]

Is this already fixed?

[kevinstephano]

@zasdfgbnm is this still an issue or just an old bug? I see a bunch of PR merges.