`cudaMemcpy`: Host-Device Data Transfer

The Bridge Between Host and Device: cudaMemcpy

Just as you would use memcpy in C to copy blocks of memory on the host, CUDA provides an analogous function, cudaMemcpy, to move data between the host's and the device's distinct memory spaces. Understanding its structure is straightforward, especially given your background in C-style APIs.

Let's start with a video that introduces the function and its parameters.

As the video highlighted, the function signature is:

cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, cudaMemcpyKind kind);

Let's break this down:

1. void* dst: A pointer to the destination memory address.
2. const void* src: A pointer to the source memory address.
3. size_t count: The number of bytes to copy.
4. cudaMemcpyKind kind: This special enumerator tells CUDA the direction of the transfer.

The kind parameter is what makes cudaMemcpy so versatile. For our purposes, the two most important values are:

• cudaMemcpyHostToDevice: Copies data from the host (CPU RAM) to the device (GPU VRAM).
• cudaMemcpyDeviceToHost: Copies data from the device back to the host.

Other kinds exist, such as cudaMemcpyDeviceToDevice for on-GPU copies and cudaMemcpyDefault which lets the driver infer the direction, but we will focus on the explicit host-device transfers.

The following reading provides a detailed textual reference for cudaMemcpy and places it within the context of a complete program.

The Complete Data Management Workflow

With cudaMalloc, cudaMemcpy, and cudaFree, we now have all the tools for a complete data management cycle. A typical CUDA program follows this explicit pattern:

1. Allocate and initialize data on the host. This is your input data, created using standard C++ mechanisms like new or malloc.
2. Allocate memory on the device. Use cudaMalloc to create space on the GPU for both the input data and the eventual output.
3. Copy input data to the device. Use cudaMemcpy with cudaMemcpyHostToDevice to move your input data from the host to the device.
4. (Kernel Execution) Launch a GPU kernel to process the data residing on the device. (We will cover this in the next lesson).
5. Copy results back to the host. Use cudaMemcpy with cudaMemcpyDeviceToHost to retrieve the computed results.
6. Clean up. Free the memory on both the device (cudaFree) and the host (free or delete[]).

Let's look at a complete code example that implements this pattern. We won't launch a kernel yet, but we will perform a "round trip": copy data to the device and then immediately copy it back to a different host buffer to verify that the transfer was successful.

#include <iostream>
#include <vector>
#include <numeric> // For std::iota
#include <cuda_runtime.h>

// Helper function from the previous lesson
void checkCudaError(cudaError_t err, const char* msg) {
    if (err != cudaSuccess) {
        std::cerr << "CUDA Error: " << msg << " (" << cudaGetErrorString(err) << ")" << std::endl;
        exit(EXIT_FAILURE);
    }
}

int main() {
    // 1. Allocate and initialize data on the host
    const int N = 1024;
    const size_t size = N * sizeof(int);
    std::vector<int> h_input(N);
    std::vector<int> h_output(N, 0); // Initialize output vector with zeros

    // Fill host input with 0, 1, 2, ..., N-1
    std::iota(h_input.begin(), h_input.end(), 0);

    // 2. Allocate memory on the device
    int* d_data = nullptr;
    checkCudaError(cudaMalloc((void**)&d_data, size), "cudaMalloc failed");

    // 3. Copy input data to the device
    std::cout << "Copying data from Host to Device..." << std::endl;
    checkCudaError(
        cudaMemcpy(d_data, h_input.data(), size, cudaMemcpyHostToDevice),
        "cudaMemcpy Host-to-Device failed"
    );

    // --- In a real program, a kernel launch would happen here ---

    // 5. Copy data back to the host (into a separate output buffer)
    std::cout << "Copying data from Device to Host..." << std::endl;
    checkCudaError(
        cudaMemcpy(h_output.data(), d_data, size, cudaMemcpyDeviceToHost),
        "cudaMemcpy Device-to-Host failed"
    );

    // 6. Verification and Cleanup
    bool success = true;
    for (int i = 0; i < N; ++i) {
        if (h_input[i] != h_output[i]) {
            std::cerr << "Mismatch at index " << i << ": host_input=" << h_input[i] << ", host_output=" << h_output[i] << std::endl;
            success = false;
            break;
        }
    }

    if (success) {
        std::cout << "Success! Data round trip verified." << std::endl;
    }

    checkCudaError(cudaFree(d_data), "cudaFree failed");
    // h_input and h_output are std::vectors, so their memory is managed automatically

    return 0;
}

This example solidifies the entire sequence. I encourage you to compile and run this code to see the data management cycle in action.

A Note on Performance: Synchronous Transfers and Pinned Memory

An important characteristic of cudaMemcpy is that it is a synchronous (or blocking) operation with respect to the host. When your CPU executes a cudaMemcpy call, it pauses and waits until the entire data transfer is complete before moving on to the next line of code. This ensures data is where you expect it to be, but it can be a performance bottleneck. We will later explore asynchronous transfers using CUDA streams to overlap data movement with computation.

Furthermore, by default, host memory allocated with new or malloc is pageable. This means the operating system can move it around in physical memory. Before the GPU can access this data, the CUDA driver must first copy it into a special non-pageable, or pinned, memory area. This staging copy happens implicitly and adds overhead.

You can eliminate this extra copy by directly allocating pinned memory using cudaMallocHost(). Data transfers from pinned memory are significantly faster. While we will focus on this optimization in a later module, it's a key concept to be aware of as you build your mental model of performance engineering.

The following short video introduces this concept and demonstrates the performance difference.

For now, continue using standard pageable memory. Mastering the fundamental cudaMemcpy pattern is the priority. We will revisit pinned memory when we begin optimizing our kernels.

Conclusion

You have now mastered the essential functions for managing data in a CUDA application. By combining the allocation from the last lesson with the transfer operations from this one, you can now fully control the flow of data between the CPU and GPU.

Key Takeaways:

• cudaMemcpy is the primary function for transferring data between the host and device.
• The function requires four arguments: destination, source, size in bytes, and a cudaMemcpyKind to specify the direction (cudaMemcpyHostToDevice or cudaMemcpyDeviceToHost).
• The standard CUDA workflow is a cycle: allocate on host and device, copy to device, compute, copy results back to host, and free memory.
• By default, cudaMemcpy is a synchronous operation, which is simple but can impact performance.
• Using standard pageable memory incurs an extra copy overhead, which can be eliminated with pinned memory for higher performance—a topic for future optimization.

We are now fully prepared for the most exciting part. With the ability to allocate memory on the GPU and populate it with data, our next lesson will focus on writing and launching your first CUDA kernel to perform a parallel computation on that data.