Milestones V2

docker-compose.test.yml

@@ -0,0 +1,32 @@
+services:
+  postgres-test:
+    image: postgres:15
+    container_name: postgres_db_test
+    environment:
+      POSTGRES_USER: postgres
+      POSTGRES_PASSWORD: postgres
+      POSTGRES_DB: clusterdev
+    ports:
+      - "5433:5432"
+    tmpfs:
+      - /var/lib/postgresql/data
+    healthcheck:
+      test: ["CMD-SHELL", "pg_isready -U postgres -d clusterdev"]
+      interval: 5s
+      timeout: 3s
+      retries: 10
+
+  migrate-test:
+    image: python:3.11-slim
+    depends_on:
+      postgres-test:
+        condition: service_healthy
+    volumes:
+      - ./src/migrations:/migrations
+    command: >
+      sh -c "
+        pip install yoyo-migrations psycopg2-binary &&
+        yoyo apply -b -d 'postgresql://postgres:postgres@postgres-test:5432/clusterdev' -v /migrations/ &&
+        yoyo list -d 'postgresql://postgres:postgres@postgres-test:5432/clusterdev' /migrations/
+      "
+    restart: "no"

examples/matmul_py/task.yml

@@ -7,6 +7,19 @@ files:
   - {"name": "reference.py", "source": "reference.py"}
   - {"name": "eval.py", "source": "../eval.py"}
 
+milestones:
+  - {
+      name: "pytorch",
+      source: "submission.py",
+      description: "PyTorch reference implementation as a performance baseline for matmul"
+    }
+  - {
+      name: "triton",
+      source: "triton_ref.py",
+      description: "Triton reference implementation as a performance baseline for matmul",
+      exclude_gpus: ['T4']
+    }
+
 lang: "py"
 
 description: |

examples/matmul_py/triton_ref.py

