// 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;
}