Skip to main content

fusedtok

CI PyPI License: MIT Python 3.10+

Fused CUDA kernels for LLM inference — RMSNorm / RoPE / SwiGLU / attention decode and friends, with zero-copy torch tensor support: up to 9.3x faster than PyTorch SDPA (attention decode, RTX 3060, see Benchmarks).

中文文档请看 README_zh.md | English below.

Why

LLM inference frameworks launch many small, memory-bound operators per token. Each launch round-trips through global memory. fusedtok fuses them into single kernels to cut memory traffic and launch overhead.

Operators

Status Kernel Notes
RMSNorm (+residual) LLaMA/Qwen style, fused residual add
LayerNorm with affine
RoPE interleaved and NeoX layouts, kv-cache pos_offset
SwiGLU fused MLP activation
Softmax (row-wise) numerically stable
SiLU / GeLU / GeLU-tanh / ReLU / Tanh / Sigmoid elementwise
add / mul elementwise binary (fused add+residual pattern)
top-k / top-p (nucleus) arrival-ticket radix + early-exit compaction, replayed from a cached CUDA graph; deterministic ties (1.5x vs torch/CUB @131k k=50, parity-to-winning across the whole k range on both test GPUs)
argmax / temperature greedy decoding helpers
sample_topp fused nucleus sampling: softmax -> top-p -> seeded draw, global-mass threshold
sample_topk fused top-k sampling: softmax -> top-k -> renormalize within the window -> seeded draw (2.1x / 1.9x vs the topk+multinomial composite @131k)
repetition penalty CTRL-style, applied to sampled token ids
decode_step the whole decode step fused: penalty -> temperature -> nucleus sample, one call, one readback
quantize_int8 / dequantize_int8 / qadd_int8 symmetric per-tensor INT8, fused dequant-add-requant
qgemm INT8 matmul, int32-exact: cp.async double-buffered pipelined IMMA GEMM with runtime tile tuning (64x64 / 128x128) + warp-per-row GEMV (M=1 decode; 2x vs fp16 projection)
qgemm_perchannel the W8A8 layout real INT8 inference uses: per-output-channel weight scales fused into the same kernel's epilogue at zero cost
attention_decode single-token causal attention with GQA over a contiguous kv-cache: online softmax, flash-decoding split over long caches, per-sequence lengths
attention_prefill fresh-sequence attention over S query rows (causal / bidirectional); convenience path - heavyweight prefill stays SDPA/flash territory (honest ~0.45x)

Install

pip install fusedtok

Prebuilt wheels on PyPI (built with CUDA 12.4): Linux x86_64 (manylinux, cp310-cp313) and Windows x86_64 (cp311-cp313). On other platforms or Python versions pip builds from source automatically:

git clone https://github.com/Hai-Wenxiang/fusedtok.git
cd fusedtok
pip install .

Requirements:

  • NVIDIA GPU of RTX 30 series (Ampere) or newer — e.g. RTX 3060/3090, RTX 4080, RTX 5090, A100, H100
  • CUDA Toolkit >= 12.0
  • A C++17 compiler (MSVC on Windows, GCC/Clang on Linux); Python 3.10+
What is "compute capability"? (click to expand)

Compute capability is NVIDIA's version number for a GPU architecture generation — not a performance score. CUDA code must be compiled for a specific architecture to run on it. The wheel builds native cubins for compute capability 8.0 (A100) and 8.6 (RTX 30) plus a compute_86 PTX fallback, so Ampere runs natively and newer architectures (RTX 40/50, ...) JIT the PTX with their driver.

Compute capability Architecture Example GPUs
7.5 Turing GTX 16xx, RTX 20xx (not supported)
8.0 / 8.6 Ampere A100, RTX 30xx
8.9 Ada RTX 40xx (via PTX)
9.0 Hopper H100 (via PTX)
12.0 Blackwell RTX 50xx (via PTX)

Check yours: run nvidia-smi to see your GPU model, then look it up at https://developer.nvidia.com/cuda-gpus

Usage

numpy in / numpy out, or torch in / torch out — including zero-copy CUDA: kernels read and write torch device buffers directly via data_ptr(), with no staging copies and no host synchronization.

import numpy as np
import torch
import fusedtok

x = np.random.randn(4, 1024).astype(np.float32)
w = np.random.rand(1024).astype(np.float32)

