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