# Problem with coalesced memory access

**URL:** <https://forums.developer.nvidia.com/t/problem-with-coalesced-memory-access/4262>\
**Category:** CUDA Programming and Performance\
**Created:** [June 23, 2008, 10:32am UTC](https://forums.developer.nvidia.com/t/problem-with-coalesced-memory-access/4262 "2008-06-23T10:32:17Z")\
**Posts on this page:** 3\
**Page:** 1

<div class="post-metadata">

**Author:** ![Fabixel](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Fabixel](https://forums.developer.nvidia.com/u/Fabixel)\
**Post date:** [June 23, 2008, 10:32am UTC](https://forums.developer.nvidia.com/t/problem-with-coalesced-memory-access/4262/1 "2008-06-23T10:32:17Z")

</div>

Hi!

I’m trying to get a simple kernel running with coalesced memory reads & writes:

```auto
 // copy assignments

  int* to   = (int*)(ASSIGNMENT(&assign_in,scan_result[blockIdx.x]));

  int* from = (int*)(ASSIGNMENT(&assign_out,blockIdx.x));

  

  for(int i=threadIdx.x; i<variables+1; i+=THREADS_PER_BLOCK)

  	to[i] = from[i];

```

That’s the complete kernel.

If I print out the addresses and the addresses modulo 64 I get:

> [@](#):
>
> &nbsp; &nbsp; &nbsp; &nbsp;
> 
> 3058540800&nbsp; 3057225984&nbsp; 0&nbsp; 0
> 
> 3058541120&nbsp; 3057226304&nbsp; 0&nbsp; 0
> 
> 3058541440&nbsp; 3057226624&nbsp; 0&nbsp; 0
> 
> 3058540800&nbsp; 3057225984&nbsp; 0&nbsp; 0
> 
> 3058541120&nbsp; 3057226304&nbsp; 0&nbsp; 0
> 
> 3058541440&nbsp; 3057226624&nbsp; 0&nbsp; 0

So all the starting addresses are multiples of 64 and the array elements are integer. The size of THREADS\_PER\_BLOCK is 192.

This should make both the reads and the writes coalesced… but the CUDA profiler tells me that incoherent accesses far outweigh the coherent ones.

So where did I make my mistake …?

---

<div class="post-metadata">

**Author:** ![JHHPC](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@JHHPC](https://forums.developer.nvidia.com/u/JHHPC)\
**Post date:** [June 23, 2008, 11:25am UTC](https://forums.developer.nvidia.com/t/problem-with-coalesced-memory-access/4262/2 "2008-06-23T11:25:43Z")

</div>

> [@](#):
>
> ```auto
>  // copy assignments
> 
>  int* to   = (int*)(ASSIGNMENT(&assign_in,scan_result[blockIdx.x]));
> 
>  int* from = (int*)(ASSIGNMENT(&assign_out,blockIdx.x));
> 
>  
> 
>  for(int i=threadIdx.x; i<variables+1; i+=THREADS_PER_BLOCK)
> 
>  	to[i] = from[i];
> 
> ```
> 
> [snapback]398378[/snapback]

Why are you checking modulo 64 and not 32 (however should produce the same in your case, so just curiousity)?

Tells the profiler uncoherent loads or stores? Or both?

In my opinion your code is alright. Try perhaps 64 or 128 threads per block.

---

<div class="post-metadata">

**Author:** ![Fabixel](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Fabixel](https://forums.developer.nvidia.com/u/Fabixel)\
**Post date:** [June 23, 2008, 11:40am UTC](https://forums.developer.nvidia.com/t/problem-with-coalesced-memory-access/4262/3 "2008-06-23T11:40:28Z")

</div>

The check for mod64 has no real reason, other than trying to be extra safe ;)

The profiler counts too many global loads, for example:

gld\_incoherent: 4511

gld\_coherent: 179

gst\_incoherent: 0

gst\_coherent: 700

The storing seems to work fine… which confuses me a bit, since loading and storing both use the same mechanism.

128 threads per block did not work.

The first two lines also contain one global memory access, so I tried to make this access only once per threadBlock:

```auto
 // copy assignments

 Â __shared__ int* to;

 Â __shared__ int* from;

 Â 

 Â if(threadIdx.x==0){

     to  = (int*)(ASSIGNMENT(&assign_in,scan_result[blockIdx.x]));

 Â Â Â Â from = (int*)(ASSIGNMENT(&assign_out,blockIdx.x));

 Â } 

 Â 

 Â __syncthreads();

 Â 

 Â for(int i=threadIdx.x; i<variables+1; i+=THREADS_PER_BLOCK)

      to[i] = from[i];

```

But besides a compiler warning, nothing really changed.

(“Advisory: Cannot tell what pointer points to, assuming global memory space”)

\*EDIT

I think I may have found the problem. Not quoted was another if(…) which also contained one global memory read. Using a shared value like above, this value is now only read once to a shared variable and global incoherent loads went down by a few thousand :)

But now I’d like to get rid of this compiler warning [External Image](http://hqnveipbwb20/public/style_emoticons/<#EMO_DIR#>/blarg.gif "Image hosted on another site. Click to open in a new tab.")