# CPU reference implementation (ground truth, runs anywhere)
y = fusedtok.rmsnorm(x, w, eps=1e-6)

# staged CUDA: copies to GPU, runs kernel, copies back
y = fusedtok.rmsnorm(x, w, cuda=True)

# zero-copy CUDA with torch tensors: kernels run in torch's own buffers,
# stream-ordered with other torch operations
xt, wt = torch.from_numpy(x).cuda(), torch.from_numpy(w).cuda()
yt = fusedtok.rmsnorm(xt, wt)          # -> CUDA torch tensor

# RoPE with kv-cache position offset, NeoX (LLaMA-HF) layout
q = torch.randn(1, 4096, device="cuda")          # new token only
q_rot, k_rot = fusedtok.rope(q, k=None, pos_offset=1023, neox=True)

# attention over a GQA kv-cache: one call per decode step, no score
# materialization, variable-length batches share one cache tensor
out = fusedtok.attention_decode(
    q_heads,                                    # [B, Hq, D] new token
    k_cache, v_cache,                           # [B, Hkv, T, D]
    lens=torch.tensor([1023, 512], dtype=torch.int32, device="cuda"))
# fresh-sequence prefill (causal by default; convenience path)
ctx = fusedtok.attention_prefill(q_all, k_all, v_all, causal=True)

# sampling side: the whole decode step in one fused call
token = fusedtok.decode_step(logits, sampled_ids, penalty=1.1,
                             p=0.9, temperature=0.8, seed=step)
# or step by step:
logits = fusedtok.repetition_penalty(logits, sampled_ids, penalty=1.1)
token = fusedtok.sample_topp(logits, p=0.9, temperature=0.8, seed=step)
# top-k sampling variant (renormalizes within the k survivors)
token = fusedtok.sample_topk(logits, k=50, temperature=0.8, seed=step)

A minimal per-token sampling loop:

import torch, fusedtok as ft

h = torch.zeros(1, 4096, device="cuda")            # decoder state
w = torch.load("rms_weight.pt").cuda()             # float32 weights
wq, wscale = ft.quantize_int8(weight_f32.ravel())  # int8 weights
generated = []
for step in range(256):
    h = ft.rmsnorm(h, w, residual=h)               # fused add + norm
    q = ft.rope(q, k=None, pos_offset=step, neox=True)
    logits = model_output(h)                       # your model
    tok = ft.decode_step(logits, generated, penalty=1.1,
                         p=0.9, temperature=0.8, seed=step)
    generated.append(int(tok))

Every function accepts float32 numpy arrays or torch tensors (other dtypes are converted with a copy) and returns float32 outputs of the same family. CUDA torch tensors may also be bfloat16 - the kernels compute in float32 and convert at the load/store boundary (norm weights are upcast to float32 automatically; sampling/selection ops stay float32). CUDA torch tensors select the zero-copy path automatically.

See examples/demo.py for a runnable tour of every operator.

Correctness

Every kernel ships with a CPU reference implementation and element-wise parity tests (pytest). Tests run on machines without a GPU (CUDA cases skip automatically).

API stability

1.0 freezes the public surface: the names in fusedtok.__all__ (30 operators + helpers) keep their signatures across the 1.x series. Type stubs (__init__.pyi, PEP 561 py.typed) ship with the package. New operators arrive in minor releases; breaking changes require a new major version and a deprecation window. Determinism promises: selection ties resolve to the earliest index; sampling is deterministic per seed.

Benchmarks

RTX 3060 (sm_86), float32, zero-copy torch tensors, CUDA-event timing over 3 independent rounds (means below; per-round values in the JSON), vs the equivalent PyTorch reference (composite eager expressions; attention references use pre-expanded heads - repeat_interleave outside the timed region). Largest shape per op; full data: docs/benchmark_rtx3060.json, reproduce with python benchmarks/bench.py:

