# Memory Access **URL:** \ **Category:** CUDA Programming and Performance\ **Created:** [June 22, 2009, 10:01am UTC](https://forums.developer.nvidia.com/t/memory-access/10305 "2009-06-22T10:01:02Z")\ **Posts on this page:** 3\ **Page:** 1
**Author:** ![cuda\_micha](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@cuda\_micha](https://forums.developer.nvidia.com/u/cuda_micha)\ **Post date:** [June 22, 2009, 10:01am UTC](https://forums.developer.nvidia.com/t/memory-access/10305/1 "2009-06-22T10:01:02Z")
Hello, I have a problem accessing shared memory. if I allocate the memory by the following way: ```auto __shared__ char shared_mem[8192]; ``` and access it (the first 4 Bytes) that way: ```auto unsigned int test1 = *(unsigned int *)(shared_mem); ``` Ill get the nvcc compiler error: “Error: Unaligned memory accesses not supported” I dont know why this error occours, because the access is aligned to 4 Byte and starts at the total offset of the allocated memory… I figured out that I can handle this problem allocating the shared memory as an array of 4 Byte values: ```auto __shared__ unsigned int shared_mem[8192/4]; ``` as a result, I dont get the compile error and for the first test everything seems to work fine… BUT! Now I have another problem… in contrast the the upper solution with the char array, the profiler now tells me that cuda has to serialize the wrap threads for memory acces… But it should be a broadcast because all threads access the same area… Is there another way or workaround to solve the problem with the “Unaligned memory accesses” ??? Thank you! ---
**Author:** ![mikemoik](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@mikemoik](https://forums.developer.nvidia.com/u/mikemoik)\ **Post date:** [June 22, 2009, 10:50am UTC](https://forums.developer.nvidia.com/t/memory-access/10305/2 "2009-06-22T10:50:58Z")
hi micha, ```auto unsigned int test1 = *(unsigned int *)(shared_mem); ``` as nvcc says this access is definitly not 4byte aligned, because shared\_mem % 4 != 0. > [@](#): > > I figured out that I can handle this problem allocating the shared memory > > as an array of 4 Byte values: > > ```auto > __shared__ unsigned int shared_mem[8192/4]; > > ``` yes, now shared\_mem is 4byte-aligened as every int. > [@](#): > > Now I have another problem… > > in contrast the the upper solution with the char array, the profiler now tells > > me that cuda has to serialize the wrap threads for memory acces… > > But it should be a broadcast because all threads access the same area… hmm depends on how u access memory right now … but serialization seems quite propable with 8bit-access. how could u compare to former code, when nvcc didnt compile? ;) yours, moik ---
**Author:** ![cuda\_micha](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@cuda\_micha](https://forums.developer.nvidia.com/u/cuda_micha)\ **Post date:** [June 22, 2009, 11:31am UTC](https://forums.developer.nvidia.com/t/memory-access/10305/3 "2009-06-22T11:31:03Z")
Hi mikemoik, thanks… seems like you can only handle with 4 byte aligned arrays in shmem… is there no chance to handle with char arrays and access them only in 4 byte aligned way?