Fix arange when step is negative #1942

naoyam · 2022-08-29T00:36:42Z

The device version of ceilDiv assumes positive inputs, so when step is negative, it gives an incorrect result. For example, I see FusionStandAloneArange results in a write error with compute-sanitizer when start = 0, stop = -1, step = -1.5 and dtype = kLong. Here's sanitizer output:

========= Invalid __global__ write of size 8 bytes
=========     at 0x310 in CudaCodeGen::kernel5(long long, long long, long long, CudaCodeGen::Tensor<long long, (int)1>)
=========     by thread (1,0,0) in block (0,0,0)
=========     Address 0x7f2fbc600808 is out of bounds
=========     and is 1 bytes after the nearest allocation at 0x7f2fbc600800 of size 8 bytes
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame: [0x21740c]

We could fix the device ceilDiv, but I'm feeling we don't want to introduce the conditional code if possible.

negative, it gives an incorrect result. For example, I see FusionStandAloneArange results in a write error with compute-sanitizer when start = 0, stop = -1, step = -1.5 and dtype = kLong.

naoyam · 2022-08-29T01:38:08Z

torch/csrc/jit/codegen/cuda/arith.cpp

-  auto size = castOp(DataType::Int, ceilDiv(sub(end, start), step));
+  // Make sure no negative value is passed to ceilDiv as the device
+  // implementation of ceilDiv assumes positive inputs
+  auto size = castOp(DataType::Int, ceilDiv(abs(sub(end, start)), abs(step)));


Actually, I'm not really sure what the right behavior should be when the signs of end - start and step are different, which I think is invalid. Do we need to do error checking? Doing so in device code could be very costly, so we probably would want to do on the host side before launching kernels.

In any case, unless we want to check it on device, the above workaround should be fine.

I think we do need to check it on the host side, but I am not sure how. Would it be a good idea to have both a BinaryOpType::CeilDiv and a BinaryOpType::CeilDivMaybeNegative, and we use CeilDivMaybeNegative for arange, and use CeilDiv for others?

I don't know if we would need additional ceilDiv expr types. I think it should be possible to determine if it's guaranteed to be safe or not by looking at Expr nodes, and we could just do host-side checks of those ceilDiv exprs that are not determined to be safe.

That said, I don't think it's important at this point. At least for now, I think it'd be just reasonable if nvFuser works if a given program is correct but could result in undefined behavior if an unsound program is given. We would definitely want to leave a note about the lack of the check.

But, if it should be validated anyway, I think we would want to extend KernelIrScanner to mark potentially unsafe ceilDiv exprs. As I mentioned above, I think we should be able to just analyze the input vals of those exprs and see if they are composed of just non-negative values such as domain extents.

A note is fine for now. We can revisit in the future if needed.

zasdfgbnm

LGTM, thanks for fixing!

zasdfgbnm · 2022-08-29T02:48:54Z

torch/csrc/jit/codegen/cuda/arith.cpp

-  auto size = castOp(DataType::Int, ceilDiv(sub(end, start), step));
+  // Make sure no negative value is passed to ceilDiv as the device
+  // implementation of ceilDiv assumes positive inputs
+  auto size = castOp(DataType::Int, ceilDiv(abs(sub(end, start)), abs(step)));


I think we do need to check it on the host side, but I am not sure how. Would it be a good idea to have both a BinaryOpType::CeilDiv and a BinaryOpType::CeilDivMaybeNegative, and we use CeilDivMaybeNegative for arange, and use CeilDiv for others?

zasdfgbnm · 2022-08-29T02:50:10Z

torch/csrc/jit/codegen/cuda/dynamic_type.h