Op Shape fusedtok PyTorch reference Speedup
attention_decode (GQA) T=16384, D=128 853 µs 7614 µs (SDPA) 8.92x
RoPE NeoX (q+k) [8192×4096] 1641 µs 10061 µs 6.13x
RMSNorm (+residual) [4096×4096] 614 µs 2061 µs 3.36x
SwiGLU [4096×4096] 614 µs 1025 µs 1.67x
top-k (k=50) [131072] 79 µs 137 µs 1.75x
top-k (k=4096, mid-k) [131072] 113 µs 127 µs 1.12x
LayerNorm [4096×4096] 446 µs 616 µs 1.38x
Softmax [4096×4096] 414 µs 432 µs 1.04x
SiLU / GeLU / add [4096×4096] ~412 µs ~411 µs ~1.0x
sample_topk k=50 [131072] 133 µs 282 µs (topk+multinomial) 2.13x
argmax [131072] 65 µs 45 µs 0.69x (incl. host readback)
int8 qgemm (IMMA) [4096×4096×4096] 3554 µs (38.7 TOPS) 1634 µs (cuBLASLt) 0.46x (honest)
int8 qgemm pc (W8A8) [4096×4096×4096] 3553 µs (38.7 TOPS) 2046 µs (cuBLASLt + broadcast) 0.58x (honest)
attention_prefill (causal) S=1024, D=128 5732 µs 2560 µs (SDPA flash) 0.45x (honest)

Row-wise kernels (norms, softmax) autotune their thread-block size per shape at first call (v0.4.1); the table reflects the tuned choices.

fusedtok vs PyTorch reference

RTX 5060 Ti (Blackwell, sm_120) — same suite, largest shape per op (full data: docs/benchmark_rtx5060ti.json):

Op Shape fusedtok PyTorch reference Speedup
RoPE NeoX (q+k) [8192×4096] 1384 µs 8368 µs 6.04x
attention_decode (GQA) T=16384, D=128 575 µs 2682 µs (SDPA) 4.67x
RMSNorm (+residual) [4096×4096] 504 µs 1657 µs 3.29x
SwiGLU [4096×4096] 504 µs 858 µs 1.70x
top-k (k=50) [131072] 27 µs 41 µs (CUB) 1.50x
top-k (k=4096, mid-k) [131072] 50 µs 54 µs (CUB) 1.09x
LayerNorm / Softmax [4096×4096] ~345 µs ~348 µs 1.0x
sample_topk k=50 [131072] 49 µs 93 µs (topk+multinomial) 1.91x
argmax [131072] 17 µs 14 µs 0.83x (incl. host readback)
int8 qgemm (IMMA) [4096×4096×4096] 2063 µs (66.6 TOPS) 800 µs (cuBLASLt) 0.39x (honest)
int8 qgemm pc (W8A8) [4096×4096×4096] 2079 µs (66.1 TOPS) 1142 µs (cuBLASLt + broadcast) 0.55x (honest)
attention_prefill (causal) S=1024, D=128 3291 µs 1421 µs (SDPA flash) 0.43x (honest)

On smaller shapes the Blackwell card shows bigger wins (softmax 2.5x, RMSNorm 3.2x at 256 rows, attention decode 3.8x at T=4096 running 235 GB/s) - the launch-overhead share shrinks as shapes grow; full sweep in the JSON.

fusedtok vs PyTorch reference (RTX 5060 Ti)

The PyPI wheel ships sm_80/sm_86 cubins plus a compute_86 PTX fallback — verified to JIT and run correctly on Blackwell (sm_120) drivers.

Fusions win big (RoPE / RMSNorm / SwiGLU) because eager mode round-trips intermediate tensors through global memory. The v0.4 selection pipeline (arrival-ticket radix rounds + early-exit compaction, replayed from a cached CUDA graph) beats torch's CUB radix select at small k on both GPUs; the v1.0 retune (in-block-sort threshold and sort chunk both dropped 2048 -> 1024 - a single block bitonic-sorting 2048 keys was the whole mid-k regression) brings the mid-k window to parity-or-winning as well (k=4096 @131k: 1.12x / 1.09x). attention_decode wins big at decode (one launch streams the GQA cache once at up to ~157 GB/s effective while SDPA pays head expansion or small-query inefficiency); attention_prefill is the honest convenience path at ~0.45x of SDPA's flash backend — no tensor cores by design, so heavyweight prefill stays with SDPA/FlashAttention. The INT8 decode GEMV moves half the bytes of an fp16 projection and runs at full memory bandwidth (2x); the pipelined IMMA GEMM (v1.0 rework: cp.async double-buffered slabs, runtime-tuned 64x64 / 128x128 tiles) reaches ~39 TOPS on a 3060 and ~67 TOPS on a 5060 Ti — 2x-4x the v0.4 kernel — but cuBLASLt (torch._int_mm) still holds a ~2.2-2.6x lead: its tiles pipeline deeper and its epilogue is tuned per-arch. For now qgemm is the exact / graph-capturable / zero-copy INT8 path, not the fastest one; honest numbers, a CUTLASS-class schedule stays future work. The per-channel variant (qgemm_perchannel, the W8A8 layout INT8 inference actually uses) fuses the per-output-channel scale multiply into the same epilogue at zero kernel cost — the composite torch reference pays for that broadcast separately, which is where its 0.55-0.58x comes from.

