3outeille

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():

3outeille

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

3outeille

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

# Model
device = "cpu"

device_map = {

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

3outeille

• 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
• 2. Use cooperative groups: https://numba.readthedocs.io/en/stable/cuda/cooperative_groups.html

3outeille

# 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)

3outeille

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

Table of Contents

• Features of OOP
• Smart Pointers
• Name Mangling and Externs
• Virtual Functions
  • Pure Virtual Functions
  • Virtual Destructors

• Virtual Table

3outeille

 /* Loop of 16x16 blocks */
for (y = 0; y < heightComp[compCtr]; y += 16)
{
    ...
    for (x = 0; x < widthComp[compCtr]; x += 16)
    {
        ...
        for (blkId = 0; blkId < 4; blkId++)
        {
            yOffset8x8 = (blkId >> 1) * 8;

3outeille

int16_t fg_compute_block_avg(int16_t *dstSampleBlk8, uint32_t widthComp, uint16_t *pNumSamples,
                      uint8_t ySize, uint8_t xSize, uint8_t bitDepth)
{
  uint32_t blockAvg   = 0;
  uint16_t numSamples = 0;
  uint8_t  k, l;

  for (k = 0; k < ySize; k++)
  {
    for (l = 0; l < xSize; l++)

3outeille

int16_t fg_compute_block_avg_sse4(int16_t *dstSampleBlk8, uint32_t widthComp, uint16_t *pNumSamples,
                      uint8_t ySize, uint8_t xSize, uint8_t bitDepth)
{
  uint16_t blockAvg   = 0;
  uint16_t numSamples = 0;

  __m128i acc = _mm_setzero_si128();

  for (int i = 0; i < ySize; i+=1, numSamples+=8)
  {