[TorchInductor] Add support for Pascal GPUs (P100, GTX 1080, etc)

[jansel]

According to @ptillet supporting Pascal GPUs with Triton should be possible if we:

1. Disable float16 support (and give a good error message saying fp16 requires a newer GPU)
2. Use aten backend for matmul/conv
3. May need to remove an assert on the Triton side?
4. Test it to make sure things work

This might be a good first issue if someone has a Pascal GPU handy to test on.

cc @ezyang @soumith @msaroufim @wconstab @ngimel @bdhirsh

[ngimel]

Would atomics work on Pascal? I thought they are also Volta+

[ptillet]

Hmm IIRC global memory atomics should be fine (shared atomics are more recent), but we should double check :)

[msaroufim]

I'm interested in this if it's not too crazy, the only assert I could find was https://github.com/openai/triton/blob/master/lib/driver/llvm.cc#L141

But more of a high level question if we're looking at something like Google Colab

• It supports K80 Kepler which is old but still popular
• T4 which is Turing and works just fine with Triton when I tried it

Are P100 widespread? My perception is the most popular GPUs on the market are K80, T4, V100, and A100. Regardless I have an old dusty 1070 in my closet

[soumith]

My perception is the most popular GPUs on the market are K80, T4, V100, and A100

Spoken like a true Google Colab user haha.
Yea, P100 is still pretty popular.

[shintaro-iwasaki]

I am curious if someone has tried it (or is working on it) on P100. As far as I tried it on P100 Google cloud, many tests pass without any changes (278 failed, 3983 passed, 517 skipped). If no one is working on it right now, I am happy to fix the remaining failures in Triton.

Environment Details (collapsed)

VM details (Google Cloud)

https://cloud.google.com/

# You need a paid account (free trial does not cover GPUs)
Google Cloud -> New Project -> Compute-Engine -> VM Instance
Machine:
GPU: NVIDIA Tesla P100 x 1
CPU: 2 vCPUs, 7.5GB memory
Boot disk:
  OS: Ubuntu 18.04 LTS
  Disk: 40GB (cannot build Triton on the default 10GB disk)
- When I tried, about $1.2 per hour.
- US instances were full when I tried.  I used Asia or Australia.
- Needed a paid account (GPU is not covered by free trial)
- Needed quota request for any GPU instance (by default, no GPU instance is allowed).  Needed to wait an hour for approval

Reproducer

## 1. Install CUDA and a driver
# Update the apt key (https://developer.nvidia.com/blog/updating-the-cuda-linux-gpg-repository-key/)
sudo apt-key del 7fa2af80
wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu1804/x86_64/cuda-keyring_1.0-1_all.deb
# Download CUDA as instructed
wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu1804/x86_64/cuda-ubuntu1804.pin
sudo mv cuda-ubuntu1804.pin /etc/apt/preferences.d/cuda-repository-pin-600
sudo apt-key adv --fetch-keys https://developer.download.nvidia.com/compute/cuda/repos/ubuntu1804/x86_64/7fa2af80.pub
sudo add-apt-repository "deb https://developer.download.nvidia.com/compute/cuda/repos/ubuntu1804/x86_64/ /"
sudo apt-get update
sudo apt-get -y install cuda
# Are you using P100?
nvidia-smi | grep "Tesla P100"

## 2. Setup the build environment
sudo apt update
sudo apt install -y build-essential wget git libz-dev
wget https://repo.anaconda.com/archive/Anaconda3-2022.05-Linux-x86_64.sh
bash Anaconda3-2022.05-Linux-x86_64.sh -b -p $(pwd)/anaconda3
eval "$($(pwd)/anaconda3/bin/conda shell.bash hook)"
conda create -y --name triton_base
conda activate triton_base
conda install -y cmake setuptools

## 3. Build Triton
git clone https://github.com/openai/triton.git
cd triton/python
pip3 install -e '.[tests]'

## 4. Test
pytest test/unit

Environment

$ nvidia-smi
+-----------------------------------------------------------------------------+
| NVIDIA-SMI 520.61.05    Driver Version: 520.61.05    CUDA Version: 11.8     |
|-------------------------------+----------------------+----------------------+
| GPU  Name        Persistence-M| Bus-Id        Disp.A | Volatile Uncorr. ECC |
| Fan  Temp  Perf  Pwr:Usage/Cap|         Memory-Usage | GPU-Util  Compute M. |
|                               |                      |               MIG M. |
|===============================+======================+======================|
|   0  Tesla P100-PCIE...  On   | 00000000:00:04.0 Off |                    0 |
| N/A   36C    P0    25W / 250W |      0MiB / 16384MiB |      0%      Default |
|                               |                      |                  N/A |
+-------------------------------+----------------------+----------------------+

Results

$ pytest test/unit
...
=== 278 failed, 3983 passed, 517 skipped in 525.26s (0:08:45) ===

[ngimel]

Currently has_triton requires voltas+ https://github.com/pytorch/torchdynamo/blob/main/torchinductor/utils.py#L40, so most tests probably don't attempt to generate kernels

[shintaro-iwasaki]

Oh, I ran Triton's test. In this sense, not all tests of OSS Triton pass on P100 now. Definitely another patch should be needed for TorchInductor. If we don't plan to change the OSS Triton side (or if Triton doesn't want any patch for P100 -- I believe we won't ask official P100 support for OpenAI), please let me know.

[ptillet]

Yeah, I suspect the failing test are all related to FP16/TF32 handling in matmuls. Nothing major if you guys plan to disable triton for matmuls on pre-Volta GPUs

[shintaro-iwasaki]

Thanks. @ptillet Regardless of whether FP16/TF32 related handling in matmuls (i.e., "tl.dot()" in general) will be skipped or fixed somehow, would you mind if I would create a patch to Triton so that the Triton's test results on P100 get green?

[ptillet]

Yep, I think that's a great patch idea :)

[msaroufim]

As an extra datapoint Kaggle notebooks now support P100 or T4 https://www.kaggle.com/product-feedback/361104

[shintaro-iwasaki]

Sorry for getting late. triton-lang/triton#827 is a PR to (sort-of) improve the Pascal support of OSS Triton.

Findings are:

• tl.dot() does not work on Pascal.
• Some atomic operations are not supported.
• Kernel invocation fails if the generated kernel needs shared memory larger than 64KB. 3 unit tests are failing because of this.
• Other than these, all the unit tests pass (though tests excluded on Volta are also skipped on Pascal too).

[ngimel]

@shintaro-iwasaki so fp16 atomics don't work on Pascal, and that's fine. I'm wondering if we should disable fp16 atomics universally because they might have bad perf, unless triton generates fast (not cas-based) code for them. Can you please check perf and generated code of fp16 atomic ops on newer hardware, and if it's cas-based, we should ban it everywhere.

[shintaro-iwasaki]

@ngimel Thanks for your comment! Would you mind teaching me the context and the priority a bit more to obtain what we really want to understand?

1. Performance depends on how dominant that atomic part is in a kernel (for example, the size of K in matmul with SPLIT_K >= 2 ... if K is extremely large, atomic part is negligible). If so, what would be the concern regarding (possibly) CAS-based atomic operations?
2. Even if its performance is worse to some extent because of CAS or whatever, will we ban it?

If we want to use this information in an architecture-specific cost model deployed by TorchInductor (e.g., performance on "newer hardware" like V100 or A100), it might need some nontrivial work, so I would like to first know the background (and perhaps it should be tracked by a different GitHub issue).