Development

See CONTRIBUTING.md for the full guide (test rules, error contract, determinism invariants). Quick start:

# Windows: run inside a VS developer prompt (vcvars64)
cmake -S . -B build -G Ninja -DCMAKE_BUILD_TYPE=Release
cmake --build build
# from repo root: PYTHONPATH picks up the built module, conftest.py adds python/
$env:PYTHONPATH = "$PWD/build"        # Windows
PYTHONPATH=$PWD/build                 # Linux
python -m pytest tests -q
python benchmarks/bench.py            # GPU benchmark + chart

Windows / Linux. Windows uses MSVC via nvcc; CI builds and runs the CPU test suite on every push.

Roadmap

  • v0.2 (done): bf16 zero-copy, radix-select top-k/top-p, fused nucleus sampling, single-read softmax, CUDA-graph verified
  • v0.3 (done): chunk-merge selection sort + parallel nucleus count, bf16x4/x8 vectorized elementwise, INT8 quantize/dequantize utilities
  • v0.4 (done): arrival-ticket selection pipeline (no cooperative launch, early-exit compaction, cached CUDA graphs), stream-aware launchers everywhere (real CUDA-graph capture), INT8 compute path (IMMA qgemm + decode GEMV), fused decode_step sampling
  • v0.4.1 (done): runtime block-size autotuning for the row-wise kernels (norms/softmax pick 128..1024 threads per shape at first call)
  • v0.5 (done): attention - GQA decode attention over a contiguous kv-cache (flash-decoding split over long caches, per-sequence lengths) and a tiled prefill path (honest ~0.45x of SDPA flash - the convenience path); single-chart-per-GPU benchmarks; Windows wheels in the PyPI publish pipeline
  • 1.0 (in development): pipelined tensor-core INT8 GEMM (cp.async double-buffering, runtime tile tuning; 17 -> 39 TOPS on a 3060) with per-channel weight scales (W8A8), fused top-k sampling (2.1x vs the topk+multinomial composite), top-k mid-range-k parity, text hygiene gate, wheel matrix expansion, API freeze

Community

License

MIT — see LICENSE. Third-party notices: NOTICES.md.

Download files

Download the file for your platform. If you're not sure which to choose, learn more about installing packages.

Source Distribution

fusedtok-1.0.0.tar.gz (427.1 kB view details)

Uploaded Source

Built Distributions

If you're not sure about the file name format, learn more about wheel file names.

fusedtok-1.0.0-cp313-cp313-win_amd64.whl (1.1 MB view details)

Uploaded CPython 3.13Windows x86-64

fusedtok-1.0.0-cp313-cp313-manylinux_2_34_x86_64.whl (1.3 MB view details)

Uploaded CPython 3.13manylinux: glibc 2.34+ x86-64

fusedtok-1.0.0-cp312-cp312-win_amd64.whl (1.1 MB view details)

Uploaded CPython 3.12Windows x86-64

fusedtok-1.0.0-cp312-cp312-manylinux_2_34_x86_64.whl (1.3 MB view details)

Uploaded CPython 3.12manylinux: glibc 2.34+ x86-64

fusedtok-1.0.0-cp311-cp311-win_amd64.whl (1.1 MB view details)

Uploaded CPython 3.11Windows x86-64

fusedtok-1.0.0-cp311-cp311-manylinux_2_34_x86_64.whl (1.3 MB view details)

Uploaded CPython 3.11manylinux: glibc 2.34+ x86-64

fusedtok-1.0.0-cp310-cp310-manylinux_2_34_x86_64.whl (1.3 MB view details)

Uploaded CPython 3.10manylinux: glibc 2.34+ x86-64

File details

Details for the file fusedtok-1.0.0.tar.gz.

File metadata

  • Download URL: fusedtok-1.0.0.tar.gz
  • Upload date:
  • Size: 427.1 kB
  • Tags: Source
  • Uploaded using Trusted Publishing? Yes
  • Uploaded via: twine/7.0.0 CPython/3.13.14

