# APIs for splitting shared memory and L1 cache

**URL:** <https://forums.developer.nvidia.com/t/apis-for-splitting-shared-memory-and-l1-cache/294291>\
**Category:** CUDA Programming and Performance\
**Created:** [May 28, 2024, 6:08am UTC](https://forums.developer.nvidia.com/t/apis-for-splitting-shared-memory-and-l1-cache/294291 "2024-05-28T06:08:44Z")\
**Posts on this page:** 4\
**Page:** 1

<div class="post-metadata">

**Author:** ![qwerty00](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@qwerty00](https://forums.developer.nvidia.com/u/qwerty00)\
**Post date:** [May 28, 2024, 6:08am UTC](https://forums.developer.nvidia.com/t/apis-for-splitting-shared-memory-and-l1-cache/294291/1 "2024-05-28T06:08:44Z")

</div>

Hi, I recently noticed that there are 3 APIs for the partition of shared memory and L1:

```auto
__global__ void ptr_kernel() {
    //...
}

cudaFuncSetAttribute(ptr_kernel, cudaFuncAttributePreferredSharedMemoryCarveout, cudaSharedmemCarveoutMaxShared);
cudaFuncSetCacheConfig(ptr_kernel, cudaFuncCachePreferShared);
cudaDeviceSetCacheConfig(cudaFuncCachePreferShared);

```

I know that `cudaDeviceSetCacheConfig` is used for device level partition, while the other two are used for a specific function that could override the global cache setting when being launched. But it seems that the `cudaFuncSetAttribute` and `cudaFuncSetCacheConfig` have the same functionality.

Another question is: does these APIs imply device synchronization, i.e they need to wait for all kernels done before changing the cache setting?

---

<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:** [May 28, 2024, 2:07pm UTC](https://forums.developer.nvidia.com/t/apis-for-splitting-shared-memory-and-l1-cache/294291/2 "2024-05-28T14:07:21Z")

</div>

> [@qwerty00](#):
>
> ```auto
> cudaFuncSetCacheConfig(ptr_kernel, cudaFuncCachePreferShared);
> cudaDeviceSetCacheConfig(cudaFuncCachePreferShared);
> 
> ```

The `cudaFuncCachePreferShared` is/was [applicable to devices that had a hardware resource that combined](https://docs.nvidia.com/cuda/cuda-runtime-api/group __CUDART__ EXECUTION.html#group __CUDART__ EXECUTION_1g6699ca1943ac2655effa0d571b2f4f15) shared memory and L1 cache. Since at least Maxwell, this has [not been true](https://stackoverflow.com/questions/52305331/how-to-call-cudadevicesetsharedmemconfig-and-cudadevicesetcacheconfig).

Here is my understanding of the progression:

- In eg. Fermi generation, the L1 cache and the shared memory were part of the same hardware block. The texture unit was separate.
- After Fermi, through Pascal, the L1 and Tex units were combined into the same hardware resource, and shared was a separate entity.
- In the Volta timeframe, all 3 units were combined into a single hardware resource: L1/Tex and share.

AFAIK, the cudaFuncCachePreferShared setting was usable in the Fermi generation. I’m not sure it has any applicability any more.

The carveout setting did not exist in the fermi era, and is applicable for volta and beyond.

~~`cudaFuncAttributePreferredSharedMemoryCarveout` has the [same](https://docs.nvidia.com/cuda/cuda-runtime-api/group __CUDART__ EXECUTION.html) applicability.~~

---

<div class="post-metadata">

**Author:** ![qwerty00](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@qwerty00](https://forums.developer.nvidia.com/u/qwerty00)\
**Post date:** [May 29, 2024, 4:48am UTC](https://forums.developer.nvidia.com/t/apis-for-splitting-shared-memory-and-l1-cache/294291/3 "2024-05-29T04:48:03Z")

</div>

Thank you. But in [NVIDIA Ampere GPU Architecture Tuning Guide](https://docs.nvidia.com/cuda/ampere-tuning-guide/index.html) Section 1.4.2.3 it says that

> In the NVIDIA Ampere GPU architecture, the portion of the L1 cache dedicated to shared memory (known as the _carveout_ ) can be selected at runtime as in previous architectures such as Volta, using `cudaFuncSetAttribute()` with the attribute `cudaFuncAttributePreferredSharedMemoryCarveout` . The NVIDIA A100 GPU supports shared memory capacity of 0, 8, 16, 32, 64, 100, 132 or 164 KB per SM. GPUs with compute capability 8.6 support shared memory capacity of 0, 8, 16, 32, 64 or 100 KB per SM.

If we can allocate size for shared memory and L1, do the shared memory and L1 shared the same chip? Or there exist some other hardware mechanisms?

---

<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:** [May 29, 2024, 2:18pm UTC](https://forums.developer.nvidia.com/t/apis-for-splitting-shared-memory-and-l1-cache/294291/4 "2024-05-29T14:18:05Z")

</div>

I’ve modified my previous response. Yes, for volta and beyond, the L1 and shared are combined. This is mentioned in the volta white paper.