NVIDIA / NVIDIA/cutlass

[BUG] Segmentation fault with high N, C, K values in Conv2dFprop

Open
#2,175 2 comments 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

? - Needs Triage bug inactive-30d inactive-90d
Dominant language
C++
Stars
10.5k
Forks
2.1k
Avg merge
3d 11h
Merged PRs (30d)
7

Description

Describe the bug
Hello, I get a segmentation fault while executing example 16 with default configuration (plus cuda11.8, ubuntu 18, with a Tesla V100 volta) and high value of batch size (n), channels (c) and filters (k):

16_ampere_tensorop_conv2dfprop --n=32 --h=512 --w=512 --c=1024 --k=1024 --r=3 --s=3

The error has it s origin in the following:

  cutlass::reference::host::TensorFillRandomUniform(
      tensor_b.host_view(),
      1,
      ElementInputB(7),
      ElementInputB(-8),
      0);

and the backtrace is the following:

Thread 1 "16_ampere_tenso" received signal SIGSEGV, Segmentation fault.
cutlass::reference::host::detail::TensorFillRandomUniformFunc<cutlass::half_t, cutlass::layout::TensorNHWC>::operator() (this=0x7fffffffaa10, coord=...)
    at /mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/tools/util/include/cutlass/util/reference/host/tensor_fill.h:855
855         view.at(coord) = func();
(gdb) backtrace
#0  cutlass::reference::host::detail::TensorFillRandomUniformFunc<cutlass::half_t, cutlass::layout::TensorNHWC>::operator() (this=0x7fffffffaa10, coord=...)
    at /mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/tools/util/include/cutlass/util/reference/host/tensor_fill.h:855
#1  0x000055555556da20 in cutlass::reference::host::detail::TensorForEachHelper<cutlass::reference::host::detail::TensorFillRandomUniformFunc<cutlass::half_t, cutlass::layout::TensorNHWC>, 4, 0>::TensorForEachHelper (this=0x7fffffff943b, func=..., extent=..., coord=...)
    at /mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/tools/util/include/cutlass/util/reference/host/tensor_foreach.h:80
#2  0x000055555556d1cc in cutlass::reference::host::detail::TensorForEachHelper<cutlass::reference::host::detail::TensorFillRandomUniformFunc<cutlass::half_t, cutlass::layout::TensorNHWC>, 4, 1>::TensorForEachHelper (this=0x7fffffff949b, func=..., extent=..., coord=...)
    at /mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/tools/util/include/cutlass/util/reference/host/tensor_foreach.h:60
#3  0x000055555556bf8a in cutlass::reference::host::detail::TensorForEachHelper<cutlass::reference::host::detail::TensorFillRandomUniformFunc<cutlass::half_t, cutlass::layout::TensorNHWC>, 4, 2>::TensorForEachHelper (this=0x7fffffff94fb, func=..., extent=..., coord=...)
    at /mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/tools/util/include/cutlass/util/reference/host/tensor_foreach.h:60
#4  0x000055555556a43c in cutlass::reference::host::detail::TensorForEachHelper<cutlass::reference::host::detail::TensorFillRandomUniformFunc<cutlass::half_t, cutlass::layout::TensorNHWC>, 4, 3>::TensorForEachHelper (this=0x7fffffff9557, func=..., extent=..., coord=...)
    at /mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/tools/util/include/cutlass/util/reference/host/tensor_foreach.h:60
#5  0x000055555556883a in cutlass::reference::host::TensorForEach<cutlass::reference::host::detail::TensorFillRandomUniformFunc<cutlass::half_t, cutlass::layout::TensorNHWC>, 4>
    (extent=..., func=...) at /mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/tools/util/include/cutlass/util/reference/host/tensor_foreach.h:95
#6  0x000055555555eec7 in cutlass::reference::host::TensorFillRandomUniform<cutlass::half_t, cutlass::layout::TensorNHWC> (dst=..., seed=1, max=7, min=-8, bits=0, pnan=0, 
    exclude_zero=false) at /mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/tools/util/include/cutlass/util/reference/host/tensor_fill.h:980
#7  0x00005555555586c2 in profile_convolution (options=...)
    at /mnt/beegfs/gap/izcagal@upvnet.upv.es/cutlass/examples/16_ampere_tensorop_conv2dfprop/ampere_tensorop_conv2dfprop.cu:552
#8  0x000055555555b233 in main (argc=8, args=0x7fffffffd418)

With other configurations or even high batch size but lower C or K I get no error:

--n=256 --h=256 --w=256 --c=8 --k=8 --r=3 --s=3 -> ok
--n=256 --h=256 --w=256 --c=512 --k=512 --r=3 --s=3 -> error
--n=1 --h=256 --w=256 --c=512 --k=512 --r=3 --s=3 -> ok
--n=256 --h=256 --w=256 --c=128 --k=128 --r=3 --s=3 -> error:

terminate called after throwing an instance of 'std::length_error'
  what():  vector::_M_default_append

--n=256 --h=256 --w=256 --c=64 --k=64 --r=3 --s=3 -> error:

terminate called after throwing an instance of 'cutlass::cuda_exception'
  what():  std::exception
Aborted (core dumped)

--n=32 --h=256 --w=256 --c=64 --k=64 --r=3 --s=3 -> ok
--n=32 --h=256 --w=256 --c=128 --k=128 --r=3 --s=3 -> ok

Can anyone explain the boundaries for each parameter? Which is the highest n, k or c number I can execute? I do understand that h and w are not the limiting factors.

Thank you.

Contributor guide

No contributing guide indexed for this repository

First steps

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. Open a pull request that references the issue number.

Research direction

Start by running examples/16_ampere_tensorop_conv2dfprop/ampere_tensorop_conv2dfprop.cu with the reported high N, C, and K values. Trace the allocation and fill path through tools/util/include/cutlass/util/reference/host/tensor_fill.h and tensor_foreach.h. Done means establishing the parameter boundaries and explaining or addressing the failure modes without the segmentation fault.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
machine-learning, performance
Issue type
Bug
Difficulty
4/5
Estimated time
3-5 days
Activity status
Stale
Clarity
Mostly clear
Newbie friendliness
35/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.