@@ -186,6 +186,14 @@ class TORCH_CUDA_CU_API IntOrDouble {
  explicit operator int64_t() const;
  explicit operator size_t() const;
  explicit operator int() const;
+
+  IntOrDouble abs() const {


Should we make this a function inside namespace IntOrDouble_functions in this file to be consistent with other functions?

Syncing nvfuser devel branch to upstream master. https://github.com/csarofeen/pytorch/ Codegen changes include: * codegen improvement: i. allow non-root trivial reductions, allow empty/no-op fusion ii. fixes vectorization checks and size calculation iii. bank conflict handle improvement iv. enables transpose scheduler * misc: i. CI tests failure fixes ii. cpp tests file clean up iii. trivial forwarding supports added in codegen runtime iv. added factory methods support in codegen Commits that's in this PR from the devel branch: ``` 7117a7e patching nvfuser conv cudnn test numerics mismatch (#2048) 65af1a4 Inserting sync for redundant parallel types is already done at the (#2023) 6ac74d1 Fix sync map (#2047) f5bca33 Bank conflict checker improvements (#2032) d2ca7e3 Minor update on cp.async code generation. (#1901) d36cf61 Test file cleanup (#2040) 0b8e83f Allow non-root trivial reductions (#2037) a2dfe40 Fix vectorize size calculation (#2035) e040676 Use withPredicate to replace setPredicate to maintain Exprs immutable (#2025) 197221b removing ci workflow (#2034) 40e2703 Reduction rand like patch (#2031) bc77266 Add utility for checking bank conflict of shared memory (#2029) ddd1cf7 Add back FusionReductionWithTrivialReduction_CUDA (#2030) fbd97e5 Revert "Cleanup trivial reduction workarounds (#2006)" (#2024) bca20c1 Cleanup trivial reduction workarounds (#2006) e4b6585 Trivial forwarding (#1995) 1a0e355 Fix contiguity analysis of predicates to match updated contiguity. (#1991) a4effa6 Enable output allocation cache (#2010) 35440b7 Patching bn inference (#2016) 0f9f0b4 Add matmul benchmark (#2007) 45045cd Enable tests previously disabled due to an aliasing bug (#2005) 967aa77 Contiguous indexing for View operations (#1990) a43cb20 Make inlining even more modular (#2004) dc45835 Test util cleanup (#2003) 3ca21eb More strict validation (#2000) a7a7d57 Fix build problem (#1999) fc235b0 Just fixes comments (#1998) 482386c cleanup (#1997) 4cbe0db Improve divisible split detection (#1970) 42ccc52 Minor build fix. (#1996) fcf8c09 Cleanup of lower_utils.cpp: Isolate out GpuLower usage (#1989) 15f2f6d Move ConcretizedBroadcastDomains to shared_ptr in GpuLower. (#1988) 8f1c7f5 Minor cleanup lower_unroll.cpp (#1994) 1d9858c Minor cleanup (#1992) f262d9c Add support for uniform RNG (#1986) eb1dad1 Remove non-const functions, remove GpuLower instance on build, pass in ca_map. (#1987) 634820c Add support for some empty fusion (#1981) eabe8d8 Segment self mapping fusions (#1954) e96aacf Enable Transpose operation (#1882) 425dce2 Add a null scheduler that helps segmenting away no-op schedules (#1835) 306d4a6 Fix canScheduleCompileTime check of transpose scheduler (#1969) b1bd32c Minor fix (#1967) bd93578 Enable transpose scheduler (#1927) b7a206e Move scheduler vectorize utilities into their own file (#1959) d9420e4 View scheduling (#1928) c668e13 Upstream push ci fixes (#1965) c40202b Fix dump effective bandwidth (#1962) 93505bc WAR on index mapping when exact and permissive maps differ (#1960) 45e95fd Allow splitting inner-most ID to create virtual innermost ID in transpose scheduler (#1930) a3ecb33 Improve the comments at the beginning of index_compute.h (#1946) f7bc341 Remove unused variables (#1955) df3393a Some cleanup (#1957) 7d1d7c8 TVDomainGuard factory (#1953) 357ba22 Fill allocation with nan on tests (#1956) 8eafc54 Fix detection of unmappable root domains (#1952) 90a51f2 Some indexing cleanups, Add eye support (#1940) ddc01e4 Exclude unsupported data types (#1951) 992e17c test the groups the same order as they are merged (#1949) 208262b Move detection of self mapping IDs to IterDomainGraph from (#1941) ac4de38 Merge pull request #1945 from csarofeen/master_merge_0828 6310948 Add full, full_like, zeros, zeros_like, ones, ones_like (#1943) aab10bc Merge remote-tracking branch 'upstream/viable/strict' into HEAD 4c254c0 Fix arange when step is negative (#1942) 89330aa Tensor factories must set the output shape as its input (#1939) ``` RUN_TORCHBENCH: nvfuser Differential Revision: [D40869846](https://our.internmc.facebook.com/intern/diff/D40869846) Pull Request resolved: pytorch#87779 Approved by: https://github.com/davidberard98

The device version of ceilDiv assumes positive inputs, so when step is

4d2fa7c

negative, it gives an incorrect result. For example, I see FusionStandAloneArange results in a write error with compute-sanitizer when start = 0, stop = -1, step = -1.5 and dtype = kLong.

naoyam requested a review from zasdfgbnm August 29, 2022 00:38

naoyam commented Aug 29, 2022

View reviewed changes

zasdfgbnm approved these changes Aug 29, 2022

View reviewed changes

naoyam added 2 commits August 28, 2022 21:38

PR feedback

e37ef6e

Add a comment on potential unsafe behavior

88c5181

naoyam merged commit 4c254c0 into devel Aug 29, 2022

zasdfgbnm deleted the fix_arange_with_negative_step_2 branch August 29, 2022 06:24

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

Fix arange when step is negative #1942

Fix arange when step is negative #1942

Uh oh!

naoyam commented Aug 29, 2022 •

edited

Loading

Uh oh!

naoyam Aug 29, 2022

Uh oh!

zasdfgbnm Aug 29, 2022

Uh oh!

naoyam Aug 29, 2022 •

edited

Loading

Uh oh!

naoyam Aug 29, 2022

Uh oh!

zasdfgbnm Aug 29, 2022

Uh oh!

zasdfgbnm left a comment

Uh oh!

zasdfgbnm Aug 29, 2022

Uh oh!

zasdfgbnm Aug 29, 2022

Uh oh!

Uh oh!

Fix arange when step is negative #1942

Fix arange when step is negative #1942

Uh oh!

Conversation

naoyam commented Aug 29, 2022 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

naoyam Aug 29, 2022

Choose a reason for hiding this comment

Uh oh!

zasdfgbnm Aug 29, 2022

Choose a reason for hiding this comment

Uh oh!

naoyam Aug 29, 2022 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

naoyam Aug 29, 2022

Choose a reason for hiding this comment

Uh oh!

zasdfgbnm Aug 29, 2022

Choose a reason for hiding this comment

Uh oh!

zasdfgbnm left a comment

Choose a reason for hiding this comment

Uh oh!

zasdfgbnm Aug 29, 2022

Choose a reason for hiding this comment

Uh oh!

zasdfgbnm Aug 29, 2022

Choose a reason for hiding this comment

Uh oh!

Uh oh!

naoyam commented Aug 29, 2022 •

edited

Loading

naoyam Aug 29, 2022 •

edited

Loading