igneum-network/igneum
1
0
Fork 0
Code Issues Pull requests Releases Activity
igneum/proto-opencl/family-probe.c

410 lines
29 KiB
C
Raw Permalink Blame History

/* family-probe (OpenCL): the step cost of every reserve candidate family of spec 1.13.2 on AMD (and any OpenCL GPU),
* standalone (no pack, no lottery kernel). Counter ASIC 3.0 item 6 (docs/plans/counter-asic-3-reserve.md), the AMD
* port of proto-cuda/family-probe.cu and proto-metal/family-probe.swift, 6 October 2026. PC 1 job on the RX 9070 XT;
* the Mac runs it on Apple OpenCL for the C-form and emulated rows only (every vendor builtin fails to build there,
* which is reported as a row, not an error).
*
* Same method as the CUDA probe: a dependent chain of one op per step per lane, 1,048,576 lanes x 4,096 steps, best
* of N, device event time, bit-exact against a CPU reference on two whole 32-lane groups (lanes 0..31 and the last
* 32). Every chain has the dot4 probe's glue: acc = OP(acc, x, y); x = x * K + acc; y = rotl(y, 7) ^ (acc + s).
*
* A family is measured through every form AMD's OpenCL C might give it, each built on its own (a variant the
* platform cannot compile prints one build=failed row with the first line of the build log and the run goes on):
* path=native an explicit intrinsic or builtin: __builtin_amdgcn_ds_bpermute / ds_swizzle (the clang builtins the
* LC compiler accepts, the same path the 5 October dot4 row took through __builtin_amdgcn_sudot4),
* sub_group_shuffle[_xor] (cl_khr_subgroup_shuffle), intel_sub_group_shuffle[_xor]
* (cl_intel_subgroups), amd_bfe / amd_perm (cl_amd_media_ops2), __builtin_amdgcn_sudot4, the
* cl_khr_integer_dot_product dot(), and the WMMA builtins for mm8
* path=sequence the plain OpenCL C form (rotate, shifts, masks, popcount, clz, the ternary select): whatever
* instruction sequence the compiler emits; one instruction or several is not read from the ISA here
* (no disassembler in the job), it is inferred from the step cost against rotr
* path=emulated a form that is certainly several instructions: the byte permute by shifts and masks, the lane
* shuffles through __local memory and two barriers, the dot4 by four byte products
* The families (the names of counter-asic-3-reserve.md and the CUDA probe): alu (reference chain, 5 ops a step, no
* acc), rotr, shflx (lane XOR 8), shl, shr, bfe (bits 7..19 of y), andn, perm (bytes y.b1, y.b3, y.b0, y.b2), popc,
* clz (clz(0) = 32), sel (bit 5 of y ? x : acc), shfla (lane + 3 mod 32), dot4 (signed, dp4a.s32.s32 semantics),
* mm8 (one 16x16x16 int8 WMMA per step per lane: the fragment layout of RDNA 4's WMMA is not verified against a CPU
* reference here, so that row is exact=unverified and OWED for exactness; its step cost stands if it builds).
*
* Output: one line per variant
* RESULT FAMILY name=<family> ms=<best> gsteps=<G steps/s> ratio=<best / alu best> exact=yes|no|unverified
* path=native|sequence|emulated variant=<name> ns=<ns per step> vendor=<v> device="<name>"
* RESULT FAMILY name=<family> variant=<name> path=<p> build=failed log="<first line>"
* and, per family, the row the status file takes (the fastest exact variant, native before sequence before emulated):
* RESULT FAMILYBEST name=<family> ms=<best> gsteps=<x> ratio=<r> exact=<e> path=<p> variant=<name>
*
* Build (Mac, Apple OpenCL): cc -std=c99 -O2 -Wno-deprecated-declarations -o family-probe-cl family-probe.c -framework OpenCL
* Build (Windows, mingw, no SDK): x86_64-w64-mingw32-gcc -std=c99 -O2 -static -DIGNEUM_CL_DYNAMIC -DCL_TARGET_OPENCL_VERSION=120 \
* -I <redist>/include -o family-probe-cl.exe family-probe.c (OpenCL.dll loaded at run time)
* Run: family-probe-cl [--list] --device-name <substring> | --device N [--lanes N] [--steps N] [--reps N] [--only name[,name]]
* (--device-name picks the card by name on its newest platform and is the job's way; a bare
* ordinal is never the default)
*/
#define CL_TARGET_OPENCL_VERSION 120
#ifdef __APPLE__
#include <OpenCL/cl.h>
#else
#include <CL/cl.h>
#endif
#ifdef IGNEUM_CL_DYNAMIC
#include "cl_dynamic.h"
#endif
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <stdint.h>
/* ---- the kernel text: one COMMON block and one body per variant, every variant its own program ---- */
static const char* COMMON =
"static inline uint pm_mix(uint x) { x ^= x >> 16; x *= 0x7feb352du; x ^= x >> 15; x *= 0x846ca68bu; x ^= x >> 16; return x; }\n"
"static inline uint rotr_var(uint x, uint n) { n &= 31u; return (x >> n) | (x << ((32u - n) & 31u)); }\n"
"#define K32 0x9E3779B1u\n"
"#define CHAIN_HEAD uint g = (uint)get_global_id(0); uint lid = (uint)get_local_id(0); uint lane = lid & 31u; (void)lane; uint x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u; uint acc = pm_mix(x);\n"
"#define CHAIN_TAIL x = x * K32 + acc; y = rotate(y, 7u) ^ (acc + s);\n"
"#define CHAIN_OUT out[g] = acc ^ x ^ y;\n"
"#define CHAIN(OP) __kernel void probe(uint steps, uint seed, __global uint* out) { CHAIN_HEAD for (uint s = 0u; s < steps; ++s) { OP; CHAIN_TAIL } CHAIN_OUT }\n"
/* the lane id inside the wave on AMD (mbcnt of the full mask), so a ds_bpermute address stays inside the lane's
own 32-lane group whether the compiler chose wave32 or wave64 */
"#define AMD_WAVE_LANE (__builtin_amdgcn_mbcnt_hi(~0u, __builtin_amdgcn_mbcnt_lo(~0u, 0u)))\n";
static const char* K_ALU =
"__kernel void probe(uint steps, uint seed, __global uint* out) {\n"
" uint g = (uint)get_global_id(0); uint x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u;\n"
" for (uint s = 0u; s < steps; ++s) { x = x * K32 + rotate(y, 7u); y = (y ^ x) + s; }\n"
" out[g] = x ^ y;\n"
"}\n";
static const char* K_ROTR = "CHAIN(acc = rotr_var(y, x))\n";
static const char* K_SHL = "CHAIN(acc = y << (x & 31u))\n";
static const char* K_SHR = "CHAIN(acc = y >> (x & 31u))\n";
static const char* K_ANDN = "CHAIN(acc = y & ~x)\n";
static const char* K_POPC = "CHAIN(acc = acc + popcount(x))\n";
static const char* K_CLZ = "CHAIN(acc = acc + clz(x))\n";
static const char* K_SEL = "CHAIN(acc = ((y >> 5u) & 1u) ? x : acc)\n";
static const char* K_BFE_C = "CHAIN(acc = (y >> 7u) & 0x1fffu)\n";
static const char* K_BFE_AMD =
"#pragma OPENCL EXTENSION cl_amd_media_ops2 : enable\n"
"CHAIN(acc = amd_bfe(y, 7u, 13u))\n";
static const char* K_PERM_C =
"CHAIN(acc = ((y >> 8) & 0xffu) | (((y >> 24) & 0xffu) << 8) | ((y & 0xffu) << 16) | (((y >> 16) & 0xffu) << 24))\n";
/* v_perm_b32 through cl_amd_media_ops2: both sources are y, so the selector bytes 1, 3, 0, 2 pick y's bytes whichever
source the extension puts in the low half; the CPU check settles the byte order */
static const char* K_PERM_AMD =
"#pragma OPENCL EXTENSION cl_amd_media_ops2 : enable\n"
"CHAIN(acc = amd_perm(y, y, 0x02000301u))\n";
/* the lane shuffles */
static const char* K_SHFLX_BPERM =
"CHAIN(uint wl = AMD_WAVE_LANE; acc = acc ^ (uint)__builtin_amdgcn_ds_bpermute((int)((wl ^ 8u) << 2), (int)x))\n";
static const char* K_SHFLX_SWZ =
/* ds_swizzle bitmask mode: and_mask 0x1f, or_mask 0, xor_mask 8 -> offset 0x201f (lanes inside each 32-group) */
"CHAIN(acc = acc ^ (uint)__builtin_amdgcn_ds_swizzle((int)x, 0x201f))\n";
static const char* K_SHFLX_KHR =
"#pragma OPENCL EXTENSION cl_khr_subgroup_shuffle : enable\n"
"CHAIN(acc = acc ^ sub_group_shuffle_xor(x, 8u))\n";
static const char* K_SHFLX_INTEL =
"#pragma OPENCL EXTENSION cl_intel_subgroups : enable\n"
"CHAIN(acc = acc ^ intel_sub_group_shuffle_xor(x, 8u))\n";
static const char* K_SHFLX_LOCAL =
"__kernel void probe(uint steps, uint seed, __global uint* out) {\n"
" __local uint buf[256];\n"
" CHAIN_HEAD\n"
" for (uint s = 0u; s < steps; ++s) { buf[lid] = x; barrier(CLK_LOCAL_MEM_FENCE); uint v = buf[lid ^ 8u]; barrier(CLK_LOCAL_MEM_FENCE); acc = acc ^ v; CHAIN_TAIL }\n"
" CHAIN_OUT\n"
"}\n";
static const char* K_SHFLA_BPERM =
"CHAIN(uint wl = AMD_WAVE_LANE; acc = acc ^ (uint)__builtin_amdgcn_ds_bpermute((int)(((wl & ~31u) | ((wl + 3u) & 31u)) << 2), (int)x))\n";
static const char* K_SHFLA_KHR =
"#pragma OPENCL EXTENSION cl_khr_subgroup_shuffle : enable\n"
"CHAIN(uint sl = get_sub_group_local_id(); acc = acc ^ sub_group_shuffle(x, (sl & ~31u) | ((sl + 3u) & 31u)))\n";
static const char* K_SHFLA_INTEL =
"#pragma OPENCL EXTENSION cl_intel_subgroups : enable\n"
"CHAIN(uint sl = get_sub_group_local_id(); acc = acc ^ intel_sub_group_shuffle(x, (sl & ~31u) | ((sl + 3u) & 31u)))\n";
static const char* K_SHFLA_LOCAL =
"__kernel void probe(uint steps, uint seed, __global uint* out) {\n"
" __local uint buf[256];\n"
" CHAIN_HEAD\n"
" for (uint s = 0u; s < steps; ++s) { buf[lid] = x; barrier(CLK_LOCAL_MEM_FENCE); uint v = buf[(lid & ~31u) | ((lid + 3u) & 31u)]; barrier(CLK_LOCAL_MEM_FENCE); acc = acc ^ v; CHAIN_TAIL }\n"
" CHAIN_OUT\n"
"}\n";
/* dot4, signed bytes, wrapping accumulate (dp4a.s32.s32 semantics, the CUDA probe's dot4i row) */
static const char* K_DOT4_C =
"static inline uint dot4e(uint a, uint b, uint acc) { int4 va = convert_int4(as_char4(a)); int4 vb = convert_int4(as_char4(b)); return acc + (uint)(va.x * vb.x + va.y * vb.y + va.z * vb.z + va.w * vb.w); }\n"
"CHAIN(acc = dot4e(x, y, acc))\n";
static const char* K_DOT4_AMD =
"CHAIN(acc = (uint)__builtin_amdgcn_sudot4(true, (int)x, true, (int)y, (int)acc, false))\n";
static const char* K_DOT4_KHR =
"#pragma OPENCL EXTENSION cl_khr_integer_dot_product : enable\n"
"CHAIN(acc = acc + (uint)dot(as_char4(x), as_char4(y)))\n";
/* mm8: one 16x16x16 int8 WMMA per step per lane through the clang builtins (gfx12 first, the gfx11 shape second).
8 bytes of A and B per lane on gfx12 (int2), 16 (int4) on gfx11; 8 accumulators per lane (int8); the chain carries
c.s0 forward as acc. No CPU reference for the fragment layout here: exact=unverified. */
static const char* K_MM8_GFX12 =
"__kernel void probe(uint steps, uint seed, __global uint* out) {\n"
" CHAIN_HEAD int8 c = (int8)((int)acc, (int)pm_mix(acc), 0, 0, 0, 0, 0, 0);\n"
" for (uint s = 0u; s < steps; ++s) { int2 a = (int2)((int)x, (int)y); int2 b = (int2)((int)(y ^ s), (int)x);\n"
" c = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12(true, a, true, b, c, false); acc = (uint)c.s0; CHAIN_TAIL }\n"
" CHAIN_OUT\n"
"}\n";
static const char* K_MM8_GFX11 =
"__kernel void probe(uint steps, uint seed, __global uint* out) {\n"
" CHAIN_HEAD int8 c = (int8)((int)acc, (int)pm_mix(acc), 0, 0, 0, 0, 0, 0);\n"
" for (uint s = 0u; s < steps; ++s) { int4 a = (int4)((int)x, (int)y, (int)(x ^ s), (int)(y + s)); int4 b = (int4)((int)(y ^ s), (int)x, (int)y, (int)x);\n"
" c = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32(true, a, true, b, c, false); acc = (uint)c.s0; CHAIN_TAIL }\n"
" CHAIN_OUT\n"
"}\n";
typedef struct { const char* family; const char* variant; const char* path; const char* body; int vendorOnly; } Variant;
/* vendorOnly: 0 = every platform, 1 = AMD builtins (clang), 2 = NVIDIA only (none here) */
static const Variant VARIANTS[] = {
{ "alu", "alu", "sequence", NULL, 0 },
{ "rotr", "rotr_c", "sequence", NULL, 0 },
{ "shflx", "shflx_bperm", "native", NULL, 1 },
{ "shflx", "shflx_swz", "native", NULL, 1 },
{ "shflx", "shflx_khr", "native", NULL, 0 },
{ "shflx", "shflx_intel", "native", NULL, 0 },
{ "shflx", "shflx_local", "emulated", NULL, 0 },
{ "shl", "shl_c", "sequence", NULL, 0 },
{ "shr", "shr_c", "sequence", NULL, 0 },
{ "bfe", "bfe_amd", "native", NULL, 0 },
{ "bfe", "bfe_c", "sequence", NULL, 0 },
{ "andn", "andn_c", "sequence", NULL, 0 },
{ "perm", "perm_amd", "native", NULL, 0 },
{ "perm", "perm_c", "emulated", NULL, 0 },
{ "popc", "popc_c", "sequence", NULL, 0 },
{ "clz", "clz_c", "sequence", NULL, 0 },
{ "sel", "sel_c", "sequence", NULL, 0 },
{ "shfla", "shfla_bperm", "native", NULL, 1 },
{ "shfla", "shfla_khr", "native", NULL, 0 },
{ "shfla", "shfla_intel", "native", NULL, 0 },
{ "shfla", "shfla_local", "emulated", NULL, 0 },
{ "dot4", "dot4_amd", "native", NULL, 1 },
{ "dot4", "dot4_khr", "native", NULL, 0 },
{ "dot4", "dot4_c", "emulated", NULL, 0 },
{ "mm8", "mm8_gfx12", "native", NULL, 1 },
{ "mm8", "mm8_gfx11", "native", NULL, 1 },
};
static const int NVARIANTS = (int)(sizeof(VARIANTS) / sizeof(VARIANTS[0]));
static const char* bodyOf(const char* variant) {
if (!strcmp(variant, "alu")) return K_ALU;
if (!strcmp(variant, "rotr_c")) return K_ROTR;
if (!strcmp(variant, "shflx_bperm")) return K_SHFLX_BPERM;
if (!strcmp(variant, "shflx_swz")) return K_SHFLX_SWZ;
if (!strcmp(variant, "shflx_khr")) return K_SHFLX_KHR;
if (!strcmp(variant, "shflx_intel")) return K_SHFLX_INTEL;
if (!strcmp(variant, "shflx_local")) return K_SHFLX_LOCAL;
if (!strcmp(variant, "shl_c")) return K_SHL;
if (!strcmp(variant, "shr_c")) return K_SHR;
if (!strcmp(variant, "bfe_amd")) return K_BFE_AMD;
if (!strcmp(variant, "bfe_c")) return K_BFE_C;
if (!strcmp(variant, "andn_c")) return K_ANDN;
if (!strcmp(variant, "perm_amd")) return K_PERM_AMD;
if (!strcmp(variant, "perm_c")) return K_PERM_C;
if (!strcmp(variant, "popc_c")) return K_POPC;
if (!strcmp(variant, "clz_c")) return K_CLZ;
if (!strcmp(variant, "sel_c")) return K_SEL;
if (!strcmp(variant, "shfla_bperm")) return K_SHFLA_BPERM;
if (!strcmp(variant, "shfla_khr")) return K_SHFLA_KHR;
if (!strcmp(variant, "shfla_intel")) return K_SHFLA_INTEL;
if (!strcmp(variant, "shfla_local")) return K_SHFLA_LOCAL;
if (!strcmp(variant, "dot4_amd")) return K_DOT4_AMD;
if (!strcmp(variant, "dot4_khr")) return K_DOT4_KHR;
if (!strcmp(variant, "dot4_c")) return K_DOT4_C;
if (!strcmp(variant, "mm8_gfx12")) return K_MM8_GFX12;
if (!strcmp(variant, "mm8_gfx11")) return K_MM8_GFX11;
return NULL;
}
/* ---- CPU reference: one whole 32-lane group (lanes g0 .. g0+31), the CUDA probe's warp_ref without mm8 ---- */
static uint32_t pm_mix(uint32_t x) { x ^= x >> 16; x *= 0x7feb352du; x ^= x >> 15; x *= 0x846ca68bu; x ^= x >> 16; return x; }
static uint32_t rotl32(uint32_t v, uint32_t n) { return (v << n) | (v >> (32u - n)); }
static uint32_t rotr_var(uint32_t x, uint32_t n) { n &= 31u; return (x >> n) | (x << ((32u - n) & 31u)); }
static uint32_t clz32(uint32_t x) { uint32_t n = 0; if (!x) return 32; 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; int i;
for (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;
}
#define K32 0x9E3779B1u
static int group_ref(const char* family, uint32_t g0, uint32_t seed, uint32_t steps, uint32_t* res) {
uint32_t x[32], y[32], acc[32], xs[32], s; int l;
for (l = 0; l < 32; ++l) { x[l] = pm_mix((g0 + (uint32_t)l) ^ seed); y[l] = x[l] ^ 0x5bd1e995u; acc[l] = pm_mix(x[l]); }
if (!strcmp(family, "alu")) {
for (s = 0; s < steps; ++s) for (l = 0; l < 32; ++l) { x[l] = x[l] * K32 + rotl32(y[l], 7u); y[l] = (y[l] ^ x[l]) + s; }
for (l = 0; l < 32; ++l) res[l] = x[l] ^ y[l];
return 1;
}
if (!strcmp(family, "mm8")) return 0; /* no reference: the WMMA fragment layout is not modelled here */
for (s = 0; s < steps; ++s) {
memcpy(xs, x, sizeof xs);
for (l = 0; l < 32; ++l) {
uint32_t xv = xs[l], yv = y[l];
if (!strcmp(family, "rotr")) acc[l] = rotr_var(yv, xv);
else if (!strcmp(family, "shflx")) acc[l] = acc[l] ^ xs[l ^ 8];
else if (!strcmp(family, "shl")) acc[l] = yv << (xv & 31u);
else if (!strcmp(family, "shr")) acc[l] = yv >> (xv & 31u);
else if (!strcmp(family, "bfe")) acc[l] = (yv >> 7u) & 0x1fffu;
else if (!strcmp(family, "andn")) acc[l] = yv & ~xv;
else if (!strcmp(family, "perm")) acc[l] = perm_ref(yv);
else if (!strcmp(family, "popc")) acc[l] = acc[l] + popc32(xv);
else if (!strcmp(family, "clz")) acc[l] = acc[l] + clz32(xv);
else if (!strcmp(family, "sel")) acc[l] = ((yv >> 5u) & 1u) ? xv : acc[l];
else if (!strcmp(family, "shfla")) acc[l] = acc[l] ^ xs[(l + 3) & 31];
else if (!strcmp(family, "dot4")) acc[l] = dot4s_ref(xv, yv, acc[l]);
else { printf("no reference for %s\n", family); exit(3); }
x[l] = xv * K32 + acc[l];
y[l] = rotl32(yv, 7u) ^ (acc[l] + s);
}
}
for (l = 0; l < 32; ++l) res[l] = acc[l] ^ x[l] ^ y[l];
return 1;
}
/* ---- devices ---- */
typedef struct { cl_platform_id p; cl_device_id d; char pname[128], dname[128], driver[64], ver[64]; cl_uint cus; } Dev;
static Dev devs[32]; static int ndevs = 0;
static void enumerate(void) {
cl_platform_id ps[8]; cl_uint np = 0, i;
if (clGetPlatformIDs(8, ps, &np) != CL_SUCCESS) return;
for (i = 0; i < np; ++i) {
cl_device_id ds[8]; cl_uint nd = 0, j;
if (clGetDeviceIDs(ps[i], CL_DEVICE_TYPE_GPU, 8, ds, &nd) != CL_SUCCESS) continue;
for (j = 0; j < nd && ndevs < 32; ++j) {
Dev* v = &devs[ndevs++]; v->p = ps[i]; v->d = ds[j];
clGetPlatformInfo(ps[i], CL_PLATFORM_NAME, sizeof v->pname, v->pname, NULL);
clGetDeviceInfo(ds[j], CL_DEVICE_NAME, sizeof v->dname, v->dname, NULL);
clGetDeviceInfo(ds[j], CL_DRIVER_VERSION, sizeof v->driver, v->driver, NULL);
clGetDeviceInfo(ds[j], CL_DEVICE_VERSION, sizeof v->ver, v->ver, NULL);
v->cus = 0; clGetDeviceInfo(ds[j], CL_DEVICE_MAX_COMPUTE_UNITS, sizeof v->cus, &v->cus, NULL);
}
}
}
typedef struct { double bestMs; int built; int exact; } Run; /* exact: 1 yes, 0 no, -1 unverified */
static const Dev* gDev = NULL; /* the device under test, for the RESULT lines */
static Run run_variant(cl_context ctx, cl_command_queue q, cl_device_id dev, const Variant* v, cl_uint lanes, cl_uint steps, int reps,
cl_mem out, uint32_t* host, const char* vendor, const char* dname, double aluBest) {
const char* srcs[2]; cl_int err; cl_program prog; cl_kernel k; int r, exact = 1, hasRef = 1; double best = -1; Run res = { -1, 0, 0 };
srcs[0] = COMMON; srcs[1] = v->body;
prog = clCreateProgramWithSource(ctx, 2, srcs, NULL, &err);
if (err != CL_SUCCESS) { printf("RESULT FAMILY name=%s variant=%s path=%s build=failed log=\"clCreateProgramWithSource %d\"\n", v->family, v->variant, v->path, (int)err); return res; }
err = clBuildProgram(prog, 1, &dev, "-cl-std=CL1.2", NULL, NULL);
if (err != CL_SUCCESS) {
size_t n = 0; char* log; char* p; char* nl;
clGetProgramBuildInfo(prog, dev, CL_PROGRAM_BUILD_LOG, 0, NULL, &n); log = (char*)calloc(n + 2, 1);
if (n) clGetProgramBuildInfo(prog, dev, CL_PROGRAM_BUILD_LOG, n, log, NULL);
p = log; while (*p == '\n' || *p == '\r' || *p == ' ') ++p;
nl = strpbrk(p, "\r\n"); if (nl) *nl = 0;
for (nl = p; *nl; ++nl) if (*nl == '"') *nl = '\'';
printf("RESULT FAMILY name=%s variant=%s path=%s build=failed log=\"%.160s\"\n", v->family, v->variant, v->path, p);
free(log); clReleaseProgram(prog); return res;
}
k = clCreateKernel(prog, "probe", &err);
if (err != CL_SUCCESS) { printf("RESULT FAMILY name=%s variant=%s path=%s build=failed log=\"clCreateKernel %d\"\n", v->family, v->variant, v->path, (int)err); clReleaseProgram(prog); return res; }
for (r = 0; r < reps; ++r) {
cl_uint seed = 0x2468aceu + (cl_uint)r * 0x9E3779B9u; size_t global = lanes, local = 256; cl_event ev; cl_ulong t0, t1; double ms; uint32_t g0s[2]; int j;
clSetKernelArg(k, 0, sizeof(cl_uint), &steps); clSetKernelArg(k, 1, sizeof(cl_uint), &seed); clSetKernelArg(k, 2, sizeof(cl_mem), &out);
err = clEnqueueNDRangeKernel(q, k, 1, NULL, &global, &local, 0, NULL, &ev);
if (err != CL_SUCCESS) { printf("RESULT FAMILY name=%s variant=%s path=%s build=failed log=\"launch failed %d\"\n", v->family, v->variant, v->path, (int)err); clReleaseKernel(k); clReleaseProgram(prog); return res; }
clWaitForEvents(1, &ev);
clGetEventProfilingInfo(ev, CL_PROFILING_COMMAND_START, sizeof t0, &t0, NULL); clGetEventProfilingInfo(ev, CL_PROFILING_COMMAND_END, sizeof t1, &t1, NULL);
ms = (double)(t1 - t0) / 1e6; clReleaseEvent(ev);
if (best < 0 || ms < best) best = ms;
clEnqueueReadBuffer(q, out, CL_TRUE, 0, (size_t)lanes * 4, host, 0, NULL, NULL);
g0s[0] = 0; g0s[1] = lanes - 32u;
for (j = 0; j < 2; ++j) {
uint32_t want[32]; int l;
if (!group_ref(v->family, g0s[j], seed, steps, want)) { hasRef = 0; break; }
for (l = 0; l < 32; ++l) if (host[g0s[j] + (uint32_t)l] != want[l]) { exact = 0; printf("MISMATCH %s lane %u: gpu %08x cpu %08x\n", v->variant, g0s[j] + (uint32_t)l, host[g0s[j] + (uint32_t)l], want[l]); break; }
}
}
{
double sps = (double)lanes * (double)steps / (best / 1000.0);
double ratio = aluBest > 0 ? best / aluBest : 0;
const char* ex = hasRef ? (exact ? "yes" : "no") : "unverified";
printf("RESULT FAMILY name=%s ms=%.3f gsteps=%.2f ratio=%.3f exact=%s path=%s variant=%s ns=%.3f lanes=%u steps=%u vendor=%s device=\"%s\" cus=%u platform=\"%s\"\n",
v->family, best, sps / 1e9, ratio, ex, v->path, v->variant, best * 1e6 / (double)steps, lanes, steps, vendor, dname, gDev->cus, gDev->pname);
res.bestMs = best; res.built = 1; res.exact = hasRef ? exact : -1;
}
clReleaseKernel(k); clReleaseProgram(prog); return res;
}
static int pathRank(const char* p) { return !strcmp(p, "native") ? 0 : !strcmp(p, "sequence") ? 1 : 2; }
int main(int argc, char** argv) {
cl_uint lanes = 1u << 20, steps = 4096u; int reps = 3, device = -1, list = 0, i; Dev* v; cl_int err; cl_context ctx; cl_command_queue q; cl_mem out; uint32_t* host; const char* vendor; const char* only = NULL; const char* devName = NULL;
Run runs[64]; double aluBest = 0; int isAmd;
for (i = 1; i < argc; ++i) {
if (!strcmp(argv[i], "--list")) list = 1;
else if (!strcmp(argv[i], "--device") && i + 1 < argc) device = atoi(argv[++i]);
else if (!strcmp(argv[i], "--lanes") && i + 1 < argc) lanes = (cl_uint)strtoul(argv[++i], 0, 10);
else if (!strcmp(argv[i], "--steps") && i + 1 < argc) steps = (cl_uint)strtoul(argv[++i], 0, 10);
else if (!strcmp(argv[i], "--reps") && i + 1 < argc) reps = atoi(argv[++i]);
else if (!strcmp(argv[i], "--only") && i + 1 < argc) only = argv[++i];
else if (!strcmp(argv[i], "--device-name") && i + 1 < argc) devName = argv[++i];
else { printf("unknown argument %s\n", argv[i]); return 2; }
}
if (lanes < 64 || (lanes & 255u)) { printf("--lanes must be a multiple of 256\n"); return 2; }
#ifdef IGNEUM_CL_DYNAMIC
if (!ig_cl_load()) { printf("%s\n", ig_cl_error); return 1; }
#endif
enumerate();
if (list || ndevs == 0) { for (i = 0; i < ndevs; ++i) printf("[%d] %s | %s | driver %s | %s | %u CUs\n", i, devs[i].dname, devs[i].pname, devs[i].driver, devs[i].ver, devs[i].cus); if (ndevs == 0) printf("no OpenCL GPU devices\n"); return ndevs ? 0 : 1; }
/* --device-name <substring>: the device whose name holds it, on the NEWEST platform when two platforms list the same
card (PC 1 lists the 9070 XT on AMD-APP 3683.0 and again on the older 3652.0: the kit worker hides the older one as
dup; here the highest driver version string wins). A bare ordinal is never the default: 6 October 2026, the first
family job ran ordinal 0, PC 1's integrated gfx1036, and reported one-CU figures as the 9070 XT's. */
if (devName) {
int best = -1;
for (i = 0; i < ndevs; ++i) if (strstr(devs[i].dname, devName) && (best < 0 || strcmp(devs[i].driver, devs[best].driver) > 0)) best = i;
if (best < 0) { printf("RESULT error no OpenCL GPU device whose name holds \"%s\" (listed: ", devName); for (i = 0; i < ndevs; ++i) printf("%s[%d] %s on %s", i ? "; " : "", i, devs[i].dname, devs[i].pname); printf(")\n"); return 2; }
device = best;
printf("RESULT device_choice name=\"%s\" index=%d platform=\"%s\" driver=%s cus=%u (of %d devices; the newest platform for that name)\n", devs[best].dname, best, devs[best].pname, devs[best].driver, devs[best].cus, ndevs);
}
if (device < 0) { printf("RESULT error no device chosen: pass --device-name <substring> (the card by name) or --device N\n"); return 2; }
if (device >= ndevs) { printf("RESULT error no device %d (have %d)\n", device, ndevs); return 2; }
v = &devs[device]; gDev = v;
vendor = strstr(v->pname, "NVIDIA") ? "nvidia" : (strstr(v->pname, "AMD") ? "amd" : (strstr(v->pname, "Apple") ? "apple" : "other"));
isAmd = !strcmp(vendor, "amd");
ctx = clCreateContext(NULL, 1, &v->d, NULL, NULL, &err); if (err != CL_SUCCESS) { printf("clCreateContext %d\n", (int)err); return 1; }
q = clCreateCommandQueue(ctx, v->d, CL_QUEUE_PROFILING_ENABLE, &err); if (err != CL_SUCCESS) { printf("clCreateCommandQueue %d\n", (int)err); return 1; }
out = clCreateBuffer(ctx, CL_MEM_READ_WRITE, (size_t)lanes * 4, NULL, &err); if (err != CL_SUCCESS) { printf("clCreateBuffer %d\n", (int)err); return 1; }
host = (uint32_t*)malloc((size_t)lanes * 4);
printf("family-probe (OpenCL) on [%d] %s | %s | driver %s | %s | %u CUs, lanes %u, steps %u, best of %d, device event time\n", device, v->dname, v->pname, v->driver, v->ver, v->cus, lanes, steps, reps);
{
char ext[16384]; ext[0] = 0; clGetDeviceInfo(v->d, CL_DEVICE_EXTENSIONS, sizeof ext, ext, NULL);
printf("RESULT EXT device=\"%s\" khr_subgroups=%s khr_subgroup_shuffle=%s intel_subgroups=%s amd_media_ops2=%s khr_integer_dot_product=%s\n", v->dname,
strstr(ext, "cl_khr_subgroups") ? "yes" : "no", strstr(ext, "cl_khr_subgroup_shuffle") ? "yes" : "no", strstr(ext, "cl_intel_subgroups") ? "yes" : "no",
strstr(ext, "cl_amd_media_ops2") ? "yes" : "no", strstr(ext, "cl_khr_integer_dot_product") ? "yes" : "no");
}
for (i = 0; i < NVARIANTS; ++i) {
Variant vv = VARIANTS[i]; vv.body = bodyOf(vv.variant);
runs[i].bestMs = -1; runs[i].built = 0; runs[i].exact = 0;
if (!vv.body) continue;
if (only && !strstr(only, vv.family)) continue;
if (vv.vendorOnly == 1 && !isAmd) { printf("RESULT FAMILY name=%s variant=%s path=%s build=skipped log=\"AMD clang builtin, platform is %s\"\n", vv.family, vv.variant, vv.path, vendor); continue; }
runs[i] = run_variant(ctx, q, v->d, &vv, lanes, steps, reps, out, host, vendor, v->dname, aluBest);
if (i == 0 && runs[i].built) aluBest = runs[i].bestMs;
}
/* the row per family: the fastest EXACT variant, native before sequence before emulated at equal exactness; a
family with no exact variant takes its fastest unverified one and says so */
for (i = 0; i < NVARIANTS; ++i) {
int j, bestIdx = -1, bestRank = 9; double bestMs = 1e30;
if (i > 0 && !strcmp(VARIANTS[i].family, VARIANTS[i - 1].family)) continue; /* one row per family, at its first variant */
if (only && !strstr(only, VARIANTS[i].family)) continue;
for (j = i; j < NVARIANTS && !strcmp(VARIANTS[j].family, VARIANTS[i].family); ++j) {
int rank;
if (!runs[j].built) continue;
rank = (runs[j].exact == 1 ? 0 : runs[j].exact == -1 ? 3 : 6) + pathRank(VARIANTS[j].path);
if (rank < bestRank || (rank == bestRank && runs[j].bestMs < bestMs)) { bestRank = rank; bestMs = runs[j].bestMs; bestIdx = j; }
}
if (bestIdx < 0) { printf("RESULT FAMILYBEST name=%s built=none\n", VARIANTS[i].family); continue; }
printf("RESULT FAMILYBEST name=%s ms=%.3f gsteps=%.2f ratio=%.3f exact=%s path=%s variant=%s device=\"%s\" cus=%u platform=\"%s\"\n", VARIANTS[i].family, runs[bestIdx].bestMs,
(double)lanes * (double)steps / (runs[bestIdx].bestMs / 1000.0) / 1e9, aluBest > 0 ? runs[bestIdx].bestMs / aluBest : 0,
runs[bestIdx].exact == 1 ? "yes" : runs[bestIdx].exact == -1 ? "unverified" : "no", VARIANTS[bestIdx].path, VARIANTS[bestIdx].variant, v->dname, v->cus, v->pname);
}
clReleaseMemObject(out); clReleaseCommandQueue(q); clReleaseContext(ctx); free(host);
printf("family-probe: done\n");
return 0;
}
Reference in a new issue View git blame Copy permalink