Device memory#

A plain CPU heap, but tracked so the host/device split is real: device pointers come from cudaMalloc and are meant to travel through cudaMemcpy.

Defined in code/simulator.h.

This page documents 6 symbols: 6 functions.

Functions#

Signature

Description

Availability

Location

Example

inline std::vector<Alloc>& allocs()

sim only

simulator.h:1957

inline cudaError_t cudaFree(void* devPtr)

Free a pointer previously returned by cudaMalloc; errors if it wasn’t one.

sim only

simulator.h:1973

View

inline cudaError_t cudaMalloc(T** devPtr, size_t size)

Allocate size bytes of “device” memory (a tracked host malloc).

sim only

simulator.h:1963

View

inline cudaError_t cudaMemcpy(void* dst, const void* src, size_t size, cudaMemcpyKind )

Copy size bytes. Host and device share one address space here, so the direction (kind) is accepted for portability but not otherwise needed.

sim only

simulator.h:1989

View

inline cudaError_t cudaMemcpyFromSymbol(void* dst, const void* symbol, size_t count, size_t offset = 0, cudaMemcpyKind = cudaMemcpyDeviceToHost)

sim only

simulator.h:2002

View

inline cudaError_t cudaMemcpyToSymbol(void* symbol, const void* src, size_t count, size_t offset = 0, cudaMemcpyKind = cudaMemcpyHostToDevice)

__constant__/__device__ symbols are ordinary globals in the sim; the symbol decays to its address (works for the common array-symbol case).

sim only

simulator.h:1997

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.

cudaFree (test_atomics.cu:103)
 93static T* devFrom(const T& init) {
 94    T* d = nullptr;
 95    CUDA_CHECK(cudaMalloc(&d, sizeof(T)));
 96    CUDA_CHECK(cudaMemcpy(d, &init, sizeof(T), cudaMemcpyHostToDevice));
 97    return d;
 98}
 99template <class T>
100static T devRead(T* d) {
101    T h{};
102    CUDA_CHECK(cudaMemcpy(&h, d, sizeof(T), cudaMemcpyDeviceToHost));
103    CUDA_CHECK(cudaFree(d));
104    return h;
105}
106
107static bool report(const char* name, bool ok, unsigned seed) {
108    std::printf("  %-28s seed=%-5u %s\n", name, seed, ok ? "PASSED" : "FAILED");
109    return ok;
110}
111
112// Runs the full battery once under a given schedule seed.
113static bool runAll(unsigned seed) {
114#ifndef __CUDACC__
cudaMalloc (test_atomics.cu:95)
 85__global__ void kSub(int* c) {
 86    if (blockIdx.x * blockDim.x + threadIdx.x < N) atomicSub(c, 1);
 87}
 88
 89// ---- Host driver -----------------------------------------------------------
 90
 91template <class T>
 92static T* devFrom(const T& init) {
 93    T* d = nullptr;
 94    CUDA_CHECK(cudaMalloc(&d, sizeof(T)));
 95    CUDA_CHECK(cudaMemcpy(d, &init, sizeof(T), cudaMemcpyHostToDevice));
 96    return d;
 97}
 98template <class T>
 99static T devRead(T* d) {
100    T h{};
101    CUDA_CHECK(cudaMemcpy(&h, d, sizeof(T), cudaMemcpyDeviceToHost));
102    CUDA_CHECK(cudaFree(d));
103    return h;
104}
cudaMemcpy (test_atomics.cu:96)
 86__global__ void kSub(int* c) {
 87    if (blockIdx.x * blockDim.x + threadIdx.x < N) atomicSub(c, 1);
 88}
 89
 90// ---- Host driver -----------------------------------------------------------
 91
 92template <class T>
 93static T* devFrom(const T& init) {
 94    T* d = nullptr;
 95    CUDA_CHECK(cudaMalloc(&d, sizeof(T)));
 96    CUDA_CHECK(cudaMemcpy(d, &init, sizeof(T), cudaMemcpyHostToDevice));
 97    return d;
 98}
 99template <class T>
100static T devRead(T* d) {
101    T h{};
102    CUDA_CHECK(cudaMemcpy(&h, d, sizeof(T), cudaMemcpyDeviceToHost));
103    CUDA_CHECK(cudaFree(d));
104    return h;
105}
106
107static bool report(const char* name, bool ok, unsigned seed) {
cudaMemcpyFromSymbol (test_portability.cu:99)
 89        float* dOut = nullptr; CUDA_CHECK(cudaMalloc(&dOut, N * sizeof(float)));
 90        LAUNCH(kScale, blocks, BLOCK, dIn, dOut, N);
 91        std::vector<float> got(N); CUDA_CHECK(cudaMemcpy(got.data(), dOut, N * sizeof(float), cudaMemcpyDeviceToHost));
 92        bool ok = true;
 93        for (int i = 0; i < N; ++i) if (got[i] != in[i] * scale[i & 3]) ok = false;
 94        all &= report("__constant__ / MemcpyToSymbol", ok);
 95        CUDA_CHECK(cudaFree(dIn)); CUDA_CHECK(cudaFree(dOut));
 96
 97        // Round-trip the symbol back out.
 98        float back[4] = {0, 0, 0, 0};
 99        CUDA_CHECK(cudaMemcpyFromSymbol(back, cScale, sizeof(back)));
100        bool rok = back[0] == 0.5f && back[1] == 1.0f && back[2] == 2.0f && back[3] == 4.0f;
101        all &= report("cudaMemcpyFromSymbol", rok);
102    }
103
104    // float4 vector-type SAXPY.
105    {
106        const int n4 = N / 4;
107        const float a = 3.0f;
108        std::vector<float4> x(n4), y(n4);
109        for (int i = 0; i < n4; ++i) {
110            x[i] = make_float4((float)i, (float)(i + 1), (float)(i + 2), (float)(i + 3));
cudaMemcpyToSymbol (test_portability.cu:84)
74    cudaDeviceProp prop{};
75    CUDA_CHECK(cudaGetDeviceProperties(&prop, 0));
76    std::printf("Device 0: %s (compute %d.%d)\n\n", prop.name, prop.major, prop.minor);
77
78    bool all = true;
79    const int blocks = N / BLOCK;
80
81    // __constant__ + cudaMemcpyToSymbol.
82    {
83        const float scale[4] = {0.5f, 1.0f, 2.0f, 4.0f};
84        CUDA_CHECK(cudaMemcpyToSymbol(cScale, scale, sizeof(scale)));
85
86        std::vector<float> in(N);
87        for (int i = 0; i < N; ++i) in[i] = static_cast<float>(i % 11);
88        float* dIn = upload(in);
89        float* dOut = nullptr; CUDA_CHECK(cudaMalloc(&dOut, N * sizeof(float)));
90        LAUNCH(kScale, blocks, BLOCK, dIn, dOut, N);
91        std::vector<float> got(N); CUDA_CHECK(cudaMemcpy(got.data(), dOut, N * sizeof(float), cudaMemcpyDeviceToHost));
92        bool ok = true;
93        for (int i = 0; i < N; ++i) if (got[i] != in[i] * scale[i & 3]) ok = false;
94        all &= report("__constant__ / MemcpyToSymbol", ok);
95        CUDA_CHECK(cudaFree(dIn)); CUDA_CHECK(cudaFree(dOut));
On this page