# changing L1 cache configuration using “cudaFuncSetCacheConfig" not working

**URL:** <https://forums.developer.nvidia.com/t/changing-l1-cache-configuration-using-a-oecudafuncsetcacheconfig-not-working/21097>\
**Category:** CUDA Programming and Performance\
**Created:** [February 17, 2011, 6:43am UTC](https://forums.developer.nvidia.com/t/changing-l1-cache-configuration-using-a-oecudafuncsetcacheconfig-not-working/21097 "2011-02-17T06:43:04Z")\
**Posts on this page:** 7\
**Page:** 1

<div class="post-metadata">

**Author:** ![embeddeduser](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@embeddeduser](https://forums.developer.nvidia.com/u/embeddeduser)\
**Post date:** [February 17, 2011, 6:43am UTC](https://forums.developer.nvidia.com/t/changing-l1-cache-configuration-using-a-oecudafuncsetcacheconfig-not-working/21097/1 "2011-02-17T06:43:04Z")

</div>

Hi,

I am trying to change L1 cache configuration using the Runtime API call “cudaFuncSetCacheConfig" as specified in the Programming Guide.

However, it does not seem to be working. My Dummy code is below

```auto
// Dummy code

//Cuda Device Properties

cudaDeviceProp deviceProp;

__global__ void MyKernel(....) {

....

}

void host_func() {

CUDA_SAFE_CALL_NO_SYNC(cudaGetDeviceProperties(&deviceProp, dev));

// Printing Device properties before changing L1 cache configuration

	printf(" Total amount of global memory: %llu bytes\n", (unsigned long long) deviceProp.totalGlobalMem);

        printf(" Total amount of constant memory: %u bytes\n", deviceProp.totalConstMem);

        printf(" Total amount of shared memory per block: %u bytes\n", deviceProp.sharedMemPerBlock);

// Changing L1 cache configuration to 48KB - L1 and 16KB - Shared Memory

	cudaFuncSetCacheConfig(MyKernel, cudaFuncCachePreferL1);

        printf("/nMSG - Configured L1 cache to 48KB size/n");

// Printing Device properties after changing L1 cache configuration. This should print SM = 16 KB

	printf(" Total amount of global memory: %llu bytes\n", (unsigned long long) deviceProp.totalGlobalMem);

        printf(" Total amount of constant memory: %u bytes\n", deviceProp.totalConstMem);

        printf(" Total amount of shared memory per block: %u bytes\n", deviceProp.sharedMemPerBlock);

}

```

The output of the above code which I get -

```auto
Using device 0: GeForce GTX 480

  Total amount of global memory: 1576468480 bytes

  Total amount of constant memory: 65536 bytes

  Total amount of shared memory per block: 49152 bytes

MSG - Configured L1 cache to 48KB size

Using device 0: GeForce GTX 480

  Total amount of global memory: 1576468480 bytes

  Total amount of constant memory: 65536 bytes

  Total amount of shared memory per block: 49152 bytes

```

I don’t understand where I am going wrong.

Thanks and Regards,

Ninad B Kothari

---

<div class="post-metadata">

**Author:** ![tera](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@tera](https://forums.developer.nvidia.com/u/tera)\
**Post date:** [February 17, 2011, 3:36pm UTC](https://forums.developer.nvidia.com/t/changing-l1-cache-configuration-using-a-oecudafuncsetcacheconfig-not-working/21097/2 "2011-02-17T15:36:08Z")

</div>

At the very minimum you need to call cudaGetDeviceProperties() again to actually see the result of the change…

I wouldn’t however be surprised if the change weren’t immediately reflected before the cache configuration has actually changed on the GPU, or even if it wouldn’t change at all.

---

<div class="post-metadata">

**Author:** ![embeddeduser](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@embeddeduser](https://forums.developer.nvidia.com/u/embeddeduser)\
**Post date:** [February 17, 2011, 5:31pm UTC](https://forums.developer.nvidia.com/t/changing-l1-cache-configuration-using-a-oecudafuncsetcacheconfig-not-working/21097/3 "2011-02-17T17:31:28Z")

</div>

> [@](#):
>
> At the very minimum you need to call cudaGetDeviceProperties() again to actually see the result of the change…
> 
> I wouldn’t however be surprised if the change weren’t immediately reflected before the cache configuration has actually changed on the GPU, or even if it wouldn’t change at all.

Hi tera you are right … and I apologize for not posting my dummy code correctly. Let me paste the correct code below

```auto
// Dummy code

//Cuda Device Properties

cudaDeviceProp deviceProp;

__global__ void MyKernel(....) {

....

}

void host_func() {

queryDevice();

// Changing L1 cache configuration to 48KB - L1 and 16KB - Shared Memory

        cudaFuncSetCacheConfig(MyKernel, cudaFuncCachePreferL1);

        printf("/nMSG - Configured L1 cache to 48KB size/n");

// Printing Device properties after changing L1 cache configuration. This should print SM = 16 KB

queryDevice();

// Call Kernel

dim3 dimBlock(16, 16); // just an example

dim3 dimGrid(4, 4); // just an example

MyKernel<<<dimGrid, dimBlock>>>();

printf("/nMSG - Check if L1 configuration has been taken after running kernel/n");

// querying Device to check if L1 configuration has been taken

queryDevice();

}

void queryDevice() {

int dev = 0;

CUDA_SAFE_CALL_NO_SYNC(cudaGetDeviceProperties(&deviceProp, dev));

printf(" Total amount of global memory: %llu bytes\n", (unsigned long long) deviceProp.totalGlobalMem);

        printf(" Total amount of constant memory: %u bytes\n", deviceProp.totalConstMem);

        printf(" Total amount of shared memory per block: %u bytes\n", deviceProp.sharedMemPerBlock);

}

```

The results are below … please excuse any typos

```auto
Using device 0: GeForce GTX 480

  Total amount of global memory: 1576468480 bytes

  Total amount of constant memory: 65536 bytes

  Total amount of shared memory per block: 49152 bytes

MSG - Configured L1 cache to 48KB size

Using device 0: GeForce GTX 480

  Total amount of global memory: 1576468480 bytes

  Total amount of constant memory: 65536 bytes

  Total amount of shared memory per block: 49152 bytes

MSG - Check if L1 configuration has been taken after running kernel

Using device 0: GeForce GTX 480

  Total amount of global memory: 1576468480 bytes

  Total amount of constant memory: 65536 bytes

  Total amount of shared memory per block: 49152 bytes

```

However, in all three cases it does not show the changed L1 configuration. Am I still missing something? Could it be that my GTX 480 does not support this? How can I be sure?

Thanks and Regards,

embeddeduser

---

<div class="post-metadata">

**Author:** ![njuffa](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@njuffa](https://forums.developer.nvidia.com/u/njuffa)\
**Post date:** [February 17, 2011, 8:45pm UTC](https://forums.developer.nvidia.com/t/changing-l1-cache-configuration-using-a-oecudafuncsetcacheconfig-not-working/21097/4 "2011-02-17T20:45:48Z")

</div>

cudaFuncSetCacheConfig() configures the L1 / shared memory split for a particular **global** function, i.e. a compute kernel. In other words, it configures a per-function property that can override a per-device property. Note that the request for large L1 may be ignored by the driver if it would prevent the kernel from executing due to lack of shared memory.

The online documentation for cudaFuncSetCacheConfig() states (emphasis mine):

“On devices where the L1 cache and shared memory use the same hardware resources, this sets through cacheConfig the preferred cache configuration for the function specified via func. This is only a preference. The runtime will use the requested configuration if possible, but it is free to choose a different configuration if required to execute func.”

[later]

The online documentation for cudaGetDeviceProperties() explains that the property sharedMemPerBlock refers to the maximum amount available, meaning this value is fixed for a given device (emphasis again mine):

“sharedMemPerBlock is the maximum amount of shared memory available to a thread block in bytes; this amount is shared by all thread blocks simultaneously resident on a multiprocessor;”

---

<div class="post-metadata">

**Author:** ![Kiran\_CUDA](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Kiran\_CUDA](https://forums.developer.nvidia.com/u/Kiran_CUDA)\
**Post date:** [January 30, 2012, 6:52am UTC](https://forums.developer.nvidia.com/t/changing-l1-cache-configuration-using-a-oecudafuncsetcacheconfig-not-working/21097/5 "2012-01-30T06:52:00Z")

</div>

> [@](#):
>
> Hi tera you are right … and I apologize for not posting my dummy code correctly. Let me paste the correct code below
> 
> ```auto
> // Dummy code
> 
> //Cuda Device Properties
> 
> cudaDeviceProp deviceProp;
> 
> __global__ void MyKernel(....) {
> 
> ....
> 
> }
> 
> void host_func() {
> 
> queryDevice();
> 
> // Changing L1 cache configuration to 48KB - L1 and 16KB - Shared Memory
> 
> cudaFuncSetCacheConfig(MyKernel, cudaFuncCachePreferL1);
> 
> printf("/nMSG - Configured L1 cache to 48KB size/n");
> 
> // Printing Device properties after changing L1 cache configuration. This should print SM = 16 KB
> 
> queryDevice();
> 
> // Call Kernel
> 
> dim3 dimBlock(16, 16); // just an example
> 
> dim3 dimGrid(4, 4); // just an example
> 
> MyKernel<<<dimGrid, dimBlock>>>();
> 
> printf("/nMSG - Check if L1 configuration has been taken after running kernel/n");
> 
> // querying Device to check if L1 configuration has been taken
> 
> queryDevice();
> 
> }
> 
> void queryDevice() {
> 
> int dev = 0;
> 
> CUDA_SAFE_CALL_NO_SYNC(cudaGetDeviceProperties(&deviceProp, dev));
> 
> printf(" Total amount of global memory: %llu bytes\n", (unsigned long long) deviceProp.totalGlobalMem);
> 
> printf(" Total amount of constant memory: %u bytes\n", deviceProp.totalConstMem);
> 
> printf(" Total amount of shared memory per block: %u bytes\n", deviceProp.sharedMemPerBlock);
> 
> }
> 
> ```
> 
> The results are below … please excuse any typos
> 
> ```auto
> Using device 0: GeForce GTX 480
> 
> Total amount of global memory: 1576468480 bytes
> 
> Total amount of constant memory: 65536 bytes
> 
> Total amount of shared memory per block: 49152 bytes
> 
> MSG - Configured L1 cache to 48KB size
> 
> Using device 0: GeForce GTX 480
> 
> Total amount of global memory: 1576468480 bytes
> 
> Total amount of constant memory: 65536 bytes
> 
> Total amount of shared memory per block: 49152 bytes
> 
> MSG - Check if L1 configuration has been taken after running kernel
> 
> Using device 0: GeForce GTX 480
> 
> Total amount of global memory: 1576468480 bytes
> 
> Total amount of constant memory: 65536 bytes
> 
> Total amount of shared memory per block: 49152 bytes
> 
> ```
> 
> However, in all three cases it does not show the changed L1 configuration. Am I still missing something? Could it be that my GTX 480 does not support this? How can I be sure?
> 
> Thanks and Regards,
> 
> embeddeduser

Any update on this issue? I am also facing the same issue

---

<div class="post-metadata">

**Author:** ![tera](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@tera](https://forums.developer.nvidia.com/u/tera)\
**Post date:** [January 30, 2012, 11:48am UTC](https://forums.developer.nvidia.com/t/changing-l1-cache-configuration-using-a-oecudafuncsetcacheconfig-not-working/21097/6 "2012-01-30T11:48:43Z")

</div>

[Yes](http://forums.nvidia.com/index.php?showtopic=171878&#entry1076542). cudaGetDeviceProperties() only reflects the maximum cache available per SM, not the current setting.

---

<div class="post-metadata">

**Author:** ![DrAnderson42](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@DrAnderson42](https://forums.developer.nvidia.com/u/DrAnderson42)\
**Post date:** [February 3, 2012, 1:04pm UTC](https://forums.developer.nvidia.com/t/changing-l1-cache-configuration-using-a-oecudafuncsetcacheconfig-not-working/21097/7 "2012-02-03T13:04:50Z")

</div>

Execute your application in the nvvp profiler. It reports the requested and actual cache configuration of every kernel launch.