cuda::barrier / async copies

cuda::barrier / async copies (sim only) The sim copies synchronously and has no memory pipeline, so an async copy is a plain memcpy and a barrier is the block barrier. This models the RESULT of the <cuda/barrier> + cooperative_groups::memcpy_async idiom, not its latency hiding.

Defined in code/simulator.h.

This page documents 7 symbols: 1 enums, 6 functions.

Enums

Signature	Description	Availability	Location	Example
enum thread_scope		sim only	simulator.h:3593

Functions

Signature	Description	Availability	Location	Example
arrival_token arrive(std::ptrdiff_t = 1)		sim only	simulator.h:3598	View
void arrive_and_drop()		sim only	simulator.h:3601
void arrive_and_wait()		sim only	simulator.h:3600	View
inline void init(barrier<Scope>* , std::ptrdiff_t )		sim only	simulator.h:3603	View
inline void memcpy_async(void* dst, const void* src, size_t n, B& )		sim only	simulator.h:3604	View
void wait(arrival_token&&) const		sim only	simulator.h:3599	View

Examples

Code from the repository’s example corpus (code\*.cu); the highlighted line is the call site. These blocks are what the View links above open.

arrive (test_web_cs_awbarrier.cu:72)

    SHARED_EXTERN(double, tmp);        // [SIM] was: extern __shared__ double tmp[];

#pragma unroll
    for (int offset = tile32.size() / 2; offset > 0; offset /= 2) {
        threadSum += tile32.shfl_down(threadSum, offset);
    }
    if (tile32.thread_rank() == 0) {
        tmp[tile32.meta_group_rank()] = threadSum;
    }

    auto token = barrier.arrive();

    barrier.wait(std::move(token));

    // The warp 0 will perform last round of reduction
    if (tile32.meta_group_rank() == 0) {
        double beta = tile32.thread_rank() < tile32.meta_group_size()
                          ? tmp[tile32.thread_rank()]
                          : 0.0;

#pragma unroll
        for (int offset = tile32.size() / 2; offset > 0; offset /= 2) {

arrive_and_wait (test_web_cs_globaltoshmem.cu:123)

            float4* const       A4s = reinterpret_cast<float4*>(&As[threadIdx.y][t4x]);
            float4* const       B4s = reinterpret_cast<float4*>(&Bs[threadIdx.y][t4x]);
            const float4* const A4  = reinterpret_cast<const float4*>(&A[a + wA * threadIdx.y + t4x]);
            const float4* const B4  = reinterpret_cast<const float4*>(&B[a + wA * threadIdx.y + t4x]);

            cuda::memcpy_async(A4s, A4, sizeof(float4), bar);
            cuda::memcpy_async(B4s, B4, sizeof(float4), bar);
        }

        // Synchronize to make sure the matrices are loaded
        bar.arrive_and_wait();

#pragma unroll
        for (int k = 0; k < BLOCK_SIZE; ++k) {
            Csub += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        }

        // Synchronize before the next iteration overwrites the tiles
        bar.arrive_and_wait();
    }

    int c                                 = wB * BLOCK_SIZE * blockIdx.y + BLOCK_SIZE * blockIdx.x;

init (test_web_cs_awbarrier.cu:108)

                                             double* partialResults, int size) {
#if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 700
    cg::thread_block cta  = cg::this_thread_block();
    cg::grid_group   grid = cg::this_grid();

    cg::thread_block_tile<32> tile32 = cg::tiled_partition<32>(cta);

    __shared__ cuda::barrier<cuda::thread_scope_block> barrier;

    if (threadIdx.x == 0) {
        init(&barrier, blockDim.x);
    }

    cg::sync(cta);

    double threadSum = 0.0;
    for (int i = grid.thread_rank(); i < size; i += grid.size()) {
        threadSum += (double)(vecA[i] * vecB[i]);
    }

    // Each thread block performs reduction of partial dotProducts and writes to
    // global mem.

memcpy_async (test_web_cs_globaltoshmem.cu:118)

    const int t4x = threadIdx.x * 4;

    for (int a = aBegin, b = bBegin; a <= aEnd; a += aStep, b += bStep) {
        // Now, one fourth of the threads load four elements of each matrix
        if (t4x < BLOCK_SIZE) {
            float4* const       A4s = reinterpret_cast<float4*>(&As[threadIdx.y][t4x]);
            float4* const       B4s = reinterpret_cast<float4*>(&Bs[threadIdx.y][t4x]);
            const float4* const A4  = reinterpret_cast<const float4*>(&A[a + wA * threadIdx.y + t4x]);
            const float4* const B4  = reinterpret_cast<const float4*>(&B[a + wA * threadIdx.y + t4x]);

            cuda::memcpy_async(A4s, A4, sizeof(float4), bar);
            cuda::memcpy_async(B4s, B4, sizeof(float4), bar);
        }

        // Synchronize to make sure the matrices are loaded
        bar.arrive_and_wait();

#pragma unroll
        for (int k = 0; k < BLOCK_SIZE; ++k) {
            Csub += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        }

wait (test_web_cs_awbarrier.cu:74)

#pragma unroll
    for (int offset = tile32.size() / 2; offset > 0; offset /= 2) {
        threadSum += tile32.shfl_down(threadSum, offset);
    }
    if (tile32.thread_rank() == 0) {
        tmp[tile32.meta_group_rank()] = threadSum;
    }

    auto token = barrier.arrive();

    barrier.wait(std::move(token));

    // The warp 0 will perform last round of reduction
    if (tile32.meta_group_rank() == 0) {
        double beta = tile32.thread_rank() < tile32.meta_group_size()
                          ? tmp[tile32.thread_rank()]
                          : 0.0;

#pragma unroll
        for (int offset = tile32.size() / 2; offset > 0; offset /= 2) {
            beta += tile32.shfl_down(beta, offset);
        }