L2 not writing back to device memory

[Vectorizer]

To familiarize myself with Nsight Compute, I am incrementally writing a matrix transpose kernel, first step being a “copy kernel” to establish some global memory throughput numbers.

I do verify that the kernel output is valid, but Nsight Compute seems to be missing the write back from L2 to Global memory.

__global__ void copyKernel(float* __restrict__ dataIn, float* __restrict__ dataOut, uint32_t xDimension, uint32_t yDimension) {
	__shared__ float cache[TILE_DIMENSION_X][TILE_DIMENSION_Y];
#pragma unroll 
	for(uint32_t yIndex{ (blockIdx.y * blockDim.y) + threadIdx.y }; yIndex < yDimension; yIndex += (gridDim.y * blockDim.y))
		for(uint32_t xIndex{ (blockIdx.x * blockDim.x) + threadIdx.x }; xIndex < xDimension; xIndex += (gridDim.x * blockDim.x)) {
			const uint32_t index{ xIndex + (xDimension * yIndex) };
			cache[threadIdx.y][threadIdx.x] = dataIn[index];
			dataOut[index] = cache[threadIdx.y][threadIdx.x];
		}
}

[Robert_Crovella]

On these forums, please do not post pictures of code. Please post code as properly formatted text. (Thanks for the edit.)

My first guess would be that the L2, being a write-back cache, found no need to write data to device memory while your kernel was running because the 4MB of data you are writing fit in the cache.

It’s hard to confirm or refute this without information such as the total extent of what you are doing (not just the kernel code) and also the device you are running on. You could try running the test with 4GB copy instead of 4MB copy.

For example, my L4 GPU has a peak bandwidth of around 300GB/s, so your reported 247GB/s at ~88% of peak might be consistent with that. The L4 GPU has a L2 cache that is considerably larger than 8MB, so it would not trigger write backs for 4MB of data.