@@ -0,0 +1,127 @@
+#!POPCORN leaderboard matmul_py
+import triton
+import triton.language as tl
+import torch
+from task import input_t, output_t
+
+
+@triton.jit
+def matmul_kernel(
+    # Pointers to matrices
+    a_ptr, b_ptr, c_ptr,
+    # Matrix dimensions
+    M, N, K,
+    # The stride variables represent how much to increase the ptr by when moving by 1
+    # element in a particular dimension. E.g. `stride_am` is how much to increase `a_ptr`
+    # by to get the element one row down (A has M rows).
+    stride_am, stride_ak,
+    stride_bk, stride_bn,
+    stride_cm, stride_cn,
+    # Meta-parameters
+    BLOCK_SIZE_M: tl.constexpr, BLOCK_SIZE_N: tl.constexpr, BLOCK_SIZE_K: tl.constexpr,
+    GROUP_SIZE_M: tl.constexpr,
+):
+    """Kernel for computing the matmul C = A x B.
+    A has shape (M, K), B has shape (K, N) and C has shape (M, N)
+    """
+    # -----------------------------------------------------------
+    # Map program ids `pid` to the block of C it should compute.
+    # This is done in a grouped ordering to promote L2 cache hit rates.
+    # See above `L2 Cache Optimizations` section for details.
+    pid = tl.program_id(axis=0)
+    num_pid_m = tl.cdiv(M, BLOCK_SIZE_M)
+    num_pid_n = tl.cdiv(N, BLOCK_SIZE_N)
+    num_pid_in_group = GROUP_SIZE_M * num_pid_n
+    group_id = pid // num_pid_in_group
+    first_pid_m = group_id * GROUP_SIZE_M
+    group_size_m = min(num_pid_m - first_pid_m, GROUP_SIZE_M)
+    pid_m = first_pid_m + (pid % group_size_m)
+    pid_n = (pid % num_pid_in_group) // group_size_m
+
+    # ----------------------------------------------------------
+    # Create pointers for the first blocks of A and B.
+    # We will advance this pointer as we move in the K direction
+    # and accumulate
+    # `a_ptrs` is a block of [BLOCK_SIZE_M, BLOCK_SIZE_K] pointers
+    # `b_ptrs` is a block of [BLOCK_SIZE_K, BLOCK_SIZE_N] pointers
+    # See above `Pointer Arithmetic` section for details
+    offs_am = (pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)) % M
+    offs_bn = (pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)) % N
+    offs_k = tl.arange(0, BLOCK_SIZE_K)
+    a_ptrs = a_ptr + (offs_am[:, None] * stride_am + offs_k[None, :] * stride_ak)
+    b_ptrs = b_ptr + (offs_k[:, None] * stride_bk + offs_bn[None, :] * stride_bn)
+
+    # -----------------------------------------------------------
+    # Iterate to compute a block of the C matrix.
+    # We accumulate into a `[BLOCK_SIZE_M, BLOCK_SIZE_N]` block
+    # of fp32 values for higher precision.
+    # `accumulator` will be converted back to fp16 after the loop.
+    accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)
+    for k in range(0, tl.cdiv(K, BLOCK_SIZE_K)):
+        # Load the next block of A and B, generate a mask by checking the K dimension.
+        # If it is out of bounds, set it to 0.
+        a = tl.load(a_ptrs, mask=offs_k[None, :] < K - k * BLOCK_SIZE_K, other=0.0)
+        b = tl.load(b_ptrs, mask=offs_k[:, None] < K - k * BLOCK_SIZE_K, other=0.0)
+        # We accumulate along the K dimension.
+        accumulator += tl.dot(a, b)
+        # Advance the ptrs to the next K block.
+        a_ptrs += BLOCK_SIZE_K * stride_ak
+        b_ptrs += BLOCK_SIZE_K * stride_bk
+    # You can fuse arbitrary activation functions here
+    # while the accumulator is still in FP32!
+    c = accumulator.to(tl.float16)
+
+    # -----------------------------------------------------------
+    # Write back the block of the output matrix C with masks.
+    offs_cm = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)
+    offs_cn = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
+    c_ptrs = c_ptr + stride_cm * offs_cm[:, None] + stride_cn * offs_cn[None, :]
+    c_mask = (offs_cm[:, None] < M) & (offs_cn[None, :] < N)
+    tl.store(c_ptrs, c, mask=c_mask)
+
+
+def triton_matmul(a, b):
+    # Check constraints.
+    assert a.shape[1] == b.shape[0], "Incompatible dimensions"
+    assert a.is_contiguous(), "Matrix A must be contiguous"
+    assert b.is_contiguous(), "Matrix B must be contiguous"
+    M, K = a.shape
+    K, N = b.shape
+    # Allocate output.
+    c = torch.empty((M, N), device=a.device, dtype=a.dtype)
+    # 1D launch kernel where each block gets its own program.
+    grid = lambda META: (triton.cdiv(M, META['BLOCK_SIZE_M']) * triton.cdiv(N, META['BLOCK_SIZE_N']), )
+    matmul_kernel[grid](
+        a, b, c,
+        M, N, K,
+        a.stride(0), a.stride(1),
+        b.stride(0), b.stride(1),
+        c.stride(0), c.stride(1),
+        BLOCK_SIZE_M=128, BLOCK_SIZE_N=128, BLOCK_SIZE_K=32,
+        GROUP_SIZE_M=8,
+    )
+    return c
+
+
+def custom_kernel(data: input_t) -> output_t:
+    a, b = data
+    # Convert to torch tensors if they aren't already
+    if not isinstance(a, torch.Tensor):
+        a = torch.tensor(a, dtype=torch.float16).cuda()
+    if not isinstance(b, torch.Tensor):
+        b = torch.tensor(b, dtype=torch.float16).cuda()
+
+    # Ensure tensors are on GPU and contiguous
+    if not a.is_cuda:
+        a = a.cuda()
+    if not b.is_cuda:
+        b = b.cuda()
+
+    a = a.contiguous()
+    b = b.contiguous()
+
+    # Use our custom Triton matmul
+    result = triton_matmul(a, b)
+
+    # Convert back to the expected output format
+    return result

pyproject.toml

@@ -44,6 +44,8 @@ relative_files = true
 exclude_lines = [
    "pragma: no cover",
    "raise NotImplementedError",
+    # For now, don't require coverage of db errors
+   "except psycopg2.Error"
 ]
 
 [tool.pytest.ini_options]