# Reconfiguring the cache / shared memory on a Fermi understanding the cudaFuncSetCacheConfig command

**URL:** <https://forums.developer.nvidia.com/t/reconfiguring-the-cache-shared-memory-on-a-fermi-understanding-the-cudafuncsetcacheconfig-command/16425>\
**Category:** CUDA Programming and Performance\
**Created:** [May 3, 2010, 8:24pm UTC](https://forums.developer.nvidia.com/t/reconfiguring-the-cache-shared-memory-on-a-fermi-understanding-the-cudafuncsetcacheconfig-command/16425 "2010-05-03T20:24:15Z")\
**Posts on this page:** 20\
**Page:** 1

<div class="post-metadata">

**Author:** ![Maddy\_Scientist](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Maddy\_Scientist](https://forums.developer.nvidia.com/u/Maddy_Scientist)\
**Post date:** [May 3, 2010, 8:24pm UTC](https://forums.developer.nvidia.com/t/reconfiguring-the-cache-shared-memory-on-a-fermi-understanding-the-cudafuncsetcacheconfig-command/16425/1 "2010-05-03T20:24:15Z")

</div>

Hi,

I’ve just installed my GTX 480 and trying to benchmark it. The application I’m using required about 124 registers on the GT200 generation in double precision. Since NVIDIA decided to limit the number registers of Fermi to only 64 registers per thread, the end result is that Fermi is slower than GT200 at double precision for my code :-(

This is using the default cache configuration of 48k / 16k for shared / L1. Switching to 16k / 48 k may help me since the registers will be spilled to the large L1, instead of out to L2 or device memory. Looking at the programming guide, to do this one uses the cudaFuncSetCacheConfig command. From the programming guide:

[codebox]// Device code

**global** void MyKernel() { … }

// Host code

// Runtime API

// cudaFuncCachePreferShared: shared memory is 48 KB

// cudaFuncCachePreferL1: shared memory is 16 KB

// cudaFuncCachePreferNone: no preference

cudaFuncSetCacheConfig(MyKernel, cudaFuncCachePreferShared)

// Driver API

// CU\_FUNC\_CACHE\_PREFER\_SHARED: shared memory is 48 KB

// CU\_FUNC\_CACHE\_PREFER\_L1: shared memory is 16 KB

// CU\_FUNC\_CACHE\_PREFER\_NONE: no preference

CUfunction myKernel;

cuFuncSetCacheConfig(myKernel, CU\_FUNC\_CACHE\_PREFER\_SHARED)[/codebox]

However, this does not compile, and instead I have to surround MyKernel by quotes, i.e., “MyKernel”, to get it to compile. When running though, I get the following error:

[codebox]

error: no instance of overloaded function “cudaFuncSetCacheConfig” matches the argument list argument types are: (, cudaFuncCache)[/codebox]

Can someone point out what I’m doing wrong here?

Thanks.

---

<div class="post-metadata">

**Author:** ![Maddy\_Scientist](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Maddy\_Scientist](https://forums.developer.nvidia.com/u/Maddy_Scientist)\
**Post date:** [May 5, 2010, 3:35pm UTC](https://forums.developer.nvidia.com/t/reconfiguring-the-cache-shared-memory-on-a-fermi-understanding-the-cudafuncsetcacheconfig-command/16425/2 "2010-05-05T15:35:38Z")

</div>

I take it no-one no the forums used this feature then? I thought the re-configurable cache was one of the selling points of the Fermi, and so knowing how to successfully do this is something we should know.

Thanks.

---

<div class="post-metadata">

**Author:** ![Lev](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Lev](https://forums.developer.nvidia.com/u/Lev)\
**Post date:** [May 5, 2010, 3:50pm UTC](https://forums.developer.nvidia.com/t/reconfiguring-the-cache-shared-memory-on-a-fermi-understanding-the-cudafuncsetcacheconfig-command/16425/3 "2010-05-05T15:50:33Z")

</div>

Do you use driver API?

---

<div class="post-metadata">

**Author:** ![MisterAnderson42](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@MisterAnderson42](https://forums.developer.nvidia.com/u/MisterAnderson42)\
**Post date:** [May 5, 2010, 6:00pm UTC](https://forums.developer.nvidia.com/t/reconfiguring-the-cache-shared-memory-on-a-fermi-understanding-the-cudafuncsetcacheconfig-command/16425/4 "2010-05-05T18:00:48Z")

</div>

> [@](#):
>
> I take it no-one no the forums used this feature then? I thought the re-configurable cache was one of the selling points of the Fermi, and so knowing how to successfully do this is something we should know.

I’ve played with it a little bit, but haven’t had the time to fully explore the new cache and update my applications yet. The problem with having a dozen things to work on all at once…

I also don’t get the error that you do:

```auto
# test.cu

__global__ void MyKernel()

	{

	}

int main()

	{

	cudaFuncSetCacheConfig(MyKernel, cudaFuncCachePreferShared);

	return 0;

	}

```

```auto
$ nvcc -c test.cu

$

```

This is with CUDA 3.0 on an iMac, but I’ve done the same with CUDA 3.0 in linux at home. The setting defiantly works, the few tests I have run do have performance differences with different cache preference settings. It even works with templated kernels, you just have to list the fully template instantiation in the call to cudaFuncSetCacheConfig, like this:

```auto
cudaFuncSetCacheConfig(MyKernel<int, float, 5>, cudaFuncCachePreferShared);

```

---

<div class="post-metadata">

**Author:** ![Lev](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Lev](https://forums.developer.nvidia.com/u/Lev)\
**Post date:** [May 5, 2010, 6:56pm UTC](https://forums.developer.nvidia.com/t/reconfiguring-the-cache-shared-memory-on-a-fermi-understanding-the-cudafuncsetcacheconfig-command/16425/5 "2010-05-05T18:56:37Z")

</div>

Strange, this function is declared with (const char\*) parameter.

---

<div class="post-metadata">

**Author:** ![Maddy\_Scientist](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Maddy\_Scientist](https://forums.developer.nvidia.com/u/Maddy_Scientist)\
**Post date:** [May 5, 2010, 7:08pm UTC](https://forums.developer.nvidia.com/t/reconfiguring-the-cache-shared-memory-on-a-fermi-understanding-the-cudafuncsetcacheconfig-command/16425/6 "2010-05-05T19:08:47Z")

</div>

> [@](#):
>
> This is with CUDA 3.0 on an iMac, but I’ve done the same with CUDA 3.0 in linux at home. The setting defiantly works, the few tests I have run do have performance differences with different cache preference settings. It even works with templated kernels, you just have to list the fully template instantiation in the call to cudaFuncSetCacheConfig, like this:
> 
> ```auto
> cudaFuncSetCacheConfig(MyKernel<int, float, 5>, cudaFuncCachePreferShared);
> 
> ```

The kernel for which I am trying to do this uses polymorphism but not templates. I wonder if the polymorphism is causing my problems. This is unlikely though, I would have assumed that when using polymorphism, cache settings would be set uniformly between all kernels with a common name, but that you would still have the ability to change the cache settings.

---

<div class="post-metadata">

**Author:** ![MisterAnderson42](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@MisterAnderson42](https://forums.developer.nvidia.com/u/MisterAnderson42)\
**Post date:** [May 5, 2010, 7:37pm UTC](https://forums.developer.nvidia.com/t/reconfiguring-the-cache-shared-memory-on-a-fermi-understanding-the-cudafuncsetcacheconfig-command/16425/7 "2010-05-05T19:37:12Z")

</div>

> [@](#):
>
> The kernel for which I am trying to do this uses polymorphism but not templates. I wonder if the polymorphism is causing my problems. This is unlikely though, I would have assumed that when using polymorphism, cache settings would be set uniformly between all kernels with a common name, but that you would still have the ability to change the cache settings.

The documentation for the const char\* version states that you need to specify the fully mangled c++ name of the function you wish to change the setting for. That indicates you would not get the behavior you wish for.

I have no idea what magic the templated version that takes the function pointer does… Perhaps this is the cause of your difficulty. If you have two different functions with the same name, perhaps the function pointer cannot be resolved to either one of them.

---

<div class="post-metadata">

**Author:** ![Sarnath](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/sarnath/32/9742_2.png) [@Sarnath](https://forums.developer.nvidia.com/u/Sarnath)\
**Post date:** [May 5, 2010, 10:44pm UTC](https://forums.developer.nvidia.com/t/reconfiguring-the-cache-shared-memory-on-a-fermi-understanding-the-cudafuncsetcacheconfig-command/16425/8 "2010-05-05T22:44:53Z")

</div>

16384/124 ~= 132 – hmm… may not b gr8  
8192/124 ~= 64 – not so gr8

---

<div class="post-metadata">

**Author:** ![Maddy\_Scientist](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Maddy\_Scientist](https://forums.developer.nvidia.com/u/Maddy_Scientist)\
**Post date:** [May 5, 2010, 10:55pm UTC](https://forums.developer.nvidia.com/t/reconfiguring-the-cache-shared-memory-on-a-fermi-understanding-the-cudafuncsetcacheconfig-command/16425/9 "2010-05-05T22:55:12Z")

</div>

I haven’t studied all the Fermi documentation, so the answer the following may be in there. Do you know if there is a way to reconfigure the cache for all kernels universally, rather than on a kernel by kernel basis? I would imagine I could set a default flag upon device initialisation that would be set for all kernels, unless cudaFuncSetCacheConfig was set. This seem like my best hope at the moment, given the trouble I’m having with my polymorphic kernels.

---

<div class="post-metadata">

**Author:** ![MisterAnderson42](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@MisterAnderson42](https://forums.developer.nvidia.com/u/MisterAnderson42)\
**Post date:** [May 6, 2010, 11:59am UTC](https://forums.developer.nvidia.com/t/reconfiguring-the-cache-shared-memory-on-a-fermi-understanding-the-cudafuncsetcacheconfig-command/16425/10 "2010-05-06T11:59:10Z")

</div>

Tim stated in another thread that if you only set a single kernel to prefer L1, that setting will stick permanently. You could create a dummy empty kernel, set it to L1, and call it at the beginning of your program.

---

<div class="post-metadata">

**Author:** ![Lev](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Lev](https://forums.developer.nvidia.com/u/Lev)\
**Post date:** [May 6, 2010, 12:03pm UTC](https://forums.developer.nvidia.com/t/reconfiguring-the-cache-shared-memory-on-a-fermi-understanding-the-cudafuncsetcacheconfig-command/16425/11 "2010-05-06T12:03:54Z")

</div>

Btw, can I obtain current status of configuration?

---

<div class="post-metadata">

**Author:** ![Maddy\_Scientist](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Maddy\_Scientist](https://forums.developer.nvidia.com/u/Maddy_Scientist)\
**Post date:** [May 7, 2010, 4:06am UTC](https://forums.developer.nvidia.com/t/reconfiguring-the-cache-shared-memory-on-a-fermi-understanding-the-cudafuncsetcacheconfig-command/16425/12 "2010-05-07T04:06:49Z")

</div>

> [@](#):
>
> Tim stated in another thread that if you only set a single kernel to prefer L1, that setting will stick permanently. You could create a dummy empty kernel, set it to L1, and call it at the beginning of your program.

Thanks for the suggestion. This works. It hasn’t helped my performance woes, but that’s another story…

---

<div class="post-metadata">

**Author:** ![dmsmall](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@dmsmall](https://forums.developer.nvidia.com/u/dmsmall)\
**Post date:** [May 7, 2010, 7:44pm UTC](https://forums.developer.nvidia.com/t/reconfiguring-the-cache-shared-memory-on-a-fermi-understanding-the-cudafuncsetcacheconfig-command/16425/13 "2010-05-07T19:44:26Z")

</div>

Are you expecting to get better double precision performance with fermi? Apparently, Nvidia has chosen not to give us the extra double precision performance on GeForce products but only on the Tesla products. There is another thread floating around that talks about this. Wanted to make sure you knew before spending too much time trying to get your 480 to run better than your old 280.

---

<div class="post-metadata">

**Author:** ![SPWorley](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@SPWorley](https://forums.developer.nvidia.com/u/SPWorley)\
**Post date:** [May 8, 2010, 12:36am UTC](https://forums.developer.nvidia.com/t/reconfiguring-the-cache-shared-memory-on-a-fermi-understanding-the-cudafuncsetcacheconfig-command/16425/14 "2010-05-08T00:36:48Z")

</div>

> [@](#):
>
> Are you expecting to get better double precision performance with fermi? Apparently, Nvidia has chosen not to give us the extra double precision performance on GeForce products but only on the Tesla products. There is another thread floating around that talks about this. Wanted to make sure you knew before spending too much time trying to get your 480 to run better than your old 280.

The GTX480 has massively better DP throughput than the GTX285… almost double.

GTX285: 88.5 Gflops

GTX480: 168 Gflops

S2050: 515 Gflops

---

<div class="post-metadata">

**Author:** ![gshi](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@gshi](https://forums.developer.nvidia.com/u/gshi)\
**Post date:** [May 28, 2010, 2:51am UTC](https://forums.developer.nvidia.com/t/reconfiguring-the-cache-shared-memory-on-a-fermi-understanding-the-cudafuncsetcacheconfig-command/16425/15 "2010-05-28T02:51:38Z")

</div>

> [@](#):
>
> Thanks for the suggestion. This works. It hasn’t helped my performance woes, but that’s another story…

It improved the performance for double precision in my case.

The key is you need to run the dummy kernel after you set the cache preference.

---

<div class="post-metadata">

**Author:** ![wanderine](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@wanderine](https://forums.developer.nvidia.com/u/wanderine)\
**Post date:** [June 4, 2010, 2:06pm UTC](https://forums.developer.nvidia.com/t/reconfiguring-the-cache-shared-memory-on-a-fermi-understanding-the-cudafuncsetcacheconfig-command/16425/16 "2010-06-04T14:06:30Z")

</div>

I only get this error

error: identifier “cudaFuncCachePreferShared” is undefined

error: identifier “cudaFuncSetCacheConfig” is undefined

?

---

<div class="post-metadata">

**Author:** ![Deus](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Deus](https://forums.developer.nvidia.com/u/Deus)\
**Post date:** [June 7, 2010, 11:59am UTC](https://forums.developer.nvidia.com/t/reconfiguring-the-cache-shared-memory-on-a-fermi-understanding-the-cudafuncsetcacheconfig-command/16425/17 "2010-06-07T11:59:09Z")

</div>

I have the same Problem with cudaFuncCachePreferL1.

I define a kernel

```auto
__global__ void dummy(){}

```

and set

```auto
cudaFuncSetCacheConfig("dummy", cudaFuncCachePreferL1);

```

but the compiler return an error

```auto
cudaErrorInvalidDeviceFunction

```

What is wrong on this code?

---

<div class="post-metadata">

**Author:** ![alexish](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@alexish](https://forums.developer.nvidia.com/u/alexish)\
**Post date:** [June 7, 2010, 12:56pm UTC](https://forums.developer.nvidia.com/t/reconfiguring-the-cache-shared-memory-on-a-fermi-understanding-the-cudafuncsetcacheconfig-command/16425/18 "2010-06-07T12:56:23Z")

</div>

> [@](#):
>
> I have the same Problem with cudaFuncCachePreferL1.
> 
> I define a kernel
> 
> ```auto
> __global__ void dummy(){}
> 
> ```
> 
> and set
> 
> ```auto
> cudaFuncSetCacheConfig("dummy", cudaFuncCachePreferL1);
> 
> ```
> 
> but the compiler return an error
> 
> ```auto
> cudaErrorInvalidDeviceFunction
> 
> ```
> 
> What is wrong on this code?

I have done some experiments with cache configuration and found that two syntax worked:

- use the complete mangled name of your kernel as a C string

=\> you have something like that cudaFuncSetCacheConfig(“\_Z11\_dummy…”, cudaFuncCachePreferL1);

- use directly the name of the kernel, but not as a C string

=\> cudaFuncSetCacheConfig(dummy, cudaFuncCachePreferL1);

I also confirm that the cache configuration is persistent.

---

<div class="post-metadata">

**Author:** ![Deus](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Deus](https://forums.developer.nvidia.com/u/Deus)\
**Post date:** [June 7, 2010, 1:07pm UTC](https://forums.developer.nvidia.com/t/reconfiguring-the-cache-shared-memory-on-a-fermi-understanding-the-cudafuncsetcacheconfig-command/16425/19 "2010-06-07T13:07:57Z")

</div>

Thank you for the quick response! Your solution work fine…

---

<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:** [June 7, 2010, 4:35pm UTC](https://forums.developer.nvidia.com/t/reconfiguring-the-cache-shared-memory-on-a-fermi-understanding-the-cudafuncsetcacheconfig-command/16425/20 "2010-06-07T16:35:46Z")

</div>

> [@](#):
>
> The application I’m using required about 124 registers on the GT200 generation in double precision. Since NVIDIA decided to limit the number registers of Fermi to only 64 registers per thread, the end result is that Fermi is slower than GT200 at double precision for my code :-(
> 
> This is using the default cache configuration of 48k / 16k for shared / L1. Switching to 16k / 48 k may help me since the registers will be spilled to the large L1, instead of out to L2 or device memory.

Instead of reconfiguring shared memory to become L1 cache and catch more of the spilled vars, maybe you can declare some of your local vars as (tid-indexed arrays in) shared memory to explicitly place them in the on-chip memory.