Artikelen

CUDA Shared Memory Swizzling: Voorkomen van Bank Conflicts

19-08-2026 12:54 • AI • https://leimao.github.io/blog/CUDA-Shared-Memor…

Wanneer we CUDA-kernels schrijven die gebruikmaken van shared memory, moeten we voorzichtig zijn met shared memory bank conflicts. Ernstige bank-conflicten kunnen leiden tot een aanzienlijke prestatiebeperking.

CUDA Shared Memory Swizzling

Introductie

Een eenvoudige manier om deze conflicten aan te pakken is door middel van padding. Padding kan echter shared memory verspillen en heeft andere nadelen.

In dit artikel bespreken we hoe we shared memory bank-conflicten kunnen oplossen met swizzling. Swizzling is een complexere techniek die kan worden gebruikt om bank-conflicten te vermijden zonder shared memory te verspillen.

CUDA Shared Memory Swizzling

Voorbeeld van Swizzling

Wanneer we CUDA shared memory gebruiken om data te cachen zonder padding te gebruiken, komt het vaak voor dat het lezen van of schrijven naar het shared memory door een warp bank-conflicten veroorzaakt. Swizzling is een techniek die de mapping van de shared memory-index herarrangeert om deze conflicten te voorkomen. Matrix-transpositie is een perfect voorbeeld van een operatie die bank-conflicten kan veroorzaken als de implementatie geen padding of swizzling gebruikt.

In een scenario waarbij het shared memory een 2D-array van floats is met een grootte van 32 × 16, leest elke warp een rij van 32 waarden uit het global memory en schrijft deze met swizzling naar het shared memory. Bij het schrijven naar het shared memory zullen er geen bank-conflicten optreden.

Om de matrix-transpositie uit te voeren, leest elke warp twee "geswizzlede" kolommen van 32 waarden uit het shared memory en schrijft deze naar het global memory. Op deze manier is er slechts één shared memory bank-conflict bij het lezen. Zonder swizzling zouden er 16 (2-way) bank-conflicten optreden. Uiteraard, als het shared memory een 2D-array van floats met een grootte van 32 × 32 zou zijn, zouden er geen bank-conflicten optreden bij het schrijven of lezen.

Swizzling-formule

Gegeven een array T array[][NX] in het shared memory, definiëren we NX × sizeof(T) == SWIZZLESIZE. De toegestane waarden voor SWIZZLESIZE zijn machten van 2 die groter zijn dan of gelijk zijn aan 32 (bijv. 32, 64, 128, 256, etc.).

Gegeven de index [y][x] in T array[][NX], kunnen we de geswizzlede index x_swz als volgt berekenen:

Bereken de index van het TC-byte blok binnen het SWIZZLE_SIZE-byte segment:

i_chunk = (y × NX + x) × sizeof(T) / sizeof(TC)
ychunk = i / (SWIZZLESIZE / sizeof(TC))
xchunk = i % (SWIZZLESIZE / sizeof(TC))

Bereken de geswizzlede index van het TC-byte blok met een XOR-operatie:

xchunkswz = ychunk ^ xchunk

Bereken de uiteindelijke geswizzlede index:

xswz = xchunk_swz × sizeof(TC) / sizeof(T) % NX + x % (sizeof(TC) / sizeof(T))

Eigenschappen van Swizzling

Deze swizzling-formule heeft de volgende eigenschappen:

De index voor en na swizzling moet een één-op-één mapping zijn.
NX moet een macht van 2 zijn.
Voor elke x en elke set {y, y+1, y+2, ..., y+31}, moet het aantal unieke geswizzlede indices x_swz gemaximaliseerd worden.

Eigenschap 1 garandeert dat er geen dataverlies optreedt tijdens het swizzlen. Eigenschap 2 garandeert dat de mapping consistent blijft.

Bewijs

Hieronder volgt een informeel wiskundig bewijs voor deze eigenschappen.

Bewijs voor Eigenschap 1: $$ \begin{align} x{\text{chunk}} &= i{\text{chunk}} \% (\text{SWIZZLESIZE} / \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} $$ Hieruit volgt een equivalente formule voor $x{\text{chunk}}$. Let op dat $\text{sizeof}(\text{T}) / \text{sizeof}(\text{TC})$ een bit-shift operatie is wanneer $\text{sizeof}(\text{TC}) / \text{sizeof}(\text{T})$ een macht van 2 is.

Voor $y{\text{chunk}}$: $$ \begin{align} y{\text{chunk}} &= i{\text{chunk}} / (\text{SWIZZLESIZE} / \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} $$

Voor $x{\text{chunk\swz}}$: $$ \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} $$ Aangezien $\oplus$ de bitwise XOR-operatie is, is de mapping een één-op-één mapping als ofwel $y{\text{chunk}}$ of $x{\text{chunk}}$ een constante is.

Voor de uiteindelijke index $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})) \\ \end{align} $$ Omdat een reeks van $\text{sizeof}(\text{TC}) / \text{sizeof}(\text{T})$ $x$-waarden wordt gemapt naar één unieke chunk-index $x{\text{chunk}}$, en de mapping tussen $x{\text{chunk}}$ en $x{\text{chunk\swz}}$ één-op-één is, zal één $x{\text{chunk\swz}}$ waarde mappen naar één unieke $x{\text{chunk\swz}} × \text{sizeof}(\text{TC}) / \text{sizeof}(\text{T}) \% \text{NX}$ waarde. Door de offset $x \% (\text{sizeof}(\text{TC}) / \text{sizeof}(\text{T}))$ toe te voegen, ontstaat er een één-op-één mapping tussen de geswizzlede index $x_{\text{swz}}$ en de originele index $x$.

