diff options
25 files changed, 1924 insertions, 0 deletions
diff --git a/tools/cuda-ab/memcpy_roundtrip.cu b/tools/cuda-ab/memcpy_roundtrip.cu new file mode 100644 index 0000000..679181f --- /dev/null +++ b/tools/cuda-ab/memcpy_roundtrip.cu @@ -0,0 +1,123 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * memcpy_roundtrip.cu - 8 MB H2D + kernel + D2H roundtrip plus memset check. + * + * Two data paths are exercised over an 8 MB (2M float) device buffer: + * + * 1. host -> device copy, a kernel (out[i] = in[i] + 0.5), device -> host + * copy; the host verifies the full chain element by element and sums the + * result as a checksum. + * 2. cudaMemset of a second 8 MB buffer with a known byte pattern, copied + * back and verified byte-for-byte. + * + * Output lines are stable and greppable: + * + * DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes> + * ROUNDTRIP bytes=<n> bad=<mismatches> csum=<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 -o memcpy_roundtrip memcpy_roundtrip.cu + * + * Then, from the guest: /tmp/memcpy_roundtrip + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#include <cstring> +#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 MemsetByte 0x5A + +__global__ void bump(const float *in, float *out, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + out[i] = in[i] + 0.5f; +} + +int main(void) +{ + cudaDeviceProp prop; + const size_t bytes = 8u << 20; /* 8 MB */ + const int n = (int)(bytes / sizeof(float)); + unsigned char *h_m; + float *h_in, *h_out; + float *d_in, *d_out; + unsigned char *d_m; + double csum = 0.0; + int bad = 0, i; + const int grid = (n + 255) / 256; + + 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); + + h_in = (float *)malloc(bytes); + h_out = (float *)malloc(bytes); + h_m = (unsigned char *)malloc(bytes); + if (!h_in || !h_out || !h_m) + return 1; + + for (i = 0; i < n; i++) + h_in[i] = (float)(i % 1024) * 0.25f; + + CHECK(cudaMalloc(&d_in, bytes)); + CHECK(cudaMalloc(&d_out, bytes)); + CHECK(cudaMalloc(&d_m, bytes)); + + /* H2D -> kernel -> D2H roundtrip. */ + CHECK(cudaMemcpy(d_in, h_in, bytes, cudaMemcpyHostToDevice)); + bump<<<grid, 256>>>(d_in, d_out, n); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_out, d_out, bytes, cudaMemcpyDeviceToHost)); + + for (i = 0; i < n; i++) { + if (h_out[i] != h_in[i] + 0.5f) + bad++; + csum += (double)h_out[i]; + } + + /* cudaMemset pattern, copied back and verified byte-for-byte. */ + CHECK(cudaMemset(d_m, MemsetByte, bytes)); + CHECK(cudaMemcpy(h_m, d_m, bytes, cudaMemcpyDeviceToHost)); + for (i = 0; i < (int)bytes; i++) { + if (h_m[i] != (unsigned char)MemsetByte) + bad++; + } + + printf("ROUNDTRIP bytes=%zu bad=%d csum=%.6f\n", bytes, bad, csum); + + if (bad != 0) { + printf("STATUS FAIL\n"); + return 1; + } + + printf("STATUS OK\n"); + CHECK(cudaFree(d_in)); + CHECK(cudaFree(d_out)); + CHECK(cudaFree(d_m)); + free(h_in); + free(h_out); + free(h_m); + return 0; +} diff --git a/tools/cuda-ab/multikernel.cu b/tools/cuda-ab/multikernel.cu new file mode 100644 index 0000000..8e5f5ed --- /dev/null +++ b/tools/cuda-ab/multikernel.cu @@ -0,0 +1,191 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * multikernel.cu - dependent three-kernel sequence over 1M floats. + * + * Three different kernels run one after another with a data dependency, each + * result verified on the host before the next stage starts: + * + * stage 1 add c[i] = a[i] + b[i] + * stage 2 scale c[i] = c[i] * 3 + * stage 3 saxpy out[i] = 1.5 * c[i] + d[i] + * + * All inputs are exact binary fractions, so the host model is bit-exact and + * any mismatch is a real device-side divergence. + * + * Output lines are stable and greppable: + * + * DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes> + * STAGE idx=<n> name=<add|scale|saxpy> bad=<mismatches> + * MULTI stages=3 bad=<mismatches> + * 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 -o multikernel multikernel.cu + * + * Then, from the guest: /tmp/multikernel + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#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) + +__global__ void k_add(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]; +} + +__global__ void k_scale(float *c, float s, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = c[i] * s; +} + +__global__ void k_saxpy(const float *x, const float *y, float *out, + float alpha, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + out[i] = alpha * x[i] + y[i]; +} + +static int verify_stage(const float *got, const float *a, const float *b, + const float *d, float scale, float alpha, int n, + int stage) +{ + float *exp = (float *)malloc((size_t)n * sizeof(float)); + int i, bad = 0; + const char *name = stage == 1 ? "add" : stage == 2 ? "scale" : "saxpy"; + + if (!exp) + return -1; + + /* Host reference model, matching the kernel chain exactly. */ + for (i = 0; i < n; i++) { + float v = a[i] + b[i]; + + if (stage >= 2) + v = v * scale; + if (stage >= 3) + v = alpha * v + d[i]; + exp[i] = v; + } + for (i = 0; i < n; i++) { + if (got[i] != exp[i]) + bad++; + } + printf("STAGE idx=%d name=%s bad=%d\n", stage, name, bad); + free(exp); + return bad; +} + +int main(void) +{ + cudaDeviceProp prop; + const int n = 1 << 20; /* 1M floats */ + const float scale = 3.0f, alpha = 1.5f; + size_t bytes = (size_t)n * sizeof(float); + float *h_a, *h_b, *h_c, *h_d, *h_out; + float *d_a, *d_b, *d_c, *d_d, *d_out; + int bad = 0, i, s; + const int grid = (n + 255) / 256; + + 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); + + h_a = (float *)malloc(bytes); + h_b = (float *)malloc(bytes); + h_c = (float *)malloc(bytes); + h_d = (float *)malloc(bytes); + h_out = (float *)malloc(bytes); + if (!h_a || !h_b || !h_c || !h_d || !h_out) + return 1; + + for (i = 0; i < n; i++) { + h_a[i] = (float)(i % 1000) * 0.5f; + h_b[i] = (float)(i % 997) * 0.25f; + h_d[i] = (float)(i % 89) * 0.125f; + } + + CHECK(cudaMalloc(&d_a, bytes)); + CHECK(cudaMalloc(&d_b, bytes)); + CHECK(cudaMalloc(&d_c, bytes)); + CHECK(cudaMalloc(&d_d, bytes)); + CHECK(cudaMalloc(&d_out, bytes)); + CHECK(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(d_d, h_d, bytes, cudaMemcpyHostToDevice)); + + /* stage 1: add */ + k_add<<<grid, 256>>>(d_a, d_b, d_c, n); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost)); + s = verify_stage(h_c, h_a, h_b, h_d, scale, alpha, n, 1); + if (s < 0) + return 1; + bad += s; + + /* stage 2: scale (depends on stage 1's c) */ + k_scale<<<grid, 256>>>(d_c, scale, n); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost)); + s = verify_stage(h_c, h_a, h_b, h_d, scale, alpha, n, 2); + if (s < 0) + return 1; + bad += s; + + /* stage 3: saxpy out = alpha * c + d (depends on stage 2's c) */ + k_saxpy<<<grid, 256>>>(d_c, d_d, d_out, alpha, n); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_out, d_out, bytes, cudaMemcpyDeviceToHost)); + s = verify_stage(h_out, h_a, h_b, h_d, scale, alpha, n, 3); + if (s < 0) + return 1; + bad += s; + + printf("MULTI stages=3 bad=%d\n", bad); + + if (bad != 0) { + printf("STATUS FAIL\n"); + return 1; + } + + printf("STATUS OK\n"); + CHECK(cudaFree(d_a)); + CHECK(cudaFree(d_b)); + CHECK(cudaFree(d_c)); + CHECK(cudaFree(d_d)); + CHECK(cudaFree(d_out)); + free(h_a); + free(h_b); + free(h_c); + free(h_d); + free(h_out); + return 0; +} diff --git a/tools/cuda-ab/run/memcpy_roundtrip b/tools/cuda-ab/run/memcpy_roundtrip Binary files differnew file mode 100755 index 0000000..a074ab0 --- /dev/null +++ b/tools/cuda-ab/run/memcpy_roundtrip diff --git a/tools/cuda-ab/run/multikernel b/tools/cuda-ab/run/multikernel Binary files differnew file mode 100755 index 0000000..02b52a9 --- /dev/null +++ b/tools/cuda-ab/run/multikernel diff --git a/tools/cuda-ab/run/streams2 b/tools/cuda-ab/run/streams2 Binary files differnew file mode 100755 index 0000000..7414b76 --- /dev/null +++ b/tools/cuda-ab/run/streams2 diff --git a/tools/cuda-ab/run/vecadd_loop b/tools/cuda-ab/run/vecadd_loop Binary files differnew file mode 100755 index 0000000..4ee5332 --- /dev/null +++ b/tools/cuda-ab/run/vecadd_loop diff --git a/tools/cuda-ab/run/victory_atomics b/tools/cuda-ab/run/victory_atomics Binary files differnew file mode 100755 index 0000000..6049773 --- /dev/null +++ b/tools/cuda-ab/run/victory_atomics diff --git a/tools/cuda-ab/run/victory_math b/tools/cuda-ab/run/victory_math Binary files differnew file mode 100755 index 0000000..5f99902 --- /dev/null +++ b/tools/cuda-ab/run/victory_math diff --git a/tools/cuda-ab/run/victory_matmul b/tools/cuda-ab/run/victory_matmul Binary files differnew file mode 100755 index 0000000..9d201a2 --- /dev/null +++ b/tools/cuda-ab/run/victory_matmul diff --git a/tools/cuda-ab/run/victory_memset2d b/tools/cuda-ab/run/victory_memset2d Binary files differnew file mode 100755 index 0000000..b463e69 --- /dev/null +++ b/tools/cuda-ab/run/victory_memset2d diff --git a/tools/cuda-ab/run/victory_reduce b/tools/cuda-ab/run/victory_reduce Binary files differnew file mode 100755 index 0000000..5a2d8eb --- /dev/null +++ b/tools/cuda-ab/run/victory_reduce diff --git a/tools/cuda-ab/run/victory_saxpy b/tools/cuda-ab/run/victory_saxpy Binary files differnew file mode 100755 index 0000000..08e9c73 --- /dev/null +++ b/tools/cuda-ab/run/victory_saxpy diff --git a/tools/cuda-ab/run/victory_streams_events b/tools/cuda-ab/run/victory_streams_events Binary files differnew file mode 100755 index 0000000..f97cac6 --- /dev/null +++ b/tools/cuda-ab/run/victory_streams_events diff --git a/tools/cuda-ab/run/victory_strided b/tools/cuda-ab/run/victory_strided Binary files differnew file mode 100755 index 0000000..5b5f77f --- /dev/null +++ b/tools/cuda-ab/run/victory_strided diff --git a/tools/cuda-ab/soak.sh b/tools/cuda-ab/soak.sh new file mode 100755 index 0000000..51b542c --- /dev/null +++ b/tools/cuda-ab/soak.sh @@ -0,0 +1,216 @@ +#!/bin/bash +# SPDX-License-Identifier: BSD-3-Clause +# +# soak.sh - CUDA client soak runner for the vGPU guest. +# +# Usage: +# tools/cuda-ab/soak.sh <seconds> [guest_dir] # soak for <seconds> +# tools/cuda-ab/soak.sh --seq [guest_dir] # one pass over all clients +# +# The client set (victory_mail, smiley, vecadd_loop, streams2, +# memcpy_roundtrip, multikernel plus the victory series: victory_saxpy, +# victory_reduce, victory_matmul, victory_strided, victory_memset2d, +# victory_atomics, victory_streams_events, victory_math) runs in the guest +# one at a time via scripts/ssh_guest.sh. Clients whose binary is missing +# from <guest_dir> (default /tmp) are skipped. A client that does not print +# STATUS OK within SOAK_TIMEOUT seconds (default 240) stops the soak with a +# SOAK FAILED line. +# +# Every run is recorded as a timestamped line per run in +# /tmp/opencode/soak-<stamp>/results.log; the full client output of each run +# is kept next to it as out-r<round>-<client>.txt. +# +# Environment overrides: +# SOAK_TIMEOUT per-client seconds in the guest (default 240) +# SSH_CMD guest ssh wrapper (default <repo>/scripts/ssh_guest.sh) +# +# Examples: +# tools/cuda-ab/soak.sh 3600 # 1 h continuous battery +# tools/cuda-ab/soak.sh --seq # single pass, quick check +set -u + +HERE="$(cd "$(dirname "$0")" && pwd)" +REPO="$(cd "$HERE/../.." && pwd)" +SSH="${SSH_CMD:-$REPO/scripts/ssh_guest.sh}" +TMO="${SOAK_TIMEOUT:-240}" +CLIENTS="victory_mail smiley vecadd_loop streams2 memcpy_roundtrip multikernel victory_saxpy victory_reduce victory_matmul victory_strided victory_memset2d victory_atomics victory_streams_events victory_math" + +MODE=soak +if [ "${1:-}" = "--seq" ]; then + MODE=seq + SECS=0 + GUEST_DIR="${2:-/tmp}" +else + SECS="${1:-}" + GUEST_DIR="${2:-/tmp}" +fi + +usage() { + echo "usage: $0 <seconds> [guest_dir]" >&2 + echo " $0 --seq [guest_dir]" >&2 + exit 1 +} + +[ -x "$SSH" ] || { echo "ssh wrapper not found: $SSH" >&2; exit 1; } +if [ "$MODE" = soak ]; then + case "$SECS" in + ''|*[!0-9]*) usage ;; + esac + [ "$SECS" -gt 0 ] || usage +fi + +STAMP="$(date +%Y%m%d-%H%M%S)" +LOGDIR="/tmp/opencode/soak-$STAMP" +mkdir -p "$LOGDIR" +RESULTS="$LOGDIR/results.log" +START=$(date +%s) +DEADLINE=$((START + SECS)) + +log() { + printf '%s %s\n' "$(date '+%Y-%m-%dT%H:%M:%S%z')" "$*" | tee -a "$RESULTS" +} + +# gssh_raw - one guest command, no retry (safe only for idempotent commands). +gssh_raw() { + timeout 60 "$SSH" "$1" 2>/dev/null +} + +# gssh - idempotent guest command, retried for flaky ssh. +gssh() { + local i out + + for i in 1 2 3; do + out="$(gssh_raw "$1")" && { printf '%s' "$out"; return 0; } + sleep 2 + done + return 1 +} + +# start_client <name> <round> - launch under remote `timeout`, output to a +# per-round file ending in EXIT=<rc>. The first attempt launches +# unconditionally; retries only re-launch if the output file is absent, so a +# dropped ssh stream cannot cause a double run. +start_client() { + local name="$1" round="$2" bin="$GUEST_DIR/$1" + local out="$GUEST_DIR/.soak-$1-r$round.out" i r + + for i in 1 2 3; do + if [ "$i" = 1 ]; then + r="$(gssh_raw "rm -f $out; nohup sh -c 'timeout $TMO $bin > $out 2>&1; echo EXIT=\$? >> $out' >/dev/null 2>&1 & echo STARTED")" || r="" + else + r="$(gssh_raw "if [ -f $out ]; then echo STARTED; else rm -f $out; nohup sh -c 'timeout $TMO $bin > $out 2>&1; echo EXIT=\$? >> $out' >/dev/null 2>&1 & echo STARTED; fi")" || r="" + fi + case "$r" in + *STARTED*) return 0 ;; + esac + sleep 3 + done + return 1 +} + +# client_alive <bin> - pgrep without self-match ([/] in the pattern keeps the +# remote check shell - whose argv contains the pattern text - from matching). +client_alive() { + local pat="[/]${1#/}" + + gssh "pgrep -f '$pat' >/dev/null && echo ALIVE" | grep -q ALIVE +} + +# run_one <name> <round> - returns 0 when STATUS OK was seen, 1 otherwise. +run_one() { + local name="$1" round="$2" bin="$GUEST_DIR/$1" + local out="$GUEST_DIR/.soak-$1-r$round.out" + local text="" r="" i t0 tnow dur rc reason + + t0=$(date +%s) + log "RUN client=$name round=$round timeout=${TMO}s" + + if ! start_client "$name" "$round"; then + log "FAILED client=$name round=$round reason=ssh-launch-failed" + return 1 + fi + + while :; do + sleep 10 + r="$(gssh "cat $out 2>/dev/null")" || true + [ -n "$r" ] && text="$r" + printf '%s' "$text" | grep -q 'EXIT=' && break + tnow=$(date +%s) + if [ $((tnow - t0)) -gt $((TMO + 120)) ]; then + if client_alive "$bin"; then + log "FAILED client=$name round=$round reason=still-running-after-$((tnow - t0))s" + else + log "FAILED client=$name round=$round reason=process-gone-without-exit-marker" + fi + return 1 + fi + done + + tnow=$(date +%s) + dur=$((tnow - t0)) + + # Flaky-ssh guard: retry the final status read a few times. + for i in 1 2 3; do + r="$(gssh "cat $out 2>/dev/null")" || true + if printf '%s' "$r" | grep -q 'EXIT='; then + text="$r" + break + fi + sleep 2 + done + + rc=$(printf '%s' "$text" | sed -n 's/.*EXIT=\([0-9][0-9]*\).*/\1/p' | tail -1) + [ -n "$rc" ] || rc=? + printf '%s\n' "$text" > "$LOGDIR/out-r${round}-${name}.txt" + + if printf '%s' "$text" | grep -q 'STATUS OK'; then + log "OK client=$name rc=$rc dur=${dur}s round=$round" + return 0 + fi + reason="$(printf '%s' "$text" | grep -E 'STATUS FAIL|CUDA_ERROR' | head -1)" + [ -n "$reason" ] || reason="no STATUS OK in output" + log "FAILED client=$name rc=$rc dur=${dur}s round=$round reason=$reason" + return 1 +} + +log "SOAK START mode=$MODE seconds=$SECS guest_dir=$GUEST_DIR timeout=${TMO}s logdir=$LOGDIR" + +# Which clients are present in the guest? Retry the listing so a flaky ssh +# 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="" + [ -n "$found" ] && break + sleep 3 +done +if [ -z "$found" ]; then + log "SOAK FAILED reason=no-client-binaries-in-$GUEST_DIR" + exit 1 +fi +log "CLIENTS $(printf '%s' "$found" | tr '\n' ' ')" + +round=0 +runs=0 +oks=0 +while :; do + [ "$MODE" = soak ] && [ "$(date +%s)" -ge "$DEADLINE" ] && break + round=$((round + 1)) + stopped=0 + for c in $found; do + if [ "$MODE" = soak ] && [ "$(date +%s)" -ge "$DEADLINE" ]; then + stopped=1 + break + fi + if ! run_one "$c" "$round"; then + log "SOAK FAILED client=$c round=$round runs=$runs ok=$oks" + exit 1 + fi + runs=$((runs + 1)) + oks=$((oks + 1)) + done + [ "$MODE" = seq ] && break + [ "$stopped" = 1 ] && break +done + +log "SOAK DONE mode=$MODE rounds=$round runs=$runs ok=$oks elapsed=$(( $(date +%s) - START ))s" +exit 0 diff --git a/tools/cuda-ab/streams2.cu b/tools/cuda-ab/streams2.cu new file mode 100644 index 0000000..a5dc5c2 --- /dev/null +++ b/tools/cuda-ab/streams2.cu @@ -0,0 +1,175 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * streams2.cu - two non-default streams, concurrent kernels, event ordering. + * + * Two independent vectorAdds run concurrently on two non-default streams. + * A cross-stream dependency is then imposed with events: stream 2's in-place + * scale waits on stream 1's add via cudaEventRecord/cudaStreamWaitEvent, and + * stream 1's copy-back waits on stream 2's scale the same way. Both streams + * are joined with cudaStreamSynchronize before the host verifies: + * + * c1 == a1 + b1 + * c2 == (a2 + b2) * 2 + * + * Output lines are stable and greppable: + * + * DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes> + * STREAMS n=<count> bad=<mismatches> + * TIMING ms=<elapsed> + * 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 -o streams2 streams2.cu + * + * Then, from the guest: /tmp/streams2 + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#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) + +__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]; +} + +__global__ void scale(float *c, float s, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = c[i] * s; +} + +int main(void) +{ + cudaDeviceProp prop; + const int n = 1 << 15; + const float s2 = 2.0f; + size_t bytes = (size_t)n * sizeof(float); + float *h_a1, *h_b1, *h_c1, *h_a2, *h_b2, *h_c2; + float *d_a1, *d_b1, *d_c1, *d_a2, *d_b2, *d_c2; + cudaStream_t s1, s2s; + cudaEvent_t ev_add1, ev_scale2, t0, t1; + float ms = 0.0f; + int bad = 0, i; + const int grid = (n + 255) / 256; + + 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); + + h_a1 = (float *)malloc(bytes); + h_b1 = (float *)malloc(bytes); + h_c1 = (float *)malloc(bytes); + h_a2 = (float *)malloc(bytes); + h_b2 = (float *)malloc(bytes); + h_c2 = (float *)malloc(bytes); + if (!h_a1 || !h_b1 || !h_c1 || !h_a2 || !h_b2 || !h_c2) + return 1; + + for (i = 0; i < n; i++) { + h_a1[i] = (float)(i % 1000) * 0.5f; + h_b1[i] = (float)(i % 997) * 0.25f; + h_a2[i] = (float)(i % 991) * 0.125f; + h_b2[i] = (float)(i % 983) * 0.0625f; + } + + CHECK(cudaMalloc(&d_a1, bytes)); + CHECK(cudaMalloc(&d_b1, bytes)); + CHECK(cudaMalloc(&d_c1, bytes)); + CHECK(cudaMalloc(&d_a2, bytes)); + CHECK(cudaMalloc(&d_b2, bytes)); + CHECK(cudaMalloc(&d_c2, bytes)); + + CHECK(cudaStreamCreate(&s1)); + CHECK(cudaStreamCreate(&s2s)); + CHECK(cudaEventCreate(&ev_add1)); + CHECK(cudaEventCreate(&ev_scale2)); + CHECK(cudaEventCreate(&t0)); + CHECK(cudaEventCreate(&t1)); + + CHECK(cudaEventRecord(t0)); + + /* Both input copies and both add kernels run concurrently. */ + CHECK(cudaMemcpyAsync(d_a1, h_a1, bytes, cudaMemcpyHostToDevice, s1)); + CHECK(cudaMemcpyAsync(d_b1, h_b1, bytes, cudaMemcpyHostToDevice, s1)); + CHECK(cudaMemcpyAsync(d_a2, h_a2, bytes, cudaMemcpyHostToDevice, s2s)); + CHECK(cudaMemcpyAsync(d_b2, h_b2, bytes, cudaMemcpyHostToDevice, s2s)); + vadd<<<grid, 256, 0, s1>>>(d_a1, d_b1, d_c1, n); + vadd<<<grid, 256, 0, s2s>>>(d_a2, d_b2, d_c2, n); + CHECK(cudaGetLastError()); + + /* stream 2's scale waits for stream 1's add... */ + CHECK(cudaEventRecord(ev_add1, s1)); + CHECK(cudaStreamWaitEvent(s2s, ev_add1, 0)); + scale<<<grid, 256, 0, s2s>>>(d_c2, s2, n); + CHECK(cudaGetLastError()); + + /* ...and stream 1's copy-back waits for stream 2's scale. */ + CHECK(cudaEventRecord(ev_scale2, s2s)); + CHECK(cudaStreamWaitEvent(s1, ev_scale2, 0)); + CHECK(cudaMemcpyAsync(h_c1, d_c1, bytes, cudaMemcpyDeviceToHost, s1)); + CHECK(cudaMemcpyAsync(h_c2, d_c2, bytes, cudaMemcpyDeviceToHost, s2s)); + CHECK(cudaEventRecord(t1, s1)); + + CHECK(cudaStreamSynchronize(s1)); + CHECK(cudaStreamSynchronize(s2s)); + CHECK(cudaEventElapsedTime(&ms, t0, t1)); + + for (i = 0; i < n; i++) { + if (h_c1[i] != h_a1[i] + h_b1[i]) + bad++; + if (h_c2[i] != (h_a2[i] + h_b2[i]) * s2) + bad++; + } + + printf("STREAMS n=%d bad=%d\n", n, bad); + printf("TIMING ms=%.3f\n", ms); + + if (bad != 0) { + printf("STATUS FAIL\n"); + return 1; + } + + printf("STATUS OK\n"); + CHECK(cudaStreamDestroy(s1)); + CHECK(cudaStreamDestroy(s2s)); + CHECK(cudaEventDestroy(ev_add1)); + CHECK(cudaEventDestroy(ev_scale2)); + CHECK(cudaEventDestroy(t0)); + CHECK(cudaEventDestroy(t1)); + CHECK(cudaFree(d_a1)); + CHECK(cudaFree(d_b1)); + CHECK(cudaFree(d_c1)); + CHECK(cudaFree(d_a2)); + CHECK(cudaFree(d_b2)); + CHECK(cudaFree(d_c2)); + free(h_a1); + free(h_b1); + free(h_c1); + free(h_a2); + free(h_b2); + free(h_c2); + return 0; +} diff --git a/tools/cuda-ab/vecadd_loop.cu b/tools/cuda-ab/vecadd_loop.cu new file mode 100644 index 0000000..bde3d05 --- /dev/null +++ b/tools/cuda-ab/vecadd_loop.cu @@ -0,0 +1,121 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * vecadd_loop.cu - 100 repeated vectorAdd launches, soak-style. + * + * The same three device buffers are reused for 100 identical vadd launches: + * the point is repetition and stability, not new data. Every 10th launch the + * result is copied back and verified element by element against the + * host-computed a+b, so a silent corruption that only appears after many + * launches is caught. + * + * Output lines are stable and greppable: + * + * DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes> + * VERIFY launch=<n> bad=<mismatches> + * LOOP launches=100 bad=<mismatches> + * 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 -o vecadd_loop vecadd_loop.cu + * + * Then, from the guest: /tmp/vecadd_loop + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#include <cuda_runtime.h> + +#define NLAUNCH 100 +#define VERIFY_EVERY 10 + +#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) + +__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]; +} + +int main(void) +{ + cudaDeviceProp prop; + const int n = 1 << 15; + size_t bytes = (size_t)n * sizeof(float); + float *h_a, *h_b, *h_c; + float *d_a, *d_b, *d_c; + int bad = 0, i, launch; + + 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); + + h_a = (float *)malloc(bytes); + h_b = (float *)malloc(bytes); + h_c = (float *)malloc(bytes); + if (!h_a || !h_b || !h_c) + return 1; + + for (i = 0; i < n; i++) { + h_a[i] = (float)(i % 1000) * 0.5f; + h_b[i] = (float)(i % 997) * 0.25f; + } + + CHECK(cudaMalloc(&d_a, bytes)); + CHECK(cudaMalloc(&d_b, bytes)); + CHECK(cudaMalloc(&d_c, bytes)); + CHECK(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice)); + + for (launch = 1; launch <= NLAUNCH; launch++) { + vadd<<<(n + 255) / 256, 256>>>(d_a, d_b, d_c, n); + CHECK(cudaGetLastError()); + if (launch % VERIFY_EVERY != 0) + continue; + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost)); + { + int vbad = 0; + + for (i = 0; i < n; i++) { + if (h_c[i] != h_a[i] + h_b[i]) + vbad++; + } + printf("VERIFY launch=%d bad=%d\n", launch, vbad); + bad += vbad; + } + } + CHECK(cudaDeviceSynchronize()); + + printf("LOOP launches=%d bad=%d\n", NLAUNCH, bad); + + if (bad != 0) { + printf("STATUS FAIL\n"); + return 1; + } + + printf("STATUS OK\n"); + CHECK(cudaFree(d_a)); + CHECK(cudaFree(d_b)); + CHECK(cudaFree(d_c)); + free(h_a); + free(h_b); + free(h_c); + return 0; +} diff --git a/tools/cuda-ab/victory_atomics.cu b/tools/cuda-ab/victory_atomics.cu new file mode 100644 index 0000000..7e5eedf --- /dev/null +++ b/tools/cuda-ab/victory_atomics.cu @@ -0,0 +1,129 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * victory_atomics.cu - global atomicAdd/atomicMax histogram over 4 MB. + * + * A deterministic 4M-byte buffer (every value 0..255) is histogrammed on the + * device into a 256-bin counter array with atomicAdd, while a global maximum + * is tracked with atomicMax. The host recomputes the exact histogram and + * maximum and verifies every bin, the max, and the total count. + * + * Output lines are stable and greppable: + * + * DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes> + * ATOMICS n=<count> bins=256 max=<value> bad=<mismatches> total=<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 -o victory_atomics victory_atomics.cu + * + * Then, from the guest: /tmp/victory_atomics + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#include <cstring> +#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 BINS 256 + +__global__ void histogram(const unsigned char *data, unsigned int *hist, + unsigned int *gmax, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) { + unsigned int v = data[i]; + + atomicAdd(&hist[v], 1u); + atomicMax(gmax, v); + } +} + +int main(void) +{ + cudaDeviceProp prop; + const int n = 1 << 22; /* 4M bytes */ + unsigned char *h_data; + unsigned int *h_hist; + unsigned int h_max = 0, h_total = 0; + unsigned char *d_data; + unsigned int *d_hist, *d_max; + unsigned int ref[BINS]; + unsigned int ref_max = 0; + int bad = 0, i; + + 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); + + h_data = (unsigned char *)malloc(n); + h_hist = (unsigned int *)malloc(BINS * sizeof(unsigned int)); + if (!h_data || !h_hist) + return 1; + + for (i = 0; i < n; i++) + h_data[i] = (unsigned char)(((long long)i * 37 + i / 256) & 0xFF); + + memset(ref, 0, sizeof(ref)); + for (i = 0; i < n; i++) { + ref[h_data[i]]++; + if (h_data[i] > ref_max) + ref_max = h_data[i]; + } + + CHECK(cudaMalloc(&d_data, n)); + CHECK(cudaMalloc(&d_hist, BINS * sizeof(unsigned int))); + CHECK(cudaMalloc(&d_max, sizeof(unsigned int))); + CHECK(cudaMemcpy(d_data, h_data, n, cudaMemcpyHostToDevice)); + CHECK(cudaMemset(d_hist, 0, BINS * sizeof(unsigned int))); + CHECK(cudaMemset(d_max, 0, sizeof(unsigned int))); + + histogram<<<(n + 255) / 256, 256>>>(d_data, d_hist, d_max, n); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_hist, d_hist, BINS * sizeof(unsigned int), + cudaMemcpyDeviceToHost)); + CHECK(cudaMemcpy(&h_max, d_max, sizeof(unsigned int), + cudaMemcpyDeviceToHost)); + + for (i = 0; i < BINS; i++) { + if (h_hist[i] != ref[i]) + bad++; + h_total += h_hist[i]; + } + if (h_max != ref_max) + bad++; + + printf("ATOMICS n=%d bins=256 max=%u bad=%d total=%u\n", n, h_max, bad, + h_total); + + if (bad != 0 || h_total != (unsigned int)n) { + printf("STATUS FAIL\n"); + return 1; + } + + printf("STATUS OK\n"); + CHECK(cudaFree(d_data)); + CHECK(cudaFree(d_hist)); + CHECK(cudaFree(d_max)); + free(h_data); + free(h_hist); + return 0; +} diff --git a/tools/cuda-ab/victory_math.cu b/tools/cuda-ab/victory_math.cu new file mode 100644 index 0000000..f9e5807 --- /dev/null +++ b/tools/cuda-ab/victory_math.cu @@ -0,0 +1,163 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * victory_math.cu - mixed int/float/double math over 64K elements. + * + * Each thread computes a float lane (sqrtf + fmaf), a double lane (sqrt + + * multiply), and an integer lane (wrapping multiply/add, xor, shift, popc). + * The host recomputes all three lanes and verifies the float/double lanes + * within a relative tolerance and the integer lane exactly. + * + * Output lines are stable and greppable: + * + * DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes> + * MATH n=<count> badf=<float_mismatches> badd=<double_mismatches> badi=<int_mismatches> csumf=<sum> csumd=<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 -o victory_math victory_math.cu + * + * Then, from the guest: /tmp/victory_math + */ + +#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) + +/* Portable popcount: device intrinsic under nvcc, host builtin elsewhere. */ +__host__ __device__ static inline unsigned int popc(unsigned int v) +{ +#if defined(__CUDA_ARCH__) + return (unsigned int)__popc(v); +#else + return (unsigned int)__builtin_popcount(v); +#endif +} + +/* Integer lane, replicated bit-for-bit on the host. */ +__host__ __device__ static inline int int_lane(unsigned int u) +{ + unsigned int w = u * 31u + 7u; + + return (int)((w ^ (u >> 3)) + popc(u & 0xFFu)); +} + +__global__ void mathmix(const unsigned int *in, float *of, double *od, + int *oi, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) { + unsigned int u = in[i]; + + of[i] = sqrtf((float)u) + + fmaf((float)(u % 7u), 1.5f, -2.0f); + od[i] = sqrt((double)u) * 1.0000001; + oi[i] = int_lane(u); + } +} + +int main(void) +{ + cudaDeviceProp prop; + const int n = 1 << 16; /* 64K elements */ + size_t ib = (size_t)n * sizeof(unsigned int); + size_t fb = (size_t)n * sizeof(float); + size_t db = (size_t)n * sizeof(double); + unsigned int *h_in; + float *h_of, *h_rf; + double *h_od, *h_rd; + int *h_oi, *h_ri; + unsigned int *d_in; + float *d_of; + double *d_od; + int *d_oi; + double csumf = 0.0, csumd = 0.0; + int badf = 0, badd = 0, badi = 0, i; + + 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); + + h_in = (unsigned int *)malloc(ib); + h_of = (float *)malloc(fb); + h_rf = (float *)malloc(fb); + h_od = (double *)malloc(db); + h_rd = (double *)malloc(db); + h_oi = (int *)malloc(ib); + h_ri = (int *)malloc(ib); + if (!h_in || !h_of || !h_rf || !h_od || !h_rd || !h_oi || !h_ri) + return 1; + + for (i = 0; i < n; i++) + h_in[i] = (unsigned int)(i % 100000) + 1u; + + CHECK(cudaMalloc(&d_in, ib)); + CHECK(cudaMalloc(&d_of, fb)); + CHECK(cudaMalloc(&d_od, db)); + CHECK(cudaMalloc(&d_oi, ib)); + CHECK(cudaMemcpy(d_in, h_in, ib, cudaMemcpyHostToDevice)); + + mathmix<<<(n + 255) / 256, 256>>>(d_in, d_of, d_od, d_oi, n); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_of, d_of, fb, cudaMemcpyDeviceToHost)); + CHECK(cudaMemcpy(h_od, d_od, db, cudaMemcpyDeviceToHost)); + CHECK(cudaMemcpy(h_oi, d_oi, ib, cudaMemcpyDeviceToHost)); + + for (i = 0; i < n; i++) { + unsigned int u = h_in[i]; + + h_rf[i] = sqrtf((float)u) + fmaf((float)(u % 7u), 1.5f, -2.0f); + h_rd[i] = sqrt((double)u) * 1.0000001; + h_ri[i] = int_lane(u); + + if (fabsf(h_of[i] - h_rf[i]) > 1e-4f * (1.0f + fabsf(h_rf[i]))) + badf++; + if (fabs(h_od[i] - h_rd[i]) > 1e-12 * (1.0 + fabs(h_rd[i]))) + badd++; + if (h_oi[i] != h_ri[i]) + badi++; + csumf += (double)h_of[i]; + csumd += h_od[i]; + } + + printf("MATH n=%d badf=%d badd=%d badi=%d csumf=%.6f csumd=%.6f\n", n, + badf, badd, badi, csumf, csumd); + + if (badf != 0 || badd != 0 || badi != 0) { + printf("STATUS FAIL\n"); + return 1; + } + + printf("STATUS OK\n"); + CHECK(cudaFree(d_in)); + CHECK(cudaFree(d_of)); + CHECK(cudaFree(d_od)); + CHECK(cudaFree(d_oi)); + free(h_in); + free(h_of); + free(h_rf); + free(h_od); + free(h_rd); + free(h_oi); + free(h_ri); + return 0; +} diff --git a/tools/cuda-ab/victory_matmul.cu b/tools/cuda-ab/victory_matmul.cu new file mode 100644 index 0000000..78ab59c --- /dev/null +++ b/tools/cuda-ab/victory_matmul.cu @@ -0,0 +1,135 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * victory_matmul.cu - 256x256 tiled matrix multiply with shared memory. + * + * C = A * B for three 256x256 float matrices of small integer values. The + * inputs are integers in [0, 10), so every product and every partial sum is + * exactly representable in float (< 2^24); whatever order the GPU accumulates + * in, the result is exact. The host replays the multiply with 64-bit integers + * and verifies every element exactly, plus a checksum over C. + * + * Output lines are stable and greppable: + * + * DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes> + * MATMUL n=256 bad=<mismatches> csum=<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 -o victory_matmul victory_matmul.cu + * + * Then, from the guest: /tmp/victory_matmul + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#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 TILE 16 +#define N 256 + +__global__ void matmul_tiled(const float *A, const float *B, float *C, int n) +{ + __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 < n / TILE; t++) { + As[threadIdx.y][threadIdx.x] = A[row * n + 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; +} + +int main(void) +{ + cudaDeviceProp prop; + const int n = N; + const size_t bytes = (size_t)n * n * sizeof(float); + float *h_a, *h_b, *h_c; + float *d_a, *d_b, *d_c; + double csum = 0.0; + int bad = 0, i, j, k; + dim3 grid(n / TILE, n / TILE); + dim3 block(TILE, TILE); + + 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); + + h_a = (float *)malloc(bytes); + h_b = (float *)malloc(bytes); + h_c = (float *)malloc(bytes); + if (!h_a || !h_b || !h_c) + return 1; + + for (i = 0; i < n * n; i++) { + h_a[i] = (float)(i % 10); + h_b[i] = (float)((i * 3 + 1) % 10); + } + + CHECK(cudaMalloc(&d_a, bytes)); + CHECK(cudaMalloc(&d_b, bytes)); + CHECK(cudaMalloc(&d_c, bytes)); + CHECK(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice)); + + matmul_tiled<<<grid, block>>>(d_a, d_b, d_c, n); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost)); + + for (i = 0; i < n; i++) { + for (j = 0; j < n; j++) { + long long ref = 0; + + for (k = 0; k < n; k++) + ref += (long long)h_a[i * n + k] * + (long long)h_b[k * n + j]; + if (h_c[i * n + j] != (float)ref) + bad++; + csum += (double)h_c[i * n + j]; + } + } + + printf("MATMUL n=%d bad=%d csum=%.6f\n", n, bad, csum); + + if (bad != 0) { + printf("STATUS FAIL\n"); + return 1; + } + + printf("STATUS OK\n"); + CHECK(cudaFree(d_a)); + CHECK(cudaFree(d_b)); + CHECK(cudaFree(d_c)); + free(h_a); + free(h_b); + free(h_c); + return 0; +} diff --git a/tools/cuda-ab/victory_memset2d.cu b/tools/cuda-ab/victory_memset2d.cu new file mode 100644 index 0000000..25a163a --- /dev/null +++ b/tools/cuda-ab/victory_memset2d.cu @@ -0,0 +1,125 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * victory_memset2d.cu - cudaMemset + pitched cudaMemcpy2D round trip. + * + * A 4096-column x 1024-row byte image is allocated with cudaMallocPitch + * (row pitch >= width), then exercised three ways: + * + * 1. cudaMemset the whole allocation (including pitch padding) to 0xA5 and + * read it back row by row with cudaMemcpy2D; every byte must match. + * 2. H2D cudaMemcpy2D of a host pattern, D2H cudaMemcpy2D of the same + * image; every byte must match the host pattern. + * 3. cudaMemset2D of the upper-left quadrant to 0x3C; read back and verify + * that exactly that rectangle changed. + * + * Output lines are stable and greppable: + * + * DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes> + * MEMSET2D w=<width> h=<height> pitch=<pitch> bad=<mismatches> + * 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 -o victory_memset2d victory_memset2d.cu + * + * Then, from the guest: /tmp/victory_memset2d + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#include <cstring> +#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 IMG_W 4096 +#define IMG_H 1024 +#define FILL_BYTE 0xA5 +#define QUAD_BYTE 0x3C + +int main(void) +{ + cudaDeviceProp prop; + const size_t w = IMG_W, h = IMG_H; + unsigned char *h_img, *h_back; + unsigned char *d_img; + size_t pitch = 0; + int bad = 0, i, row, col; + + 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); + + h_img = (unsigned char *)malloc(w * h); + h_back = (unsigned char *)malloc(w * h); + if (!h_img || !h_back) + return 1; + + CHECK(cudaMallocPitch(&d_img, &pitch, w, h)); + + /* 1. Full cudaMemset over the pitched allocation, read back with 2D. */ + CHECK(cudaMemset(d_img, FILL_BYTE, pitch * h)); + CHECK(cudaMemcpy2D(h_back, w, d_img, pitch, w, h, + cudaMemcpyDeviceToHost)); + for (i = 0; i < (int)(w * h); i++) { + if (h_back[i] != (unsigned char)FILL_BYTE) + bad++; + } + + /* 2. Pitched H2D -> D2H round trip of a host pattern. */ + for (row = 0; row < (int)h; row++) + for (col = 0; col < (int)w; col++) + h_img[row * w + col] = (unsigned char)((row + col) & 0xFF); + CHECK(cudaMemcpy2D(d_img, pitch, h_img, w, w, h, + cudaMemcpyHostToDevice)); + memset(h_back, 0, w * h); + CHECK(cudaMemcpy2D(h_back, w, d_img, pitch, w, h, + cudaMemcpyDeviceToHost)); + for (i = 0; i < (int)(w * h); i++) { + if (h_back[i] != h_img[i]) + bad++; + } + + /* 3. cudaMemset2D of the upper-left quadrant only. */ + CHECK(cudaMemset2D(d_img, pitch, QUAD_BYTE, w / 2, h / 2)); + memset(h_back, 0, w * h); + CHECK(cudaMemcpy2D(h_back, w, d_img, pitch, w, h, + cudaMemcpyDeviceToHost)); + for (row = 0; row < (int)h; row++) { + for (col = 0; col < (int)w; col++) { + unsigned char expect = h_img[row * w + col]; + + if (row < (int)(h / 2) && col < (int)(w / 2)) + expect = QUAD_BYTE; + if (h_back[row * w + col] != expect) + bad++; + } + } + + printf("MEMSET2D w=%zu h=%zu pitch=%zu bad=%d\n", w, h, pitch, bad); + + if (bad != 0) { + printf("STATUS FAIL\n"); + return 1; + } + + printf("STATUS OK\n"); + CHECK(cudaFree(d_img)); + free(h_img); + free(h_back); + return 0; +} diff --git a/tools/cuda-ab/victory_reduce.cu b/tools/cuda-ab/victory_reduce.cu new file mode 100644 index 0000000..a18188f --- /dev/null +++ b/tools/cuda-ab/victory_reduce.cu @@ -0,0 +1,112 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * victory_reduce.cu - block reduction to shared memory + global atomics. + * + * Every block reduces its own slice of the values 1..N with a shared-memory + * tree reduction; each block's partial sum is then folded into one global + * counter with 64-bit atomicAdd. The host verifies the exact closed form + * N*(N+1)/2, so any lost or double-counted block fails the client. + * + * Output lines are stable and greppable: + * + * DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes> + * REDUCE n=<count> sum=<gpu_total> expected=<closed_form> bad=<0|1> + * 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 -o victory_reduce victory_reduce.cu + * + * Then, from the guest: /tmp/victory_reduce + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#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 NTHREADS 256 + +__global__ void reduce_sum(const unsigned int *in, unsigned long long *out, + int n) +{ + __shared__ unsigned long long s[NTHREADS]; + int i = blockIdx.x * blockDim.x + threadIdx.x; + + s[threadIdx.x] = (i < n) ? (unsigned long long)in[i] : 0ULL; + __syncthreads(); + + for (int off = NTHREADS / 2; off > 0; off >>= 1) { + if (threadIdx.x < off) + s[threadIdx.x] += s[threadIdx.x + off]; + __syncthreads(); + } + + if (threadIdx.x == 0) + atomicAdd(out, s[0]); +} + +int main(void) +{ + cudaDeviceProp prop; + const int n = 1 << 20; /* 1M values */ + size_t bytes = (size_t)n * sizeof(unsigned int); + unsigned int *h_in; + unsigned long long h_sum = 0; + unsigned int *d_in; + unsigned long long *d_sum; + unsigned long long expected, bad; + int i; + + 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); + + h_in = (unsigned int *)malloc(bytes); + if (!h_in) + return 1; + for (i = 0; i < n; i++) + h_in[i] = (unsigned int)(i + 1); + + CHECK(cudaMalloc(&d_in, bytes)); + CHECK(cudaMalloc(&d_sum, sizeof(unsigned long long))); + CHECK(cudaMemcpy(d_in, h_in, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemset(d_sum, 0, sizeof(unsigned long long))); + + reduce_sum<<<(n + NTHREADS - 1) / NTHREADS, NTHREADS>>>(d_in, d_sum, n); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(&h_sum, d_sum, sizeof(h_sum), cudaMemcpyDeviceToHost)); + + expected = (unsigned long long)n * (unsigned long long)(n + 1) / 2ULL; + bad = (h_sum == expected) ? 0 : 1; + + printf("REDUCE n=%d sum=%llu expected=%llu bad=%llu\n", n, h_sum, + expected, bad); + + if (bad != 0) { + printf("STATUS FAIL\n"); + return 1; + } + + printf("STATUS OK\n"); + CHECK(cudaFree(d_in)); + CHECK(cudaFree(d_sum)); + free(h_in); + return 0; +} diff --git a/tools/cuda-ab/victory_saxpy.cu b/tools/cuda-ab/victory_saxpy.cu new file mode 100644 index 0000000..387eb6c --- /dev/null +++ b/tools/cuda-ab/victory_saxpy.cu @@ -0,0 +1,124 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * victory_saxpy.cu - SAXPY (y = a*x + y) over 1M floats, two alphas. + * + * Two back-to-back launches over a 4 MB buffer: + * + * y = 3.0 * x + y + * y = 0.25 * x + y + * + * The host replays the same two steps from the same inputs and verifies the + * device result element by element with a small tolerance (the device may + * contract multiply-add to FMA; the host may not). + * + * Output lines are stable and greppable: + * + * DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes> + * SAXPY n=<count> launches=2 bad=<mismatches> csum=<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 -o victory_saxpy victory_saxpy.cu + * + * Then, from the guest: /tmp/victory_saxpy + */ + +#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) + +__global__ void saxpy(const float *x, float *y, float alpha, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + y[i] = alpha * x[i] + y[i]; +} + +int main(void) +{ + cudaDeviceProp prop; + const int n = 1 << 20; /* 1M floats = 4 MB */ + const float alpha1 = 3.0f, alpha2 = 0.25f; + size_t bytes = (size_t)n * sizeof(float); + float *h_x, *h_y0, *h_ref, *h_y; + float *d_x, *d_y; + double csum = 0.0; + int bad = 0, i; + + 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); + + h_x = (float *)malloc(bytes); + h_y0 = (float *)malloc(bytes); + h_ref = (float *)malloc(bytes); + h_y = (float *)malloc(bytes); + if (!h_x || !h_y0 || !h_ref || !h_y) + return 1; + + for (i = 0; i < n; i++) { + h_x[i] = (float)(i % 1000) * 0.5f; + h_y0[i] = (float)(i % 997) * 0.25f; + h_ref[i] = h_y0[i]; + } + + /* Host reference: the same two SAXPY steps, in the same order. */ + for (i = 0; i < n; i++) + h_ref[i] = alpha1 * h_x[i] + h_ref[i]; + for (i = 0; i < n; i++) + h_ref[i] = alpha2 * h_x[i] + h_ref[i]; + + CHECK(cudaMalloc(&d_x, bytes)); + CHECK(cudaMalloc(&d_y, bytes)); + CHECK(cudaMemcpy(d_x, h_x, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(d_y, h_y0, bytes, cudaMemcpyHostToDevice)); + + saxpy<<<(n + 255) / 256, 256>>>(d_x, d_y, alpha1, n); + saxpy<<<(n + 255) / 256, 256>>>(d_x, d_y, alpha2, n); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_y, d_y, bytes, cudaMemcpyDeviceToHost)); + + for (i = 0; i < n; i++) { + float tol = 1e-3f + 1e-5f * fabsf(h_ref[i]); + + if (fabsf(h_y[i] - h_ref[i]) > tol) + bad++; + csum += (double)h_y[i]; + } + + printf("SAXPY n=%d launches=2 bad=%d csum=%.6f\n", n, bad, csum); + + if (bad != 0) { + printf("STATUS FAIL\n"); + return 1; + } + + printf("STATUS OK\n"); + CHECK(cudaFree(d_x)); + CHECK(cudaFree(d_y)); + free(h_x); + free(h_y0); + free(h_ref); + free(h_y); + return 0; +} diff --git a/tools/cuda-ab/victory_streams_events.cu b/tools/cuda-ab/victory_streams_events.cu new file mode 100644 index 0000000..f54bd85 --- /dev/null +++ b/tools/cuda-ab/victory_streams_events.cu @@ -0,0 +1,172 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * victory_streams_events.cu - 3 streams, overlapped copies, event ordering. + * + * Three non-default streams each copy their own inputs, run a vectorAdd, and + * are then chained with events: + * + * stream 0: vadd c0 = a0 + b0 + * stream 1: waits event e0 from stream 0, then c1 += c0 + * stream 2: waits event e1 from stream 1, then c2 += c1 + * + * Every stream also does its own async D2H copy; the whole chain is timed + * with cudaEventRecord/cudaEventElapsedTime. The host verifies all three + * outputs; inputs are multiples of 0.25 so the arithmetic is exact in float. + * + * Output lines are stable and greppable: + * + * DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes> + * STREAMS3 n=<count> bad=<mismatches> + * TIMING ms=<elapsed> + * 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 -o victory_streams_events victory_streams_events.cu + * + * Then, from the guest: /tmp/victory_streams_events + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#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 NSTREAMS 3 + +__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]; +} + +__global__ void vadd_into(float *c, const float *s, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = c[i] + s[i]; +} + +int main(void) +{ + cudaDeviceProp prop; + const int n = 1 << 16; /* 256 KB per stream */ + size_t bytes = (size_t)n * sizeof(float); + float *h_a[NSTREAMS], *h_b[NSTREAMS], *h_c[NSTREAMS]; + float *d_a[NSTREAMS], *d_b[NSTREAMS], *d_c[NSTREAMS]; + cudaStream_t st[NSTREAMS]; + cudaEvent_t ev[NSTREAMS], t0, t1; + float ms = 0.0f; + int bad = 0, i, s; + const int grid = (n + 255) / 256; + + 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); + + for (s = 0; s < NSTREAMS; s++) { + h_a[s] = (float *)malloc(bytes); + h_b[s] = (float *)malloc(bytes); + h_c[s] = (float *)malloc(bytes); + if (!h_a[s] || !h_b[s] || !h_c[s]) + return 1; + for (i = 0; i < n; i++) { + h_a[s][i] = (float)((i + s) % 1000) * 0.5f; + h_b[s][i] = (float)((i + 2 * s) % 997) * 0.25f; + } + CHECK(cudaMalloc(&d_a[s], bytes)); + CHECK(cudaMalloc(&d_b[s], bytes)); + CHECK(cudaMalloc(&d_c[s], bytes)); + CHECK(cudaStreamCreate(&st[s])); + CHECK(cudaEventCreate(&ev[s])); + } + CHECK(cudaEventCreate(&t0)); + CHECK(cudaEventCreate(&t1)); + + CHECK(cudaEventRecord(t0)); + + for (s = 0; s < NSTREAMS; s++) { + CHECK(cudaMemcpyAsync(d_a[s], h_a[s], bytes, + cudaMemcpyHostToDevice, st[s])); + CHECK(cudaMemcpyAsync(d_b[s], h_b[s], bytes, + cudaMemcpyHostToDevice, st[s])); + vadd<<<grid, 256, 0, st[s]>>>(d_a[s], d_b[s], d_c[s], n); + CHECK(cudaGetLastError()); + CHECK(cudaEventRecord(ev[s], st[s])); + } + + /* Cross-stream chain: 1 waits on 0, 2 waits on 1. */ + CHECK(cudaStreamWaitEvent(st[1], ev[0], 0)); + vadd_into<<<grid, 256, 0, st[1]>>>(d_c[1], d_c[0], n); + CHECK(cudaGetLastError()); + CHECK(cudaEventRecord(ev[1], st[1])); + + CHECK(cudaStreamWaitEvent(st[2], ev[1], 0)); + vadd_into<<<grid, 256, 0, st[2]>>>(d_c[2], d_c[1], n); + CHECK(cudaGetLastError()); + + /* Copies back overlap across the three streams. */ + for (s = 0; s < NSTREAMS; s++) + CHECK(cudaMemcpyAsync(h_c[s], d_c[s], bytes, + cudaMemcpyDeviceToHost, st[s])); + CHECK(cudaEventRecord(t1, st[2])); + + for (s = 0; s < NSTREAMS; s++) + CHECK(cudaStreamSynchronize(st[s])); + CHECK(cudaEventElapsedTime(&ms, t0, t1)); + + /* c0 = a0+b0; c1 = (a1+b1)+c0; c2 = (a2+b2)+c1 (all exact floats). */ + for (i = 0; i < n; i++) { + float c0 = h_a[0][i] + h_b[0][i]; + float c1 = h_a[1][i] + h_b[1][i] + c0; + float c2 = h_a[2][i] + h_b[2][i] + c1; + + if (h_c[0][i] != c0) + bad++; + if (h_c[1][i] != c1) + bad++; + if (h_c[2][i] != c2) + bad++; + } + + printf("STREAMS3 n=%d bad=%d\n", n, bad); + printf("TIMING ms=%.3f\n", ms); + + if (bad != 0) { + printf("STATUS FAIL\n"); + return 1; + } + + printf("STATUS OK\n"); + for (s = 0; s < NSTREAMS; s++) { + CHECK(cudaEventDestroy(ev[s])); + CHECK(cudaStreamDestroy(st[s])); + CHECK(cudaFree(d_a[s])); + CHECK(cudaFree(d_b[s])); + CHECK(cudaFree(d_c[s])); + free(h_a[s]); + free(h_b[s]); + free(h_c[s]); + } + CHECK(cudaEventDestroy(t0)); + CHECK(cudaEventDestroy(t1)); + return 0; +} diff --git a/tools/cuda-ab/victory_strided.cu b/tools/cuda-ab/victory_strided.cu new file mode 100644 index 0000000..2fb8739 --- /dev/null +++ b/tools/cuda-ab/victory_strided.cu @@ -0,0 +1,138 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * victory_strided.cu - strided/offset memory access over 4 MB of floats. + * + * Two kernels permute an n-element array (n = 1M floats, 4 MB) with a + * stride co-prime to n (idx = i*7 mod n), so every element is written exactly + * once and the host can verify every single position: + * + * out1[(i*7) % n] = in[i] + 1.0f (plain permutation) + * out2[(i*7) % n] = in[off + i] * 2.0f (offset base pointer) + * + * All values are multiples of 0.25, so both operations are exact in float + * and the host comparison is exact. + * + * Output lines are stable and greppable: + * + * DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes> + * STRIDED n=<count> stride=7 off=<offset> bad=<mismatches> csum=<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 -o victory_strided victory_strided.cu + * + * Then, from the guest: /tmp/victory_strided + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#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 STRIDE 7 +#define OFFSET 4097 + +__global__ void permute_add(const float *in, float *out, int n, int stride) +{ + long long i = (long long)blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) { + long long idx = (i * (long long)stride) % (long long)n; + + out[idx] = in[i] + 1.0f; + } +} + +__global__ void permute_offset(const float *in, float *out, int n, int off, + int stride) +{ + long long i = (long long)blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) { + long long idx = (i * (long long)stride) % (long long)n; + + out[idx] = in[(long long)off + i] * 2.0f; + } +} + +int main(void) +{ + cudaDeviceProp prop; + const int n = 1 << 20; /* 1M floats = 4 MB */ + const long long total = (long long)n + OFFSET; + size_t bytes = (size_t)n * sizeof(float); + size_t total_bytes = (size_t)total * sizeof(float); + float *h_in, *h_out1, *h_out2; + float *d_in, *d_out1, *d_out2; + double csum = 0.0; + int bad = 0, i; + + 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); + + h_in = (float *)malloc(total_bytes); + h_out1 = (float *)malloc(bytes); + h_out2 = (float *)malloc(bytes); + if (!h_in || !h_out1 || !h_out2) + return 1; + + for (i = 0; i < (int)total; i++) + h_in[i] = (float)(i % 1000) * 0.25f; + + CHECK(cudaMalloc(&d_in, total_bytes)); + CHECK(cudaMalloc(&d_out1, bytes)); + CHECK(cudaMalloc(&d_out2, bytes)); + CHECK(cudaMemcpy(d_in, h_in, total_bytes, cudaMemcpyHostToDevice)); + + permute_add<<<(n + 255) / 256, 256>>>(d_in, d_out1, n, STRIDE); + permute_offset<<<(n + 255) / 256, 256>>>(d_in, d_out2, n, OFFSET, + STRIDE); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_out1, d_out1, bytes, cudaMemcpyDeviceToHost)); + CHECK(cudaMemcpy(h_out2, d_out2, bytes, cudaMemcpyDeviceToHost)); + + for (i = 0; i < n; i++) { + long long idx = ((long long)i * STRIDE) % (long long)n; + + if (h_out1[idx] != h_in[i] + 1.0f) + bad++; + if (h_out2[idx] != h_in[OFFSET + i] * 2.0f) + bad++; + csum += (double)h_out1[idx]; + } + + printf("STRIDED n=%d stride=%d off=%d bad=%d csum=%.6f\n", n, STRIDE, + OFFSET, bad, csum); + + if (bad != 0) { + printf("STATUS FAIL\n"); + return 1; + } + + printf("STATUS OK\n"); + CHECK(cudaFree(d_in)); + CHECK(cudaFree(d_out1)); + CHECK(cudaFree(d_out2)); + free(h_in); + free(h_out1); + free(h_out2); + return 0; +} |