PyTorch Developer Podcast

Compiler collectives

wookim@fb.com (PyTorch) — Sun, 4 Aug 2024 01:47:46 +0000

Compiler collectives are a PT2 feature where by compiler instances across multiple ranks use NCCL collectives to communicate information to other instances. This is used to ensure we consistently decide if inputs or static or dynamic across all ranks. See also PR at https://github.com/pytorch/pytorch/pull/130935

TORCH_TRACE and tlparse

wookim@fb.com (PyTorch) — Mon, 29 Apr 2024 00:01:28 +0000

TORCH_TRACE and tlparse are a structured log and log parser for PyTorch 2. It gives useful information about what code was compiled and what the intermediate build products look like.

Higher order operators

wookim@fb.com (PyTorch) — Sun, 21 Apr 2024 19:28:57 +0000

Higher order operators are a special form of operators in torch.ops which have relaxed input argument requirements: in particular, they can accept any form of argument, including Python callables. Their name is based off of their most common use case, which is to represent higher order functions like control flow operators. However, they are also used to implement other variants of basic operators and can also be used to smuggle in Python data that is quite unusual. They are implemented using a Python dispatcher.

Inductor - Post-grad FX passes

wookim@fb.com (PyTorch) — Fri, 12 Apr 2024 07:00:00 +0000

The post-grad FX passes in Inductor run after AOTAutograd has functionalized and normalized the input program into separate forward/backward graphs. As such, they generally can assume that the graph in question is functionalized, except for some mutations to inputs at the end of the graph. At the end of post-grad passes, there are special passes that reintroduce mutation into the graph before going into the rest of Inductor lowering which is generally aware of passes. The post-grad FX passes are varied but are typically domain specific passes making local changes to specific parts of the graph.

CUDA graph trees

wookim@fb.com (PyTorch) — Sun, 24 Mar 2024 07:00:00 +0000

CUDA graph trees are the internal implementation of CUDA graphs used in PT2 when you say mode="reduce-overhead". Their primary innovation is that they allow the reuse of memory across multiple CUDA graphs, as long as they form a tree structure of potential paths you can go down with the CUDA graph. This greatly reduced the memory usage of CUDA graphs in PT2. There are some operational implications to using CUDA graphs which are described in the podcast.

Min-cut partitioner

wookim@fb.com (PyTorch) — Sun, 17 Mar 2024 07:00:00 +0000

The min-cut partitioner makes decisions about what to save for backwards when splitting the forward and backwards graph from the joint graph traced by AOTAutograd. Crucially, it doesn't actually do a "split"; instead, it is deciding how much of the joint graph should be used for backwards. I also talk about the backward retracing problem.

AOTInductor

wookim@fb.com (PyTorch) — Sat, 2 Mar 2024 08:00:00 +0000

AOTInductor is a feature in PyTorch that lets you export an inference model into a self-contained dynamic library, which can subsequently be loaded and used to run optimized inference. It is aimed primarily at CUDA and CPU inference applications, for situations when your model export once to be exported once while your runtime may still get continuous updates. One of the big underlying organizing principles is a limited ABI which does not include libtorch, which allows these libraries to stay stable over updates to the runtime. There are many export-like use cases you might be interested in using AOTInductor for, and some of the pieces should be useful, but AOTInductor does not necessarily solve them.

Tensor subclasses and PT2

wookim@fb.com (PyTorch) — Sat, 24 Feb 2024 08:00:00 +0000

Tensor subclasses allow you to add extend PyTorch with new types of tensors without having to write any C++. They have been used to implement DTensor, FP8, Nested Jagged Tensor and Complex Tensor. Recent work by Brian Hirsh means that we can compile tensor subclasses in PT2, eliminating their overhead. The basic mechanism by which this compilation works is a desugaring process in AOTAutograd. There are some complications involving views, dynamic shapes and tangent metadata mismatch.

Compiled autograd

wookim@fb.com (PyTorch) — Mon, 19 Feb 2024 08:00:00 +0000

Compiled autograd is an extension to PT2 that permits compiling the entirety of a backward() call in PyTorch. This allows us to fuse accumulate grad nodes as well as trace through arbitrarily complicated Python backward hooks. Compiled autograd is an important part of our plans for compiled DDP/FSDP as well as for whole-graph compilation.

PT2 extension points

wookim@fb.com (PyTorch) — Mon, 5 Feb 2024 09:00:00 +0000

We discuss some extension points for customizing PT2 behavior across Dynamo, AOTAutograd and Inductor.

Inductor - Define-by-run IR

wookim@fb.com (PyTorch) — Wed, 24 Jan 2024 08:00:00 +0000

Define-by-run IR is how Inductor defines the internal compute of a pointwise/reduction operation. It is characterized by a function that calls a number of functions in the 'ops' namespace, where these ops can be overridden by different handlers depending on what kind of semantic analysis you need to do. The ops Inductor supports include regular arithmetic operators, but also memory load/store, indirect indexing, masking and collective operations like reductions.

Unsigned integers

wookim@fb.com (PyTorch) — Wed, 17 Jan 2024 14:00:00 +0000

Traditionally, unsigned integer support in PyTorch was not great; we only support uint8. Recently, we added support for uint16, uint32 and uint64. Bare bones functionality works, but I'm entreating the community to help us build out the rest. In particular, for most operations, we plan to use PT2 to build anything else. But if you have an eager kernel you really need, send us a PR and we'll put it in. While most of the implementation was straightforward, there are some weirdnesses related to type promotion inconsistencies with numpy and dealing with the upper range of uint64. There is also upcoming support for sub-byte dtypes uint1-7, and these will exclusively be implemented via PT2.

Inductor - IR

wookim@fb.com (PyTorch) — Tue, 16 Jan 2024 09:00:00 +0000

Inductor IR is an intermediate representation that lives between ATen FX graphs and the final Triton code generated by Inductor. It was designed to faithfully represent PyTorch semantics and accordingly models views, mutation and striding. When you write a lowering from ATen operators to Inductor IR, you get a TensorBox for each Tensor argument which contains a reference to the underlying IR (via StorageBox, and then a Buffer/ComputedBuffer) that says how the Tensor was computed. The inner computation is represented via define-by-run, which allows for compact definition of IR representation, while still allowing you to extract an FX graph out if you desire. Scheduling then takes buffers of inductor IR and decides what can be fused. Inductor IR may have too many nodes, this would be a good thing to refactor in the future.

Dynamo - VariableTracker

wookim@fb.com (PyTorch) — Fri, 12 Jan 2024 17:40:41 +0000

I talk about VariableTracker in Dynamo. VariableTracker is Dynamo's representation of the Python. I talk about some recent changes, namely eager guards and mutable VT. I also tell you how to find the functionality you care about in VariableTracker (https://docs.google.com/document/d/1XDPNK3iNNShg07jRXDOrMk2V_i66u1hEbPltcsxE-3E/edit#heading=h.i6v7gqw5byv6).

Unbacked SymInts

wookim@fb.com (PyTorch) — Tue, 21 Feb 2023 08:00:00 +0000

This podcast goes over the basics of unbacked SymInts. You might want to listen to this one before listening to https://pytorch-dev-podcast.simplecast.com/episodes/zero-one-specialization Some questions we answer (h/t from Gregory Chanan):