Eigenschap 2 is triviaal aan te tonen. Voor Eigenschap 3 is het volgende inzicht nuttig: $$ \begin{align} x{\text{swz}} &= \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} $$ Voor elke $x$ en elke set $\{y, y+1, y+2, \dots, y+\text{NX}-1\}$, is het aantal unieke geswizzlede indices $x{\text{swz}}$ gelijk aan $\text{NX}$, wat het maximum is.

Voorbeelden

Matrix Transpositie

In dit voorbeeld is de matrix-transpositie in CUDA geïmplementeerd op drie verschillende manieren:

Transpositie met shared memory bank-conflicten.
Transpositie zonder bank-conflicten via padding.
Transpositie zonder bank-conflicten via swizzling.

Implementatie (`transpose.cu`)

#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)
{
    // Verspil wat shared memory om bank-conflicten te vermijden als BLOCK_TILE_SKEW_SIZE_X != 0.
    __shared__ T shm[BLOCK_TILE_SIZE_Y][BLOCK_TILE_SIZE_X + BLOCK_TILE_SKEW_SIZE_X];

    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. Geen shared memory bank conflict.
        shm[shm_to_idx_y][shm_to_idx_x] = input_matrix[input_matrix_from_idx];
    }

    __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. Geen bank conflict als 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];

    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))
    {
        shm[shm_to_idx_y][shm_to_idx_x_swizzled] = input_matrix[input_matrix_from_idx];
    }

    __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))
    {
        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)
{
    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);

    for (size_t i{0}; i < M; ++i)
    {
        for (size_t j{0}; j < N; ++j)
        {
            matrix_transposed_reference[j * M + i] = matrix[i * N + j];
        }
    }

    T *d_matrix, *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, *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()
{
    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));
        }
    }

    size_t const M{8192};
    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;
}

Resultaten

Het programma is gebouwd en uitgevoerd op een platform met een Intel i9-9900K CPU en een NVIDIA RTX 3090 GPU.

$ nvcc transpose.cu -o transpose
$ ./transpose
8192 x 8192 Matrix
Transpose with Shared Memory Bank Conflict: 1.10 ms
Transpose without Shared Memory Bank Conflict via Padding: 0.92 ms
Transpose without Shared Memory Bank Conflict via Swizzling: 0.92 ms

We zien dat de transpose-kernel met shared memory bank-conflicten de hoogste latentie heeft. De kernels die gebruikmaken van padding en swizzling hebben dezelfde latentie en zijn in dit geval 20% sneller.

Let op dat deze implementatie ongeveer 65% van de piek-geheugenbandbreedte van een RTX 3090 GPU behaalt. De prestaties kunnen aanzienlijk worden verbeterd door gebruik te maken van gevectoriseerde geheugentoegang, mits wordt aangenomen dat de matrix altijd is gepadded (meestal toegewezen via cudaMallocPitch), zodat elke rij aan de coalescing-vereisten voldoet.

Swizzling vs. Padding

Swizzling en padding zijn twee veelgebruikte technieken om shared memory bank-conflicten op te lossen.

Swizzling

Voordeel: Verspilt geen ruimte in het shared memory.
Nadeel: Complexer om te implementeren en te begrijpen, omdat de index-mapping niet lineair is.

Padding

Voordeel: Eenvoudig te implementeren en te begrijpen.
Nadeel: Verspilt shared memory-ruimte. Daarnaast kan het de adres-uitlijning van de data doorbreken als de padding-grootte niet zorgvuldig is gekozen. Dit gebeurt vaak bij gevectoriseerde geheugentoegang op 2D-gepadde arrays, wat kan leiden tot ongedefinieerd gedrag.

Referenties

CUDA Coalesced Memory Access
CUDA Shared Memory Bank
Swizzling - Wikipedia
Advanced Performance Optimization in CUDA - NVIDIA GTC 2024
Using TMA to Transfer Multi-Dimensional Arrays - CUDA Programming Guide
CUtensorMapSwizzle - CUDA Driver API

article_shortcut	Samenvatting laten zien
receipt_long	Volledige tekst laten zien
favorite	Artikel opslaan als favoriet
keep	Artikel vastzetten in linker kantlijn
settings	Categorie-instellingen beheren

Tech-update

CUDA Shared Memory Swizzling: Voorkomen van Bank Conflicts

CUDA Shared Memory Swizzling

Introductie

CUDA Shared Memory Swizzling

Voorbeeld van Swizzling

Swizzling-formule

Eigenschappen van Swizzling

Bewijs

Voorbeelden

Matrix Transpositie

Implementatie (`transpose.cu`)

Resultaten

Swizzling vs. Padding

Referenties

CUDA Shared Memory Swizzling

Introductie

CUDA Shared Memory Swizzling

Voorbeeld van Swizzling

Swizzling-formule

Eigenschappen van Swizzling

Bewijs

Voorbeelden

Matrix Transpositie

Implementatie (`transpose.cu`)

Resultaten

Swizzling vs. Padding

Referenties

Zichtbare categorieën

Over deze site

CUDA Shared Memory Swizzling: Voorkomen van Bank Conflicts

CUDA Shared Memory Swizzling

Introductie

CUDA Shared Memory Swizzling

Voorbeeld van Swizzling

Swizzling-formule

Eigenschappen van Swizzling

Bewijs

Voorbeelden

Matrix Transpositie

Implementatie (transpose.cu)

Resultaten

Swizzling vs. Padding

Referenties

CUDA Shared Memory Swizzling

Introductie

CUDA Shared Memory Swizzling

Voorbeeld van Swizzling

Swizzling-formule

Eigenschappen van Swizzling

Bewijs

Voorbeelden

Matrix Transpositie

Implementatie (transpose.cu)

Resultaten

Swizzling vs. Padding

Referenties

Implementatie (`transpose.cu`)

Implementatie (`transpose.cu`)