1 files changed, 40 insertions, 0 deletions
diff --git a/tools/cuda-bridge/kernels.cu b/tools/cuda-bridge/kernels.cu
new file mode 100644
index 0000000..508c94c
--- /dev/null
+++ b/tools/cuda-bridge/kernels.cu
@@ -0,0 +1,40 @@
+// SPDX-License-Identifier: BSD-3-Clause
+//
+// kernels.cu - CUDA kernels for the host-side CUDA bridge.
+//
+// Exactly mirrors the victory_mail semantics: mail_kernel writes every
+// character of the message from device memory, vadd computes c = a + b.
+// vadd_mutated is the falsifiability variant (adds 1.0f); it exists only to
+// prove the bridge executes the real requested kernel, never a canned answer.
+//
+// Compiled to PTX at build time by nvcc (see Makefile); the PTX is embedded
+// into cuda_bridge and loaded with cuModuleLoadData.
+
+#define MAIL "the little gpu man says you've got mail"
+
+__device__ const char kMail[] = MAIL;
+
+extern "C" __global__ void mail_kernel(char *out, int len)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < len)
+		out[i] = kMail[i];
+}
+
+extern "C" __global__ void vadd(const float *a, const float *b, float *c, int n)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < n)
+		c[i] = a[i] + b[i];
+}
+
+extern "C" __global__ void vadd_mutated(const float *a, const float *b,
+					float *c, int n)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < n)
+		c[i] = a[i] + b[i] + 1.0f;
+}