Manual Memory Management

Before unified memory was available — or whenever full control over data movement is wanted — memory is allocated on the device explicitly with cudaMalloc (C/C++) or the device attribute (CUDA Fortran), and data is copied between host and device with cudaMemcpy or, in CUDA Fortran, plain array assignment. This episode introduces this classical approach, the direction constants of cudaMemcpy and the device memory management API.

Explicit copies also change the synchronisation picture: cudaMemcpy blocks the host until the transfer is complete, so it acts as a built-in synchronisation point. Two exercises put this into practice — the manual-memory counterpart of the previous episode’s vector-scalar multiplication, and a vector addition whose execution configuration you tune for a large vector.

Questions

• What do you gain by allocating and copying GPU memory yourself, when managed memory already works?

• Which parts of an existing host program have to change to move its data to the device?

• How does a copy know which direction it is going in?

Objectives

• Perform explicit memory allocation and data transfers with cudaMalloc/cudaMemcpy or equivalent in Fortran

Instructor note

• 30 min teaching

• 30 min exercises

Classical (manual) memory management

Before unified memory was available (or when full control over data movement is desired), memory must be explicitly allocated on the device via cudaMalloc or __device__ specifier in C/C++ or device attribute in Fortran and data explicitly copied via cudaMemcpy in C/C++/Fortran or regular assignment in Fortran. And also freed again by calling cudaFree in C/C++ or deallocate for Fortran allocatable arrays if not deallocated automatically.

C/C++Fortran

double *Y_d{}, *X_d{}, *Y{}, *X{};
Y = (double*)malloc(1024 * sizeof(double));
X = (double*)malloc(1024 * sizeof(double));
// allocate 1024 * sizeof(double) bytes on device, return device pointer in first argument
cudaMalloc((void**)&Y_d, 1024 * sizeof(double));
cudaMalloc((void**)&X_d, 1024 * sizeof(double));

// Copy 1024 * sizeof(double) bytes from X at host → device X_d
cudaMemcpy(X_d, X, 1024 * sizeof(double), cudaMemcpyDefault);
cudaMemcpy(Y_d, Y, 1024 * sizeof(double), cudaMemcpyDefault);

add_kernel<<<4, 256>>>(X_d, Y_d, 1024);
// Explicit sychronisation; Not necessary in this case, see comment below.
cudaDeviceSynchronize();

// Copy device → host
cudaMemcpy(X, X_d, 1024 * sizeof(double), cudaMemcpyDefault);

cudaFree(Y_d); free(Y);
cudaFree(X_d); free(X);

The last argument of cudaMemcpy specifies the transfer direction:

Direction	Constant
Host → Host	cudaMemcpyHostToHost
Host → Device	cudaMemcpyHostToDevice
Device → Host	cudaMemcpyDeviceToHost
Device → Device	cudaMemcpyDeviceToDevice
Auto-detect (requires UVA)	cudaMemcpyDefault

Tip

With cudaMemcpyDefault (requires UVA) the runtime detects the transfer direction from the two pointers, as in the example above; the explicit constants document the intended direction.

Advantages of manual management: explicit control over when data moves; cudaMemcpy provides a synchronisation point, i.e. cudaMemcpy blocks the host if data has to be transferred from host to device or vice versa. One can think of it as having a built-in cudaDeviceSynchronize in this case. Analog for corresponding assignments in CUDA Fortran.

Disadvantage: the programmer must manage all transfers.

Warning

Segmentation faults are the typical symptom of incorrect memory handling: a host pointer passed to a kernel, or a device pointer dereferenced in host code, leads to incorrect results or a crash. Keep track of which pointer refers to which memory (see the pointer rules in the last episode).

Device memory management API

C/C++Fortran

cudaError_t cudaMalloc(void** devPtr, size_t size);
// Allocate size bytes of device memory

cudaError_t cudaFree(void* devPtr);

cudaError_t cudaMemcpy(void* dst, const void* src,
                        size_t count, cudaMemcpyKind kind);
// Copy count bytes from src to dst

cudaError_t cudaMemset(void* devPtr, int value, size_t count);
// Set count bytes to pattern value