- Are unbacked symints only for export? Because otherwise I could just break / wait for the actual size. But maybe I can save some retracing / graph breaks perf if I have them too? So the correct statement is "primarily" for export?

- Why am I looking into the broadcasting code at all? Naively, I would expect the export graph to be just a list of ATen ops strung together. Why do I recurse that far down? Why can't I annotate DONT_TRACE_ME_BRO?

- How does 0/1 specialization fit into this? I understand we may want to 0/1 specialize in a dynamic shape regime in "eager" mode (is there a better term?), but that doesn't seem to matter for export?

- So far we've mainly been talking about how to handle our own library code. There is a worry about pushing complicated constraints downstream, similar to torchscript. What constraints does this actually push?

Zero-one specialization

wookim@fb.com (PyTorch) — Mon, 20 Feb 2023 08:00:00 +0000

Mikey Dagistes joins me to ask some questions about the recent recent composability sync https://www.youtube.com/watch?v=NJV7YFbtoR4 where we discussed 0/1 specialization and its implications on export in PT2. What's the fuss all about? What do I need to understand about PT2 to understand why 0/1 specialization is a thing?

torchdynamo

wookim@fb.com (PyTorch) — Tue, 6 Dec 2022 08:00:00 +0000

What is torchdynamo? From a bird's eye view, what exactly does it do? What are some important things to know about it? How does it differ from other graph capture mechanisms?

For more reading, check out https://docs.google.com/document/d/13K03JN4gkbr40UMiW4nbZYtsw8NngQwrTRnL3knetGM/edit#

PyTorch 2.0

wookim@fb.com (PyTorch) — Sun, 4 Dec 2022 08:00:00 +0000

Soumith's keynote on PT2.0: https://youtu.be/vbtGZL7IrAw?t=1037
PT2 Manifesto: https://docs.google.com/document/d/1tlgPcR2YmC3PcQuYDPUORFmEaBPQEmo8dsh4eUjnlyI/edit#
PT2 Architecture: https://docs.google.com/document/d/1wpv8D2iwGkKjWyKof9gFdTf8ISszKbq1tsMVm-3hSuU/edit#

History of functorch

wookim@fb.com (PyTorch) — Mon, 7 Nov 2022 23:14:31 +0000

Join me with Richard Zou to talk about the history of functorch. What was the thought process behind the creation of functorch? How did it get started? JAX’s API and model is fairly different from PyTorch’s, how did we validate that it would work in PyTorch? Where did functorch go after the early user studies? Where is it going next?

Learning rate schedulers

wookim@fb.com (PyTorch) — Mon, 13 Jun 2022 16:00:00 +0000

What’s a learning rate? Why might you want to schedule it? How does the LR scheduler API in PyTorch work? What the heck is up with the formula implementation? Why is everything terrible?

Weak references

wookim@fb.com (PyTorch) — Mon, 6 Jun 2022 16:00:00 +0000

What are they good for? (Caches. Private fields.) C++ side support, how it’s implemented / release resources. Python side support, how it’s implemented. Weak ref tensor hazard due to resurrection. Downsides of weak references in C++. Scott Wolchok’s release resources optimization.

Other episodes to listen to first: https://pytorch-dev-podcast.simplecast.com/episodes/reference-counting https://pytorch-dev-podcast.simplecast.com/episodes/pyobject-preservation

Strides

wookim@fb.com (PyTorch) — Mon, 30 May 2022 16:00:00 +0000

