cuda::pipeline#

The multi-stage async-copy API. The sim copies synchronously and has no memory pipeline, so acquire/commit/wait/release are bookkeeping no-ops and a staged copy simply completes before consumer_wait() returns. That models the RESULT of double/multi-buffering exactly; it cannot model the latency hiding, which is the whole point of the API on real hardware.

Defined in code/simulator.h.

This page documents 19 symbols: 1 enums, 18 functions.

Enums#

Signature

Description

Availability

Location

Example

enum pipeline_role

The multi-stage async-copy API. The sim copies synchronously and has no memory pipeline, so acquire/commit/wait/release are bookkeeping no-ops and a staged copy simply completes before …

sim only

simulator.h:3614

Functions#

Signature

Description

Availability

Location

Example

explicit aligned_size_t(size_t s) : v_(s)

sim only

simulator.h:3644

void consumer_release()

sim only

simulator.h:3622

void consumer_wait()

sim only

simulator.h:3621

bool consumer_wait_for(const Rep&, const Period&)

sim only

simulator.h:3623

inline pipeline<thread_scope_thread> make_pipeline()

sim only

simulator.h:3631

inline pipeline<Scope> make_pipeline(const Group&, pipeline_shared_state<Scope, N>*)

sim only

simulator.h:3632

inline pipeline<Scope> make_pipeline(const Group&, pipeline_shared_state<Scope, N>*, pipeline_role)

sim only

simulator.h:3634

inline void memcpy_async(void* dst, const void* src, Shape size, pipeline<Scope>&)

The pipeline-taking memcpy_async overloads (the barrier-taking one is above).

sim only

simulator.h:3649

View

inline void memcpy_async(const Group&, void* dst, const void* src, Shape size, pipeline<Scope>&)

sim only

simulator.h:3653

View

inline void memcpy_async(const Group&, T* dst, const T* src, Shape size, pipeline<Scope>&)

sim only

simulator.h:3657

View

inline void memcpy_async(const G& , T* dst, const T* src, size_t bytes)

sim only

simulator.h:3664

View

inline void memcpy_async(const G& , T* dst, const T* src, size_t bytes, B& )

sim only

simulator.h:3667

View

void producer_acquire()

sim only

simulator.h:3619

void producer_commit()

sim only

simulator.h:3620

bool quit()

sim only

simulator.h:3624

operator size_t() const

sim only

simulator.h:3645

View

inline void wait(const G& g)

sim only

simulator.h:3670

View

inline void wait_prior(const G& g)

sim only

simulator.h:3671

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.

