blob: 508c94c9efcbf19a9dcd1dbb39fa38e0e736f575 (
plain)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
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;
}
|
