Unified-memory stream attachment (advisory in one address space)

Defined in code/simulator.h.

This page documents 4 symbols: 3 macros, 1 functions.

Macros

Signature	Description	Availability	Location	Example
cudaMemAttachGlobal	Expands to 0x01u	sim only	simulator.h:2465
cudaMemAttachHost	Expands to 0x02u	sim only	simulator.h:2468
cudaMemAttachSingle	Expands to 0x04u	sim only	simulator.h:2471

Functions

Signature	Description	Availability	Location	Example
inline cudaError_t cudaStreamAttachMemAsync(cudaStream_t , void* , size_t = 0, unsigned = cudaMemAttachSingle)		sim only	simulator.h:2473	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.

cudaStreamAttachMemAsync (test_web_cs_um_streams.cu:84)

        result[i] *= beta;
        for (int j = 0; j < n; j++) result[i] += alpha * A[i * n + j] * x[j];
    }
}

// Run one task: small ones on the host, large ones through cuBLAS on a stream.
template <typename T>
static void execute(Task<T>& t, cublasHandle_t* handle, cudaStream_t* stream, int tid) {
    if (t.size < 100) {
        std::printf("Task [%u], thread [%d] executing on host (%u)\n", t.id, tid, t.size);
        check(cudaStreamAttachMemAsync(stream[0], t.data,   0, cudaMemAttachHost), "attach data");
        check(cudaStreamAttachMemAsync(stream[0], t.vector, 0, cudaMemAttachHost), "attach vector");
        check(cudaStreamAttachMemAsync(stream[0], t.result, 0, cudaMemAttachHost), "attach result");
        check(cudaStreamSynchronize(stream[0]), "cudaStreamSynchronize");
        gemv(int(t.size), int(t.size), T(1), t.data, t.vector, T(0), t.result);
    } else {
        std::printf("Task [%u], thread [%d] executing on device (%u)\n", t.id, tid, t.size);
        const double one = 1.0, zero = 0.0;
        cublasSetStream(handle[tid + 1], stream[tid + 1]);
        check(cudaStreamAttachMemAsync(stream[tid + 1], t.data,   0, cudaMemAttachSingle), "attach data");
        check(cudaStreamAttachMemAsync(stream[tid + 1], t.vector, 0, cudaMemAttachSingle), "attach vector");
        check(cudaStreamAttachMemAsync(stream[tid + 1], t.result, 0, cudaMemAttachSingle), "attach result");