memcpy_async (test_web_cs_globaltoshmem.cu:118)
108    const int t4x = threadIdx.x * 4;
109
110    for (int a = aBegin, b = bBegin; a <= aEnd; a += aStep, b += bStep) {
111        // Now, one fourth of the threads load four elements of each matrix
112        if (t4x < BLOCK_SIZE) {
113            float4* const       A4s = reinterpret_cast<float4*>(&As[threadIdx.y][t4x]);
114            float4* const       B4s = reinterpret_cast<float4*>(&Bs[threadIdx.y][t4x]);
115            const float4* const A4  = reinterpret_cast<const float4*>(&A[a + wA * threadIdx.y + t4x]);
116            const float4* const B4  = reinterpret_cast<const float4*>(&B[a + wA * threadIdx.y + t4x]);
117
118            cuda::memcpy_async(A4s, A4, sizeof(float4), bar);
119            cuda::memcpy_async(B4s, B4, sizeof(float4), bar);
120        }
121
122        // Synchronize to make sure the matrices are loaded
123        bar.arrive_and_wait();
124
125#pragma unroll
126        for (int k = 0; k < BLOCK_SIZE; ++k) {
127            Csub += As[threadIdx.y][k] * Bs[k][threadIdx.x];
128        }
memcpy_async (test_web_cs_globaltoshmem.cu:118)
108    const int t4x = threadIdx.x * 4;
109
110    for (int a = aBegin, b = bBegin; a <= aEnd; a += aStep, b += bStep) {
111        // Now, one fourth of the threads load four elements of each matrix
112        if (t4x < BLOCK_SIZE) {
113            float4* const       A4s = reinterpret_cast<float4*>(&As[threadIdx.y][t4x]);
114            float4* const       B4s = reinterpret_cast<float4*>(&Bs[threadIdx.y][t4x]);
115            const float4* const A4  = reinterpret_cast<const float4*>(&A[a + wA * threadIdx.y + t4x]);
116            const float4* const B4  = reinterpret_cast<const float4*>(&B[a + wA * threadIdx.y + t4x]);
117
118            cuda::memcpy_async(A4s, A4, sizeof(float4), bar);
119            cuda::memcpy_async(B4s, B4, sizeof(float4), bar);
120        }
121
122        // Synchronize to make sure the matrices are loaded
123        bar.arrive_and_wait();
124
125#pragma unroll
126        for (int k = 0; k < BLOCK_SIZE; ++k) {
127            Csub += As[threadIdx.y][k] * Bs[k][threadIdx.x];
128        }
memcpy_async (test_web_cs_globaltoshmem.cu:118)
108    const int t4x = threadIdx.x * 4;
109
110    for (int a = aBegin, b = bBegin; a <= aEnd; a += aStep, b += bStep) {
111        // Now, one fourth of the threads load four elements of each matrix
112        if (t4x < BLOCK_SIZE) {
113            float4* const       A4s = reinterpret_cast<float4*>(&As[threadIdx.y][t4x]);
114            float4* const       B4s = reinterpret_cast<float4*>(&Bs[threadIdx.y][t4x]);
115            const float4* const A4  = reinterpret_cast<const float4*>(&A[a + wA * threadIdx.y + t4x]);
116            const float4* const B4  = reinterpret_cast<const float4*>(&B[a + wA * threadIdx.y + t4x]);
117
118            cuda::memcpy_async(A4s, A4, sizeof(float4), bar);
119            cuda::memcpy_async(B4s, B4, sizeof(float4), bar);
120        }
121
122        // Synchronize to make sure the matrices are loaded
123        bar.arrive_and_wait();
124
125#pragma unroll
126        for (int k = 0; k < BLOCK_SIZE; ++k) {
127            Csub += As[threadIdx.y][k] * Bs[k][threadIdx.x];
128        }
memcpy_async (test_web_cs_globaltoshmem.cu:118)
108    const int t4x = threadIdx.x * 4;
109
110    for (int a = aBegin, b = bBegin; a <= aEnd; a += aStep, b += bStep) {
111        // Now, one fourth of the threads load four elements of each matrix
112        if (t4x < BLOCK_SIZE) {
113            float4* const       A4s = reinterpret_cast<float4*>(&As[threadIdx.y][t4x]);
114            float4* const       B4s = reinterpret_cast<float4*>(&Bs[threadIdx.y][t4x]);
115            const float4* const A4  = reinterpret_cast<const float4*>(&A[a + wA * threadIdx.y + t4x]);
116            const float4* const B4  = reinterpret_cast<const float4*>(&B[a + wA * threadIdx.y + t4x]);
117
118            cuda::memcpy_async(A4s, A4, sizeof(float4), bar);
119            cuda::memcpy_async(B4s, B4, sizeof(float4), bar);
120        }
121
122        // Synchronize to make sure the matrices are loaded
123        bar.arrive_and_wait();
124
125#pragma unroll
126        for (int k = 0; k < BLOCK_SIZE; ++k) {
127            Csub += As[threadIdx.y][k] * Bs[k][threadIdx.x];
128        }
memcpy_async (test_web_cs_globaltoshmem.cu:118)
108    const int t4x = threadIdx.x * 4;
109
110    for (int a = aBegin, b = bBegin; a <= aEnd; a += aStep, b += bStep) {
111        // Now, one fourth of the threads load four elements of each matrix
112        if (t4x < BLOCK_SIZE) {
113            float4* const       A4s = reinterpret_cast<float4*>(&As[threadIdx.y][t4x]);
114            float4* const       B4s = reinterpret_cast<float4*>(&Bs[threadIdx.y][t4x]);
115            const float4* const A4  = reinterpret_cast<const float4*>(&A[a + wA * threadIdx.y + t4x]);
116            const float4* const B4  = reinterpret_cast<const float4*>(&B[a + wA * threadIdx.y + t4x]);
117
118            cuda::memcpy_async(A4s, A4, sizeof(float4), bar);
119            cuda::memcpy_async(B4s, B4, sizeof(float4), bar);
120        }
121
122        // Synchronize to make sure the matrices are loaded
123        bar.arrive_and_wait();
124
125#pragma unroll
126        for (int k = 0; k < BLOCK_SIZE; ++k) {
127            Csub += As[threadIdx.y][k] * Bs[k][threadIdx.x];
128        }
size_t (test_web_bf_oversub2.cu:47)
37    if (err != cudaSuccess || deviceCount == 0) {
38        std::printf("No CUDA device available (error %d: %s).\n"
39                    "Compilation and toolkit are fine; kernels need an NVIDIA GPU.\n",
40                    static_cast<int>(err), cudaGetErrorString(err));
41        return 0;
42    }
43    cudaDeviceProp prop{}; cudaGetDeviceProperties(&prop, 0);
44    std::printf("Device 0: %s (compute %d.%d)\n", prop.name, prop.major, prop.minor);
45
46    // [SIM] source oversubscribes with n GB (n=1..30); here a single modest in-bounds size.
47    const size_t memory_size = size_t(8) * 1024 * 1024;   // 8 MB managed
48    const size_t data_size = memory_size / sizeof(uint64_t);
49
50    uint64_t *x;
51    CHECK(cudaMallocManaged(&x, memory_size));
52
53    // [SIM] pre-fill with a non-zero sentinel so "all zeros" actually proves the kernel ran.
54    for (size_t i = 0; i < data_size; ++i)
55        x[i] = 0xDEADBEEFull;
56
57    LAUNCH(gpu_touch, (unsigned)((data_size - 1) / 1024 + 1), 1024, x, data_size);   // [SIM]
58    CHECK(cudaGetLastError());
wait (test_web_cs_awbarrier.cu:74)
64#pragma unroll
65    for (int offset = tile32.size() / 2; offset > 0; offset /= 2) {
66        threadSum += tile32.shfl_down(threadSum, offset);
67    }
68    if (tile32.thread_rank() == 0) {
69        tmp[tile32.meta_group_rank()] = threadSum;
70    }
71
72    auto token = barrier.arrive();
73
74    barrier.wait(std::move(token));
75
76    // The warp 0 will perform last round of reduction
77    if (tile32.meta_group_rank() == 0) {
78        double beta = tile32.thread_rank() < tile32.meta_group_size()
79                          ? tmp[tile32.thread_rank()]
80                          : 0.0;
81
82#pragma unroll
83        for (int offset = tile32.size() / 2; offset > 0; offset /= 2) {
84            beta += tile32.shfl_down(beta, offset);
85        }
On this page