# Why has beginner's experiment to demonstrate L1 cache incoherency failed.

**URL:** <https://forums.developer.nvidia.com/t/why-has-beginners-experiment-to-demonstrate-l1-cache-incoherency-failed/37318>\
**Category:** CUDA Programming and Performance\
**Created:** [March 23, 2015, 11:57am UTC](https://forums.developer.nvidia.com/t/why-has-beginners-experiment-to-demonstrate-l1-cache-incoherency-failed/37318 "2015-03-23T11:57:45Z")\
**Posts on this page:** 5\
**Page:** 1

<div class="post-metadata">

**Author:** ![RobertE](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@RobertE](https://forums.developer.nvidia.com/u/RobertE)\
**Post date:** [March 23, 2015, 11:57am UTC](https://forums.developer.nvidia.com/t/why-has-beginners-experiment-to-demonstrate-l1-cache-incoherency-failed/37318/1 "2015-03-23T11:57:45Z")

</div>

L1 Cache coherency of CUDA blocks

Device global memory is cached in the same L1 cache for all threads in a block, but we are warned that the L1 caches for different blocks can become incoherent.

Putting this together with the fact that data is transferred in 128 Byte “chunks”, I wondered what would  
happen if two blocks wrote to different parts of the same chunk of global memory. I expected that  
the two L1 caches would become incoherent and the results eventually written to RAM would be indeterminate.

However, my first experiment shows no evidence of this behaviour:

```auto
////////////////////// EXPERIMENTS KERNELS START ///////////////////////////////

//assume dimBlock(WARPSIZE, 4), dimGrid(32)
__global__ void Experiment_On_4KBytes_Kernel_A(uint8_t * __restrict__ pBytes)
{
	unsigned offset = 128 * threadIdx.x + 4 * blockIdx.x + threadIdx.y;
	*(pBytes + offset) = offset % 43; //prime number hash
}

//assume 1 thread
__global__ void Experiment_On_4KBytes_Kernel_B(uint8_t * __restrict__ pBytes)
{
	unsigned errs = 0;

	for (unsigned chunk = 0; chunk < 32; ++chunk)
		for (unsigned block = 0; block < 32; ++block)
			for (unsigned word = 0; word < 4; ++word)
			{
				unsigned offset = 128 * chunk + 4 * block + word;
				errs += ( (unsigned) pBytes[offset] != offset % 43);
			}

	printf("EXPERIMENT RESULT: %u\n",errs); //the result was 0, even though 32 different blocks wrote in each chunk of 128 bytes! 
}

////////////////////// EXPERIMENTS KERNELS END ///////////////////////////////

.
.
.
////////////////////// EXPERIMENTS START ///////////////////////////////
	{
		dim3 dimBlock(WARPSIZE, 4), dimGrid(32);

		Experiment_On_4KBytes_Kernel_A<<<dimGrid, dimBlock, 0, stream>>>(pBytestream);

		Experiment_On_4KBytes_Kernel_B<<<1, 1, 0, stream>>>(pBytestream);
	}
////////////////////// EXPERIMENTS END ///////////////////////////////

```

This suggests to me that either my experiment is not valid or there is a feature in Invidia GPUs to prevent  
WAW hazards, perhaps using the following idea I found on the internet:

- Write-back big trick: keep track of whether other caches also contain a cached line. If not, a cache has an “exclusive” on the line, and can read and write the line as if it were the only CPU.
 reference: inst.eecs.berkeley.edu/~cs194-6/fa08/ppt/lec10.ppt

Any ideas on the experiment or on the GPU operation would be very helpful.

---

<div class="post-metadata">

**Author:** ![Robert\_Crovella](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/robert_crovella/32/14043_2.png) [@Robert\_Crovella](https://forums.developer.nvidia.com/u/Robert_Crovella)\
**Post date:** [March 23, 2015, 2:39pm UTC](https://forums.developer.nvidia.com/t/why-has-beginners-experiment-to-demonstrate-l1-cache-incoherency-failed/37318/2 "2015-03-23T14:39:02Z")

</div>

What sort of GPU are you running on?

---

<div class="post-metadata">

**Author:** ![RobertE](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@RobertE](https://forums.developer.nvidia.com/u/RobertE)\
**Post date:** [March 24, 2015, 9:54am UTC](https://forums.developer.nvidia.com/t/why-has-beginners-experiment-to-demonstrate-l1-cache-incoherency-failed/37318/3 "2015-03-24T09:54:02Z")

</div>

The gpu is a GeForce GTX 7800 Ti.

---

<div class="post-metadata">

**Author:** ![Robert\_Crovella](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/robert_crovella/32/14043_2.png) [@Robert\_Crovella](https://forums.developer.nvidia.com/u/Robert_Crovella)\
**Post date:** [March 24, 2015, 1:58pm UTC](https://forums.developer.nvidia.com/t/why-has-beginners-experiment-to-demonstrate-l1-cache-incoherency-failed/37318/4 "2015-03-24T13:58:41Z")

</div>

I don’t know what GeForce GTX 7800 Ti is.  
GTX 780 Ti is a Kepler GPU that has L1 disabled (for global loads)

[url][https://docs.nvidia.com/cuda/kepler-tuning-guide/index.html#l1-cache[/url]](https://docs.nvidia.com/cuda/kepler-tuning-guide/index.html#l1-cache%5B/url%5D)

---

<div class="post-metadata">

**Author:** ![RobertE](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@RobertE](https://forums.developer.nvidia.com/u/RobertE)\
**Post date:** [March 25, 2015, 1:23pm UTC](https://forums.developer.nvidia.com/t/why-has-beginners-experiment-to-demonstrate-l1-cache-incoherency-failed/37318/5 "2015-03-25T13:23:20Z")

</div>

Thank you. I did mean 780 Ti, so your posting and the link about L1 being disabled have answered my question.