5 files changed, 1076 insertions, 1 deletions

diff --git a/tools/cuda-ab/run/transformer_async b/tools/cuda-ab/run/transformer_async
new file mode 100755
index 0000000..2911b5c
--- /dev/null
+++ b/tools/cuda-ab/run/transformer_async
diff --git a/tools/cuda-ab/run/transformer_stages b/tools/cuda-ab/run/transformer_stages
new file mode 100755
index 0000000..1461a3e
--- /dev/null
+++ b/tools/cuda-ab/run/transformer_stages
diff --git a/tools/cuda-ab/soak.sh b/tools/cuda-ab/soak.sh
index c150692..4f1430f 100755
--- a/tools/cuda-ab/soak.sh
+++ b/tools/cuda-ab/soak.sh
@@ -192,7 +192,7 @@ log "SOAK START mode=$MODE seconds=$SECS guest_dir=$GUEST_DIR timeout=${TMO}s lo
 # cannot silently shrink the battery.
 found=""
 for i in 1 2 3; do
-	found="$(gssh "for b in $CLIENTS; do [ -x $GUEST_DIR/\$b ] && echo \$b; done")" || found=""
+	found="$(gssh "for b in $CLIENTS; do [ -x $GUEST_DIR/\$b ] && echo \$b; done; true")" || found=""
 	[ -n "$found" ] && break
 	sleep 3
 done