File hashes

Hashes for fusedtok-1.0.0.tar.gz
Algorithm Hash digest
SHA256 599a72c8269a41a4e7b0ac0e67db2ba317319c7123b2deaa128d6e5048d4e6ab
MD5 e0a1bc2c98d369ba20fe3d473ee5d4ff
BLAKE2b-256 6de95bfd055de995a97d15d1a72786a41c5ada48dd25f3ea9118c9fc9d112205

See more details on using hashes here.

Provenance

The following attestation bundles were made for fusedtok-1.0.0.tar.gz:

Publisher: publish.yml on Hai-Wenxiang/fusedtok

Attestations: Values shown here reflect the state when the release was signed and may no longer be current.

File details

Details for the file fusedtok-1.0.0-cp313-cp313-win_amd64.whl.

File metadata

  • Download URL: fusedtok-1.0.0-cp313-cp313-win_amd64.whl
  • Upload date:
  • Size: 1.1 MB
  • Tags: CPython 3.13, Windows x86-64
  • Uploaded using Trusted Publishing? Yes
  • Uploaded via: twine/7.0.0 CPython/3.13.14

File hashes

Hashes for fusedtok-1.0.0-cp313-cp313-win_amd64.whl
Algorithm Hash digest
SHA256 330d93d6c461fd36685ac66e6344731c03db24dcd6b722867305eb45e44befd8
MD5 f7bb1cd7176a4aaaad2c7269588eb3e3
BLAKE2b-256 d57adc94ab5f7bd26795f453af474e1e038e6f8fdfe60273a5a62ac5e022c322

See more details on using hashes here.

Provenance

The following attestation bundles were made for fusedtok-1.0.0-cp313-cp313-win_amd64.whl:

Publisher: publish.yml on Hai-Wenxiang/fusedtok

Attestations: Values shown here reflect the state when the release was signed and may no longer be current.

File details

Details for the file fusedtok-1.0.0-cp313-cp313-manylinux_2_34_x86_64.whl.

File metadata

File hashes

Hashes for fusedtok-1.0.0-cp313-cp313-manylinux_2_34_x86_64.whl
Algorithm Hash digest
SHA256 e16dbc7fbd1ab6e24acce4330fd28a68f2b5bab4c10c6450d5f11d1cb9168fd1
MD5 b795a0e30f34ee35e0855bbde58bb0aa
BLAKE2b-256 c7fe44ab76897796215d207c7b59aa04164a7c1474eff7afb43bb2af2972ab55

See more details on using hashes here.

Provenance

The following attestation bundles were made for fusedtok-1.0.0-cp313-cp313-manylinux_2_34_x86_64.whl:

Publisher: publish.yml on Hai-Wenxiang/fusedtok

Attestations: Values shown here reflect the state when the release was signed and may no longer be current.

File details

Details for the file fusedtok-1.0.0-cp312-cp312-win_amd64.whl.

File metadata

  • Download URL: fusedtok-1.0.0-cp312-cp312-win_amd64.whl
  • Upload date:
  • Size: 1.1 MB
  • Tags: CPython 3.12, Windows x86-64
  • Uploaded using Trusted Publishing? Yes
  • Uploaded via: twine/7.0.0 CPython/3.13.14

File hashes

Hashes for fusedtok-1.0.0-cp312-cp312-win_amd64.whl
Algorithm Hash digest
SHA256 b55df5dcd679f8f4491645e6117503b99c05f52eec70c9e5087f45aaaef1e368
MD5 06203406b532bb9d413ed4b5ebad5ac3
BLAKE2b-256 f4a3e8065751882810b26e7d9f67df5be64654c2bbc811398aa9e6138dc9b58b

See more details on using hashes here.

Provenance

The following attestation bundles were made for fusedtok-1.0.0-cp312-cp312-win_amd64.whl:

Publisher: publish.yml on Hai-Wenxiang/fusedtok

Attestations: Values shown here reflect the state when the release was signed and may no longer be current.

File details

Details for the file fusedtok-1.0.0-cp312-cp312-manylinux_2_34_x86_64.whl.

File metadata

File hashes

