187 lines
12 KiB
C
187 lines
12 KiB
C
/* dot4-probe (OpenCL): dp4a-class throughput on AMD and NVIDIA through OpenCL, standalone (no pack, no lottery kernel).
|
|
* Counter ASIC 2.0 layer 7 (docs/analysis/int8-matrix-family.md), 5 October 2026. PC job; the Mac can run it on Apple
|
|
* OpenCL for the scalar rows only.
|
|
*
|
|
* Dependent chains, same shape as --memprobe's ALU chain (host.c, probe_alu) and the Metal and CUDA probes:
|
|
* 1,048,576 lanes x 4,096 steps, best of 3, device event time.
|
|
* alu x = x * K + rotate(y, 7); y = (y ^ x) + s the card's integer baseline
|
|
* dot4e scalar emulation: 4 sign-extended byte products summed into a wrapping int accumulator
|
|
* dot4_amd __builtin_amdgcn_sudot4(a, 1, b, 1, acc, 0): V_DOT4_I32_IU8 with both operands signed, no clamp (RDNA 3;
|
|
* the same builtin on RDNA 4 is the thing to check: LLVM gates it on the dot7-insts / dot8-insts features)
|
|
* dot4_khr cl_khr_integer_dot_product: acc + dot(as_char4(a), as_char4(b)) (OpenCL C 3.0 extension, section
|
|
* 6.2.2.16 of the OpenCL C specification; listed by a platform or not)
|
|
* dot4_nv inline PTX "dp4a.s32.s32" (NVIDIA's OpenCL compiler accepts inline PTX asm; PTX ISA 9.7.1.24)
|
|
* Each variant is 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. Every variant that runs is checked bit for bit against the CPU reference
|
|
* on two lanes, so an intrinsic with different semantics (a saturating accumulate, say) shows as NO in the ok column.
|
|
*
|
|
* Build (Mac, Apple OpenCL): cc -std=c99 -O2 -o dot4-probe-cl dot4-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 dot4-probe-cl.exe dot4-probe.c (OpenCL.dll loaded at run time)
|
|
* Run: dot4-probe-cl [--list] [--device N] [--lanes N] [--steps N] [--reps N]
|
|
*/
|
|
#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>
|
|
|
|
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"
|
|
"#define CHAIN_HEAD uint g = (uint)get_global_id(0); uint x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u; int acc = (int)pm_mix(x);\n"
|
|
"#define CHAIN_TAIL x = x * 0x9E3779B1u + (uint)acc; y = rotate(y, 7u) ^ ((uint)acc + s);\n"
|
|
"#define CHAIN_OUT out[g] = (uint)acc ^ x ^ y;\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 * 0x9E3779B1u + rotate(y, 7u); y = (y ^ x) + s; }\n"
|
|
" out[g] = x ^ y;\n"
|
|
"}\n";
|
|
static const char* K_EMUL =
|
|
"static inline int dot4e(uint a, uint b, int acc) {\n"
|
|
" int4 va = convert_int4(as_char4(a)); int4 vb = convert_int4(as_char4(b));\n"
|
|
" return acc + va.x * vb.x + va.y * vb.y + va.z * vb.z + va.w * vb.w;\n"
|
|
"}\n"
|
|
"__kernel void probe(uint steps, uint seed, __global uint* out) {\n"
|
|
" CHAIN_HEAD\n"
|
|
" for (uint s = 0u; s < steps; ++s) { acc = dot4e(x, y, acc); CHAIN_TAIL }\n"
|
|
" CHAIN_OUT\n"
|
|
"}\n";
|
|
static const char* K_AMD =
|
|
"__kernel void probe(uint steps, uint seed, __global uint* out) {\n"
|
|
" CHAIN_HEAD\n"
|
|
" for (uint s = 0u; s < steps; ++s) { acc = __builtin_amdgcn_sudot4(true, (int)x, true, (int)y, acc, false); CHAIN_TAIL }\n"
|
|
" CHAIN_OUT\n"
|
|
"}\n";
|
|
static const char* K_KHR =
|
|
"#pragma OPENCL EXTENSION cl_khr_integer_dot_product : enable\n"
|
|
"__kernel void probe(uint steps, uint seed, __global uint* out) {\n"
|
|
" CHAIN_HEAD\n"
|
|
" for (uint s = 0u; s < steps; ++s) { acc = acc + dot(as_char4(x), as_char4(y)); CHAIN_TAIL }\n"
|
|
" CHAIN_OUT\n"
|
|
"}\n";
|
|
static const char* K_NV =
|
|
"static inline int dot4nv(uint a, uint b, int acc) { int d; asm(\"dp4a.s32.s32 %0, %1, %2, %3;\" : \"=r\"(d) : \"r\"(a), \"r\"(b), \"r\"(acc)); return d; }\n"
|
|
"__kernel void probe(uint steps, uint seed, __global uint* out) {\n"
|
|
" CHAIN_HEAD\n"
|
|
" for (uint s = 0u; s < steps; ++s) { acc = dot4nv(x, y, acc); CHAIN_TAIL }\n"
|
|
" CHAIN_OUT\n"
|
|
"}\n";
|
|
|
|
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 int32_t dot4_ref(uint32_t a, uint32_t b, int32_t acc) {
|
|
int32_t r = acc; int i;
|
|
for (i = 0; i < 4; ++i) { int32_t ba = (int8_t)((a >> (8 * i)) & 0xffu), bb = (int8_t)((b >> (8 * i)) & 0xffu); r = (int32_t)((uint32_t)r + (uint32_t)(ba * bb)); }
|
|
return r;
|
|
}
|
|
static uint32_t lane_ref(int alu, uint32_t g, uint32_t seed, uint32_t steps) {
|
|
uint32_t x = pm_mix(g ^ seed), y = x ^ 0x5bd1e995u, s; int32_t acc;
|
|
if (alu) { for (s = 0; s < steps; ++s) { x = x * 0x9E3779B1u + rotl32(y, 7u); y = (y ^ x) + s; } return x ^ y; }
|
|
acc = (int32_t)pm_mix(x);
|
|
for (s = 0; s < steps; ++s) { acc = dot4_ref(x, y, acc); x = x * 0x9E3779B1u + (uint32_t)acc; y = rotl32(y, 7u) ^ ((uint32_t)acc + s); }
|
|
return (uint32_t)acc ^ x ^ y;
|
|
}
|
|
|
|
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);
|
|
}
|
|
}
|
|
}
|
|
|
|
static int run_variant(cl_context ctx, cl_command_queue q, cl_device_id dev, const char* name, const char* body, int alu, cl_uint lanes, cl_uint steps, int reps, cl_mem out, uint32_t* host, const char* vendor, const char* dname) {
|
|
const char* srcs[2] = { COMMON, body }; cl_int err; cl_program prog; cl_kernel k; int r, ok = 1; double best = -1;
|
|
prog = clCreateProgramWithSource(ctx, 2, srcs, NULL, &err);
|
|
if (err != CL_SUCCESS) { printf("| %s | build failed | clCreateProgramWithSource %d | | | | |\n", name, (int)err); return 0; }
|
|
err = clBuildProgram(prog, 1, &dev, "-cl-std=CL1.2", NULL, NULL);
|
|
if (err != CL_SUCCESS) {
|
|
size_t n = 0; char* log; char* nl;
|
|
clGetProgramBuildInfo(prog, dev, CL_PROGRAM_BUILD_LOG, 0, NULL, &n); log = (char*)calloc(n + 1, 1);
|
|
if (n) clGetProgramBuildInfo(prog, dev, CL_PROGRAM_BUILD_LOG, n, log, NULL);
|
|
while (*log == '\n' || *log == '\r') ++log;
|
|
nl = strpbrk(log, "\r\n"); if (nl) *nl = 0;
|
|
printf("| %s | build failed | %.160s | | | | |\n", name, log);
|
|
printf("RESULT DOT4 vendor=%s device=\"%s\" kernel=%s build=failed\n", vendor, dname, name);
|
|
clReleaseProgram(prog); return 0;
|
|
}
|
|
k = clCreateKernel(prog, "probe", &err);
|
|
if (err != CL_SUCCESS) { printf("| %s | build failed | clCreateKernel %d | | | | |\n", name, (int)err); clReleaseProgram(prog); return 0; }
|
|
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 gs[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("| %s | launch failed | %d | | | | |\n", name, (int)err); clReleaseKernel(k); clReleaseProgram(prog); return 0; }
|
|
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);
|
|
gs[0] = 0; gs[1] = lanes - 1;
|
|
for (j = 0; j < 2; ++j) { uint32_t want = lane_ref(alu, gs[j], seed, steps); if (host[gs[j]] != want) { ok = 0; printf("MISMATCH %s lane %u: gpu %08x cpu %08x\n", name, gs[j], host[gs[j]], want); } }
|
|
}
|
|
{
|
|
double sps = (double)lanes * (double)steps / (best / 1000.0);
|
|
printf("| %s | %u | %u | %.3f | %.2f | %.3f | %s |\n", name, lanes, steps, best, sps / 1e9, best * 1e6 / (double)steps, ok ? "yes" : "NO");
|
|
printf("RESULT DOT4 vendor=%s device=\"%s\" kernel=%s lanes=%u steps=%u best_ms=%.3f gsteps_per_s=%.2f ok=%d\n", vendor, dname, name, lanes, steps, best, sps / 1e9, ok);
|
|
}
|
|
clReleaseKernel(k); clReleaseProgram(prog); return 1;
|
|
}
|
|
|
|
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;
|
|
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 { printf("unknown argument %s\n", argv[i]); 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"));
|
|
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("dot4-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[8192]; ext[0] = 0; clGetDeviceInfo(v->d, CL_DEVICE_EXTENSIONS, sizeof ext, ext, NULL);
|
|
printf("cl_khr_integer_dot_product listed: %s\n", strstr(ext, "cl_khr_integer_dot_product") ? "yes" : "no");
|
|
}
|
|
printf("| kernel | lanes | steps | best ms | G steps/s (= G dot4/s for dot4 rows) | ns per dependent step | lanes 0 and last ok |\n|---|---|---|---|---|---|---|\n");
|
|
run_variant(ctx, q, v->d, "alu", K_ALU, 1, lanes, steps, reps, out, host, vendor, v->dname);
|
|
run_variant(ctx, q, v->d, "dot4e", K_EMUL, 0, lanes, steps, reps, out, host, vendor, v->dname);
|
|
run_variant(ctx, q, v->d, "dot4_khr", K_KHR, 0, lanes, steps, reps, out, host, vendor, v->dname);
|
|
if (!strcmp(vendor, "amd")) run_variant(ctx, q, v->d, "dot4_amd", K_AMD, 0, lanes, steps, reps, out, host, vendor, v->dname);
|
|
if (!strcmp(vendor, "nvidia")) run_variant(ctx, q, v->d, "dot4_nv", K_NV, 0, lanes, steps, reps, out, host, vendor, v->dname);
|
|
clReleaseMemObject(out); clReleaseCommandQueue(q); clReleaseContext(ctx); free(host);
|
|
printf("dot4-probe: done\n");
|
|
return 0;
|
|
}
|