Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[WIP: Do not merge] Fix cat while keeping the memory layout #286

Open
wants to merge 1,922 commits into
base: master
Choose a base branch
from

Conversation

shuhand0
Copy link
Collaborator

@shuhand0 shuhand0 commented Feb 3, 2023

Fixes #ISSUE_NUMBER

YJ Shi and others added 30 commits December 9, 2022 00:09
… helper func (pytorch#90214)

Fixes pytorch#90213. Also a torch.tensor to tvm.nd.array helper function is added to avoid data copy with dlpack.

@jansel @Chillee

Pull Request resolved: pytorch#90214
Approved by: https://github.com/wconstab
Summary: botnet26t_256 shows random tolerance failure on CI. The root
cause of this randomness is still to-be-invesitgated, but let's use a
larger tolerance for now.

Pull Request resolved: pytorch#90383
Approved by: https://github.com/ezyang
…from original FQN to flat_param (pytorch#89899)

**Motivation:**
Add a helper to map from the FQN to the corresponding flat_param. The helper will directly get flat_param from fsdp_state and flat_handler as flat_param is not registered to the module if `use_orig_params` is True.

Pull Request resolved: pytorch#89899
Approved by: https://github.com/awgu
Signed-off-by: Edward Z. Yang <ezyang@fb.com>

Pull Request resolved: pytorch#89867
Approved by: https://github.com/albanD
Fixes pytorch#90375

Signed-off-by: Edward Z. Yang <ezyang@fb.com>

Pull Request resolved: pytorch#90392
Approved by: https://github.com/voznesenskym
Test Plan: CI

Differential Revision: D41666330

Pull Request resolved: pytorch#90023
Approved by: https://github.com/SherlockNoMad
…ytorch#90357)

Summary:
Copying QInt8 and QInt32 from cpu to vulkan:
 - Added shader nchw_to_image_int8
 - Added shader nchw_to_image_int32

Copying QInt8 and QInt32 from vulkan to cpu
Note: This functionality is currently disabled until issues on Android are resolved.
- Added shader image_to_nchw_int32
- QInt8 works with the same existing image_to_nchw_quantized shaders

Added multiple tests for each supported dtype:
- cpu_to_vulkan_and_dequantize:
These tests check the correctness of copying quantized cpu tensor to vulkan by comparing the output of the following:
  - cpu float tensor -> quantize -> to vulkan -> dequantize -> to cpu
  - cpu float tensor -> quantize -> dequantize
- cpu_to_vulkan_and_vulkan_to_cpu
(currently disabled until copying vulkan quantized to cpu is enabled):
These tests check the correctness of copying from cpu to vulkan and from vulkan to cpu by creating a random cpu float tensor, quantizing it, then copying it to vulkan, then back to cpu and comparing the output tensor to the original quantized tensor.
- quantize_per_tensor_and_vulkan_to_cpu
(currently disabled until copying vulkan quantized to cpu is enabled):
These tests check the correctness of copying quantized tensor from vulkan to cpu by comparing the output of the following:
  - cpu float tensor -> to vulkan -> quantize -> to cpu
  - cpu float tensor -> quantize

Test Plan:
On Mac
```
cd ~/fbsource
buck1 run -c pt.vulkan_full_precision=1 //xplat/caffe2:pt_vulkan_quantized_api_test_binAppleMac\#macosx-arm64
```

On Android
```
cd ~/fbsource
buck1 build -c ndk.custom_libcxx=false -c pt.enable_qpl=0 -c pt.vulkan_full_precision=1 //xplat/caffe2:pt_vulkan_quantized_api_test_binAndroid\#android-arm64 --show-output
adb push buck-out/gen/xplat/caffe2/pt_vulkan_quantized_api_test_binAndroid\#android-arm64 /data/local/tmp/vulkan_quantized_api_test
adb shell "/data/local/tmp/vulkan_quantized_api_test"
```

Reviewed By: kimishpatel

Differential Revision: D41654287

Pull Request resolved: pytorch#90357
Approved by: https://github.com/SS-JIA
This is an automated pull request to update the first-party submodule for [pytorch/FBGEMM](https://github.com/pytorch/FBGEMM).

New submodule commit: pytorch/FBGEMM@f99e161

Test Plan: Ensure that CI jobs succeed on GitHub before landing.
Pull Request resolved: pytorch#74729
Approved by: https://github.com/malfet
If PyTorch is package into a wheel with [nvidia-cublas-cu11](https://pypi.org/project/nvidia-cublas-cu11/), which is designated as PureLib, but `torch` wheel is not, can cause a torch_globals loading problem.

Fix that by searching for `nvidia/cublas/lib/libcublas.so.11` an `nvidia/cudnn/lib/libcudnn.so.8` across all `sys.path` folders.

Test plan:
```
docker pull amazonlinux:2
docker run --rm -t amazonlinux:2 bash -c 'yum install -y python3 python3-devel python3-distutils patch;python3 -m pip install torch==1.13.0;curl -OL https://patch-diff.githubusercontent.com/raw/pytorch/pytorch/pull/90411.diff; pushd /usr/local/lib64/python3.7/site-packages; patch -p1 </90411.diff; popd; python3 -c "import torch;print(torch.__version__, torch.cuda.is_available())"'
```

Fixes pytorch#88869

Pull Request resolved: pytorch#90411
Approved by: https://github.com/atalman
…ion, weights and bias (pytorch#90012)

Summary:
- Registered vulkan_prepack::create_qconv2d_context to the QuantizedCPU backend.
- Registered vulkan_prepack::run_qconv2d_context to the Vulkan backend.
- Added function test_quantized_conv2d, in order to test Vulkan Quantized Conv2d with QUInt8 activation, weight and bias (all QUInt8).
- Added multiples tests for vulkan quantized conv2d (regular, depthwise and pointwise). All these tests make use of the test_quantized_conv2d function.

This function tests the correctness of vulkan quantized conv2d, by comparing the following two processes:
(we start with randomly generated float cpu tensors)
- random float cpu tensors -> to vulkan -> quantize them -> apply vulkan conv2d quantized op -> dequantize result -> to cpu
- random float cpu tensors -> quantize them -> dequantize -> apply cpu floating point conv2d op on dequantized tensors -> quantize result -> dequantize

This function takes three boolean flags that modify its behavior:
- prepacking:
  - if false, then we directly call at::native::vulkan::ops::quantized_conv2d
  - if true, then we call vulkan_prepack::create_qconv2d_context and vulkan_prepack::run_qconv2d_context.
- compute_quantization_params & random_quantization_params:
  - if both are false, all quantization params are fixed (given as input)
  - if compute_quantization_params is true, all params are computed
  - if random_quantization_params is true, the input params are random and the output params are computed.
(compute_quantization_params takes precedence over random_quantization_params)

Test Plan:
On Mac
```
cd ~/fbsource
buck1 run -c pt.vulkan_full_precision=1 //xplat/caffe2:pt_vulkan_quantized_api_test_binAppleMac\#macosx-arm64
```

On Android
```
cd ~/fbsource
buck1 build -c ndk.custom_libcxx=false -c pt.enable_qpl=0 -c pt.vulkan_full_precision=1 //xplat/caffe2:pt_vulkan_quantized_api_test_binAndroid\#android-arm64 --show-output
adb push buck-out/gen/xplat/caffe2/pt_vulkan_quantized_api_test_binAndroid\#android-arm64 /data/local/tmp/vulkan_quantized_api_test
adb shell "/data/local/tmp/vulkan_quantized_api_test"
```

Reviewed By: SS-JIA

Differential Revision: D41047096

Pull Request resolved: pytorch#90012
Approved by: https://github.com/salilsdesai
Summary: cuda:: is a ambiguous namespace. Make it explicit c10::cuda

Differential Revision: D41469007
/caffe2/caffe2/core/context_gpu.cu(564): error: "caffe2::cuda" is ambiguous/caffe2/caffe2/core/context_gpu.cu(564): error: expected a ";"/caffe2/caffe2/core/context_gpu.cu(568): warning #12-D: parsing restarts here after previous syntax error
Remark: The warnings can be suppressed with "-diag-suppress <warning-number>"/caffe2/caffe2/core/context_gpu.cu(569): error: "caffe2::cuda" is ambiguous/caffe2/caffe2/core/context_gpu.cu(628): error: "caffe2::cuda" is ambiguous
4 errors detected in the compilation of "/caffe2/caffe2/core/context_gpu.cu".

Pull Request resolved: pytorch#89534
Approved by: https://github.com/malfet
Fixes pytorch/torchdynamo#1959, pytorch#90260
However, I wasn't able to make existing stride tests fail before the fix, even though I'm comparing all, not just significant strides.
Separately running refs on meta tensors produces wrong strides as shown in pytorch#90260, however, it looks like in meta tests some other way of computing meta info is used (I've been running
```
pytest -s -v test/test_meta.py -k test_meta_outplace_expand_cuda_float64
```
and verified that it has sample input that should fail, and that it indeed compares all the strides, but the produced `meta_rs` results somehow still had correct strides).

Edit: @SherlockNoMad helped me figure out how to fail the tests, and now I've set the correct ops for checking. `expand` fails for some test inputs because it special-cases 0-dim input case, correctly modeling it in prims would require a lot of changes, so skipping that for now.

Pull Request resolved: pytorch#90341
Approved by: https://github.com/SherlockNoMad
add save and load stats in memory_tracker, so that users could plot the traces in another place, rather than just inside trainer
Pull Request resolved: pytorch#90144
Approved by: https://github.com/rohan-varma
Continuation of pytorch#88207

A compile time guard was preventing ActivityType::CUDA from being available on rocm. This caused both the GPU_FALLBACK and CUDA modes to be active at the same time. So operators were being charged gpu time for the hipEventRecord ranges and the actual kernel execution times. This caused incorrect (and often negative) cuda times, in e.g. table().

Previously a cmake variable was not being propagated to a '-D', causing an issue on Windows, which uses cuda but not cupti.

Pull Request resolved: pytorch#89785
Approved by: https://github.com/jeffdaily, https://github.com/malfet
Get rid of std::iterator inheritance/references for `c10::DictIterator`, `c10::IListRefIterator` and `c10::ListIterator`

Followup after pytorch#90174

Fixes deprecation warning and extension compilation failures using VC++
that raises following errors:
```
C:\actions-runner\_work\pytorch\pytorch\build\win_tmp\build\torch\include\ATen/core/IListRef.h(517): error C4996: 'std::iterator<std::bidirectional_iterator_tag,T,ptrdiff_t,T *,T &>::value_type': warning STL4015: The std::iterator class template (used as a base class to provide typedefs) is deprecated in C++17. (The <iterator> header is NOT deprecated.) The C++ Standard has never required user-defined iterators to derive from std::iterator. To fix this warning, stop deriving from std::iterator and start providing publicly accessible typedefs named iterator_category, value_type, difference_type, pointer, and reference. Note that value_type is required to be non-const, even for constant iterators. You can define _SILENCE_CXX17_ITERATOR_BASE_CLASS_DEPRECATION_WARNING or _SILENCE_ALL_CXX17_DEPRECATION_WARNINGS to acknowledge that you have received this warning.

C:\actions-runner\_work\pytorch\pytorch\build\win_tmp\build\torch\include\ATen/core/List.h(169): error C4996: 'std::iterator<std::random_access_iterator_tag,T,ptrdiff_t,T *,T &>::difference_type': warning STL4015: The std::iterator class template (used as a base class to provide typedefs) is deprecated in C++17. (The <iterator> header is NOT deprecated.) The C++ Standard has never required user-defined iterators to derive from std::iterator. To fix this warning, stop deriving from std::iterator and start providing publicly accessible typedefs named iterator_category, value_type, difference_type, pointer, and reference. Note that value_type is required to be non-const, even for constant iterators. You can define _SILENCE_CXX17_ITERATOR_BASE_CLASS_DEPRECATION_WARNING or _SILENCE_ALL_CXX17_DEPRECATION_WARNINGS to acknowledge that you have received this warning.

```

Discovered while working on pytorch#85969
Pull Request resolved: pytorch#90379
Approved by: https://github.com/ezyang, https://github.com/dagitses
Summary:
This PR adds symmetric quant in the backend config for executorch

Test Plan:
NA, will be tested in meta internal flow

Reviewers:

Subscribers:

Tasks:

Tags:
Pull Request resolved: pytorch#90304
Approved by: https://github.com/cccclai, https://github.com/jcaip, https://github.com/andrewor14
I'm not sure how lintrunner worked without this lol.

Signed-off-by: Edward Z. Yang <ezyang@fb.com>

Pull Request resolved: pytorch#90179
Approved by: https://github.com/albanD, https://github.com/voznesenskym
Signed-off-by: Edward Z. Yang <ezyang@fb.com>

Pull Request resolved: pytorch#90185
Approved by: https://github.com/voznesenskym
…Tuple (pytorch#90186)

Signed-off-by: Edward Z. Yang <ezyang@fb.com>

Pull Request resolved: pytorch#90186
Approved by: https://github.com/voznesenskym
Signed-off-by: Edward Z. Yang <ezyang@fb.com>

Pull Request resolved: pytorch#90202
Approved by: https://github.com/voznesenskym
…h#90286)

The original implementation of cond() operator support in dynamo operated by recursively calling export() on the inner subgraph.  This is problematic for a number of reasons:

* My original motivating reason: the original implementation had to play tricks to feed real tensors to the recursive export call, which means that it doesn't work well with tracing with dynamic shapes (where we MUST stay in fake tensors to accurately track dynamic shapes across the cond invocation)
* If there are pending side effects, the recursive export() call won't see those side effects (as they are only tracked by Dynamo, not actually applied to the Python environment.) You can see an example where dynamo cond tracing does the wrong thing at pytorch#90208
* If there were side effects inside the true/false branch, these side effects were silently lost (as the export only returns the graph of tensor operations, and not any of the residual Python bytecodes necessary to reapply any side effects.) This could have substantive effects on the export of subsequent parts of the model, as those parts of the models could rely on the side effects.
* It was not possible to track NN module accesses inside the true/false branches, necessitating a hack where the NN module was explicitly passed in as an input to cond pytorch#87020 (comment) which doesn't really make any sense from a backend compilation perspective
* Guards induced from the inside of the true/false branch were not properly propagated to the top level guards; they were just silently dropped (in fact, the original implementation checked that the true/false branch produce the same guards which... is not useful? Like, I don't think that actually is even necessary for correctness)

This PR replaces the old implementation with a new implementation based on graphstate checkpointing. The basic idea is to process a cond(), we checkpoint the state of our interpreter, run the true branch, rollback to our checkpoint, run the false branch, rollback to our checkpoint and then merge the changes from both of the checkpoints. I require the true/false branches to have exactly the same side effects, but union their guards.

Some of the details:

* Dynamo is too aggressive with tracking side effects when processing closures, c.f. https://github.com/pytorch/torchdynamo/pull/233/files#r1040480078 The basic problem is whenever I define a closure, this immediately counts as a side effect, even if I didn't actually mutate anything. This triggered on the nested cond export example. To prevent this from happening, I optimistically avoid tracking side effects, but if a STORE_DEREF happens, I restart analysis with the relevant Source.name() added to `mutated_closure_cell_contents` so we start tracking on closure allocation. This is enough to fix the relevant test.
* For the most part, I assert that the graph states must be equivalent after applying the true/false branches. During debugging, I found it useful to be able to compare two graph states and give a better description about what the divergence was. You can test this using the `diff()` method I've added to a few structures.
* The implementation now supports NestedUserFunctionVariable, which is nice as it allows the true/false branches to be defined closer to the cond implementation.
* I fixed the naming of the true/false subgraphs; previously they were named `name_0`, `name_1`, now they are named `cond_true_0` and `cond_false_0`
* I added `name_to_input` to the saved graph state. I don't actually know if this is necessary, but it seemed like a good idea.
* I have to play some tricks to get the speculating execution of the true/false branch to record into a subgraph. After a careful read of OutputGraph, I found that what would work is overriding graph with a fresh Graph that we want to write things into, and manually setting up the inputs/outputs. It's a little delicate as you have to make sure you reset the Graph to its original before you restore a checkpoint, as checkpoints don't actually save graph for efficiency, and just undo changes on the graph. This capability may usefully get refactored to OutputGraph but I didn't do it in this PR for simplicity.

There are some further problems with the cond() implementation that I leave for future work. Most of these were preexisting with the original implementation.

* Not a problem per se, but if an NN module is used by both the true/false branch, it will show up in the final graph twice (since it has to be a submodule of the GraphModule that makes use of it.) I hope the export pipeline can deal with this.
* List of tensor output for cond is not supported.
* The true/false return values may not have consistent sizes/dims/etc, and we don't check them for consistency.
* If we modify fake tensors in the true/false branches, we aren't rolling them back, c.f. https://github.com/pytorch/torchdynamo/issues/1840

Signed-off-by: Edward Z. Yang <ezyang@fb.com>

Pull Request resolved: pytorch#90286
Approved by: https://github.com/voznesenskym
Summary:
In this logic, we are traversing the entries to find the module for STACK_GLOBAL entries.

According to https://github.com/python/cpython/blob/2837241f22be33a5597707b2aa723cb2cf6f3967/Lib/pickletools.py#L1799 we need to look for GET, BINGET and LONG_BINGET.

So this diff updates that. Also while testing, I found some cases of empty modules, for cases such as tanh. For this I added the option to skip processing when this is the case.

Test Plan: Tested with f392778829

Differential Revision: D41748595

Pull Request resolved: pytorch#90223
Approved by: https://github.com/PaliC
Summary:
This diff introduces a set of changes that makes it possible for the host to get assertions from CUDA devices. This includes the introduction of

**`CUDA_KERNEL_ASSERT2`**

A preprocessor macro to be used within a CUDA kernel that, upon an assertion failure, writes the assertion message, file, line number, and possibly other information to UVM (Managed memory). Once this is done, the original assertion is triggered, which places the GPU in a Bad State requiring recovery. In my tests, data written to UVM appears there before the GPU reaches the Bad State and is still accessible from the host after the GPU is in this state.

Messages are written to a multi-message buffer which can, in theory, hold many assertion failures. I've done this as a precaution in case there are several, but I don't actually know whether that is possible and a simpler design which holds only a single message may well be all that is necessary.

**`TORCH_DSA_KERNEL_ARGS`**

This preprocess macro is added as an _argument_ to a kernel function's signature. It expands to supply the standardized names of all the arguments needed by `C10_CUDA_COMMUNICATING_KERNEL_ASSERTION` to handle device-side assertions. This includes, eg, the name of the pointer to the UVM memory the assertion would be written to. This macro abstracts the arguments so there is a single point of change if the system needs to be modified.

**`c10::cuda::get_global_cuda_kernel_launch_registry()`**

This host-side function returns a singleton object that manages the host's part of the device-side assertions. Upon allocation, the singleton allocates sufficient UVM (Managed) memory to hold information about several device-side assertion failures. The singleton also provides methods for getting the current traceback (used to identify when a kernel was launched). To avoid consuming all the host's memory the singleton stores launches in a circular buffer; a unique "generation number" is used to ensure that kernel launch failures map to their actual launch points (in the case that the circular buffer wraps before the failure is detected).

**`TORCH_DSA_KERNEL_LAUNCH`**

This host-side preprocessor macro replaces the standard
```
kernel_name<<<blocks, threads, shmem, stream>>>(args)
```
invocation with
```
TORCH_DSA_KERNEL_LAUNCH(blocks, threads, shmem, stream, args);
```
Internally, it fetches the UVM (Managed) pointer and generation number from the singleton and append these to the standard argument list. It also checks to ensure the kernel launches correctly. This abstraction on kernel launches can be modified to provide additional safety/logging.

**`c10::cuda::c10_retrieve_device_side_assertion_info`**
This host-side function checks, when called, that no kernel assertions have occurred. If one has. It then raises an exception with:
1. Information (file, line number) of what kernel was launched.
2. Information (file, line number, message) about the device-side assertion
3. Information (file, line number) about where the failure was detected.

**Checking for device-side assertions**

Device-side assertions are most likely to be noticed by the host when a CUDA API call such as `cudaDeviceSynchronize` is made and fails with a `cudaError_t` indicating
> CUDA error: device-side assert triggered CUDA kernel errors

Therefore, we rewrite `C10_CUDA_CHECK()` to include a call to `c10_retrieve_device_side_assertion_info()`. To make the code cleaner, most of the logic of `C10_CUDA_CHECK()` is now contained within a new function `c10_cuda_check_implementation()` to which `C10_CUDA_CHECK` passes the preprocessor information about filenames, function names, and line numbers. (In C++20 we can use `std::source_location` to eliminate macros entirely!)

# Notes on special cases

* Multiple assertions from the same block are recorded
* Multiple assertions from different blocks are recorded
* Launching kernels from many threads on many streams seems to be handled correctly
* If two process are using the same GPU and one of the processes fails with a device-side assertion the other process continues without issue
* X Multiple assertions from separate kernels on different streams seem to be recorded, but we can't reproduce the test condition
* X Multiple assertions from separate devices should be all be shown upon exit, but we've been unable to generate a test that produces this condition

Differential Revision: D37621532

Pull Request resolved: pytorch#84609
Approved by: https://github.com/ezyang, https://github.com/malfet
DenisVieriu97 and others added 26 commits January 20, 2023 13:50
* Fix nonzero for empty tensors

* Update TestConsistency list
* Fix nn.functional.gelu
- fix gelu_out_mps key
- add calculation for gelu with tanh
- remove gelu from blocklist

* - add test_gelu_tanh test
- Enable global manual seeding via torch.manual_seed() + test case
- Add torch.mps.synchronize() to wait for MPS stream to finish + test case
- Enable the following python interfaces for MPS:
torch.mps.get_rng_state()
torch.mps.set_rng_state()
torch.mps.is_available()
torch.mps.synchronize()
torch.mps.manual_seed()
torch.mps.seed()
torch.mps.is_initialized()
torch.mps.init()
* Test MPS CI runners

* Cherry pick remaining files

* Enable lintrunner:

* Change lint  runner

* Retrigger checks

* Retrigger checks #2

* Retrigger checks #3

* Retrigger checks #4

* Retrigger checks #5

* Retrigger checks #5

* Retrigger checks #7

* Retrigger checks #8

* Retrigger checks #9

* Retrigger checks #9 (change arch to arm)

* Retrigger checks #10

* Retrigger checks #11

* Retrigger checks #12

* Retrigger checks #13

* Retrigger checks #14

* Retrigger checks #14

* Retrigger checks #15

* Retrigger checks #16

* Retrigger checks #16

* Retrigger checks #17

* Retrigger checks #19

* Retrigger checks #20

* Retrigger checks #21

* Fix lintrunner

* Fix lintrunner

* Remove lint.json
* Fix view slice reshape

* Don't use arrayViews for reshapedTensors
* Fix blocklist dict search

* Fix scatter reduce

* Skip unimplemented ops
- Also fix indentation and formatting
- add revert caculation for save_var used in backward path
- add backward test for native_batch_norm and _native_batch_norm_legit
- remove interpolate from blocklist since it works
- move interpolate_area to unimplemented list
- report error for int64
- for int16,int32 cast tensor to float than floor
- unblock floor_divide and divfloor_rounding testing
- block test_divmode int64 div floor
- add extra tests test_div_floor_int, block for now due to fast math issue
- fix num_output_dims calculation
- fix median_out_mps key
- cast tensor sent to sortWithTensor and argSortWithTensor
- note down same issue for unique
- unblock median from blocklist
- adding test_median_int16 test
- unblock nn.functional.softplus test
- unblock test_softplus test
* Calculate output shape inside nonzero op

* nonzero optimizations

* Fix lintrunner
* Add support for negative dimensions in cumsum

* Expand test with dtypes

* Fix lint issues

---------

Co-authored-by: abhipathak97 <abhipathak97@mps10.scv.apple.com>
* Fix test mps

* Address PR comments

* Add to blocklist remaining failures

* Fix comments
- remove the duplicate calculation from layer_norm_mps
- raise atol rtol for native_layer_norm
* Adding blocklist for macOS 12
- add blocklist for macOS 12
- move nn.functional.conv_transpose2d to the list

* - add allowlist for macOS 13.3
- move nn.functional.conv_transpose2d to the list
- move pow and __rpow__ to the list
- This patch fixes a regression caused by recent MPS module interface. Since we now compile torch._C.is_mps_available()
only if USE_MPS is defined, then it may cause failures on CUDA (and other devices when USE_MPS is not defined) if we upstream.
So, this patch checks if is_mps_available is implemented first and then calls it.
- Also use the unique name `default_mps_generator` to avoid conflicts with CPU default generator
* Update gradient allowlist and blocklist

* Remove prints

* Fix lintrunner
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet