179 lines
12 KiB
Text
179 lines
12 KiB
Text
// family-probe (CUDA): the step cost of every reserve candidate family of spec 1.13.2 on NVIDIA, standalone (no pack).
|
|
// Counter ASIC 3.0 item 6 (docs/plans/counter-asic-3-reserve.md), 6 October 2026. PC job; not run on the Mac. Same
|
|
// method as dot4-probe.cu (docs/analysis/int8-matrix-family.md section 4): a dependent chain of one op per step per
|
|
// lane, 1,048,576 lanes x 4,096 steps, best of N, event time, bit-exact against a CPU reference on two whole warps.
|
|
//
|
|
// Every chain has the dot4 probe's glue: acc = OP(acc, x, y); x = x * K + acc; y = rotl(y, 7) ^ (acc + s). Reference
|
|
// rows (live families): alu (add-xor-rotate, 5 ops per step counted, no acc), rotr (the live rotr_var text), shflx
|
|
// (the live shfl: dst ^= src of lane (lane XOR 8)). Candidate rows, the seven families of 1.13.2:
|
|
// shl acc = y << (x & 31) shr acc = y >> (x & 31)
|
|
// bfe acc = bfe.u32(y, 7, 13) (inline PTX) bfec acc = (y >> 7) & 0x1fff (the C form, what nvcc emits)
|
|
// andn acc = y & ~x perm acc = __byte_perm(y, y, 0x2031): bytes (y.b1, y.b3, y.b0, y.b2)
|
|
// popc acc = acc + __popc(x) clz acc = acc + __clz(x) (clz(0) = 32)
|
|
// sel acc = bit 5 of y ? x : acc shfla acc ^= __shfl_sync(x, (lane + 3) & 31)
|
|
// Comparison rows: dot4i (__dp4a, PTX dp4a.s32.s32, the dot4 probe's intrinsic row) and mm8 (one mma.sync
|
|
// m8n8k16 u8 x u8 -> s32 per step per warp, inline PTX, sm_75+; PTX ISA 9.4 section 9.7.16.5.3 fragment layout:
|
|
// lane i holds A[i/4][4(i%4)..+3] as the bytes of x, B[4(i%4)..+3][i/4] as the bytes of y, C and D [i/4][2(i%4)+j]
|
|
// as acc (j = 0) and acc2 (j = 1); the chain carries d0 forward as acc, d1 as acc2). The shifted and extracted forms
|
|
// take y (fresh every step) as the shifted register so the chain never drains to 0.
|
|
//
|
|
// Build (Linux or WSL2, CUDA Toolkit): nvcc -O2 -arch=sm_120 -o family-probe family-probe.cu
|
|
// (older toolkits: nvcc -O2 -gencode arch=compute_89,code=compute_89 ...)
|
|
// Run: family-probe [--lanes N] [--steps N] [--reps N] [--device N]
|
|
#include <cuda_runtime.h>
|
|
#include <cstdio>
|
|
#include <cstdlib>
|
|
#include <cstring>
|
|
#include <cstdint>
|
|
|
|
__host__ __device__ inline uint32_t pm_mix(uint32_t x) { x ^= x >> 16; x *= 0x7feb352du; x ^= x >> 15; x *= 0x846ca68bu; x ^= x >> 16; return x; }
|
|
__host__ __device__ inline uint32_t rotl32(uint32_t v, uint32_t n) { return (v << n) | (v >> (32u - n)); }
|
|
__host__ __device__ inline uint32_t rotr_var(uint32_t x, uint32_t n) { n &= 31u; return (x >> n) | (x << ((32u - n) & 31u)); }
|
|
|
|
#define K32 0x9E3779B1u
|
|
#define CHAIN(NAME, OP) \
|
|
__global__ void NAME(uint32_t steps, uint32_t seed, uint32_t* out) { \
|
|
uint32_t g = blockIdx.x * blockDim.x + threadIdx.x; uint32_t lane = threadIdx.x & 31u; (void)lane; \
|
|
uint32_t x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u; uint32_t acc = pm_mix(x); \
|
|
for (uint32_t s = 0; s < steps; ++s) { OP; x = x * K32 + acc; y = rotl32(y, 7u) ^ (acc + s); } \
|
|
out[g] = acc ^ x ^ y; \
|
|
}
|
|
|
|
__global__ void probe_alu(uint32_t steps, uint32_t seed, uint32_t* out) {
|
|
uint32_t g = blockIdx.x * blockDim.x + threadIdx.x;
|
|
uint32_t x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u;
|
|
for (uint32_t s = 0; s < steps; ++s) { x = x * K32 + rotl32(y, 7u); y = (y ^ x) + s; }
|
|
out[g] = x ^ y;
|
|
}
|
|
__device__ inline uint32_t bfe_ptx(uint32_t v) { uint32_t r; asm("bfe.u32 %0, %1, 7, 13;" : "=r"(r) : "r"(v)); return r; }
|
|
CHAIN(probe_rotr, acc = rotr_var(y, x))
|
|
CHAIN(probe_shflx, acc = acc ^ __shfl_xor_sync(0xffffffffu, x, 8))
|
|
CHAIN(probe_shl, acc = y << (x & 31u))
|
|
CHAIN(probe_shr, acc = y >> (x & 31u))
|
|
CHAIN(probe_bfe, acc = bfe_ptx(y))
|
|
CHAIN(probe_bfec, acc = (y >> 7u) & 0x1fffu)
|
|
CHAIN(probe_andn, acc = y & ~x)
|
|
CHAIN(probe_perm, acc = __byte_perm(y, y, 0x2031))
|
|
CHAIN(probe_popc, acc = acc + (uint32_t)__popc((int)x))
|
|
CHAIN(probe_clz, acc = acc + (uint32_t)__clz((int)x))
|
|
CHAIN(probe_sel, acc = ((y >> 5u) & 1u) ? x : acc)
|
|
CHAIN(probe_shfla, acc = acc ^ __shfl_sync(0xffffffffu, x, (int)((lane + 3u) & 31u)))
|
|
CHAIN(probe_dot4i, acc = (uint32_t)__dp4a((int)x, (int)y, (int)acc))
|
|
|
|
// one mma.sync m8n8k16 per step per warp; the chain carries d0 as acc, d1 as acc2
|
|
__global__ void probe_mm8(uint32_t steps, uint32_t seed, uint32_t* out) {
|
|
uint32_t g = blockIdx.x * blockDim.x + threadIdx.x;
|
|
uint32_t x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u; uint32_t acc = pm_mix(x), acc2 = pm_mix(acc);
|
|
for (uint32_t s = 0; s < steps; ++s) {
|
|
uint32_t d0, d1;
|
|
asm volatile("mma.sync.aligned.m8n8k16.row.col.s32.u8.u8.s32 {%0,%1}, {%2}, {%3}, {%4,%5};"
|
|
: "=r"(d0), "=r"(d1) : "r"(x), "r"(y), "r"(acc), "r"(acc2));
|
|
acc = d0; acc2 = d1;
|
|
x = x * K32 + acc; y = rotl32(y, 7u) ^ (acc + s);
|
|
}
|
|
out[g] = acc ^ acc2 ^ x ^ y;
|
|
}
|
|
|
|
// ---- CPU reference: one whole warp (lanes g0 .. g0+31, lane = g AND 31) ----
|
|
static uint32_t clz32(uint32_t x) { if (!x) return 32; uint32_t n = 0; while (!(x & 0x80000000u)) { x <<= 1; ++n; } return n; }
|
|
static uint32_t popc32(uint32_t x) { uint32_t n = 0; while (x) { n += x & 1u; x >>= 1; } return n; }
|
|
static uint32_t perm_ref(uint32_t y) { return ((y >> 8) & 0xffu) | (((y >> 24) & 0xffu) << 8) | ((y & 0xffu) << 16) | (((y >> 16) & 0xffu) << 24); }
|
|
static uint32_t dot4s_ref(uint32_t a, uint32_t b, uint32_t acc) {
|
|
int32_t r = (int32_t)acc;
|
|
for (int i = 0; i < 4; ++i) { int32_t ba = (int32_t)(int8_t)((a >> (8 * i)) & 0xffu); int32_t bb = (int32_t)(int8_t)((b >> (8 * i)) & 0xffu); r = (int32_t)((uint32_t)r + (uint32_t)(ba * bb)); }
|
|
return (uint32_t)r;
|
|
}
|
|
static void warp_ref(const char* name, uint32_t g0, uint32_t seed, uint32_t steps, uint32_t* res) {
|
|
uint32_t x[32], y[32], acc[32], acc2[32], xs[32];
|
|
for (int l = 0; l < 32; ++l) { x[l] = pm_mix((g0 + l) ^ seed); y[l] = x[l] ^ 0x5bd1e995u; acc[l] = pm_mix(x[l]); acc2[l] = pm_mix(acc[l]); }
|
|
if (!strcmp(name, "alu")) {
|
|
for (uint32_t s = 0; s < steps; ++s) for (int l = 0; l < 32; ++l) { x[l] = x[l] * K32 + rotl32(y[l], 7u); y[l] = (y[l] ^ x[l]) + s; }
|
|
for (int l = 0; l < 32; ++l) res[l] = x[l] ^ y[l];
|
|
return;
|
|
}
|
|
for (uint32_t s = 0; s < steps; ++s) {
|
|
memcpy(xs, x, sizeof xs);
|
|
if (!strcmp(name, "mm8")) {
|
|
uint32_t D[8][8];
|
|
for (int r = 0; r < 8; ++r) for (int n = 0; n < 8; ++n) {
|
|
uint32_t c = (n & 1) ? acc2[r * 4 + n / 2] : acc[r * 4 + n / 2];
|
|
for (int k = 0; k < 16; ++k) {
|
|
uint32_t a = (xs[r * 4 + k / 4] >> (8 * (k % 4))) & 0xffu;
|
|
uint32_t b = (y[n * 4 + k / 4] >> (8 * (k % 4))) & 0xffu;
|
|
c += a * b;
|
|
}
|
|
D[r][n] = c;
|
|
}
|
|
for (int l = 0; l < 32; ++l) { acc[l] = D[l / 4][2 * (l % 4)]; acc2[l] = D[l / 4][2 * (l % 4) + 1]; }
|
|
for (int l = 0; l < 32; ++l) { x[l] = xs[l] * K32 + acc[l]; y[l] = rotl32(y[l], 7u) ^ (acc[l] + s); }
|
|
continue;
|
|
}
|
|
for (int l = 0; l < 32; ++l) {
|
|
uint32_t xv = xs[l], yv = y[l];
|
|
if (!strcmp(name, "rotr")) acc[l] = rotr_var(yv, xv);
|
|
else if (!strcmp(name, "shflx")) acc[l] = acc[l] ^ xs[l ^ 8];
|
|
else if (!strcmp(name, "shl")) acc[l] = yv << (xv & 31u);
|
|
else if (!strcmp(name, "shr")) acc[l] = yv >> (xv & 31u);
|
|
else if (!strcmp(name, "bfe") || !strcmp(name, "bfec")) acc[l] = (yv >> 7u) & 0x1fffu;
|
|
else if (!strcmp(name, "andn")) acc[l] = yv & ~xv;
|
|
else if (!strcmp(name, "perm")) acc[l] = perm_ref(yv);
|
|
else if (!strcmp(name, "popc")) acc[l] = acc[l] + popc32(xv);
|
|
else if (!strcmp(name, "clz")) acc[l] = acc[l] + clz32(xv);
|
|
else if (!strcmp(name, "sel")) acc[l] = ((yv >> 5u) & 1u) ? xv : acc[l];
|
|
else if (!strcmp(name, "shfla")) acc[l] = acc[l] ^ xs[(l + 3) & 31];
|
|
else if (!strcmp(name, "dot4i")) acc[l] = dot4s_ref(xv, yv, acc[l]);
|
|
else { printf("no reference for %s\n", name); exit(3); }
|
|
x[l] = xv * K32 + acc[l];
|
|
y[l] = rotl32(yv, 7u) ^ (acc[l] + s);
|
|
}
|
|
}
|
|
for (int l = 0; l < 32; ++l) res[l] = acc[l] ^ x[l] ^ y[l] ^ (!strcmp(name, "mm8") ? acc2[l] : 0u);
|
|
}
|
|
|
|
#define CK(x) do { cudaError_t e = (x); if (e != cudaSuccess) { printf("CUDA error %s at %s:%d\n", cudaGetErrorString(e), __FILE__, __LINE__); return 1; } } while (0)
|
|
typedef void (*kernel_t)(uint32_t, uint32_t, uint32_t*);
|
|
|
|
int main(int argc, char** argv) {
|
|
uint32_t lanes = 1u << 20, steps = 4096u; int reps = 3, device = 0;
|
|
for (int i = 1; i < argc; ++i) {
|
|
if (!strcmp(argv[i], "--lanes") && i + 1 < argc) lanes = (uint32_t)strtoul(argv[++i], 0, 10);
|
|
else if (!strcmp(argv[i], "--steps") && i + 1 < argc) steps = (uint32_t)strtoul(argv[++i], 0, 10);
|
|
else if (!strcmp(argv[i], "--reps") && i + 1 < argc) reps = atoi(argv[++i]);
|
|
else if (!strcmp(argv[i], "--device") && i + 1 < argc) device = atoi(argv[++i]);
|
|
else { printf("unknown argument %s\n", argv[i]); return 2; }
|
|
}
|
|
CK(cudaSetDevice(device));
|
|
cudaDeviceProp p; CK(cudaGetDeviceProperties(&p, device));
|
|
printf("family-probe (CUDA) on %s, sm_%d%d, %d SMs, %d MHz, lanes %u, steps %u, best of %d, event time\n", p.name, p.major, p.minor, p.multiProcessorCount, p.clockRate / 1000, lanes, steps, reps);
|
|
uint32_t* d_out; CK(cudaMalloc(&d_out, (size_t)lanes * 4));
|
|
uint32_t* h_out = (uint32_t*)malloc((size_t)lanes * 4);
|
|
cudaEvent_t e0, e1; CK(cudaEventCreate(&e0)); CK(cudaEventCreate(&e1));
|
|
const char* names[] = { "alu", "rotr", "shflx", "shl", "shr", "bfe", "bfec", "andn", "perm", "popc", "clz", "sel", "shfla", "dot4i", "mm8" };
|
|
kernel_t kernels[] = { probe_alu, probe_rotr, probe_shflx, probe_shl, probe_shr, probe_bfe, probe_bfec, probe_andn, probe_perm, probe_popc, probe_clz, probe_sel, probe_shfla, probe_dot4i, probe_mm8 };
|
|
const int nk = (int)(sizeof(names) / sizeof(names[0]));
|
|
printf("| kernel | lanes | steps | best ms | G steps/s | ns per step | ratio to alu | warps 0 and last ok |\n|---|---|---|---|---|---|---|---|\n");
|
|
float alu_best = 0;
|
|
for (int k = 0; k < nk; ++k) {
|
|
float best = 1e30f; int ok = 1;
|
|
for (int r = 0; r < reps; ++r) {
|
|
uint32_t seed = 0x2468aceu + (uint32_t)r * 0x9E3779B9u;
|
|
CK(cudaEventRecord(e0));
|
|
kernels[k]<<<lanes / 256, 256>>>(steps, seed, d_out);
|
|
CK(cudaEventRecord(e1)); CK(cudaEventSynchronize(e1)); CK(cudaGetLastError());
|
|
float ms = 0; CK(cudaEventElapsedTime(&ms, e0, e1)); if (ms < best) best = ms;
|
|
CK(cudaMemcpy(h_out, d_out, (size_t)lanes * 4, cudaMemcpyDeviceToHost));
|
|
uint32_t g0s[2] = { 0u, lanes - 32u };
|
|
for (int j = 0; j < 2; ++j) {
|
|
uint32_t want[32]; warp_ref(names[k], g0s[j], seed, steps, want);
|
|
for (int l = 0; l < 32; ++l) if (h_out[g0s[j] + l] != want[l]) { ok = 0; printf("MISMATCH %s lane %u: gpu %08x cpu %08x\n", names[k], g0s[j] + l, h_out[g0s[j] + l], want[l]); break; }
|
|
}
|
|
}
|
|
if (k == 0) alu_best = best;
|
|
double sps = (double)lanes * (double)steps / (best / 1000.0);
|
|
double ratio = alu_best > 0 ? best / alu_best : 0;
|
|
printf("| %s | %u | %u | %.3f | %.2f | %.3f | %.2f | %s |\n", names[k], lanes, steps, best, sps / 1e9, best * 1e6 / (double)steps, ratio, ok ? "yes" : "NO");
|
|
printf("RESULT FAMILY vendor=nvidia device=\"%s\" kernel=%s lanes=%u steps=%u best_ms=%.3f gsteps_per_s=%.2f ns_per_step=%.3f ratio_alu=%.3f ok=%d\n", p.name, names[k], lanes, steps, best, sps / 1e9, best * 1e6 / (double)steps, ratio, ok);
|
|
}
|
|
printf("family-probe: done\n");
|
|
return 0;
|
|
}
|