Mike Ruberry has an RFC about stride-agnostic operator semantics (https://github.com/pytorch/pytorch/issues/78050), so let's talk about strides. What are they? How are they used to implement views and memory format? How do you handle them properly when writing kernels? In what sense are strides overspecified, and therefore, not worth slavishly reimplementing in a system like PrimTorch? What does Edward think we should do about them?

My blog post that covers strides along with other topics can be found at http://blog.ezyang.com/2019/05/pytorch-internals/

AOTAutograd

wookim@fb.com (PyTorch) — Mon, 9 May 2022 16:00:00 +0000

AOTAutograd is a cool new feature in functorch for capturing both forward and backward traces of PyTorch operators, letting you run them through a compiler and then drop the compiled kernels back into a normal PyTorch eager program. Today, Horace joins me to tell me how it works, what it is good to use for, and what our future plans for it are.

Dispatcher questions with Sherlock

wookim@fb.com (PyTorch) — Mon, 2 May 2022 16:00:00 +0000

Sherlock recently joined the PyTorch team, having previously worked on ONNX Runtime at Microsoft, and Sherlock’s going to ask me some questions about the dispatcher, and I’m going to answer them. We talked about the history of the dispatcher, how to override dispatching order, multiple dispatch, how to organize various dispatch keys and torch function mode. The companion video is at https://youtu.be/_qB2Ho1O3u4

New CI

wookim@fb.com (PyTorch) — Mon, 25 Apr 2022 16:00:00 +0000

PyTorch recently moved all of its CI from CircleCI to GitHub Actions. There were a lot of improvements in the process, making my old podcast about CI obsolete! Today, Eli Uriegas joins me to talk about why we moved to GitHub Actions, how the new CI system is put together, and what some cool features about our new CI.

Python exceptions

wookim@fb.com (PyTorch) — Sun, 17 Apr 2022 16:00:00 +0000

C++ has exceptions, Python has exceptions. But they’re not the same thing! How do exceptions work in CPython, how do we translate exceptions from C++ to Python (hint: it’s different for direct bindings versus pybind11), and what do warnings (which we also translate from C++ to Python) have in common with this infrastructure?

Torch vs ATen APIs

wookim@fb.com (PyTorch) — Mon, 11 Apr 2022 16:00:00 +0000

PyTorch’s torch API is the Python API everyone knows and loves, but there’s also another API, the ATen API, which most of PyTorch’s internal subsystems are built on. How to tell them apart? What implications do these have on our graph mode IR design? Also, a plug for PrimTorch, a new set of operators, not designed for eager mode, that is supposed to be even lower level than ATen.

All about NVIDIA GPUs

wookim@fb.com (Natalia Gimelshein) — Fri, 24 Sep 2021 16:00:00 +0000

PyTorch is in the business of shipping numerical software that can run fast on your CUDA-enabled NVIDIA GPU, but it turns out there is a lot of heterogeneity in NVIDIA’s physical GPU offering and when it comes to what is fast and what is slow, what specific GPU you have on hand matters quite a bit. Yet there are literally hundreds of distinct NVIDIA GPU models on the market, how do you make sense of the madness? Today, Natalia Gimelshein joins me to talk about everything that’s going on in the NVIDIA GPU market, and what, as a framework developer, you have to care about to make sense of it all.

Further reading.

NVIDIA microarchitectures on Wikipedia https://en.wikipedia.org/wiki/Category:Nvidia_microarchitectures
A slightly old post about matching SM to architecture https://arnon.dk/matching-sm-architectures-arch-and-gencode-for-various-nvidia-cards/

Tensor subclasses and Liskov substitution principle

wookim@fb.com (PyTorch) — Thu, 16 Sep 2021 13:00:00 +0000

A lot of recent work going in PyTorch is all about adding new and interesting Tensor subclasses, and this all leads up to the question of, what exactly is OK to make a tensor subclass? One answer to this question comes from an old principle from Barbara Liskov called the Liskov substitution principle, which informally can be stated as S is a subtype of T if anywhere you have T, it can be replaced with S without altering "desirable" properties of this program. In this podcast I'll talk about LSP and how it relates to the design of Tensor subclasses and a hypothetical "abstract Tensor specification" which really doesn't exist but which sort of implicitly exists in the corpus of existing PyTorch programs.

Half precision

wookim@fb.com (PyTorch) — Fri, 10 Sep 2021 13:00:00 +0000

In this episode I talk about reduced precision floating point formats float16 (aka half precision) and bfloat16. I'll discuss what floating point numbers are, how these two formats vary, and some of the practical considerations that arise when you are working with numeric code in PyTorch that also needs to work in reduced precision. Did you know that we do all CUDA computations in float32, even if the source tensors are stored as float16? Now you know!

Further reading.

The Wikipedia article on IEEE floating point is pretty great https://en.wikipedia.org/wiki/IEEE_754
How bfloat16 works out when doing training https://arxiv.org/abs/1905.12322
Definition of acc_type in PyTorch https://github.com/pytorch/pytorch/blob/master/aten/src/ATen/AccumulateType.h

DataLoader with multiple workers leaks memory

wookim@fb.com (PyTorch) — Wed, 1 Sep 2021 13:00:00 +0000

Today I'm going to talk about a famous issue in PyTorch, DataLoader with num_workers > 0 causes memory leak (https://github.com/pytorch/pytorch/issues/13246). This bug is a good opportunity to talk about DataSet/DataLoader design in PyTorch, fork and copy-on-write memory in Linux and Python reference counting; you have to know about all of these things to understand why this bug occurs, but once you do, it also explains why the workarounds help.

Further reading.

A nice summary of the full issue https://github.com/pytorch/pytorch/issues/13246#issuecomment-905703662
DataLoader architecture RFC https://github.com/pytorch/pytorch/issues/49440
Cinder Python https://github.com/facebookincubator/cinder

Batching

wookim@fb.com (PyTorch) — Wed, 18 Aug 2021 13:00:00 +0000

PyTorch operates on its input data in a batched manner, typically processing multiple batches of an input at once (rather than once at a time, as would be the case in typical programming). In this podcast, we talk a little about the implications of batching operations in this way, and then also about how PyTorch's API is structured for batching (hint: poorly) and how Numpy introduced a concept of ufunc/gufuncs to standardize over broadcasting and batching behavior. There is some overlap between this podcast and previous podcasts about TensorIterator and vmap; you may also be interested in those episodes.

Further reading.

ufuncs and gufuncs https://numpy.org/doc/stable/reference/ufuncs.html and https://numpy.org/doc/stable/reference/c-api/generalized-ufuncs.html
A brief taxonomy of PyTorch operators by shape behavior http://blog.ezyang.com/2020/05/a-brief-taxonomy-of-pytorch-operators-by-shape-behavior/
Related episodes on TensorIterator and vmap https://pytorch-dev-podcast.simplecast.com/episodes/tensoriterator and https://pytorch-dev-podcast.simplecast.com/episodes/vmap

Multiple dispatch in __torch_function__

wookim@fb.com (PyTorch) — Tue, 10 Aug 2021 13:00:00 +0000

Python is a single dispatch OO language, but there are some operations such as binary magic methods which implement a simple form of multiple dispatch. torch_function__ (through its Numpy predecessor __array_function) generalizes this mechanism so that invocations of torch.add with different subclasses work properly. This podcast describes how this mechanism works and how it can be used (in an unconventional way) to build composable subclasses ala JAX in functorch.

Further reading:

This podcast in written form https://dev-discuss.pytorch.org/t/functorch-levels-as-dynamically-allocated-classes/294
Multiple dispatch resolution rules in the RFC https://github.com/pytorch/rfcs/blob/master/RFC-0001-torch-function-for-methods.md#process-followed-during-a-functionmethod-call

Multithreading

wookim@fb.com (PyTorch) — Tue, 3 Aug 2021 13:00:00 +0000

Writing multithreading code has always been a pain, and in PyTorch there are buckets and buckets of multithreading related issues you have to be aware about and deal with when writing code that makes use of it. We'll cover how you interface with multithreading in PyTorch, what goes into implementing those interfaces (thread pools!) and also some miscellaneous stuff like TLS, forks and data structure thread safety that is also relevant.

Asynchronous versus synchronous execution

wookim@fb.com (PyTorch) — Tue, 27 Jul 2021 13:00:00 +0000

CUDA is asynchronous, CPU is synchronous. Making them play well together can be one of the more thorny and easy to get wrong aspects of the PyTorch API. I talk about why non_blocking is difficult to use correctly, a hypothetical "asynchronous CPU" device which would help smooth over some of the API problems and also why it used to be difficult to implement async CPU (but it's not hard anymore!) At the end, I also briefly talk about how async/sync impedance can also show up in unusual places, namely the CUDA caching allocator.

Further reading.

CUDA semantics which discuss non_blocking somewhat https://pytorch.org/docs/stable/notes/cuda.html
Issue requesting async cpu https://github.com/pytorch/pytorch/issues/44343

gradcheck

wookim@fb.com (PyTorch) — Fri, 23 Jul 2021 13:00:00 +0000

We talk about gradcheck, the property based testing mechanism that we use to verify the correctness of analytic gradient formulas in PyTorch. I'll talk a bit about testing in general, property based testing and why gradcheck is a particularly useful property based test. There will be some calculus, although I've tried to keep the math mostly to intuitions and pointers on what to read up on elsewhere.

Further reading.

Gradcheck mechanics, a detailed mathematical explanation of how it works https://pytorch.org/docs/stable/notes/gradcheck.html In particular, it also explains how gradcheck extends to complex numbers
JAX has a pretty good explanation about vjp and jvp at https://jax.readthedocs.io/en/latest/notebooks/autodiff_cookbook.html
Fast gradcheck tracking issue https://github.com/pytorch/pytorch/issues/53876

torch.use_deterministic_algorithms

wookim@fb.com (PyTorch) — Wed, 21 Jul 2021 13:00:00 +0000

torch.use_deterministic_algorithms lets you force PyTorch to use deterministic algorithms. It's very useful for debugging!

There are some errors in the recording: the feature is called torch.use_deterministic_algorithms, and there is not actually a capability to warn (this was in an old version of the PR but taken out), we just error if you hit nondeterministic code.

Docs: https://pytorch.org/docs/stable/generated/torch.use_deterministic_algorithms.html#torch.use_deterministic_algorithms

Reference counting

wookim@fb.com (PyTorch) — Tue, 20 Jul 2021 13:00:00 +0000

Reference counting is a common memory management technique in C++ but PyTorch does its reference counting in a slightly idiosyncratic way using intrusive_ptr. We'll talk about why intrusive_ptr exists, the reason why refcount bumps are slow in C++ (but not in Python), what's up with const Tensor& everywhere, why the const is a lie and how TensorRef lets you create a const Tensor& from a TensorImpl* without needing to bump your reference count.

Further reading.

Why you shouldn't feel bad about passing tensor by reference https://dev-discuss.pytorch.org/t/we-shouldnt-feel-bad-about-passing-tensor-by-reference/85
Const correctness in PyTorch https://github.com/zdevito/ATen/issues/27
TensorRef RFC https://github.com/pytorch/rfcs/pull/16

Memory layout

wookim@fb.com (PyTorch) — Tue, 13 Jul 2021 13:00:00 +0000

Memory layout specifies how the logical multi-dimensional tensor maps its elements onto physical linear memory. Some layouts admit more efficient implementations, e.g., NCHW versus NHWC. Memory layout makes use of striding to allow users to conveniently represent their tensors with different physical layouts without having to explicitly tell every operator what to do.

Further reading.

Tutorial https://pytorch.org/tutorials/intermediate/memory_format_tutorial.html
Memory format RFC https://github.com/pytorch/pytorch/issues/19092
Layout permutation proposal (not implemented) https://github.com/pytorch/pytorch/issues/32078

pytorch-probot

wookim@fb.com (PyTorch) — Mon, 12 Jul 2021 14:28:17 +0000

pytorch-probot is a GitHub application that we use to automate common tasks in GitHub. I talk about what it does and some design philosophy for it. Repo is at: https://github.com/pytorch/pytorch-probot

API design via lexical and dynamic scoping

wookim@fb.com (PyTorch) — Fri, 9 Jul 2021 13:00:00 +0000

Lexical and dynamic scoping are useful tools to reason about various API design choices in PyTorch, related to context managers, global flags, dynamic dispatch, and how to deal with BC-breaking changes. I'll walk through three case studies, one from Python itself (changing the meaning of division to true division), and two from PyTorch (device context managers, and torch function for factory functions).

Further reading.

Me unsuccessfully asking around if there was a way to simulate __future__ in libraries https://stackoverflow.com/questions/66927362/way-to-opt-into-bc-breaking-changes-on-methods-within-a-single-module
A very old issue asking for a way to change the default GPU device https://github.com/pytorch/pytorch/issues/260 and a global GPU flag https://github.com/pytorch/pytorch/issues/7535
A more modern issue based off the lexical module idea https://github.com/pytorch/pytorch/issues/27878
Array module NEP https://numpy.org/neps/nep-0037-array-module.html

Intro to distributed

wookim@fb.com (PyTorch) — Thu, 8 Jul 2021 13:00:00 +0000

Today, Shen Li (mrshenli) joins me to talk about distributed computation in PyTorch. What is distributed? What kinds of things go into making distributed work in PyTorch? What's up with all of the optimizations people want to do here?

Further reading.

PyTorch distributed overview https://pytorch.org/tutorials/beginner/dist_overview.html
Distributed data parallel https://pytorch.org/docs/stable/notes/ddp.html

Double backwards

wookim@fb.com (PyTorch) — Wed, 7 Jul 2021 13:00:00 +0000

Double backwards is PyTorch's way of implementing higher order differentiation. Why might you want it? How does it work? What are some of the weird things that happen when you do this?

Further reading.

Epic PR that added double backwards support for convolution initially https://github.com/pytorch/pytorch/pull/1643

Functional modules

wookim@fb.com (PyTorch) — Tue, 6 Jul 2021 13:00:00 +0000

Functional modules are a proposed mechanism to take PyTorch's existing NN module API and transform it into a functional form, where all the parameters are explicit argument. Why would you want to do this? What does functorch have to do with it? How come PyTorch's existing APIs don't seem to need this? What are the design problems?

Further reading.

Proposal in GitHub issues https://github.com/pytorch/pytorch/issues/49171
Linen design in flax https://flax.readthedocs.io/en/latest/design_notes/linen_design_principles.html

CUDA graphs

wookim@fb.com (PyTorch) — Mon, 28 Jun 2021 13:00:00 +0000

What are CUDA graphs? How are they implemented? What does it take to actually use them in PyTorch?

Further reading.

NVIDIA has docs on CUDA graphs https://developer.nvidia.com/blog/cuda-graphs/
Nuts and bolts implementation PRs from mcarilli: https://github.com/pytorch/pytorch/pull/51436 https://github.com/pytorch/pytorch/pull/46148

Default arguments

wookim@fb.com (PyTorch) — Fri, 25 Jun 2021 13:00:00 +0000

What do default arguments have to do with PyTorch design? Why are default arguments great for clients (call sites) but not for servers (implementation sites)? In what sense are default arguments a canonicalization to max arity? What problems does this canonicalization cause? Can you canonicalize to minimum arity? What are some lessons to take?

Further reading. https://github.com/pytorch/pytorch/issues/54613 stop serializing default arguments

Anatomy of a domain library

wookim@fb.com (PyTorch) — Thu, 24 Jun 2021 13:00:00 +0000

What's a domain library? Why do they exist? What do they do for you? What should you know about developing in PyTorch main library versus in a domain library? How coupled are they with PyTorch as a whole? What's cool about working on domain libraries?

Further reading.

The classic trio of domain libraries is https://pytorch.org/audio/stable/index.html https://pytorch.org/text/stable/index.html and https://pytorch.org/vision/stable/index.html

Line notes.

why do domain libraries exist? lots of domains specific gadgets,
inappropriate for PyTorch
what does a domain library do
- operator implementations (old days: pure python, not anymore)
  - with autograd support and cuda acceleration
  - esp encoding/decoding, e.g., for domain file formats
    - torchbind for custom objects
    - takes care of getting the dependencies for you
  - esp transformations, e.g., for data augmentation
- models, esp pretrained weights
- datasets
- reference scripts
- full wheel/conda packaging like pytorch
- mobile compatibility
separate repos: external contributors with direct access
- manual sync to fbcode; a lot easier to land code! less
  motion so lower risk
coupling with pytorch? CI typically runs on nightlies
- pytorch itself tests against torchvision, canary against
  extensibility mechanisms
- mostly not using internal tools (e.g., TensorIterator),
  too unstable (this would be good to fix)
closer to research side of pytorch; francesco also part of papers

TensorAccessor

wookim@fb.com (PyTorch) — Wed, 23 Jun 2021 13:00:00 +0000

What's TensorAccessor? Why not just use a raw pointer? What's PackedTensorAccessor? What are some future directions for mixing statically typed and typed erase code inside PyTorch proper?

Further reading.

TensorAccessor source code, short and sweet https://github.com/pytorch/pytorch/blob/master/aten/src/ATen/core/TensorAccessor.h
Legacy THCDeviceTensor https://github.com/pytorch/pytorch/blob/master/aten/src/THC/THCDeviceTensor.cuh

Random number generators

wookim@fb.com (PyTorch) — Tue, 22 Jun 2021 13:00:00 +0000

Why are RNGs important? What is the generator concept? How do PyTorch's CPU and CUDA RNGs differ? What are some of the reasons why Philox is a good RNG for CUDA? Why doesn't the generator class have virtual methods for getting random numbers? What's with the next normal double and what does it have to do with Box Muller transform? What's up with csprng?

Further reading.

CUDAGeneratorImpl has good notes about CUDA graph interaction and pointers to all of the rest of the stuff https://github.com/pytorch/pytorch/blob/1dee99c973fda55e1e9cac3d50b4d4982b6c6c26/aten/src/ATen/CUDAGeneratorImpl.h
Transform uniformly distributed random numbers to other distributions with https://github.com/pytorch/pytorch/blob/master/aten/src/ATen/core/TransformationHelper.h
torchcsprng https://github.com/pytorch/csprng

vmap

wookim@fb.com (PyTorch) — Mon, 21 Jun 2021 13:00:00 +0000

What is vmap? How is it implemented? How does our implementation compare to JAX's? What is a good way of understanding what vmap does? What's up with random numbers? Why are there some issues with the vmap that PyTorch currently ships?

Further reading.

Tracking issue for vmap support https://github.com/pytorch/pytorch/issues/42368
BatchedTensor source code https://github.com/pytorch/pytorch/blob/master/aten/src/ATen/BatchedTensorImpl.h , logical-physical transformation helper code https://github.com/pytorch/pytorch/blob/master/aten/src/ATen/VmapTransforms.h (well documented, worth a read)
functorch, the better, more JAX-y implementation of vmap https://github.com/facebookresearch/functorch
Autodidax https://jax.readthedocs.io/en/latest/autodidax.html which contains a super simple vmap implementation that is a good model for the internal implementation that PyTorch has

Expect tests

wookim@fb.com (PyTorch) — Fri, 18 Jun 2021 13:00:00 +0000

What's an expect test? Why should you use them? Why is inline expect test better than out of line? How to write a good expect test?

Further reading. expecttest source implementation https://github.com/pytorch/pytorch/blob/master/torch/testing/_internal/expecttest.py (only 311 lines!)

XLA

wookim@fb.com (PyTorch) — Thu, 17 Jun 2021 13:00:00 +0000

What's PyTorch XLA? Why should you care? How is it implemented? How does PyTorch XLA trade off functionality versus ease of performance debugging? What are some new developments in this space?

Further reading.

XLA's repo has lots of really good docs. Check out https://github.com/pytorch/xla/blob/master/OP_LOWERING_GUIDE.md and also the main https://github.com/pytorch/xla/blob/master/README.md
Alex Suhan's RFC about lazy core https://github.com/pytorch/rfcs/pull/18

TH

wookim@fb.com (PyTorch) — Wed, 16 Jun 2021 13:00:00 +0000

What is TH? Why might you care? What is so horrible about it? What the heck is the generic/ folder? Why are we porting everything to C++? What are some downsides of having ported all our TH code to C++?

Further reading.

The TH to ATen porting guide has lots of explanations of old school TH idioms https://github.com/pytorch/pytorch/wiki/TH-to-ATen-porting-guide
Old notes about refcounting in TH https://github.com/pytorch/pytorch/blob/master/aten/src/README.md

TorchScript

wookim@fb.com (PyTorch) — Tue, 15 Jun 2021 13:00:00 +0000

There is a really good TorchScript overview at https://github.com/pytorch/pytorch/blob/master/torch/csrc/jit/OVERVIEW.md and in this 20min podcast, I want to give you some of the highlights from this document.

CMake

wookim@fb.com (PyTorch) — Mon, 14 Jun 2021 13:00:00 +0000

Why is PyTorch's build so g-dang complicated. How to avoid having to deal with cmake at all? And if you have to deal with cmake, what are the most important things to know? And if you were going to improve our cmake, how would you go about doing it...

Further reading.

The official CMake documentation is a great help and well worth reading https://cmake.org/documentation
If you work in torch/csrc chances are you'll need to edit this file https://github.com/pytorch/pytorch/blob/master/tools/build_variables.bzl

Liner notes.

multiple build systems: cmake, buck, xplat buck, ovrsource buck, bazel
- tools/build_variables.bzl is read from cmake! append_filelist
  - but not used uniformly for all components! (ouch!)
mashed together ATen and Caffe2 build systems (e.g., main library libtorch_cpu is defined in caffe2/CMakeLists.txt)
cmake: not very much syntax, "everything is a function". This means you can look up constructs relatively easily; e.g., even if() is a command
the general cmake model: "set a bunch of variables, run a bunch of commands". cmake is VERY GREPPABLE
- but not everything is in CMakeLists.txt; check *.cmake too
- the directory structure makes no sense, you really need to grep.
  (doing a lot of set PARENT_SCOPE to propagate stuff)
- renaming a file? grep for it
- primary hazard of refactoring: need to make sure all the variables
  are setup at the new location
many directories are not recursive glob, beware of adding new directories
old school cmake: literally everything is stuffed in variables (CMAKE_CXX_FLAGS). new school cmake: attach things to targets, things propagate when you depend on targets (public/private dependencies)
add_library: the most important thing
don't randomly change things and pray: have hypotheses and test them

torchdeploy

wookim@fb.com (PyTorch) — Fri, 11 Jun 2021 13:00:00 +0000

torchdeploy is a way of running multiple Python interpreters inside the same process. It can be used to deploy Python PyTorch programs in situations where the GIL is a problem, not the CPython interpreter. How does it work, and what kind of challenges does it pose for people who want to write code that calls from C++ to Python?

Further reading.

How the torchdeploy build system works https://dev-discuss.pytorch.org/t/torch-deploy-the-build/238
Description of the single interpreter per Tensor invariant https://github.com/pytorch/pytorch/issues/57756
Recent work on making it possible to load C extensions into torchdeploy https://dev-discuss.pytorch.org/t/running-multiple-python-interpreters-via-custom-dynamic-loading/241

C++ frontend

wookim@fb.com (PyTorch) — Thu, 10 Jun 2021 13:00:00 +0000

What's the C++ frontend? Why is avoiding templates so important? Why is Tensor a reference type? How do we simulate keyword arguments in C++? Where did the nn Module support in the C++ API come from? Why did we reimplement all modules in C++? How are modules implemented in C++? What are some performance challenges of writing Python in C++, and how are we working around them?

Further reading.

C++ frontend tutorial https://pytorch.org/tutorials/advanced/cpp_frontend.html
Writing Python in C++ (a manifesto) https://github.com/pytorch/pytorch/wiki/Writing-Python-in-cpp-(a-manifesto)
MaybeOwned PR https://github.com/pytorch/pytorch/pull/53317

PyObject preservation

wookim@fb.com (PyTorch) — Wed, 9 Jun 2021 13:00:00 +0000

Given two separately refcounted objects, how can you arrange for each of them to stay live so long as the other is live? Why doesn't just having a strong-strong or strong-weak reference between the two objects work? What is object resurrection in CPython? What's a finalizer and why does it make things more complicated? How does Python GC work?

Further reading.

PyObject preservation PR https://github.com/pytorch/pytorch/pull/56017
Sam Gross's original PoC, which works fine if the two objects in question are both PyObjects https://github.com/colesbury/refcount/
PEP 442 Safe object finalization https://www.python.org/dev/peps/pep-0442/
Essential reading about Python GC https://devguide.python.org/garbage_collector/

Mobile selective build

wookim@fb.com (PyTorch) — Tue, 8 Jun 2021 13:00:00 +0000

What is mobile selective build? Why are we so obsessed with reducing binary size? How does selective build work? Why doesn't static linking just work? Why can't you just read out the ops used in a TorchScript model to determine what operators you actually need? What are the tradeoffs of statically determining the operator dependency graph versus tracing? What's up with the SELECTIVE_NAME macro? How the heck does selective build work at all when you have multiple mobile apps in a single Buck build system? What takeaways should I have as a regular PyTorch developer?

Further reading:

Official open source mobile documentation on custom selective builds https://pytorch.org/mobile/android/#custom-build
How to rebuild the op dependency yaml https://github.com/pytorch/pytorch/blob/master/tools/code_analyzer/build.sh

Liner notes:

binary size is premium; ship only what you actually need
big idea:
- get the ops your model needs -> apply this to build of pytorch
get the ops your model needs
- TorchScript ~> read it out directly from the model itself
- but what if ops use other ops?
  - need a dependency graph. done with static analysis llvm (jiakai) ~> with a (possibly inaccurate) yaml checked in for easy kickstart if you don't want to run the pass (updated by bot, not operational since Feb, recommend rebuilding from scratch if you run into trouble)
- other possibility: dynamic tracing
  - pro: no need for dependency graph, just look at what was called; works for dtypes
  - con: need representative inputs, if control flow might not cover everything
apply this to build of pytorch
- ordinarily: static linking ensures stuff that isn't used gets pruned
  - but this doesn't work with distributed operator registration based on static initializers
- how?
  - codegen - just don't generate it
  - no codegen - SELECTIVE_NAME - C++ doesn't support string in macro
- build system integration
  - buck constraint: only one library
    - therefore: generate multiple copies of glue library
  - alt: atomize library into each operator. caffe2 used to do this; each library takes a long time to build (1m) and crashes xcode because there's too many
common hiccups
- modify implementation details, some op is/isn't called anymore ~> error! usually just means some yaml needs regenerating. PyTorch Edge developers are very friendly and can help

torch.nn

wookim@fb.com (PyTorch) — Mon, 7 Jun 2021 13:00:00 +0000

What goes into the implementation of torch.nn? Why do NN modules exist in the first place? What's the function of Parameter? How do modules actually track all the parameters in question? What is all of the goop in the top level NN module class? What are some new developments in torch.nn modules? What are some open problems with our modules?

Further reading:

Implementation of nn.Module https://github.com/pytorch/pytorch/blob/master/torch/nn/modules/module.py
nn.Module is complicated and that means its sometimes a bit slow. Some analysis at https://dev-discuss.pytorch.org/t/overhead-in-nn-module-causing-massive-slowdowns-compared-to-raw-cublas-or-torchscript/110
Lazy modules PR https://github.com/pytorch/pytorch/pull/44538 and factory kwargs https://github.com/pytorch/pytorch/pull/54508

Liner notes:

python for hackability (c++ is reimplemented)
parameters
- parameter collection (for optimization)
- buffers: not considered optimizable
modules
- functorial operation (_apply)
- jit script: staged computation (init is not scripted)
- __call__ to forward (extra instrumentation)
- serialization / state_dict
new stuff: device kwarg (joel schlosser)
new stuff: lazy modules (emcastillo)
open problems: parameter initialization

Code generation

wookim@fb.com (PyTorch) — Fri, 4 Jun 2021 13:00:00 +0000

Why does PyTorch use code generation as part of its build process? Why doesn't it use C++ templates? What things is code generation used for? What are the pros/consof using code generation? What are some other ways to do the same things we currently do with code generation?

Further reading.

Top level file for the new code generation pipeline https://github.com/pytorch/pytorch/blob/master/tools/codegen/gen.py
Out of tree external backend code generation from Brian Hirsh: https://github.com/pytorch/xla/issues/2871
Documentation for native_functions.yaml https://github.com/pytorch/pytorch/blob/master/aten/src/ATen/native/README.md (have you seen this README before? Yes you've seen this README before. Imma post it again.)

Outline:

High level: reduce the amount of code in PyTorch, easier to develop
Strongly typed python
Stuff we're using codegen for
- Meta point: stuff c++ metaprogramming can't do
- C++ apis (functions, methods on classes)
  - Especially for forwarding (operator dot doko)
  - Prototypes for c++ to implement
- YAML files used by external frameworks for binding (accidental)
- Python arg parsing
- pyi generation
- Autograd classes for saving saved data
- Otherwise complicated constexpr computation (e.g., parsing JIT
  schema)
Pros
- Better surface syntax (native_functions.yaml, jit schema,
  derivatives.yaml)
- Better error messages (template messages famously bad)
- Easier to organize complicated code; esp nontrivial input
  data structure
- Easier to debug by looking at generated code
Con
- Not as portable (template can be used by anyone)
- Less good modeling for C++ type based metaprogramming (we've replicated a crappy version of C++ type system in our codegen)
Counterpoints in the design space
- C++ templates: just as efficient
- Boxed fallback: simpler, less efficient
Open question: can you have best of both worlds, e.g., with partially evaluated interpreters?

Why is autograd so complicated

wookim@fb.com (PyTorch) — Thu, 3 Jun 2021 13:00:00 +0000

Why is autograd so complicated? What are the constraints and features that go into making it complicated? What's up with it being written in C++? What's with derivatives.yaml and code generation? What's going on with views and mutation? What's up with hooks and anomaly mode? What's reentrant execution? Why is it relevant to checkpointing? What's the distributed autograd engine?

Further reading.

Autograd notes in the docs https://pytorch.org/docs/stable/notes/autograd.html
derivatives.yaml https://github.com/pytorch/pytorch/blob/master/tools/autograd/derivatives.yaml
Paper on autograd engine in PyTorch https://openreview.net/pdf/25b8eee6c373d48b84e5e9c6e10e7cbbbce4ac73.pdf

__torch_function__

wookim@fb.com (PyTorch) — Wed, 2 Jun 2021 13:00:00 +0000

What is __torch_function__? Why would I want to use it? What does it have to do with keeping extra metadata on Tensors or torch.fx? How is it implemented? Why is __torch_function__ a really popular way of extending functionality in PyTorch? What makes it different from the dispatcher extensibility mechanism? What are some downsides of it being written this way? What are we doing about it?

Further reading.

__torch_function__ RFC: https://github.com/pytorch/rfcs/blob/master/RFC-0001-torch-function-for-methods.md
One of the original GitHub issues tracking the overall design discussion https://github.com/pytorch/pytorch/issues/24015
Documentation for using __torch_function__ https://pytorch.org/docs/stable/notes/extending.html#extending-torch

TensorIterator

wookim@fb.com (PyTorch) — Tue, 1 Jun 2021 13:00:00 +0000

You walk into the whiteboard room to do a technical interview. The interviewer looks you straight in the eye and says, "OK, can you show me how to add the elements of two lists together?" Confused, you write down a simple for loop that iterates through each element and adds them together. Your interviewer rubs his hands together evilly and cackles, "OK, let's make it more complicated."

What does TensorIterator do? Why the heck is TensorIterator so complicated? What's going on with broadcasting? Type promotion? Overlap checks? Layout? Dimension coalescing? Parallelization? Vectorization?

Further reading.

PyTorch TensorIterator internals https://labs.quansight.org/blog/2020/04/pytorch-tensoriterator-internals/
Why is TensorIterator so slow https://dev-discuss.pytorch.org/t/comparing-the-performance-of-0-4-1-and-master/136
Broadcasting https://pytorch.org/docs/stable/notes/broadcasting.html and type promotion https://pytorch.org/docs/stable/tensor_attributes.html#type-promotion-doc

native_functions.yaml

wookim@fb.com (PyTorch) — Fri, 28 May 2021 13:00:00 +0000

What does native_functions.yaml have to do with the TorchScript compiler? What multiple use cases is native_functions.yaml trying to serve? What's up with the JIT schema type system? Why isn't it just Python types? What the heck is the (a!) thingy inside the schema? Why is it important that I actually annotate all of my functions accurately with this information? Why is my seemingly BC change to native_functions.yaml actually breaking people's code? Do I have to understand the entire compiler to understand how to work with these systems?

Further reading.

native_functions.yaml README https://github.com/pytorch/pytorch/blob/master/aten/src/ATen/native/README.md
Tracking issue for serializing default arguments https://github.com/pytorch/pytorch/issues/54613
Test for BC breaking changes in native_functions.yaml https://github.com/pytorch/pytorch/blob/master/test/backward_compatibility/check_backward_compatibility.py

Serialization

wookim@fb.com (PyTorch) — Thu, 27 May 2021 13:00:00 +0000

What is serialization? Why do I care about it? How is serialization done in general in Python? How does pickling work? How does PyTorch implement pickling for its objects? What are some pitfalls of pickling implementation? What does backwards compatibility and forwards compatibility mean in the context of serialization? What's the difference between directly pickling and using torch.save/load? So what the heck is up with JIT/TorchScript serialization? Why did we use zip files? What were some design principles for the serialization format? Why are there two implementations of serialization in PyTorch? Is the fact that PyTorch uses pickling for serialization mean that our serialization format is insecure?

Further reading.

TorchScript serialization design doc https://github.com/pytorch/pytorch/blob/master/torch/csrc/jit/docs/serialization.md
Evolution of serialization formats over time https://github.com/pytorch/pytorch/issues/31877
Code pointers:
- Tensor __reduce_ex__ https://github.com/pytorch/pytorch/blob/de845020a0da39e621db984515bc1cce03f526ea/torch/_tensor.py#L97-L178
- Python side serialization https://github.com/pytorch/pytorch/blob/de845020a0da39e621db984515bc1cce03f526ea/torch/serialization.py#L384-L499
- C++ side serialization https://github.com/pytorch/pytorch/tree/master/torch/csrc/jit/serialization

Continuous integration

wookim@fb.com (PyTorch) — Wed, 26 May 2021 13:00:00 +0000

How is our CI put together? What is the history of the CI? What constraints are under the CI? Why does the CI use Docker? Why are build and test split into two phases? Why are some parts of the CI so convoluted? How does the HUD work? What kinds of configurations is PyTorch tested under? How did we decide what configurations to test? What are some of the weird CI configurations? What's up with the XLA CI? What's going on with the Facebook internal builds?

Further reading.

The CI HUD for viewing the status of master https://ezyang.github.io/pytorch-ci-hud/build/pytorch-master
Structure of CI https://github.com/pytorch/pytorch/blob/master/.circleci/README.md
How to debug Windows problems on CircleCI https://github.com/pytorch/pytorch/wiki/Debugging-Windows-with-Remote-Desktop-or-CDB-(CLI-windbg)-on-CircleCI

Stacked diffs and ghstack

wookim@fb.com (PyTorch) — Tue, 25 May 2021 13:00:00 +0000

What's a stacked diff? Why might you want to do it? What does the workflow for stacked diffs with ghstack look like? How to use interactive rebase to edit earlier diffs in my stack? How can you actually submit a stacked diff to PyTorch? What are some things to be aware of when using ghstack?

Further reading.

The ghstack repository https://github.com/ezyang/ghstack/
A decent explanation of how the stacked diff workflow works on Phabricator, including how to do rebases https://kurtisnusbaum.medium.com/stacked-diffs-keeping-phabricator-diffs-small-d9964f4dcfa6

Shared memory

wookim@fb.com (PyTorch) — Mon, 24 May 2021 13:00:00 +0000

What is shared memory? How is it used in your operating system? How is it used in PyTorch? What's shared memory good for in deep learning? Why use multiple processes rather than one process on a single node? What's the point of PyTorch's shared memory manager? How are allocators for shared memory implemented? How does CUDA shared memory work? What is the difference between CUDA shared memory and CPU shared memory? How did we implement safer CUDA shared memory?

Further reading.

Implementations of vanilla shared memory allocator https://github.com/pytorch/pytorch/blob/master/aten/src/TH/THAllocator.cpp and the fancy managed allocator https://github.com/pytorch/pytorch/blob/master/torch/lib/libshm/libshm.h
Multiprocessing best practices describes some things one should be careful about when working with shared memory https://pytorch.org/docs/stable/notes/multiprocessing.html
More details on how CUDA shared memory works https://pytorch.org/docs/stable/multiprocessing.html#multiprocessing-cuda-sharing-details

Automatic mixed precision

wookim@fb.com (PyTorch) — Fri, 21 May 2021 13:00:00 +0000

What is automatic mixed precision? How is it implemented? What does it have to do with mode dispatch keys, fallthrough kernels? What are AMP policies? How is its cast caching implemented? How does torchvision also support AMP? What's up with Intel's CPU autocast implementation?

Further reading.

Autocast implementation lives at https://github.com/pytorch/pytorch/blob/master/aten/src/ATen/autocast_mode.cpp
How to add autocast implementations to custom operators that are out of tree https://pytorch.org/tutorials/advanced/dispatcher.html#autocast
CPU autocasting PR https://github.com/pytorch/pytorch/pull/57386

Conjugate views

wookim@fb.com (PyTorch) — Thu, 20 May 2021 13:00:00 +0000

What are complex numbers? What is conjugation? Why is conjugation so common in linear algebra? Why would we like conjugation to behave similarly to transposition (and why is matrix multiply with a transposed input so fast?) What is a conjugate view? How is it implemented? What's the relationship between views, laziness and call-by-name evaluation?

Further reading.

Pull request that adds conjugate views https://github.com/pytorch/pytorch/pull/54987
The idea of conjugate views originally came up when we were deciding which complex autograd convention to use in https://github.com/pytorch/pytorch/issues/41857 . PyTorch uses the conjugate Wirtinger derivative which, true to its name, involves a lot of conjugations in its formulas.
Conjugate views are a form of bidirectional lens. This nice presentation explains what the concept is https://www.cis.upenn.edu/~bcpierce/papers/lenses-etapsslides.pdf

History and constraints of Tensor

wookim@fb.com (PyTorch) — Wed, 19 May 2021 13:00:00 +0000

What historical constraints and design choices lead to the design of Tensor/Storage (and their Impl variants) as they are today? Why do we use intrusive refcounting? Why are we trying to get rid of virtual methods on TensorImpl? Why are there so many frickin' bitfields?

Further reading.

PyTorch internals blog post http://blog.ezyang.com/2019/05/pytorch-internals/
Writing Python in C++, a manifesto https://github.com/pytorch/pytorch/wiki/Writing-Python-in-cpp-(a-manifesto)
At time of writing, the breakdown of all fields on TensorImpl https://github.com/pytorch/pytorch/blob/71f4c5c1f436258adc303b710efb3f41b2d50c4e/c10/core/TensorImpl.h#L2155-L2177

How new operators are authored

wookim@fb.com (PyTorch) — Tue, 18 May 2021 13:00:00 +0000

What's the general process by which a new operator is added to PyTorch? Why is this actually something of a rare occurrence? How do you integrate an operator with the rest of PyTorch's system so it can be run end-to-end? What should I expect if I'm writing a CPU and CUDA kernel? What tools are available to me to make the job easier? How can I debug my kernels? How do I test them?

Further reading.

The README for the native/ directory, where all kernels get put https://github.com/pytorch/pytorch/blob/master/aten/src/ATen/native/README.md
A high level overview of how TensorIterator works https://labs.quansight.org/blog/2020/04/pytorch-tensoriterator-internals/
Where OpInfos live https://github.com/pytorch/pytorch/blob/master/torch/testing/_internal/common_methods_invocations.py

The life and death of Variable

wookim@fb.com (PyTorch) — Mon, 17 May 2021 13:00:00 +0000

What is a Variable? Why did it exist as a wrapper in the first place? Why did it get removed? How did we remove it? What are some of the lingering consequences of its removal?

Further reading:

The release notes of PyTorch 0.4 do a good job explaining the user visible consequences of the removal, at the time, including how we "simulate" concepts on Variable that don't make sense anymore https://pytorch.org/blog/pytorch-0_4_0-migration-guide/
Part 1: Removal of Variable wrapper in C++ https://github.com/pytorch/pytorch/pull/17072
Part 2: Merge of Variable and Tensor types in C++ https://github.com/pytorch/pytorch/pull/28620

Backend extensibility

wookim@fb.com (PyTorch) — Fri, 14 May 2021 13:00:00 +0000

What's the current state of backend extensibility? How did PyTorch evolve from being a CPU and CUDA only framework to also support AMD ROCm and XLA? What are some problems with adding an out-of-tree backend, and what's some work to make it better?

Further reading:

Script for HIPifying PyTorch's source when enabling ROCm https://github.com/pytorch/pytorch/blob/master/tools/amd_build/build_amd.py
PyTorch/XLA https://github.com/pytorch/xla/
Brian Hirsh's spec on what out-of-tree backend codegen looks like https://github.com/pytorch/xla/issues/2871

The road to structured kernels

wookim@fb.com (PyTorch) — Thu, 13 May 2021 13:00:00 +0000

Structured kernels are a new way to write kernels in PyTorch. Why did they take so long? What finally convinced us that we should do them? Why did it end up taking me the better part of a year to only be half done with them?

Further reading:

Structured kernels RFC https://github.com/pytorch/rfcs/blob/rfc-0005/RFC-0005-structured-kernel-definitions.md
Taxonomy of PyTorch operators by shape behavior http://blog.ezyang.com/2020/05/a-brief-taxonomy-of-pytorch-operators-by-shape-behavior/
Bram Wasti's lazy tensor prototype https://github.com/pytorch/pytorch/pull/25753

Functionalization

wookim@fb.com (PyTorch) — Wed, 12 May 2021 13:00:00 +0000

Functionalization is the process by which we remove mutation from autograd graphs in PyTorch, leaving us with a purely functional graph that we can execute in the normal way. Why do we need to do functionalization? What makes it not so easy to do? How do we do it? And how does it compare to mutation removal that you might see in a compiler?

Further reading:

Section 3.1 of this paper on PyTorch AD https://openreview.net/pdf/25b8eee6c373d48b84e5e9c6e10e7cbbbce4ac73.pdf predates our implementation of inplace autograd but accurately reports the subtleties and correctly predicts the implementation strategy we ended up taking
RFC to generalize the functionalization mechanism to be available to arbitrary backends https://github.com/pytorch/rfcs/pull/19
Code that handles lazily updating views when the base is updated https://github.com/pytorch/pytorch/blob/e5e095cbe4dbc5a601f98e6134dcbd59c6342d7d/torch/csrc/autograd/variable.cpp#L556-L603

Just enough CUDA to be dangerous

wookim@fb.com (PyTorch) — Tue, 11 May 2021 13:00:00 +0000

Ever wanted to learn about CUDA but not sure where to start? In this sixteen minute episode I try to jam in as much CUDA knowledge as could be reasonably expected in a podcast. You won't know how to write a kernel after this episode, but you'll know about what a GPU is, what the general CUDA programming model is, why asynchronous execution makes everything complicated, and some general principles PyTorch abides by when designing CUDA kernels.

Further reading:

PyTorch docs on CUDA semantics https://pytorch.org/docs/stable/notes/cuda.html
The book I was recommended for learning CUDA when I first showed up at PyToch: Programming Massively Parallel Processors https://www.amazon.com/Programming-Massively-Parallel-Processors-Hands/dp/0128119861
The environment variable that makes CUDA synchronous is CUDA_LAUNCH_BLOCKING=1. cuda-memcheck is also useful for debugging CUDA problems https://docs.nvidia.com/cuda/cuda-memcheck/index.html

Inference mode

wookim@fb.com (PyTorch) — Mon, 10 May 2021 13:00:00 +0000

What's inference mode? Why doesn't my code run fast if I use no_grad or make sure requires_grad=False? How come inference mode is safe but AutoNonVariableTypeMode is not?

Further reading:

Inference mode RFC https://github.com/.../rfc0011/RFC-0011-InferenceMode.md
Inference mode docs for C++ frontend users https://github.com/.../cpp/source/notes/inference_mode.rst
Tracking issue for Python frontend support https://github.com/pytorch/pytorch/issues/56608

Vectorization

wookim@fb.com (PyTorch) — Fri, 7 May 2021 13:00:00 +0000

What is vectorization? How do you use it in PyTorch? What are some of the traps and pitfalls of writing vectorized code in PyTorch?

Further reading:

native/cpu README https://github.com/.../aten/src/ATen/native/cpu/README.md
Vec256 classes https://github.com/.../tree/master/aten/src/ATen/cpu/vec256
AVX512 support tracking issue https://github.com/pytorch/pytorch/issues/56187

Dynamic library structure

wookim@fb.com (PyTorch) — Thu, 6 May 2021 13:00:00 +0000

Why is PyTorch split into so many libraries? What's the point of these splits? What do Windows, mobile and CUDA have to do with the library splits?

Further reading:

c10 folder architecture description https://github.com/.../wiki/Software-Architecture-for-c10
Implementation of the TORCH_API visibility macros https://github.com/.../blob/master/c10/macros/Export.h
An example of virtual call based hook to break library structure https://github.com/pytorch/pytorch/blob/master/c10/core/impl/DeviceGuardImplInterface.h

History and constraints of the dispatcher

wookim@fb.com (PyTorch) — Wed, 5 May 2021 05:20:33 +0000

Why is the dispatcher the way it is today? How did evolve over time, and what constraints got added so that it is the kind of complicated piece it is today?

Further reading:

How the dispatcher actually works http://blog.ezyang.com/.../lets-talk-about-the-pytorch.../
Zachary DeVito's original version of ATen, before it got merged back into PyTorch mainline https://github.com/zdevito/ATen
The multiple dispatch patch: https://github.com/pytorch/pytorch/pull/25653

Binding C++ objects to Python

wookim@fb.com (Edward Yang - PyTorch Research Engineer at Facebook AI) — Tue, 4 May 2021 01:19:10 +0000

In this episode, we will discuss how to bind a C++ object in Python? We'll try to answer the following questions: How does pybind11 do it? What's different about how we implement it for Tensor? What are some downsides of the approach?

Note from the future: I recorded and then decided I didn't like my follow up episode about how to preserve PyObjects even when they go dead in Python. Maybe some day!

Further reading:

Python bindings for Tensor in PyTorch https://github.com/.../csrc/autograd/python_variable.cpp
pybind11 hash map for maintaining object identity https://github.com/.../inc.../pybind11/detail/internals.h...
Tensor subclasses don't save their properties https://github.com/pytorch/pytorch/issues/47117
(but the situation here is more complicated than I imply in the podcast)