# global memory latency

**URL:** <https://forums.developer.nvidia.com/t/global-memory-latency/4248>\
**Category:** CUDA Programming and Performance\
**Created:** [June 21, 2008, 7:09pm UTC](https://forums.developer.nvidia.com/t/global-memory-latency/4248 "2008-06-21T19:09:03Z")\
**Posts on this page:** 5\
**Page:** 1

<div class="post-metadata">

**Author:** ![santyhyammer](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@santyhyammer](https://forums.developer.nvidia.com/u/santyhyammer)\
**Post date:** [June 21, 2008, 7:09pm UTC](https://forums.developer.nvidia.com/t/global-memory-latency/4248/1 "2008-06-21T19:09:03Z")

</div>

Does a global memory RW operation really take 400-600 cycles even when it’s coalesced? Same question for local **device** memory .

And other question… is this coalesced?

```auto
__device__ float4 d[2048];

__global__ void myKernel ()

{

    const unsigned int tidx(blockDim.x*blockIdx.x+threadIdx.x);

    d[tidx] = make_float4(1.0f,2.0f,3.0f,4.0f);

}

```

or, to be coalesced needs to be

```auto
__device__ float d[2048];

```

in 32bit size and not 128.

thx

---

<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 22, 2008, 10:36am UTC](https://forums.developer.nvidia.com/t/global-memory-latency/4248/2 "2008-06-22T10:36:42Z")

</div>

> [@](#):
>
> Does a global memory RW operation really take 400-600 cycles even when it’s coalesced? Same question for local **device** memory .
> 
> [snapback]397678[/snapback]

Yes but the system transfers more than 4 bytes per latency.

> [@](#):
>
> And other question… is this coalesced?
> 
> ```auto
> __device__ float4 d[2048];
> 
> __global__ void myKernel ()
> 
> {
> 
>    const unsigned int tidx(blockDim.x*blockIdx.x+threadIdx.x);
> 
>    d[tidx] = make_float4(1.0f,2.0f,3.0f,4.0f);
> 
> }
> 
> ```
> 
> or, to be coalesced needs to be
> 
> ```auto
> __device__ float d[2048];
> 
> ```
> 
> in 32bit size and not 128.
> 
> thx
> 
> [snapback]397678[/snapback]

Haven’t workes with e.g. float4 yet, sry. But try profiling with the cuda visual profiler and look at the coherent/incoherent counts. If float4 works like the normal float than you should have a coalesced load, as thread 0 block 0 treats element 0 and so on

Johannes

---

<div class="post-metadata">

**Author:** ![Ailleur](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Ailleur](https://forums.developer.nvidia.com/u/Ailleur)\
**Post date:** [June 22, 2008, 2:58pm UTC](https://forums.developer.nvidia.com/t/global-memory-latency/4248/3 "2008-06-22T14:58:06Z")

</div>

> [@](#):
>
> Does a global memory RW operation really take 400-600 cycles even when it’s coalesced? Same question for local **device** memory .
> 
> And other question… is this coalesced?
> 
> ```auto
> __device__ float4 d[2048];
> 
> __global__ void myKernel ()
> 
> {
> 
> �   const unsigned int tidx(blockDim.x*blockIdx.x+threadIdx.x);
> 
> �   d[tidx] = make_float4(1.0f,2.0f,3.0f,4.0f);
> 
> }
> 
> ```
> 
> or, to be coalesced needs to be
> 
> ```auto
> __device__ float d[2048];
> 
> ```
> 
> in 32bit size and not 128.
> 
> thx
> 
> [snapback]397678[/snapback]

Yep this should work with float4s. Float3 cannot be coalesced (well, not out of the box anyway) but float4s are no problem since each thread will be fetching/writing 128bits words.

---

<div class="post-metadata">

**Author:** ![santyhyammer](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@santyhyammer](https://forums.developer.nvidia.com/u/santyhyammer)\
**Post date:** [June 22, 2008, 3:02pm UTC](https://forums.developer.nvidia.com/t/global-memory-latency/4248/4 "2008-06-22T15:02:27Z")

</div>

> [@](#):
>
> Yep this should work with float4s. Float3 cannot be coalesced (well, not out of the box anyway) but float4s are no problem since each thread will be fetching/writing 128bits words.
> 
> [snapback]398034[/snapback]

I figured it… any aligned type should be coalesced, thx!

And… is the local memory faster than the global one or is it excatly the same?

ps: the CUDa profiled does not work with my app.

---

<div class="post-metadata">

**Author:** ![Ailleur](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Ailleur](https://forums.developer.nvidia.com/u/Ailleur)\
**Post date:** [June 22, 2008, 3:16pm UTC](https://forums.developer.nvidia.com/t/global-memory-latency/4248/5 "2008-06-22T15:16:58Z")

</div>

> [@](#):
>
> I figured it… any aligned type should be coalesced, thx!
> 
> And… is the local memory faster than the global one or is it excatly the same?
> 
> ps: the CUDa profiled does not work with my app.
> 
> [snapback]398036[/snapback]

Section 5.1.2.2 of the programming guide. (beta2v2)