Exercises

Vector-Scalar Multiplication with Manual Memory

This is the manual-memory counterpart to the managed-memory exercise from the previous episode. The CPU code below computes c[i] = a[i] * b for a vector.

Tasks:

1. Turn eval into a CUDA kernel (__global__ in C/C++, attributes(global) in CUDA Fortran). In CUDA Fortran, also give the scalar arguments b and size the value attribute.

2. Allocate device memory for the input and output vectors with cudaMalloc (in CUDA Fortran: declare them with device attribute).

3. Copy the input vector to the device with cudaMemcpy (or array assignment in CUDA Fortran).

4. Launch the kernel.

5. Copy the result back to the host.

6. Add an appropriate synchronisation where needed.

Downloadable templates: C/C++ — CUDA Fortran. The C/C++ version needs cuda_utils.h; place it in the same directory as the .cu file (once per working directory).

C/C++CUDA Fortran

#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include "cuda_utils.h"

// CUDA Kernel //
// TODO: make eval function a CUDA kernel
void eval(float* a, float b, float* c, int size) {
    for (int i = 0; i < size; i++) {
        c[i] = a[i] * b;
    }
}

int main() {
    srand48(time(NULL));
    const int vector_size = 10000;

    // allocate memory on cpu //
    float* vector_a_cpu;
    float* vector_c_cpu;
    // TODO: allocate host memory

    float b = drand48();
    for (int i = 0; i < vector_size; i++) vector_a_cpu[i] = drand48();
    for (int i = 0; i < vector_size; i++) vector_c_cpu[i] = 0;

    // allocate memory on gpu //
    float* vector_a_gpu;
    float* vector_c_gpu;
    // TODO: allocate device memory

    // TODO: copy input data to gpu

    // TODO: add kernel call

    // Synchronization needed?
    // cudaDeviceSynchronize();

    // TODO: copy results back to host

    // verify results, free host / device memory ...
}

Solution

The key changes are:

• Allocate device pointers with cudaMalloc (C/C++) or declare arrays with the device attribute (CUDA Fortran).

• Add a cudaMemcpy from host to device for the input and from device to host for the result (in CUDA Fortran, array assignment between host and device arrays implies a copy).

• Mark the computation function with __global__ / attributes(global).

• Launch with eval<<<1, 1>>>(...) as a starting point.

• cudaDeviceSynchronize() is not strictly required before the cudaMemcpy of the result because cudaMemcpy is itself synchronising on the default stream. Calling it makes the dependency explicit. In the vector-addition exercise below it is needed so that the timer measures the kernel and not the copy. The same holds for Fortran.

Full solutions: C/C++ — CUDA Fortran. The C/C++ version needs cuda_utils.h; place it in the same directory as the .cu file (once per working directory).

Vector Addition with Manual Memory

Same primitives as the previous exercise, on a different kernel: element-wise vector addition c[i] = a[i] + b[i].

In this template the host allocation, device allocation and copies are already filled in. The remaining task is the execution configuration:

1. Find an execution configuration (number of blocks, threads per block) that performs well for vector_size = 100 000 000.

2. Compare the reported bandwidth and kernel time to the trivial <<<4, 4>>> launch supplied in the template.

Note

The full templates are long because they include timing instrumentation; please open them via the download link rather than reading them inline.

Downloadable templates: C/C++ — CUDA Fortran. The C/C++ version needs cuda_utils.h; place it in the same directory as the .cu file (once per working directory).

Solution

A useful starting point for an A100/H100-class GPU is <<<2048, 512>>> (see the reference solution). On other hardware the optimum will be different — the point of the exercise is to compare a few configurations and observe the bandwidth.

Full solutions: C/C++ — CUDA Fortran. The C/C++ version needs cuda_utils.h; place it in the same directory as the .cu file (once per working directory).

Keypoints

• Manual memory management (cudaMalloc + cudaMemcpy or device attribute and assignments respectively) gives explicit control over data movement

See also

• Explicit memory management [3]

• Further reading for this episode: Episode 2: Manual memory management