CUDA Shared Memory Swizzling

NVIDIA's CUDA programming guide introduces a swizzling technique to eliminate shared memory bank conflicts without wasting memory, using an XOR-based formula that rearranges shared memory indices. The method, detailed in a blog post, ensures one-to-one mapping and maximizes unique swizzled indices for warps, reducing bank conflicts from 16 to 1 in a 32×16 float matrix transpose example.

CUDA Shared Memory Swizzling Introduction When we write CUDA kernels that use shared memory, we have to be careful about shared memory bank conflicts /blog/CUDA-Shared-Memory-Bank/ . Having severe shared memory bank conflicts can introduce a significant performance penalty. One simple way to deal with shared memory bank conflicts is to use padding. However, padding can waste shared memory and can have other drawbacks. In this blog post, I would like to discuss how to deal with shared memory bank conflicts using swizzling. Swizzling is a more complicated technique that can be used to avoid shared memory bank conflicts without wasting shared memory. CUDA Shared Memory Swizzling Swizzling Example When we use CUDA shared memory to cache data without using padding, it’s very common that either reading from or writing to shared memory by a warp can cause shared memory bank conflicts. Swizzling is a technique that rearranges the mapping of the shared memory index to avoid shared memory bank conflicts. Matrix transpose is a perfect example that can have shared memory bank conflicts if the implementation does not use padding or swizzling. In the above example, the shared memory is a 2D array of float with size 32 × 16 . In terms of matrix transpose, each warp reads a row of 32 values from the global memory and writes them to the shared memory with swizzling. There will be no shared memory bank conflicts when writing to the shared memory. To perform matrix transpose, each wrap reads two swizzled “columns” of 32 values from the shared memory and writes them to the global memory. For example, the swizzled column 0 and 1 are colored in yellow and cyan, respectively. In this way, there will be only one shared memory bank conflict when reading from the shared memory. Without using swizzling, there will be 16 2-way shared memory bank conflicts when reading from the shared memory. Of course, obviously, if the shared memory is a 2D array of float with size 32 × 32 , there will be no shared memory bank conflicts when writing to the shared memory and reading from the shared memory. Swizzling Formula Given an array of T array NX on shared memory, we define NX × sizeof T == SWIZZLE SIZE . The allowed values of SWIZZLE SIZE are powers of 2 that is larger than or equal to 32, such as 32, 64, 128, 256, …, etc. Given the index y x in T array NX , we can compute the swizzled index x swz as follows: Compute the index of the TC -byte chunk within SWIZZLE SIZE -byte segment: i chunk = y × NX + x × sizeof T / sizeof TC y chunk = i / SWIZZLE SIZE / sizeof TC x chunk = i % SWIZZLE SIZE / sizeof TC Compute the swizzled index of TC -byte chunk using XOR operation https://en.wikipedia.org/wiki/XOR gate : x chunk swz = y chunk ^ x chunk Compute swizzled index: x swz = x chunk swz × sizeof TC / sizeof T % NX + x % sizeof TC / sizeof T Swizzling Properties This swizzling formula has the following properties: - The index before and after swizzling must be one to one mapped. - $\text{NX}$ must be a power of 2. - Given any $x$ and any $\{y, y+1, y+2, \cdots, y+31\}$, the number of unique swizzled index $x {\text{swz}}$ should be maximized. The property 1 ensures that there will be no data loss during swizzling. The property 2 ensures that the index before and after swizzling will be one to one mapped. Here I am going to show some informal mathematical proofs for the properties from the swizzling formula. Proof We will first prove the property 1. $$ \begin{align} x {\text{chunk}} &= i {\text{chunk}} \% \text{SWIZZLE SIZE} / \text{sizeof} \text{TC} \\ &= \left \left y × \text{NX} + x\right × \text{sizeof} \text{T} / \text{sizeof} \text{TC} \right \% \text{NX} × \text{sizeof} \text{T} / \text{sizeof} \text{TC} \\ &= \left y × \text{NX} × \text{sizeof} \text{T} / \text{sizeof} \text{TC} + x × \text{sizeof} \text{T} / \text{sizeof} \text{TC} \right \% \text{NX} × \text{sizeof} \text{T} / \text{sizeof} \text{TC} \\ &= \left y × \text{NX} × \text{sizeof} \text{T} / \text{sizeof} \text{TC} \% \text{NX} × \text{sizeof} \text{T} / \text{sizeof} \text{TC} + x × \text{sizeof} \text{T} / \text{sizeof} \text{TC} \% \text{NX} × \text{sizeof} \text{T} / \text{sizeof} \text{TC} \right \% \text{NX} × \text{sizeof} \text{T} / \text{sizeof} \text{TC} \\ &= \left x × \text{sizeof} \text{T} / \text{sizeof} \text{TC} \% \text{NX} × \text{sizeof} \text{T} / \text{sizeof} \text{TC} \right \% \text{NX} × \text{sizeof} \text{T} / \text{sizeof} \text{TC} \\ &= x × \text{sizeof} \text{T} / \text{sizeof} \text{TC} \% \text{NX} × \text{sizeof} \text{T} / \text{sizeof} \text{TC} \\ &= \left x \% \text{NX} \right × \text{sizeof} \text{T} / \text{sizeof} \text{TC} \\ &= x × \text{sizeof} \text{T} / \text{sizeof} \text{TC} \\ \end{align} $$ It seems that we have derived another equivalent formula for $x {\text{chunk}}$. Note that $\text{sizeof} \text{T} / \text{sizeof} \text{TC} $ is a bit right shifting operation when $\text{sizeof} \text{TC} / \text{sizeof} \text{T} $ is a power of 2. $$ \begin{align} y {\text{chunk}} &= i {\text{chunk}} / \text{SWIZZLE SIZE} / \text{sizeof} \text{TC} \\ &= \left \left y × \text{NX} + x\right × \text{sizeof} \text{T} / \text{sizeof} \text{TC} \right / \text{NX} × \text{sizeof} \text{T} / \text{sizeof} \text{TC} \\ &= \left y × \text{NX} × \text{sizeof} \text{T} / \text{sizeof} \text{TC} + x × \text{sizeof} \text{T} / \text{sizeof} \text{TC} \right / \text{NX} × \text{sizeof} \text{T} / \text{sizeof} \text{TC} \\ &= y × \text{NX} × \text{sizeof} \text{T} / \text{sizeof} \text{TC} / \text{NX} × \text{sizeof} \text{T} / \text{sizeof} \text{TC} + x × \text{sizeof} \text{T} / \text{sizeof} \text{TC} / \text{NX} × \text{sizeof} \text{T} / \text{sizeof} \text{TC} \\ &= y + x / \text{NX} \\ &= y \\ \end{align} $$ It seems that we have also derived another equivalent formula for $y {\text{chunk}}$. $$ \begin{align} x {\text{chunk swz}} &= y {\text{chunk}} \oplus x {\text{chunk}} \\ &= y \oplus \left x × \text{sizeof} \text{T} / \text{sizeof} \text{TC} \right \\ &= y / \text{NX} × \text{sizeof} \text{T} / \text{sizeof} \text{TC} × \text{NX} × \text{sizeof} \text{T} / \text{sizeof} \text{TC} + \left y \% \text{NX} × \text{sizeof} \text{T} / \text{sizeof} \text{TC} \right \oplus \left x × \text{sizeof} \text{T} / \text{sizeof} \text{TC} \right \\ &= y / \text{NX} × \text{sizeof} \text{T} / \text{sizeof} \text{TC} × \text{NX} × \text{sizeof} \text{T} / \text{sizeof} \text{TC} + \left \left y \% \text{NX} \right × \text{sizeof} \text{T} / \text{sizeof} \text{TC} \right \oplus \left x × \text{sizeof} \text{T} / \text{sizeof} \text{TC} \right \\ &= y / \text{NX} × \text{sizeof} \text{T} / \text{sizeof} \text{TC} × \text{NX} × \text{sizeof} \text{T} / \text{sizeof} \text{TC} + \left y \% \text{NX} \right \oplus x × \text{sizeof} \text{T} / \text{sizeof} \text{TC} \\ \end{align} $$ Note that $\oplus$ is the bitwise XOR operation. If either $y {\text{chunk}}$ or $x {\text{chunk}}$ is a constant, the mapping is a one to one mapping. $$ \begin{align} x {\text{swz}} &= x {\text{chunk swz}} × \text{sizeof} \text{TC} / \text{sizeof} \text{T} \% \text{NX} + x \% \text{sizeof} \text{TC} / \text{sizeof} \text{T} \\ \end{align} $$ Here the proof becomes a little bit informal. Because a consecutive of $\text{sizeof} \text{TC} / \text{sizeof} \text{T} $ $x$ values will be mapped to one unique trunk index $x {\text{chunk}}$, the mapping between $x {\text{chunk}}$ and $x {\text{chunk swz}}$ is a one to one mapping, one $x {\text{chunk swz}}$ value will map to one unique $x {\text{chunk swz}} × \text{sizeof} \text{TC} / \text{sizeof} \text{T} \% \text{NX}$ value. To create the one to one mapping between the swizzled index $x {\text{swz}}$ and the original index $x$, the offset $x \% \text{sizeof} \text{TC} / \text{sizeof} \text{T} $ is added. Therefore, the index before and after swizzling must be one to one mapped. The property 2 is trivial to show. The property 3 might be somewhat easier to see given the following expression for $x {\text{swz}}$. $$ \begin{align} x {\text{swz}} &= x {\text{chunk swz}} × \text{sizeof} \text{TC} / \text{sizeof} \text{T} \% \text{NX} + x \% \text{sizeof} \text{TC} / \text{sizeof} \text{T} \\ &= \left y / \text{NX} × \text{sizeof} \text{T} / \text{sizeof} \text{TC} × \text{NX} × \text{sizeof} \text{T} / \text{sizeof} \text{TC} + \left y \% \text{NX} \right \oplus x × \text{sizeof} \text{T} / \text{sizeof} \text{TC} \right × \text{sizeof} \text{TC} / \text{sizeof} \text{T} \% \text{NX} + x \% \text{sizeof} \text{TC} / \text{sizeof} \text{T} \\ &= \left y \% \text{NX} \right \oplus x × \text{sizeof} \text{T} / \text{sizeof} \text{TC} × \text{sizeof} \text{TC} / \text{sizeof} \text{T} \% \text{NX} + x \% \text{sizeof} \text{TC} / \text{sizeof} \text{T} \\ &= \left y \% \text{NX} \right \oplus x × \text{sizeof} \text{T} / \text{sizeof} \text{TC} × \text{sizeof} \text{TC} / \text{sizeof} \text{T} + x \% \text{sizeof} \text{TC} / \text{sizeof} \text{T} \\ \end{align} $$ Given any $x$ and any $\{y, y+1, y+2, \cdots, y+\text{NX}-1\}$, the number of unique swizzled index $x {\text{swz}}$ is $\text{NX}$ which is maximized. Examples Matrix Transpose In this example, we implemented matrix transpose CUDA kernels using shared memory in three different ways: - Transpose with shared memory bank conflict. - Transpose without shared memory bank conflict via padding. - Transpose without shared memory bank conflict via swizzling. 123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202203204205206207208209210211212213214215216217218219220221222223224225226227228229230231232233234235236237238239240241242243244245246247248249250251252253254255256257258259260261262263264265266267268269270271272273274275276277278279280281282283284285286287288289290291292293294295296297298299300301302303304305306307308309310311312313314315316317318319320321322323324325326327328329330331332333334335336337338339340341342343344345346347348349350351352353354355356357358359360361362363364365366367 | include <algorithm include <cassert include <chrono include <cstdio include <functional include <iomanip include <iostream include <random include <vector include <cuda runtime.h define CHECK CUDA ERROR val check val , val, FILE , LINE void check cudaError t err, char const func, char const file, int line { if err = cudaSuccess { std::cerr << "CUDA Runtime Error at: " << file << ":" << line << std::endl; std::cerr << cudaGetErrorString err << " " << func << std::endl; std::exit EXIT FAILURE ; }} define CHECK LAST CUDA ERROR check last FILE , LINE void check last char const file, int line { cudaError t const err{cudaGetLastError }; if err = cudaSuccess { std::cerr << "CUDA Runtime Error at: " << file << ":" << line << std::endl; std::cerr << cudaGetErrorString err << std::endl; std::exit EXIT FAILURE ; }}template <class T float measure performance std::function<T cudaStream t bound function, cudaStream t stream, size t num repeats = 10, size t num warmups = 10 { cudaEvent t start, stop; float time; CHECK CUDA ERROR cudaEventCreate &start ; CHECK CUDA ERROR cudaEventCreate &stop ; for size t i{0}; i < num warmups; ++i { bound function stream ; } CHECK CUDA ERROR cudaStreamSynchronize stream ; CHECK CUDA ERROR cudaEventRecord start, stream ; for size t i{0}; i < num repeats; ++i { bound function stream ; } CHECK CUDA ERROR cudaEventRecord stop, stream ; CHECK CUDA ERROR cudaEventSynchronize stop ; CHECK LAST CUDA ERROR ; CHECK CUDA ERROR cudaEventElapsedTime &time, start, stop ; CHECK CUDA ERROR cudaEventDestroy start ; CHECK CUDA ERROR cudaEventDestroy stop ; float const latency{time / num repeats}; return latency;}constexpr size t div up size t a, size t b { return a + b - 1 / b; }template <typename T, size t BLOCK TILE SIZE X = 32, size t BLOCK TILE SIZE Y = 32, size t BLOCK TILE SKEW SIZE X = 0 global void transpose T output matrix, T const input matrix, size t M, size t N { // Waste some shared memory to avoid bank conflicts if // BLOCK TILE SKEW SIZE X = 0. shared T shm BLOCK TILE SIZE Y BLOCK TILE SIZE X + BLOCK TILE SKEW SIZE X ; // In some algorithms, such as matrix multiplication, // a warp of threads have to access a column of the 2D matrix in the shared // memory. Using the conventional index mapping, if the column size is not a // multiple of the warp size, there will be bank conflicts. size t const input matrix from idx x{threadIdx.x + blockIdx.x blockDim.x}; size t const input matrix from idx y{threadIdx.y + blockIdx.y blockDim.y}; size t const input matrix from idx{input matrix from idx x + input matrix from idx y N}; size t const shm to idx x{threadIdx.x}; size t const shm to idx y{threadIdx.y}; if input matrix from idx y < M && input matrix from idx x < N { // Coalesced global memory access. // No shared memory bank conflict. shm shm to idx y shm to idx x = input matrix input matrix from idx ; } // Make sure the buffer in a block is filled. syncthreads ; size t const block thread idx{threadIdx.x + threadIdx.y blockDim.x}; size t const shm from idx x{block thread idx / BLOCK TILE SIZE Y}; size t const shm from idx y{block thread idx % BLOCK TILE SIZE Y}; size t const output matrix to idx x{shm from idx y + blockIdx.y blockDim.y}; size t const output matrix to idx y{shm from idx x + blockIdx.x blockDim.x}; size t const output matrix to idx{output matrix to idx x + output matrix to idx y M}; if output matrix to idx y < N && output matrix to idx x < M { // Coalesced global memory access. // No shared memory bank conflict if BLOCK TILE SKEW SIZE X = 1. output matrix output matrix to idx = shm shm from idx y shm from idx x ; }}template <typename T, size t BLOCK TILE SIZE X = 32, size t BLOCK TILE SIZE Y = 32 global void transpose swizzling T output matrix, T const input matrix, size t M, size t N { shared T shm BLOCK TILE SIZE Y BLOCK TILE SIZE X ; // In some algorithms, such as matrix multiplication, // a warp of threads have to access a column of the 2D matrix in the shared // memory. Using the conventional index mapping, if the column size is not a // multiple of the warp size, there will be bank conflicts. size t const input matrix from idx x{threadIdx.x + blockIdx.x blockDim.x}; size t const input matrix from idx y{threadIdx.y + blockIdx.y blockDim.y}; size t const input matrix from idx{input matrix from idx x + input matrix from idx y N}; size t const shm to idx x{threadIdx.x}; size t const shm to idx y{threadIdx.y}; size t const shm to idx x swizzled{ shm to idx x ^ shm to idx y % BLOCK TILE SIZE X}; if input matrix from idx y < M && input matrix from idx x < N { // Coalesced global memory access. // No shared memory bank conflict. shm shm to idx y shm to idx x swizzled = input matrix input matrix from idx ; } // Make sure the buffer in a block is filled. syncthreads ; size t const block thread idx{threadIdx.x + threadIdx.y blockDim.x}; size t const shm from idx x{block thread idx / BLOCK TILE SIZE Y}; size t const shm from idx y{block thread idx % BLOCK TILE SIZE Y}; size t const shm from idx x swizzled{ shm from idx x ^ shm from idx y % BLOCK TILE SIZE X}; size t const output matrix to idx x{shm from idx y + blockIdx.y blockDim.y}; size t const output matrix to idx y{shm from idx x + blockIdx.x blockDim.x}; size t const output matrix to idx{output matrix to idx x + output matrix to idx y M}; if output matrix to idx y < N && output matrix to idx x < M { // Coalesced global memory access. // No shared memory bank conflict. output matrix output matrix to idx = shm shm from idx y shm from idx x swizzled ; }}template <typename T void launch transpose with shm bank conflict T d output matrix, T const d input matrix, size t M, size t N, cudaStream t stream { constexpr size t BLOCK TILE SIZE X{32}; constexpr size t BLOCK TILE SIZE Y{32}; constexpr size t BLOCK TILE SKEW SIZE X{0}; dim3 const block size{BLOCK TILE SIZE X, BLOCK TILE SIZE Y}; dim3 const grid size{static cast<unsigned int div up N, block size.x , static cast<unsigned int div up M, block size.y }; transpose<T, BLOCK TILE SIZE X, BLOCK TILE SIZE Y, BLOCK TILE SKEW SIZE X <<<grid size, block size, 0, stream d output matrix, d input matrix, M, N ; CHECK LAST CUDA ERROR ;}template <typename T void launch transpose without shm bank conflict via padding T d output matrix, T const d input matrix, size t M, size t N, cudaStream t stream { constexpr size t BLOCK TILE SIZE X{32}; constexpr size t BLOCK TILE SIZE Y{32}; constexpr size t BLOCK TILE SKEW SIZE X{1}; dim3 const block size{BLOCK TILE SIZE X, BLOCK TILE SIZE Y}; dim3 const grid size{static cast<unsigned int div up N, block size.x , static cast<unsigned int div up M, block size.y }; transpose<T, BLOCK TILE SIZE X, BLOCK TILE SIZE Y, BLOCK TILE SKEW SIZE X <<<grid size, block size, 0, stream d output matrix, d input matrix, M, N ; CHECK LAST CUDA ERROR ;}template <typename T void launch transpose without shm bank conflict via swizzling T d output matrix, T const d input matrix, size t M, size t N, cudaStream t stream { constexpr size t BLOCK TILE SIZE X{32}; constexpr size t BLOCK TILE SIZE Y{32}; dim3 const block size{BLOCK TILE SIZE X, BLOCK TILE SIZE Y}; dim3 const grid size{static cast<unsigned int div up N, block size.x , static cast<unsigned int div up M, block size.y }; transpose swizzling<T, BLOCK TILE SIZE X, BLOCK TILE SIZE Y <<<grid size, block size, 0, stream d output matrix, d input matrix, M, N ; CHECK LAST CUDA ERROR ;}template <typename T bool is equal T const data 1, T const data 2, size t size { for size t i{0}; i < size; ++i { if data 1 i = data 2 i { return false; } } return true;}template <typename T bool verify transpose implementation std::function<void T , T const , size t, size t, cudaStream t transpose function, size t M, size t N { // Fixed random seed for reproducibility std::mt19937 gen{0}; cudaStream t stream; size t const matrix size{M N}; std::vector<T matrix matrix size, 0.0f ; std::vector<T matrix transposed matrix size, 1.0f ; std::vector<T matrix transposed reference matrix size, 2.0f ; std::uniform real distribution<T uniform dist -256, 256 ; for size t i{0}; i < matrix size; ++i { matrix i = uniform dist gen ; } // Create the reference transposed matrix using CPU. for size t i{0}; i < M; ++i { for size t j{0}; j < N; ++j { size t const from idx{i N + j}; size t const to idx{j M + i}; matrix transposed reference to idx = matrix from idx ; } } T d matrix; T d matrix transposed; CHECK CUDA ERROR cudaMalloc &d matrix, matrix size sizeof T ; CHECK CUDA ERROR cudaMalloc &d matrix transposed, matrix size sizeof T ; CHECK CUDA ERROR cudaStreamCreate &stream ; CHECK CUDA ERROR cudaMemcpy d matrix, matrix.data , matrix size sizeof T , cudaMemcpyHostToDevice ; transpose function d matrix transposed, d matrix, M, N, stream ; CHECK CUDA ERROR cudaStreamSynchronize stream ; CHECK CUDA ERROR cudaMemcpy matrix transposed.data , d matrix transposed, matrix size sizeof T , cudaMemcpyDeviceToHost ; bool const correctness{is equal matrix transposed.data , matrix transposed reference.data , matrix size }; CHECK CUDA ERROR cudaFree d matrix ; CHECK CUDA ERROR cudaFree d matrix transposed ; CHECK CUDA ERROR cudaStreamDestroy stream ; return correctness;}template <typename T float profile transpose implementation std::function<void T , T const , size t, size t, cudaStream t transpose function, size t M, size t N { constexpr int num repeats{100}; constexpr int num warmups{10}; cudaStream t stream; size t const matrix size{M N}; T d matrix; T d matrix transposed; CHECK CUDA ERROR cudaMalloc &d matrix, matrix size sizeof T ; CHECK CUDA ERROR cudaMalloc &d matrix transposed, matrix size sizeof T ; CHECK CUDA ERROR cudaStreamCreate &stream ; std::function<void cudaStream t const transpose function wrapped{ std::bind transpose function, d matrix transposed, d matrix, M, N, std::placeholders:: 1 }; float const transpose function latency{measure performance transpose function wrapped, stream, num repeats, num warmups }; CHECK CUDA ERROR cudaFree d matrix ; CHECK CUDA ERROR cudaFree d matrix transposed ; CHECK CUDA ERROR cudaStreamDestroy stream ; return transpose function latency;}void print latencty std::string const& kernel name, float latency { std::cout << kernel name << ": " << std::fixed << std::setprecision 2 << latency << " ms" << std::endl;}int main { // Unit tests. for size t m{1}; m <= 64; ++m { for size t n{1}; n <= 64; ++n { assert verify transpose implementation<float &launch transpose with shm bank conflict<float , m, n ; assert verify transpose implementation<float &launch transpose without shm bank conflict via padding<float , m, n ; assert verify transpose implementation<float &launch transpose without shm bank conflict via swizzling< float , m, n ; } } // M: Number of rows. size t const M{8192}; // N: Number of columns. size t const N{8192}; std::cout << M << " x " << N << " Matrix" << std::endl; float const latency with shm bank conflict{ profile transpose implementation<float &launch transpose with shm bank conflict<float , M, N }; print latencty "Transpose with Shared Memory Bank Conflict", latency with shm bank conflict ; float const latency without shm bank conflict via padding{ profile transpose implementation<float &launch transpose without shm bank conflict via padding<float , M, N }; print latencty "Transpose without Shared Memory Bank Conflict via Padding", latency without shm bank conflict via padding ; float const latency without shm bank conflict via swizzling{ profile transpose implementation<float &launch transpose without shm bank conflict via swizzling<float , M, N }; print latencty "Transpose without Shared Memory Bank Conflict via Swizzling", latency without shm bank conflict via swizzling ; return 0;} | The program was built and performed on a platform that has an Intel i9-9900K CPU and an NVIDIA RTX 3090 GPU. 123456 | bash $ nvcc transpose.cu -o transpose$ ./transpose8192 x 8192 MatrixTranspose with Shared Memory Bank Conflict: 1.10 msTranspose without Shared Memory Bank Conflict via Padding: 0.92 msTranspose without Shared Memory Bank Conflict via Swizzling: 0.92 ms | We could see that the transpose kernel with shared memory bank conflict has the highest latency, while the transpose kernel without shared memory bank conflict via padding and swizzling have the same latency and run 20% faster than the kernel with shared memory bank conflict in this case. Note that this implementation achieves ~65% of the peak memory bandwidth of an RTX 3090 GPU. The performance can be further improved significantly using vectorized memory access if the implementation assumes the matrix is always padded and usually allocated using cudaMallocPitch https://docs.nvidia.com/cuda/cuda-runtime-api/group CUDART MEMORY.html group CUDART MEMORY 1g32bd7a39135594788a542ae72217775c so that each row will continue to meet the coalescing requirement. Swizzling vs Padding Swizzling and padding are two common techniques to deal with shared memory bank conflicts. The advantage of swizzling is that it does not waste shared memory space. The disadvantage of swizzling is that it is more complicated to implement and understand because the index mapping is not linear. The advantage of padding is that it is simple to implement and understand. The disadvantage of padding is that it wastes shared memory space and can break the address alignment of the data if the padding size is not selected carefully especially when we access the data via large trunks using reinterpret cast and cause undefined behaviors . This usually happens when vectorized memory access /blog/CUDA-Vectorized-Memory-Access/ is performed on 2D padded arrays which accidentally breaks the alignment of the data. References CUDA Shared Memory Swizzling