Skip to content

Instantly share code, notes, and snippets.

View csullivan's full-sized avatar

Chris Sullivan csullivan

62 followers · 21 following
  • NVIDIA
  • Portland
View GitHub Profile
@csullivan
csullivan / 2024_09_26_nsys_single_instance_wgmma_register_and_shared_layout.txt
Created September 26, 2024 16:55
Performance comparison: 5% gain using wgmma with LHS in registers vs shared. [1] https://github.com/csullivan/wgmma-intrin
Time (%) Total Time (ns) Instances Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Name
-------- --------------- --------- -------- -------- -------- -------- ----------- ----------------------------------------------------------------------------------------------------
---- 2495089 101 24703.9 24736.0 24544 27520 302.9 wgmma_f16_m64n256k16_kernel_shared_layout(__half *, __half *, __half *)
---- 2361204 101 23378.3 23423.0 23231 25600 245.6 wgmma_f16_m64n256k16_register_layout_kernel(__half *, __half *, __half *)
@csullivan
csullivan / 2024_09_26_nsys_single_instance_wgmma_register_and_shared_layout.txt
Created September 26, 2024 16:54
Performance comparison: 5% gain using wgmma with LHS in registers vs shared.
Time (%) Total Time (ns) Instances Avg (ns) Med (ns) Min (ns) Max (ns) StdDev (ns) Name
-------- --------------- --------- -------- -------- -------- -------- ----------- ----------------------------------------------------------------------------------------------------
---- 2495089 101 24703.9 24736.0 24544 27520 302.9 wgmma_f16_m64n256k16_kernel_shared_layout(__half *, __half *, __half *)
---- 2361204 101 23378.3 23423.0 23231 25600 245.6 wgmma_f16_m64n256k16_register_layout_kernel(__half *, __half *, __half *)
@csullivan
csullivan / install_python3.11.sh
Last active May 14, 2024 04:53
sudo add-apt-repository ppa:deadsnakes/ppa
sudo apt install python3.11 python3.11-distutils python3.11-venv libpython3.11-dev
curl -sS https://bootstrap.pypa.io/get-pip.py | python3.11
@csullivan
csullivan / note.md
Last active September 7, 2023 06:50
CUTLASS CMake configuration for Hopper (sm90a)
@csullivan
csullivan / sharded_decode.py
Last active September 6, 2023 23:30
Sharded decode, sharding rewrite done after FuseOpsByPattern (cublas/cutlass byoc) -- With debug tracing calls
# Ignore `tvm.save_and_copy_tensor` packed functions inserted for debugging
@R.function
def decode(input_ids1: R.Tensor((1, 1), dtype="int32"), all_seq_len: R.Shape(["n"]), kv_cache: R.Tuple(R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Object, R.Obj
@csullivan
csullivan / test_torch_distributed.py
Created July 17, 2023 05:44
Minimal NCCL torch.distributed example
import os
import torch
import torch.distributed as dist
def read_file_and_all_reduce():
# Get the rank and world size from environment variables
rank = int(os.environ['LOCAL_RANK'])
world_size = int(os.environ['WORLD_SIZE'])
# initialize the process group
@csullivan
csullivan / bradbury_if_only_we_had_taller_been.txt
Created June 2, 2023 15:40
The fence we walked between the years
Did balance us serene;
It was a place half in the sky where
In the green of leaf and promising of peach
We’d reach our hand to touch, and almost touch the sky.
If we could reach and touch, we said,
‘Twould teach us not to, never to, be dead.
We ached and almost touched that stuff;
Our reach was never quite enough.
@csullivan
csullivan / CMake_3_20_Ubuntu_18_04.md
Last active March 30, 2023 05:39 — forked from bmegli/CMake_3_20_Ubuntu_18_04.md
apt reversible source builds with checkinstall-- Example: CMake 3.20 in Ubuntu 18.04 (reversible way)

Motivatation

  • modern CMake is required for building a lot of new software
  • CMake is dependency for many packages (e.g. ROS related)
  • we don't want to remove CMake (which would remove packages that depend on it)
  • we want safe procedure to update CMake that can be reversed easily

Current version in OS

Check current version

@csullivan
csullivan / 2022-10-25_chris-sullivan_notes.md
Last active November 14, 2022 17:29

Authored-by: Eric Lunderberg

Notes summarizing discussion between @Lunderberg and @csullivan on 2022_10_25

Considerations of Pad/Crop represented separately from bijective transformations

From previous conversation, possibility of representing pad/crop separately from the layout transform. This would allow algebraic

@csullivan
csullivan / bcdedit.txt
Last active October 30, 2021 23:41
Repairing EFI/GPT Bootloader using the DISKPART and BCDEdit command
Source: https://www.retrospect.com/au/support/kb/repair_efi_bootloader
This guide details how to rebuild the BCD boot store for windows.
Use diskpart to make System partition available.
(This will allow us to use bcdedit later in this guide to repair the bootloader.)
Boot from the Windows installation media