simpleCopyKernel fails to advance the pointers.

[Artem-B]

nvbandwidth/kernels.cu

Lines 20 to 28 in 4a49bda

	__global__ void simpleCopyKernel(unsigned long long loopCount, uint4 *dst, uint4 *src) {
	for (unsigned int i = 0; i < loopCount; i++) {
	const int idx = blockIdx.x * blockDim.x + threadIdx.x;
	size_t offset = idx * sizeof(uint4);
	uint4* dst_uint4 = reinterpret_cast<uint4*>((char*)dst + offset);
	uint4* src_uint4 = reinterpret_cast<uint4*>((char*)src + offset);
	__stcg(dst_uint4, __ldcg(src_uint4));
	}
	}

It appears that offset or pointers never advance during loop iterations, and we always load/store from/to the same locations. It does not do as much copying as it's intended to.

https://godbolt.org/z/rM3dP5GsE

[Artem-B]

It may be a false alarm. Looks like the caller slices the data to give each thread 16 bytes, so the single copy is intentional, and the calculation of idx/offset inside the body loop is just misleading -- it makes things look as if we intended to make offset depend on i.