Ferdinand Mom 3outeille

Matmul benchmark of Group-ordering vs Row-major ordering on A100 => No significant improvment over row-major ordering

https://triton-lang.org/master/getting-started/tutorials/03-matrix-multiplication.html#l2-cache-optimizations

matmul-performance:
        M  group_ordering  row_major_ordering
0   256.0        3.640889            3.640889
1   384.0       11.059200           12.288000
2   512.0       23.831273           23.831273
3   640.0       39.384616           39.384616
4 768.0 58.982401 58.982401

Problem: We have blocks that are scheduled later than others which imply that we won't get the "true max value" at the time we need it.
Direction: We should find a way to wait for all threads of all blocks to finish
Solution:
- 1. Split into 2 kernels
1. Use cooperative groups: https://numba.readthedocs.io/en/stable/cuda/cooperative_groups.html

⚠️ ⚠️ This Gist Has Been Moved to This Repo ⚠️ ⚠️

	#VERBOSE=0 torchrun --nproc_per_node 3 self_contained_pp_LOC.py
	import os, random, numpy as np, torch, torch.nn as nn, torch.distributed as dist, torch.nn.functional as F
	from torch.optim import AdamW
	from torch.utils.data import DataLoader, DistributedSampler
	from datasets import load_dataset
	from transformers import AutoConfig, AutoModelForCausalLM, AutoTokenizer

	STEP, local_rank, world_size, verbose = 0, int(os.environ["LOCAL_RANK"]), int(os.environ["WORLD_SIZE"]), os.environ.get("VERBOSE", "0") == "1"

	def set_all_seed(seed):

	import torch
	from torch.nn import functional as F
	from torch import distributed as dist
	import os
	import numpy as np
	import random

	def set_random_seed(seed: int):
	torch.manual_seed(seed)
	if torch.cuda.is_available():

	from copy import deepcopy
	import torch
	from datasets import load_dataset
	from torch.optim import SGD
	from torch.utils.data import DataLoader
	from transformers import AutoModelForCausalLM, AutoTokenizer
	import random
	import os
	import numpy as np


	import torch
	from tqdm import tqdm
	from transformers import AutoModelForCausalLM, AutoTokenizer
	from datasets import load_dataset

	# Model
	device = "cpu"

	device_map = {

	# To run
	# mkdir build && cd build
	# cmake ..
	# make -j && ./bank conflict <offset> <is_debug>

	cmake_minimum_required(VERSION 3.0)

	set(CMAKE_CXX_FLAGS "-O3 -std=c++14")

	set(CUDA_NVCC_FLAGS -arch=compute_52 -code=sm_75)

Ferdinand Mom 3outeille

⚠️ ⚠️ This Gist Has Been Moved to This Repo ⚠️ ⚠️

Table of Contents