Counter ASIC 3.0 PC 1 AMD: OpenCL family probe (every 1.13.2 family through AMD's builtins, khr/intel sub-groups, media ops, C forms and local emulation; bit-exact on Apple OpenCL) and the clBuildProgram time line in the OpenCL worker
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
This commit is contained in:
igneum-labs
parent
4ae285f25e
commit
eb19dfdb18
2 changed files with 398 additions and 0 deletions
393
proto-opencl/family-probe.c
Normal file
393
proto-opencl/family-probe.c
Normal file
|
|
@ -0,0 +1,393 @@
|
|||
/* 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 N] [--lanes N] [--steps N] [--reps N] [--only name[,name]]
|
||||
*/
|
||||
#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]; } 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);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
typedef struct { double bestMs; int built; int exact; } Run; /* exact: 1 yes, 0 no, -1 unverified */
|
||||
|
||||
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\"\n",
|
||||
v->family, best, sps / 1e9, ratio, ex, v->path, v->variant, best * 1e6 / (double)steps, lanes, steps, vendor, dname);
|
||||
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 = 0, 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;
|
||||
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 { 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\n", i, devs[i].dname, devs[i].pname, devs[i].driver, devs[i].ver); if (ndevs == 0) printf("no OpenCL GPU devices\n"); return ndevs ? 0 : 1; }
|
||||
if (device < 0 || device >= ndevs) { printf("no device %d (have %d)\n", device, ndevs); return 2; }
|
||||
v = &devs[device];
|
||||
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, lanes %u, steps %u, best of %d, device event time\n", device, v->dname, v->pname, v->driver, v->ver, 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\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);
|
||||
}
|
||||
clReleaseMemObject(out); clReleaseCommandQueue(q); clReleaseContext(ctx); free(host);
|
||||
printf("family-probe: done\n");
|
||||
return 0;
|
||||
}
|
||||
|
|
@ -510,6 +510,7 @@ static const char* exchangeName(int m) { return m == 1 ? "sub_group_shuffle_xor
|
|||
// Returns 0 on success, 1 on build failure (log printed).
|
||||
static int buildProgram(Device* dv, const DeviceInfo* di, const char* src, size_t srcLen, int exchangeMode, int groupSize, const char* extra) {
|
||||
cl_int err = 0;
|
||||
double tb = wallMs(); /* the pack's compile cost (Counter ASIC 3.0 item 2, 6 October 2026): printed as one line below */
|
||||
const char* std;
|
||||
// The sub-group built-ins need OpenCL C 2.0 or 3.0. OpenCL 3.0 devices may report "OpenCL C 1.2" as the default
|
||||
// CL_DEVICE_OPENCL_C_VERSION while supporting 3.0 (the 3.0 API lists all versions; the 1.2 API cannot ask), so
|
||||
|
|
@ -534,6 +535,10 @@ static int buildProgram(Device* dv, const DeviceInfo* di, const char* src, size_
|
|||
clReleaseProgram(dv->prog); dv->prog = NULL;
|
||||
return 1;
|
||||
}
|
||||
/* The OpenCL build time of this pack's kernel text (clBuildProgram alone), the equivalent of the CUDA worker's
|
||||
* NVRTC line: a per-day item-derivation program (item 2) sits inside every hash-kernel compile, so its cost is
|
||||
* read here. Wall time, printed before the kernels are created. */
|
||||
printf("build %.1f ms clBuildProgram (exchange %d, group %d)\n", wallMs() - tb, exchangeMode, groupSize);
|
||||
dv->kHash = clCreateKernel(dv->prog, "igneum_hash", &err); CL_CHECK_ERR(err, "clCreateKernel igneum_hash");
|
||||
dv->kHashBound = clCreateKernel(dv->prog, "igneum_hash_bound", &err);
|
||||
if (err != CL_SUCCESS) dv->kHashBound = NULL; /* kernel.cl without the bound kernel: fine outside --serve */
|
||||
|
|
|
|||
Loading…
Reference in a new issue