Hashes for fusedtok-1.0.0-cp312-cp312-manylinux_2_34_x86_64.whl
Algorithm Hash digest
SHA256 0a1e97e8d4e9f2bd5e615a98ebc51c12417be18cf6c6da17b60b0fdd07963674
MD5 af26e6713f0cab241411f54911b39694
BLAKE2b-256 66317c3289715089d71c8e74a650d5cede6f8643b3e84daa7842783b000e9498

See more details on using hashes here.

Provenance

The following attestation bundles were made for fusedtok-1.0.0-cp312-cp312-manylinux_2_34_x86_64.whl:

Publisher: publish.yml on Hai-Wenxiang/fusedtok

Attestations: Values shown here reflect the state when the release was signed and may no longer be current.

File details

Details for the file fusedtok-1.0.0-cp311-cp311-win_amd64.whl.

File metadata

  • Download URL: fusedtok-1.0.0-cp311-cp311-win_amd64.whl
  • Upload date:
  • Size: 1.1 MB
  • Tags: CPython 3.11, Windows x86-64
  • Uploaded using Trusted Publishing? Yes
  • Uploaded via: twine/7.0.0 CPython/3.13.14

File hashes

Hashes for fusedtok-1.0.0-cp311-cp311-win_amd64.whl
Algorithm Hash digest
SHA256 27614a16be740b5e96a806d72bd61567a93048479110583d931489fa9718e334
MD5 95bb6a46a3fa9dc4827cccb75d1ed1ba
BLAKE2b-256 52643a7546bd80ff693241cce2b8aa192c3d43b7a6ab20858a75a36c256efb7c

See more details on using hashes here.

Provenance

The following attestation bundles were made for fusedtok-1.0.0-cp311-cp311-win_amd64.whl:

Publisher: publish.yml on Hai-Wenxiang/fusedtok

Attestations: Values shown here reflect the state when the release was signed and may no longer be current.

File details

Details for the file fusedtok-1.0.0-cp311-cp311-manylinux_2_34_x86_64.whl.

File metadata

File hashes

Hashes for fusedtok-1.0.0-cp311-cp311-manylinux_2_34_x86_64.whl
Algorithm Hash digest
SHA256 7dcfcc3e45d5364c94a5b92132d4b61878d4f8dbb204344960057533c68345c9
MD5 2ff744742749c1f2afce7b6171772064
BLAKE2b-256 20df2daad2fb8224b771faf2a778815149ed7063a77188b964b15275ef30e3ad

See more details on using hashes here.

Provenance

The following attestation bundles were made for fusedtok-1.0.0-cp311-cp311-manylinux_2_34_x86_64.whl:

Publisher: publish.yml on Hai-Wenxiang/fusedtok

Attestations: Values shown here reflect the state when the release was signed and may no longer be current.

File details

Details for the file fusedtok-1.0.0-cp310-cp310-manylinux_2_34_x86_64.whl.

File metadata

File hashes

Hashes for fusedtok-1.0.0-cp310-cp310-manylinux_2_34_x86_64.whl
Algorithm Hash digest
SHA256 2828d2830ad5551b1f36dec21fc028fa6d0b0a999d92280d75daf22cc97986d0
MD5 be2438ed93d6a964b65f4be2f023b0b5
BLAKE2b-256 bd6f0068a072858b4785e683dbf5029487a757275366c4da1eaa48b821465c42

See more details on using hashes here.

Provenance

The following attestation bundles were made for fusedtok-1.0.0-cp310-cp310-manylinux_2_34_x86_64.whl:

Publisher: publish.yml on Hai-Wenxiang/fusedtok

Attestations: Values shown here reflect the state when the release was signed and may no longer be current.

Release history Release notifications | RSS feed

1.1.0

8 files

1.0.1

8 files

This release

1.0.0 This release

8 files

0.5.1

3 files

0.5.0

3 files

0.4.1

2 files

0.4.0

2 files

0.3.1

2 files

0.3.0

2 files

0.2.1

2 files

0.2.0

2 files

0.1.2

2 files

0.1.1

2 files

0.1.0

2 files

Anthropic, PBC Visionary sponsor Bloomberg Visionary sponsor Hudson River Trading Visionary sponsor Meta Visionary sponsor NVIDIA Visionary sponsor Microsoft Sustainability sponsor Depot Continuous Integration AWS Cloud computing and Security Sponsor Datadog Monitoring Fastly CDN Google Download Analytics Sentry Error logging StatusPage Status page