cancel
Showing results for 
Search instead for 
Did you mean: 

Archives Discussions

bubu
Adept II

better control of the cache

I think it would be good to add a mechanism to control better the memory caches.

 

For instance, you could add a flag to the clCreateBuffer like CL_CACHED.

clCreateBuffer ( .....,CL_MEM_READ|CL_CACHED  ... )

 

Other option is to add a modifier keyword to the __global one:

__kernel void MyKernel ( __global __cached int *w )

{

   const int data = w[get_global_id(0)];

}

 

Sometimes I want the global memory cached but other times not.

thx.

0 Likes
3 Replies

bubu,
This will be fixed in the next release. Please use const restrict on your pointers to enable caching.
0 Likes

Originally posted by: MicahVillmow bubu,  Please use const restrict on your pointers to enable caching.


Btw, is write caching possible?

0 Likes

bubu,

Read & Write paths are different in ATI GPUs. So AFAIK write caching is not possible. Even for read  data which does not enforce coherency(const data) inside various cache levels is cached in read. Otherwise no caching is done. 

0 Likes