diff --git a/tools/cuda-ab/transformer_async.cu b/tools/cuda-ab/transformer_async.cu
new file mode 100644
index 0000000..213363a
--- /dev/null
+++ b/tools/cuda-ab/transformer_async.cu
@@ -0,0 +1,522 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * transformer.cu - single-head attention forward + 4-layer FFN mini-transformer.
+ *
+ * Two independent checks in one binary, both verified against an internal
+ * double-precision CPU reference (tolerance 1e-3 relative, 1e-6 absolute floor):
+ *
+ *  1. Attention: X[64x256], Wq/Wk/Wv/Wo[256x256] (head dim 256)
+ *       Q = X*Wq, K = X*Wk, V = X*Wv
+ *       S = Q*K^T / 16            (scale = 1/sqrt(head_dim))
+ *       A = softmax(S)            (row-wise)
+ *       Y = (A*V)*Wo
+ *  2. Mini-transformer FFN head: h <- ReLU(h*Wf[l]) for 4 layers, then
+ *       logits = h*Wh, P = softmax(logits)  (row-wise over 256)
+ *
+ * Inputs and weights are multiples of 0.25/0.125, so every GEMM is exact in
+ * fp32; only the softmax exponentials differ from the CPU reference, hence
+ * the relative tolerance.
+ *
+ * Output lines are stable and greppable:
+ *
+ *   DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes>
+ *   TRANSFORMER dh=256 seq=64 bad=<mismatches> ysum=<sum>
+ *   FFN4 layers=4 bad=<mismatches> psum=<sum>
+ *   STATUS OK|FAIL
+ *
+ * Build with a static CUDA runtime (--cudart=static) so the guest only needs
+ * the driver's libcuda.so, not a CUDA toolkit installation:
+ *
+ *   nvcc -O2 --cudart=static -arch=sm_75 -o transformer transformer.cu
+ *
+ * Then, from the guest:  /tmp/transformer
+ */
+
+#include <cstdio>
+#include <cstdlib>
+#include <cstdint>
+#include <cmath>
+#include <cuda_runtime.h>
+
+#define CHECK(x)                                                         \
+	do {                                                             \
+		cudaError_t e_ = (x);                                    \
+		if (e_ != cudaSuccess) {                                 \
+			printf("CUDA_ERROR %s at line %d\n",             \
+			       cudaGetErrorString(e_), __LINE__);        \
+			printf("STATUS FAIL\n");                         \
+			return 1;                                        \
+		}                                                        \
+	} while (0)
+
+#define SEQ 64
+#define DH 256
+#define TILE 16
+#define TOL 1e-3f
+
+__global__ void gemm_nn(const float *A, const float *B, float *C,
+			int M, int N, int K)
+{
+	__shared__ float As[TILE][TILE];
+	__shared__ float Bs[TILE][TILE];
+	int row = blockIdx.y * TILE + threadIdx.y;
+	int col = blockIdx.x * TILE + threadIdx.x;
+	float acc = 0.0f;
+
+	for (int t = 0; t < K / TILE; t++) {
+		As[threadIdx.y][threadIdx.x] = A[row * K + t * TILE + threadIdx.x];
+		Bs[threadIdx.y][threadIdx.x] =
+			B[(t * TILE + threadIdx.y) * N + col];
+		__syncthreads();
+		for (int k = 0; k < TILE; k++)
+			acc += As[threadIdx.y][k] * Bs[k][threadIdx.x];
+		__syncthreads();
+	}
+	C[row * N + col] = acc;
+}
+
+__global__ void gemm_nt(const float *A, const float *B, float *C,
+			int M, int N, int K)
+{
+	/* C[i][j] = sum_k A[i][k] * B[j][k]  (B is [N][K], row-major). */
+	__shared__ float As[TILE][TILE];
+	__shared__ float Bs[TILE][TILE];
+	int row = blockIdx.y * TILE + threadIdx.y;
+	int col = blockIdx.x * TILE + threadIdx.x;
+	float acc = 0.0f;
+
+	for (int t = 0; t < K / TILE; t++) {
+		As[threadIdx.y][threadIdx.x] = A[row * K + t * TILE + threadIdx.x];
+		Bs[threadIdx.y][threadIdx.x] =
+			B[(blockIdx.x * TILE + threadIdx.y) * K +
+			  t * TILE + threadIdx.x];
+		__syncthreads();
+		for (int k = 0; k < TILE; k++)
+			acc += As[threadIdx.y][k] * Bs[threadIdx.x][k];
+		__syncthreads();
+	}
+	C[row * N + col] = acc;
+}
+
+__global__ void softmax_rows(const float *S, float *P, int rows, int cols,
+			     float scale)
+{
+	extern __shared__ float sm[];
+	int row = blockIdx.x;
+	int t = threadIdx.x;
+	float v = S[row * cols + t] * scale;
+	int stride;
+
+	sm[t] = v;
+	__syncthreads();
+	for (stride = blockDim.x >> 1; stride > 0; stride >>= 1) {
+		if (t < stride)
+			sm[t] = fmaxf(sm[t], sm[t + stride]);
+		__syncthreads();
+	}
+	v = expf(v - sm[0]);
+	sm[t] = v;
+	__syncthreads();
+	for (stride = blockDim.x >> 1; stride > 0; stride >>= 1) {
+		if (t < stride)
+			sm[t] += sm[t + stride];
+		__syncthreads();
+	}
+	P[row * cols + t] = v / sm[0];
+}
+
+__global__ void relu_inplace(float *h, int n)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < n && h[i] < 0.0f)
+		h[i] = 0.0f;
+}
+
+/* ------------------------------------------------------------------ host */
+
+static void launch_gemm_nn(const float *A, const float *B, float *C,
+			   int M, int N, int K)
+{
+	gemm_nn<<<dim3(N / TILE, M / TILE), dim3(TILE, TILE)>>>(A, B, C, M, N, K);
+}
+
+static void launch_gemm_nt(const float *A, const float *B, float *C,
+			   int M, int N, int K)
+{
+	gemm_nt<<<dim3(N / TILE, M / TILE), dim3(TILE, TILE)>>>(A, B, C, M, N, K);
+}
+
+static void launch_softmax(const float *S, float *P, int rows, int cols,
+			   float scale)
+{
+	softmax_rows<<<rows, cols, (size_t)cols * sizeof(float)>>>(S, P, rows,
+								   cols,
+								   scale);
+}
+
+/* Deterministic inputs: multiples of 0.25 (X) and 0.125 (weights). */
+static float xval(int i)
+{
+	return (float)(((i % 7) - 3)) * 0.25f;
+}
+
+static float wval(int i, int seed)
+{
+	return (float)(((i * 13 + seed * 101 + 7) % 5) - 2) * 0.125f;
+}
+
+static void ref_gemm(const double *A, const double *B, double *C,
+		     int M, int N, int K)
+{
+	for (int i = 0; i < M; i++)
+		for (int j = 0; j < N; j++) {
+			double acc = 0.0;
+
+			for (int k = 0; k < K; k++)
+				acc += A[i * K + k] * B[k * N + j];
+			C[i * N + j] = acc;
+		}
+}
+
+static void ref_gemm_nt(const double *A, const double *B, double *C,
+			int M, int N, int K)
+{
+	for (int i = 0; i < M; i++)
+		for (int j = 0; j < N; j++) {
+			double acc = 0.0;
+
+			for (int k = 0; k < K; k++)
+				acc += A[i * K + k] * B[j * K + k];
+			C[i * N + j] = acc;
+		}
+}
+
+static void ref_softmax(const double *S, double *P, int rows, int cols,
+			double scale)
+{
+	for (int i = 0; i < rows; i++) {
+		const double *s = S + i * cols;
+		double *p = P + i * cols;
+		double m = s[0] * scale, sum = 0.0;
+
+		for (int j = 1; j < cols; j++)
+			if (s[j] * scale > m)
+				m = s[j] * scale;
+		for (int j = 0; j < cols; j++) {
+			p[j] = exp(s[j] * scale - m);
+			sum += p[j];
+		}
+		for (int j = 0; j < cols; j++)
+			p[j] /= sum;
+	}
+}
+
+static void ref_relu(double *h, int n)
+{
+	for (int i = 0; i < n; i++)
+		if (h[i] < 0.0)
+			h[i] = 0.0;
+}
+
+static int count_bad(const float *gpu, const double *ref, int n)
+{
+	int bad = 0;
+
+	for (int i = 0; i < n; i++) {
+		double g = (double)gpu[i];
+		double r = ref[i];
+
+		if (fabs(g - r) > (double)TOL * (fabs(r) + 1e-6))
+			bad++;
+	}
+	return bad;
+}
+
+static double dsum(const float *h, int n)
+{
+	double s = 0.0;
+
+	for (int i = 0; i < n; i++)
+		s += (double)h[i];
+	return s;
+}
+
+#define NF 4 /* FFN layers */
+
+int main(void)
+{
+	cudaDeviceProp prop;
+	const int seqdh = SEQ * DH, dh2 = DH * DH, seq2 = SEQ * SEQ;
+	const size_t b_seqdh = (size_t)seqdh * sizeof(float);
+	const size_t b_dh2 = (size_t)dh2 * sizeof(float);
+	const size_t b_seq2 = (size_t)seq2 * sizeof(float);
+	float *h_x, *h_y, *h_p;
+	double *r_x, *r_q, *r_k, *r_v, *r_s, *r_a, *r_av, *r_y;
+	double *r_f0, *r_f1, *r_logits, *r_p;
+	float *d_x, *d_q, *d_k, *d_v, *d_s, *d_a, *d_av, *d_y;
+	float *d_wq, *d_wk, *d_wv, *d_wo;
+	float *d_wf[NF], *d_wh, *d_f0, *d_f1, *d_logits, *d_p;
+	float *d_ping, *d_pong, *tmpf;
+	double *r_ping, *r_pong, *tmpd;
+	int bad_y, bad_p, i, l;
+
+	setvbuf(stdout, NULL, _IONBF, 0);
+
+	CHECK(cudaGetDeviceProperties(&prop, 0));
+	printf("DEVICE %s cc=%d.%d sm=%d mem=%llu\n", prop.name,
+	       prop.major, prop.minor, prop.multiProcessorCount,
+	       (unsigned long long)prop.totalGlobalMem);
+
+	/* host inputs + CPU reference workspace */
+	h_x = (float *)malloc(b_seqdh);
+	h_y = (float *)malloc(b_seqdh);
+	h_p = (float *)malloc(b_seqdh);
+	r_x = (double *)malloc((size_t)seqdh * sizeof(double));
+	r_q = (double *)malloc((size_t)seqdh * sizeof(double));
+	r_k = (double *)malloc((size_t)seqdh * sizeof(double));
+	r_v = (double *)malloc((size_t)seqdh * sizeof(double));
+	r_s = (double *)malloc((size_t)seq2 * sizeof(double));
+	r_a = (double *)malloc((size_t)seq2 * sizeof(double));
+	r_av = (double *)malloc((size_t)seqdh * sizeof(double));
+	r_y = (double *)malloc((size_t)seqdh * sizeof(double));
+	r_f0 = (double *)malloc((size_t)seqdh * sizeof(double));
+	r_f1 = (double *)malloc((size_t)seqdh * sizeof(double));
+	r_logits = (double *)malloc((size_t)seqdh * sizeof(double));
+	r_p = (double *)malloc((size_t)seqdh * sizeof(double));
+	if (!h_x || !h_y || !h_p || !r_x || !r_q || !r_k || !r_v || !r_s ||
+	    !r_a || !r_av || !r_y || !r_f0 || !r_f1 || !r_logits || !r_p)
+		return 1;
+
+	for (i = 0; i < seqdh; i++) {
+		h_x[i] = xval(i);
+		r_x[i] = (double)xval(i);
+	}
+
+	/* device buffers */
+	CHECK(cudaMalloc(&d_x, b_seqdh));
+	CHECK(cudaMalloc(&d_q, b_seqdh));
+	CHECK(cudaMalloc(&d_k, b_seqdh));
+	CHECK(cudaMalloc(&d_v, b_seqdh));
+	CHECK(cudaMalloc(&d_s, b_seq2));
+	CHECK(cudaMalloc(&d_a, b_seq2));
+	CHECK(cudaMalloc(&d_av, b_seqdh));
+	CHECK(cudaMalloc(&d_y, b_seqdh));
+	CHECK(cudaMalloc(&d_wq, b_dh2));
+	CHECK(cudaMalloc(&d_wk, b_dh2));
+	CHECK(cudaMalloc(&d_wv, b_dh2));
+	CHECK(cudaMalloc(&d_wo, b_dh2));
+	CHECK(cudaMalloc(&d_wh, b_dh2));
+	CHECK(cudaMalloc(&d_f0, b_seqdh));
+	CHECK(cudaMalloc(&d_f1, b_seqdh));
+	CHECK(cudaMalloc(&d_logits, b_seqdh));
+	CHECK(cudaMalloc(&d_p, b_seqdh));
+	for (l = 0; l < NF; l++)
+		CHECK(cudaMalloc(&d_wf[l], b_dh2));
+
+	/* weights: deterministic, uploaded once */
+	{
+		float *w = (float *)malloc(b_dh2);
+
+		if (!w)
+			return 1;
+		for (i = 0; i < dh2; i++)
+			w[i] = wval(i, 1);
+		CHECK(cudaMemcpy(d_wq, w, b_dh2, cudaMemcpyHostToDevice));
+		for (i = 0; i < dh2; i++)
+			w[i] = wval(i, 2);
+		CHECK(cudaMemcpy(d_wk, w, b_dh2, cudaMemcpyHostToDevice));
+		for (i = 0; i < dh2; i++)
+			w[i] = wval(i, 3);
+		CHECK(cudaMemcpy(d_wv, w, b_dh2, cudaMemcpyHostToDevice));
+		for (i = 0; i < dh2; i++)
+			w[i] = wval(i, 4);
+		CHECK(cudaMemcpy(d_wo, w, b_dh2, cudaMemcpyHostToDevice));
+		for (i = 0; i < dh2; i++)
+			w[i] = wval(i, 20);
+		CHECK(cudaMemcpy(d_wh, w, b_dh2, cudaMemcpyHostToDevice));
+		for (l = 0; l < NF; l++) {
+			for (i = 0; i < dh2; i++)
+				w[i] = wval(i, 10 + l);
+			CHECK(cudaMemcpy(d_wf[l], w, b_dh2,
+					 cudaMemcpyHostToDevice));
+		}
+		free(w);
+	}
+	CHECK(cudaMemcpy(d_x, h_x, b_seqdh, cudaMemcpyHostToDevice));
+
+	/* ---- check 1: attention forward (async chain, stage readbacks) ---- */
+	launch_gemm_nn(d_x, d_wq, d_q, SEQ, DH, DH);
+	launch_gemm_nn(d_x, d_wk, d_k, SEQ, DH, DH);
+	launch_gemm_nn(d_x, d_wv, d_v, SEQ, DH, DH);
+	launch_gemm_nt(d_q, d_k, d_s, SEQ, SEQ, DH);
+	launch_softmax(d_s, d_a, SEQ, SEQ, 1.0f / 16.0f);
+	launch_gemm_nn(d_a, d_v, d_av, SEQ, DH, SEQ);
+	launch_gemm_nn(d_av, d_wo, d_y, SEQ, DH, DH);
+	CHECK(cudaGetLastError());
+	CHECK(cudaDeviceSynchronize());
+	CHECK(cudaMemcpy(h_y, d_y, b_seqdh, cudaMemcpyDeviceToHost));
+	{
+		float *hs = (float *)malloc(b_seqdh);
+		float *hs2 = (float *)malloc(b_seq2);
+		double *rw = (double *)malloc((size_t)dh2 * sizeof(double));
+		double *rq = (double *)malloc((size_t)seqdh * sizeof(double));
+		double *rk = (double *)malloc((size_t)seqdh * sizeof(double));
+		double *rv = (double *)malloc((size_t)seqdh * sizeof(double));
+		double *rs = (double *)malloc((size_t)seq2 * sizeof(double));
+		double *ra = (double *)malloc((size_t)seq2 * sizeof(double));
+		double *rav = (double *)malloc((size_t)seqdh * sizeof(double));
+		double *ry = (double *)malloc((size_t)seqdh * sizeof(double));
+		int st;
+
+		for (i = 0; i < dh2; i++)
+			rw[i] = (double)wval(i, 1);
+		ref_gemm(r_x, rw, rq, SEQ, DH, DH);
+		for (i = 0; i < dh2; i++)
+			rw[i] = (double)wval(i, 2);
+		ref_gemm(r_x, rw, rk, SEQ, DH, DH);
+		for (i = 0; i < dh2; i++)
+			rw[i] = (double)wval(i, 3);
+		ref_gemm(r_x, rw, rv, SEQ, DH, DH);
+		ref_gemm_nt(rq, rk, rs, SEQ, SEQ, DH);
+		ref_softmax(rs, ra, SEQ, SEQ, 1.0 / 16.0);
+		ref_gemm(ra, rv, rav, SEQ, DH, SEQ);
+		for (i = 0; i < dh2; i++)
+			rw[i] = (double)wval(i, 4);
+		ref_gemm(rav, rw, ry, SEQ, DH, DH);
+
+		CHECK(cudaMemcpy(hs, d_q, b_seqdh, cudaMemcpyDeviceToHost));
+		st = count_bad(hs, rq, seqdh);
+		printf("ASYNC Q bad=%d\n", st);
+		CHECK(cudaMemcpy(hs, d_k, b_seqdh, cudaMemcpyDeviceToHost));
+		st = count_bad(hs, rk, seqdh);
+		printf("ASYNC K bad=%d\n", st);
+		CHECK(cudaMemcpy(hs, d_v, b_seqdh, cudaMemcpyDeviceToHost));
+		st = count_bad(hs, rv, seqdh);
+		printf("ASYNC V bad=%d\n", st);
+		CHECK(cudaMemcpy(hs2, d_s, b_seq2, cudaMemcpyDeviceToHost));
+		st = count_bad(hs2, rs, seq2);
+		printf("ASYNC S bad=%d\n", st);
+		CHECK(cudaMemcpy(hs2, d_a, b_seq2, cudaMemcpyDeviceToHost));
+		st = count_bad(hs2, ra, seq2);
+		printf("ASYNC A bad=%d\n", st);
+		CHECK(cudaMemcpy(hs, d_av, b_seqdh, cudaMemcpyDeviceToHost));
+		st = count_bad(hs, rav, seqdh);
+		printf("ASYNC AV bad=%d\n", st);
+		st = count_bad(h_y, ry, seqdh);
+		printf("ASYNC Y bad=%d\n", st);
+	}
+
+	{
+		double *r_w = (double *)malloc((size_t)dh2 * sizeof(double));
+
+		if (!r_w)
+			return 1;
+		for (i = 0; i < dh2; i++)
+			r_w[i] = (double)wval(i, 1);
+		ref_gemm(r_x, r_w, r_q, SEQ, DH, DH);
+		for (i = 0; i < dh2; i++)
+			r_w[i] = (double)wval(i, 2);
+		ref_gemm(r_x, r_w, r_k, SEQ, DH, DH);
+		for (i = 0; i < dh2; i++)
+			r_w[i] = (double)wval(i, 3);
+		ref_gemm(r_x, r_w, r_v, SEQ, DH, DH);
+		ref_gemm_nt(r_q, r_k, r_s, SEQ, SEQ, DH);
+		ref_softmax(r_s, r_a, SEQ, SEQ, 1.0 / 16.0);
+		ref_gemm(r_a, r_v, r_av, SEQ, DH, SEQ);
+		for (i = 0; i < dh2; i++)
+			r_w[i] = (double)wval(i, 4);
+		ref_gemm(r_av, r_w, r_y, SEQ, DH, DH);
+		free(r_w);
+	}
+
+	bad_y = count_bad(h_y, r_y, seqdh);
+	printf("TRANSFORMER dh=256 seq=64 bad=%d ysum=%.6f\n", bad_y,
+	       dsum(h_y, seqdh));
+
+	/* ---- check 2: 4-layer FFN stack + softmax head ---- */
+	CHECK(cudaMemcpy(d_f0, h_x, b_seqdh, cudaMemcpyHostToDevice));
+	d_ping = d_f0;
+	d_pong = d_f1;
+	for (l = 0; l < NF; l++) {
+		launch_gemm_nn(d_ping, d_wf[l], d_pong, SEQ, DH, DH);
+		relu_inplace<<<(seqdh + 255) / 256, 256>>>(d_pong, seqdh);
+		tmpf = d_ping;
+		d_ping = d_pong;
+		d_pong = tmpf;
+	}
+	launch_gemm_nn(d_ping, d_wh, d_logits, SEQ, DH, DH);
+	launch_softmax(d_logits, d_p, SEQ, DH, 1.0f);
+	CHECK(cudaGetLastError());
+	CHECK(cudaDeviceSynchronize());
+	CHECK(cudaMemcpy(h_p, d_p, b_seqdh, cudaMemcpyDeviceToHost));
+
+	{
+		double *r_w = (double *)malloc((size_t)dh2 * sizeof(double));
+
+		if (!r_w)
+			return 1;
+		memcpy(r_f0, r_x, (size_t)seqdh * sizeof(double));
+		r_ping = r_f0;
+		r_pong = r_f1;
+		for (l = 0; l < NF; l++) {
+			for (i = 0; i < dh2; i++)
+				r_w[i] = (double)wval(i, 10 + l);
+			ref_gemm(r_ping, r_w, r_pong, SEQ, DH, DH);
+			ref_relu(r_pong, seqdh);
+			tmpd = r_ping;
+			r_ping = r_pong;
+			r_pong = tmpd;
+		}
+		for (i = 0; i < dh2; i++)
+			r_w[i] = (double)wval(i, 20);
+		ref_gemm(r_ping, r_w, r_logits, SEQ, DH, DH);
+		ref_softmax(r_logits, r_p, SEQ, DH, 1.0);
+		free(r_w);
+	}
+
+	bad_p = count_bad(h_p, r_p, seqdh);
+	printf("FFN4 layers=4 bad=%d psum=%.6f\n", bad_p, dsum(h_p, seqdh));
+
+	if (bad_y != 0 || bad_p != 0) {
+		printf("STATUS FAIL\n");
+		return 1;
+	}
+
+	printf("STATUS OK\n");
+	CHECK(cudaFree(d_x));
+	CHECK(cudaFree(d_q));
+	CHECK(cudaFree(d_k));
+	CHECK(cudaFree(d_v));
+	CHECK(cudaFree(d_s));
+	CHECK(cudaFree(d_a));
+	CHECK(cudaFree(d_av));
+	CHECK(cudaFree(d_y));
+	CHECK(cudaFree(d_wq));
+	CHECK(cudaFree(d_wk));
+	CHECK(cudaFree(d_wv));
+	CHECK(cudaFree(d_wo));
+	CHECK(cudaFree(d_wh));
+	for (l = 0; l < NF; l++)
+		CHECK(cudaFree(d_wf[l]));
+	CHECK(cudaFree(d_f0));
+	CHECK(cudaFree(d_f1));
+	CHECK(cudaFree(d_logits));
+	CHECK(cudaFree(d_p));
+	free(h_x);
+	free(h_y);
+	free(h_p);
+	free(r_x);
+	free(r_q);
+	free(r_k);
+	free(r_v);
+	free(r_s);
+	free(r_a);
+	free(r_av);
+	free(r_y);
+	free(r_f0);
+	free(r_f1);
+	free(r_logits);
+	free(r_p);
+	return 0;
+}
diff --git a/tools/cuda-ab/transformer_stages.cu b/tools/cuda-ab/transformer_stages.cu
new file mode 100644
index 0000000..d7f67c0
--- /dev/null
+++ b/tools/cuda-ab/transformer_stages.cu
@@ -0,0 +1,553 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * transformer.cu - single-head attention forward + 4-layer FFN mini-transformer.
+ *
+ * Two independent checks in one binary, both verified against an internal
+ * double-precision CPU reference (tolerance 1e-3 relative, 1e-6 absolute floor):
+ *
+ *  1. Attention: X[64x256], Wq/Wk/Wv/Wo[256x256] (head dim 256)
+ *       Q = X*Wq, K = X*Wk, V = X*Wv
+ *       S = Q*K^T / 16            (scale = 1/sqrt(head_dim))
+ *       A = softmax(S)            (row-wise)
+ *       Y = (A*V)*Wo
+ *  2. Mini-transformer FFN head: h <- ReLU(h*Wf[l]) for 4 layers, then
+ *       logits = h*Wh, P = softmax(logits)  (row-wise over 256)
+ *
+ * Inputs and weights are multiples of 0.25/0.125, so every GEMM is exact in
+ * fp32; only the softmax exponentials differ from the CPU reference, hence
+ * the relative tolerance.
+ *
+ * Output lines are stable and greppable:
+ *
+ *   DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes>
+ *   TRANSFORMER dh=256 seq=64 bad=<mismatches> ysum=<sum>
+ *   FFN4 layers=4 bad=<mismatches> psum=<sum>
+ *   STATUS OK|FAIL
+ *
+ * Build with a static CUDA runtime (--cudart=static) so the guest only needs
+ * the driver's libcuda.so, not a CUDA toolkit installation:
+ *
+ *   nvcc -O2 --cudart=static -arch=sm_75 -o transformer transformer.cu
+ *
+ * Then, from the guest:  /tmp/transformer
+ */
+
+#include <cstdio>
+#include <cstdlib>
+#include <cstdint>
+#include <cmath>
+#include <cuda_runtime.h>
+
+#define CHECK(x)                                                         \
+	do {                                                             \
+		cudaError_t e_ = (x);                                    \
+		if (e_ != cudaSuccess) {                                 \
+			printf("CUDA_ERROR %s at line %d\n",             \
+			       cudaGetErrorString(e_), __LINE__);        \
+			printf("STATUS FAIL\n");                         \
+			return 1;                                        \
+		}                                                        \
+	} while (0)
+
+#define SEQ 64
+#define DH 256
+#define TILE 16
+#define TOL 1e-3f
+
+__global__ void gemm_nn(const float *A, const float *B, float *C,
+			int M, int N, int K)
+{
+	__shared__ float As[TILE][TILE];
+	__shared__ float Bs[TILE][TILE];
+	int row = blockIdx.y * TILE + threadIdx.y;
+	int col = blockIdx.x * TILE + threadIdx.x;
+	float acc = 0.0f;
+
+	for (int t = 0; t < K / TILE; t++) {
+		As[threadIdx.y][threadIdx.x] = A[row * K + t * TILE + threadIdx.x];
+		Bs[threadIdx.y][threadIdx.x] =
+			B[(t * TILE + threadIdx.y) * N + col];
+		__syncthreads();
+		for (int k = 0; k < TILE; k++)
+			acc += As[threadIdx.y][k] * Bs[k][threadIdx.x];
+		__syncthreads();
+	}
+	C[row * N + col] = acc;
+}
+
+__global__ void gemm_nt(const float *A, const float *B, float *C,
+			int M, int N, int K)
+{
+	/* C[i][j] = sum_k A[i][k] * B[j][k]  (B is [N][K], row-major). */
+	__shared__ float As[TILE][TILE];
+	__shared__ float Bs[TILE][TILE];
+	int row = blockIdx.y * TILE + threadIdx.y;
+	int col = blockIdx.x * TILE + threadIdx.x;
+	float acc = 0.0f;
+
+	for (int t = 0; t < K / TILE; t++) {
+		As[threadIdx.y][threadIdx.x] = A[row * K + t * TILE + threadIdx.x];
+		Bs[threadIdx.y][threadIdx.x] =
+			B[(blockIdx.x * TILE + threadIdx.y) * K +
+			  t * TILE + threadIdx.x];
+		__syncthreads();
+		for (int k = 0; k < TILE; k++)
+			acc += As[threadIdx.y][k] * Bs[threadIdx.x][k];
+		__syncthreads();
+	}
+	C[row * N + col] = acc;
+}
+
+__global__ void softmax_rows(const float *S, float *P, int rows, int cols,
+			     float scale)
+{
+	extern __shared__ float sm[];
+	int row = blockIdx.x;
+	int t = threadIdx.x;
+	float v = S[row * cols + t] * scale;
+	int stride;
+
+	sm[t] = v;
+	__syncthreads();
+	for (stride = blockDim.x >> 1; stride > 0; stride >>= 1) {
+		if (t < stride)
+			sm[t] = fmaxf(sm[t], sm[t + stride]);
+		__syncthreads();
+	}
+	v = expf(v - sm[0]);
+	sm[t] = v;
+	__syncthreads();
+	for (stride = blockDim.x >> 1; stride > 0; stride >>= 1) {
+		if (t < stride)
+			sm[t] += sm[t + stride];
+		__syncthreads();
+	}
+	P[row * cols + t] = v / sm[0];
+}
+
+__global__ void relu_inplace(float *h, int n)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < n && h[i] < 0.0f)
+		h[i] = 0.0f;
+}
+
+/* ------------------------------------------------------------------ host */
+
+static void launch_gemm_nn(const float *A, const float *B, float *C,
+			   int M, int N, int K)
+{
+	gemm_nn<<<dim3(N / TILE, M / TILE), dim3(TILE, TILE)>>>(A, B, C, M, N, K);
+}
+
+static void launch_gemm_nt(const float *A, const float *B, float *C,
+			   int M, int N, int K)
+{
+	gemm_nt<<<dim3(N / TILE, M / TILE), dim3(TILE, TILE)>>>(A, B, C, M, N, K);
+}
+
+static void launch_softmax(const float *S, float *P, int rows, int cols,
+			   float scale)
+{
+	softmax_rows<<<rows, cols, (size_t)cols * sizeof(float)>>>(S, P, rows,
+								   cols,
+								   scale);
+}
+
+/* Deterministic inputs: multiples of 0.25 (X) and 0.125 (weights). */
+static float xval(int i)
+{
+	return (float)(((i % 7) - 3)) * 0.25f;
+}
+
+static float wval(int i, int seed)
+{
+	return (float)(((i * 13 + seed * 101 + 7) % 5) - 2) * 0.125f;
+}
+
+static void ref_gemm(const double *A, const double *B, double *C,
+		     int M, int N, int K)
+{
+	for (int i = 0; i < M; i++)
+		for (int j = 0; j < N; j++) {
+			double acc = 0.0;
+
+			for (int k = 0; k < K; k++)
+				acc += A[i * K + k] * B[k * N + j];
+			C[i * N + j] = acc;
+		}
+}
+
+static void ref_gemm_nt(const double *A, const double *B, double *C,
+			int M, int N, int K)
+{
+	for (int i = 0; i < M; i++)
+		for (int j = 0; j < N; j++) {
+			double acc = 0.0;
+
+			for (int k = 0; k < K; k++)
+				acc += A[i * K + k] * B[j * K + k];
+			C[i * N + j] = acc;
+		}
+}
+
+static void ref_softmax(const double *S, double *P, int rows, int cols,
+			double scale)
+{
+	for (int i = 0; i < rows; i++) {
+		const double *s = S + i * cols;
+		double *p = P + i * cols;
+		double m = s[0] * scale, sum = 0.0;
+
+		for (int j = 1; j < cols; j++)
+			if (s[j] * scale > m)
+				m = s[j] * scale;
+		for (int j = 0; j < cols; j++) {
+			p[j] = exp(s[j] * scale - m);
+			sum += p[j];
+		}
+		for (int j = 0; j < cols; j++)
+			p[j] /= sum;
+	}
+}
+
+static void ref_relu(double *h, int n)
+{
+	for (int i = 0; i < n; i++)
+		if (h[i] < 0.0)
+			h[i] = 0.0;
+}
+
+static int count_bad(const float *gpu, const double *ref, int n)
+{
+	int bad = 0;
+
+	for (int i = 0; i < n; i++) {
+		double g = (double)gpu[i];
+		double r = ref[i];
+
+		if (fabs(g - r) > (double)TOL * (fabs(r) + 1e-6))
+			bad++;
+	}
+	return bad;
+}
+
+static double dsum(const float *h, int n)
+{
+	double s = 0.0;
+
+	for (int i = 0; i < n; i++)
+		s += (double)h[i];
+	return s;
+}
+
+static void dump_bad(const char *tag, const float *gpu, const double *ref,
+		     int n)
+{
+	int i, shown = 0;
+
+	for (i = 0; i < n && shown < 8; i++) {
+		double g = (double)gpu[i];
+		double r = ref[i];
+
+		if (fabs(g - r) > (double)TOL * (fabs(r) + 1e-6)) {
+			printf("BAD %s i=%d gpu=%.6f ref=%.6f (tile %d,%d row %d col %d)\n",
+			       tag, i, g, r, i / 16, (i % 256) / 16,
+			       i / 256, i % 256);
+			shown++;
+		}
+	}
+}
+
+#define NF 4 /* FFN layers */
+
+int main(void)
+{
+	cudaDeviceProp prop;
+	const int seqdh = SEQ * DH, dh2 = DH * DH, seq2 = SEQ * SEQ;
+	const size_t b_seqdh = (size_t)seqdh * sizeof(float);
+	const size_t b_dh2 = (size_t)dh2 * sizeof(float);
+	const size_t b_seq2 = (size_t)seq2 * sizeof(float);
+	float *h_x, *h_y, *h_p;
+	double *r_x, *r_q, *r_k, *r_v, *r_s, *r_a, *r_av, *r_y;
+	double *r_f0, *r_f1, *r_logits, *r_p;
+	float *d_x, *d_q, *d_k, *d_v, *d_s, *d_a, *d_av, *d_y;
+	float *d_wq, *d_wk, *d_wv, *d_wo;
+	float *d_wf[NF], *d_wh, *d_f0, *d_f1, *d_logits, *d_p;
+	float *d_ping, *d_pong, *tmpf;
+	double *r_ping, *r_pong, *tmpd;
+	int bad_y, bad_p, i, l;
+
+	setvbuf(stdout, NULL, _IONBF, 0);
+
+	CHECK(cudaGetDeviceProperties(&prop, 0));
+	printf("DEVICE %s cc=%d.%d sm=%d mem=%llu\n", prop.name,
+	       prop.major, prop.minor, prop.multiProcessorCount,
+	       (unsigned long long)prop.totalGlobalMem);
+
+	/* host inputs + CPU reference workspace */
+	h_x = (float *)malloc(b_seqdh);
+	h_y = (float *)malloc(b_seqdh);
+	h_p = (float *)malloc(b_seqdh);
+	r_x = (double *)malloc((size_t)seqdh * sizeof(double));
+	r_q = (double *)malloc((size_t)seqdh * sizeof(double));
+	r_k = (double *)malloc((size_t)seqdh * sizeof(double));
+	r_v = (double *)malloc((size_t)seqdh * sizeof(double));
+	r_s = (double *)malloc((size_t)seq2 * sizeof(double));
+	r_a = (double *)malloc((size_t)seq2 * sizeof(double));
+	r_av = (double *)malloc((size_t)seqdh * sizeof(double));
+	r_y = (double *)malloc((size_t)seqdh * sizeof(double));
+	r_f0 = (double *)malloc((size_t)seqdh * sizeof(double));
+	r_f1 = (double *)malloc((size_t)seqdh * sizeof(double));
+	r_logits = (double *)malloc((size_t)seqdh * sizeof(double));
+	r_p = (double *)malloc((size_t)seqdh * sizeof(double));
+	if (!h_x || !h_y || !h_p || !r_x || !r_q || !r_k || !r_v || !r_s ||
+	    !r_a || !r_av || !r_y || !r_f0 || !r_f1 || !r_logits || !r_p)
+		return 1;
+
+	for (i = 0; i < seqdh; i++) {
+		h_x[i] = xval(i);
+		r_x[i] = (double)xval(i);
+	}
+
+	/* device buffers */
+	CHECK(cudaMalloc(&d_x, b_seqdh));
+	CHECK(cudaMalloc(&d_q, b_seqdh));
+	CHECK(cudaMalloc(&d_k, b_seqdh));
+	CHECK(cudaMalloc(&d_v, b_seqdh));
+	CHECK(cudaMalloc(&d_s, b_seq2));
+	CHECK(cudaMalloc(&d_a, b_seq2));
+	CHECK(cudaMalloc(&d_av, b_seqdh));
+	CHECK(cudaMalloc(&d_y, b_seqdh));
+	CHECK(cudaMalloc(&d_wq, b_dh2));
+	CHECK(cudaMalloc(&d_wk, b_dh2));
+	CHECK(cudaMalloc(&d_wv, b_dh2));
+	CHECK(cudaMalloc(&d_wo, b_dh2));
+	CHECK(cudaMalloc(&d_wh, b_dh2));
+	CHECK(cudaMalloc(&d_f0, b_seqdh));
+	CHECK(cudaMalloc(&d_f1, b_seqdh));
+	CHECK(cudaMalloc(&d_logits, b_seqdh));
+	CHECK(cudaMalloc(&d_p, b_seqdh));
+	for (l = 0; l < NF; l++)
+		CHECK(cudaMalloc(&d_wf[l], b_dh2));
+
+	/* weights: deterministic, uploaded once */
+	{
+		float *w = (float *)malloc(b_dh2);
+
+		if (!w)
+			return 1;
+		for (i = 0; i < dh2; i++)
+			w[i] = wval(i, 1);
+		CHECK(cudaMemcpy(d_wq, w, b_dh2, cudaMemcpyHostToDevice));
+		for (i = 0; i < dh2; i++)
+			w[i] = wval(i, 2);
+		CHECK(cudaMemcpy(d_wk, w, b_dh2, cudaMemcpyHostToDevice));
+		for (i = 0; i < dh2; i++)
+			w[i] = wval(i, 3);
+		CHECK(cudaMemcpy(d_wv, w, b_dh2, cudaMemcpyHostToDevice));
+		for (i = 0; i < dh2; i++)
+			w[i] = wval(i, 4);
+		CHECK(cudaMemcpy(d_wo, w, b_dh2, cudaMemcpyHostToDevice));
+		for (i = 0; i < dh2; i++)
+			w[i] = wval(i, 20);
+		CHECK(cudaMemcpy(d_wh, w, b_dh2, cudaMemcpyHostToDevice));
+		for (l = 0; l < NF; l++) {
+			for (i = 0; i < dh2; i++)
+				w[i] = wval(i, 10 + l);
+			CHECK(cudaMemcpy(d_wf[l], w, b_dh2,
+					 cudaMemcpyHostToDevice));
+		}
+		free(w);
+	}
+	CHECK(cudaMemcpy(d_x, h_x, b_seqdh, cudaMemcpyHostToDevice));
+
+	/* ---- check 1: attention forward, stage-by-stage verify ---- */
+	{
+		float *h_stage = (float *)malloc(b_seqdh);
+		float *h_s2 = (float *)malloc(b_seq2);
+		double *r_stage = (double *)malloc((size_t)seqdh * sizeof(double));
+		double *r_s2 = (double *)malloc((size_t)seq2 * sizeof(double));
+		double *r_w2 = (double *)malloc((size_t)dh2 * sizeof(double));
+		int st;
+
+		if (!h_stage || !h_s2 || !r_stage || !r_s2 || !r_w2)
+			return 1;
+
+		launch_gemm_nn(d_x, d_wq, d_q, SEQ, DH, DH);
+		CHECK(cudaDeviceSynchronize());
+		CHECK(cudaMemcpy(h_stage, d_q, b_seqdh, cudaMemcpyDeviceToHost));
+		for (i = 0; i < dh2; i++)
+			r_w2[i] = (double)wval(i, 1);
+		ref_gemm(r_x, r_w2, r_stage, SEQ, DH, DH);
+		st = count_bad(h_stage, r_stage, seqdh);
+		printf("STAGE Q bad=%d\n", st);
+		dump_bad("Q", h_stage, r_stage, seqdh);
+
+		launch_gemm_nn(d_x, d_wk, d_k, SEQ, DH, DH);
+		CHECK(cudaDeviceSynchronize());
+		CHECK(cudaMemcpy(h_stage, d_k, b_seqdh, cudaMemcpyDeviceToHost));
+		for (i = 0; i < dh2; i++)
+			r_w2[i] = (double)wval(i, 2);
+		ref_gemm(r_x, r_w2, r_stage, SEQ, DH, DH);
+		st = count_bad(h_stage, r_stage, seqdh);
+		printf("STAGE K bad=%d\n", st);
+		dump_bad("K", h_stage, r_stage, seqdh);
+
+		launch_gemm_nn(d_x, d_wv, d_v, SEQ, DH, DH);
+		CHECK(cudaDeviceSynchronize());
+		CHECK(cudaMemcpy(h_stage, d_v, b_seqdh, cudaMemcpyDeviceToHost));
+		for (i = 0; i < dh2; i++)
+			r_w2[i] = (double)wval(i, 3);
+		ref_gemm(r_x, r_w2, r_stage, SEQ, DH, DH);
+		st = count_bad(h_stage, r_stage, seqdh);
+		printf("STAGE V bad=%d\n", st);
+		dump_bad("V", h_stage, r_stage, seqdh);
+
+		launch_gemm_nt(d_q, d_k, d_s, SEQ, SEQ, DH);
+		CHECK(cudaDeviceSynchronize());
+		CHECK(cudaMemcpy(h_s2, d_s, b_seq2, cudaMemcpyDeviceToHost));
+		ref_gemm_nt(r_q, r_k, r_s, SEQ, SEQ, DH);
+		st = count_bad(h_s2, r_s, seq2);
+		printf("STAGE S bad=%d\n", st);
+		dump_bad("S", h_s2, r_s, seq2);
+
+		launch_softmax(d_s, d_a, SEQ, SEQ, 1.0f / 16.0f);
+		CHECK(cudaDeviceSynchronize());
+		CHECK(cudaMemcpy(h_s2, d_a, b_seq2, cudaMemcpyDeviceToHost));
+		ref_softmax(r_s, r_a, SEQ, SEQ, 1.0 / 16.0);
+		st = count_bad(h_s2, r_a, seq2);
+		printf("STAGE A bad=%d\n", st);
+		dump_bad("A", h_s2, r_a, seq2);
+
+		launch_gemm_nn(d_a, d_v, d_av, SEQ, DH, SEQ);
+		CHECK(cudaDeviceSynchronize());
+		CHECK(cudaMemcpy(h_stage, d_av, b_seqdh, cudaMemcpyDeviceToHost));
+		ref_gemm(r_a, r_v, r_av, SEQ, DH, SEQ);
+		st = count_bad(h_stage, r_av, seqdh);
+		printf("STAGE AV bad=%d\n", st);
+
+		launch_gemm_nn(d_av, d_wo, d_y, SEQ, DH, DH);
+		CHECK(cudaDeviceSynchronize());
+		CHECK(cudaMemcpy(h_y, d_y, b_seqdh, cudaMemcpyDeviceToHost));
+		for (i = 0; i < dh2; i++)
+			r_w2[i] = (double)wval(i, 4);
+		ref_gemm(r_av, r_w2, r_y, SEQ, DH, DH);
+		st = count_bad(h_y, r_y, seqdh);
+		printf("STAGE Y bad=%d\n", st);
+	}
+
+	{
+		double *r_w = (double *)malloc((size_t)dh2 * sizeof(double));
+
+		if (!r_w)
+			return 1;
+		for (i = 0; i < dh2; i++)
+			r_w[i] = (double)wval(i, 1);
+		ref_gemm(r_x, r_w, r_q, SEQ, DH, DH);
+		for (i = 0; i < dh2; i++)
+			r_w[i] = (double)wval(i, 2);
+		ref_gemm(r_x, r_w, r_k, SEQ, DH, DH);
+		for (i = 0; i < dh2; i++)
+			r_w[i] = (double)wval(i, 3);
+		ref_gemm(r_x, r_w, r_v, SEQ, DH, DH);
+		ref_gemm_nt(r_q, r_k, r_s, SEQ, SEQ, DH);
+		ref_softmax(r_s, r_a, SEQ, SEQ, 1.0 / 16.0);
+		ref_gemm(r_a, r_v, r_av, SEQ, DH, SEQ);
+		for (i = 0; i < dh2; i++)
+			r_w[i] = (double)wval(i, 4);
+		ref_gemm(r_av, r_w, r_y, SEQ, DH, DH);
+		free(r_w);
+	}
+
+	bad_y = count_bad(h_y, r_y, seqdh);
+	printf("TRANSFORMER dh=256 seq=64 bad=%d ysum=%.6f\n", bad_y,
+	       dsum(h_y, seqdh));
+
+	/* ---- check 2: 4-layer FFN stack + softmax head ---- */
+	CHECK(cudaMemcpy(d_f0, h_x, b_seqdh, cudaMemcpyHostToDevice));
+	d_ping = d_f0;
+	d_pong = d_f1;
+	for (l = 0; l < NF; l++) {
+		launch_gemm_nn(d_ping, d_wf[l], d_pong, SEQ, DH, DH);
+		relu_inplace<<<(seqdh + 255) / 256, 256>>>(d_pong, seqdh);
+		tmpf = d_ping;
+		d_ping = d_pong;
+		d_pong = tmpf;
+	}
+	launch_gemm_nn(d_ping, d_wh, d_logits, SEQ, DH, DH);
+	launch_softmax(d_logits, d_p, SEQ, DH, 1.0f);
+	CHECK(cudaGetLastError());
+	CHECK(cudaDeviceSynchronize());
+	CHECK(cudaMemcpy(h_p, d_p, b_seqdh, cudaMemcpyDeviceToHost));
+
+	{
+		double *r_w = (double *)malloc((size_t)dh2 * sizeof(double));
+
+		if (!r_w)
+			return 1;
+		memcpy(r_f0, r_x, (size_t)seqdh * sizeof(double));
+		r_ping = r_f0;
+		r_pong = r_f1;
+		for (l = 0; l < NF; l++) {
+			for (i = 0; i < dh2; i++)
+				r_w[i] = (double)wval(i, 10 + l);
+			ref_gemm(r_ping, r_w, r_pong, SEQ, DH, DH);
+			ref_relu(r_pong, seqdh);
+			tmpd = r_ping;
+			r_ping = r_pong;
+			r_pong = tmpd;
+		}
+		for (i = 0; i < dh2; i++)
+			r_w[i] = (double)wval(i, 20);
+		ref_gemm(r_ping, r_w, r_logits, SEQ, DH, DH);
+		ref_softmax(r_logits, r_p, SEQ, DH, 1.0);
+		free(r_w);
+	}
+
+	bad_p = count_bad(h_p, r_p, seqdh);
+	printf("FFN4 layers=4 bad=%d psum=%.6f\n", bad_p, dsum(h_p, seqdh));
+
+	if (bad_y != 0 || bad_p != 0) {
+		printf("STATUS FAIL\n");
+		return 1;
+	}
+
+	printf("STATUS OK\n");
+	CHECK(cudaFree(d_x));
+	CHECK(cudaFree(d_q));
+	CHECK(cudaFree(d_k));
+	CHECK(cudaFree(d_v));
+	CHECK(cudaFree(d_s));
+	CHECK(cudaFree(d_a));
+	CHECK(cudaFree(d_av));
+	CHECK(cudaFree(d_y));
+	CHECK(cudaFree(d_wq));
+	CHECK(cudaFree(d_wk));
+	CHECK(cudaFree(d_wv));
+	CHECK(cudaFree(d_wo));
+	CHECK(cudaFree(d_wh));
+	for (l = 0; l < NF; l++)
+		CHECK(cudaFree(d_wf[l]));
+	CHECK(cudaFree(d_f0));
+	CHECK(cudaFree(d_f1));
+	CHECK(cudaFree(d_logits));
+	CHECK(cudaFree(d_p));
+	free(h_x);
+	free(h_y);
+	free(h_p);
+	free(r_x);
+	free(r_q);
+	free(r_k);
+	free(r_v);
+	free(r_s);
+	free(r_a);
+	free(r_av);
+	free(r_y);
+	free(r_f0);
+	free(r_f1);
+	free(r_logits);
+	free(r_p);
+	return 0;
+}