BraneTechnologies Documentation

HIP Execution Model #

At its core, HIP organizes computation across a hierarchy of execution units. Understanding this model is essential for writing efficient GPU code:

The CPU (host) manages the overall execution flow – allocating memory, launching kernels, and coordinating the GPU’s work. The GPU (device) handles the actual parallel execution of your code. Your computation is organized into a grid, which contains multiple blocks, each with their own cluster of threads. On AMD hardware, threads execute in groups called wavefronts, while NVIDIA uses the term warps for the same concept.

This hierarchical arrangement enables massive parallelism – potentially thousands or millions of threads working simultaneously on different pieces of your data. Each thread can execute its portion of a kernel independently, and blocks operate without synchronizing with each other unless you explicitly program such coordination through global memory operations or stream synchronization.

Think of this hierarchy like an organizational chart: the grid is the company, blocks are departments, and threads are individual workers. Each worker follows the same instructions (your kernel code) but operates on different data.

Example: Vector Addition with HIP #

#include <hip/hip_runtime.h>
#include <iostream>

// HIP kernel to perform vector addition
__global__ void vector_add(const float* A, const float* B, float* C, int N) {
    // Calculate global thread ID
    int id = blockIdx.x * blockDim.x + threadIdx.x;
    
    // Ensure we don't access beyond array bounds
    if (id < N) {
        C[id] = A[id] + B[id]; // Perform addition
    }
}

int main() {
    int N = 1024; // Number of elements
    size_t size = N * sizeof(float);
    
    // Error handling variable
    hipError_t err;

    // Allocate memory on the GPU
    float *A, *B, *C;
    err = hipMallocManaged(&A, size);
    if (err != hipSuccess) {
        std::cerr << "Failed to allocate memory for A: " << hipGetErrorString(err) << std::endl;
        return -1;
    }
    
    err = hipMallocManaged(&B, size);
    if (err != hipSuccess) {
        std::cerr << "Failed to allocate memory for B: " << hipGetErrorString(err) << std::endl;
        hipFree(A);
        return -1;
    }
    
    err = hipMallocManaged(&C, size);
    if (err != hipSuccess) {
        std::cerr << "Failed to allocate memory for C: " << hipGetErrorString(err) << std::endl;
        hipFree(A);
        hipFree(B);
        return -1;
    }

    // Initialize input data
    for (int i = 0; i < N; i++) {
        A[i] = i;
        B[i] = i * 2;
    }

    // Define grid and block dimensions
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;

    // Launch the kernel
    hipLaunchKernelGGL(vector_add, 
                       dim3(blocksPerGrid), 
                       dim3(threadsPerBlock), 
                       0,    // Shared memory bytes (none used here)
                       0,    // Stream ID (default stream)
                       A, B, C, N);
    
    // Check for kernel launch errors
    err = hipGetLastError();
    if (err != hipSuccess) {
        std::cerr << "Kernel launch failed: " << hipGetErrorString(err) << std::endl;
        hipFree(A);
        hipFree(B);
        hipFree(C);
        return -1;
    }
    
    // Wait for the kernel to finish
    hipDeviceSynchronize();

    // Print some results
    std::cout << "C[0] = " << C[0] << std::endl;
    std::cout << "C[N-1] = " << C[N-1] << std::endl;

    // Free memory
    hipFree(A);
    hipFree(B);
    hipFree(C);

    return 0;
}

This example demonstrates several key concepts. First, we define our kernel function with the __global__ qualifier, indicating it runs on the GPU but is called from the host. Inside the kernel, each thread calculates its global ID based on its position in the block and grid, then performs work only if that ID is within our data bounds.

When launching the kernel, we specify how to distribute the computation across the GPU. The blocksPerGrid value determines how many blocks to create, while threadsPerBlock specifies the number of threads per block. This example uses one-dimensional organization, but HIP also supports 2D and 3D arrangements for problems with spatial locality.

Notice how we handle error checking throughout the code. GPU programming can be unforgiving, and proper error handling helps identify issues early in development. The hipDeviceSynchronize() call ensures all GPU operations complete before we access the results, an important synchronization point between host and device.

Example: Different Memory Allocation Strategies #

Here’s how these different strategies look in practice:

// Device memory with explicit transfers
float* deviceArray;
hipMalloc(&deviceArray, size);
hipMemcpy(deviceArray, hostArray, size, hipMemcpyHostToDevice);
// Launch kernel with deviceArray
hipMemcpy(hostArray, deviceArray, size, hipMemcpyDeviceToHost);
hipFree(deviceArray);

// Managed memory (automatic transfers)
float* managedArray;
hipMallocManaged(&managedArray, size);
// Initialize and use directly from host
// Launch kernel with managedArray
// Access results directly from host
hipFree(managedArray);

// Pinned host memory for faster transfers
float* pinnedArray;
hipHostMalloc(&pinnedArray, size);
// Use pinnedArray from host
hipMemcpyAsync(deviceArray, pinnedArray, size, hipMemcpyHostToDevice, stream);
// Launch kernel
hipMemcpyAsync(pinnedArray, deviceArray, size, hipMemcpyDeviceToHost, stream);
hipHostFree(pinnedArray);

Your choice of memory management strategy should be guided by your application’s specific needs. During early development, managed memory often provides the best balance of simplicity and performance. For production code with well-understood memory access patterns, explicit memory management might yield better performance.

Using Shared Memory #

Shared memory can significantly improve performance by reducing global memory accesses:

__global__ void sharedMemExample(float* input, float* output, int n) {
    // Declare shared memory array
    __shared__ float sharedData[256];
    
    int tid = threadIdx.x;
    int gid = blockIdx.x * blockDim.x + threadIdx.x;
    
    // Load data from global to shared memory
    if (gid < n) {
        sharedData[tid] = input[gid];
    }
    
    // Ensure all threads in block have loaded their data
    __syncthreads();
    
    // Process data in shared memory (example: simple transformation)
    if (gid < n) {
        // All threads in block can access any element in shared memory
        float result = 0;
        for (int i = 0; i < 3; i++) {
            int idx = tid + i - 1;
            if (idx >= 0 && idx < blockDim.x) {
                result += sharedData[idx];
            }
        }
        output[gid] = result / 3.0f;  // Average of element and neighbors
    }
}

In this example, each thread loads one element from global memory into shared memory. After synchronizing with __syncthreads() (ensuring all data is loaded), threads can quickly access both their own data and neighboring elements. This pattern is especially powerful for algorithms that require access to adjacent data elements, like image filters or stencil operations.

Example: Stream-based Overlapping Execution #

For complex applications, asynchronous operations allow overlapping of computation with data transfers. This stream-based approach keeps the GPU busy while data moves between host and device:

#include <hip/hip_runtime.h>

int main() {
    const int segmentSize = 1024 * 1024;
    const int segments = 4;
    const int totalSize = segmentSize * segments;
    
    // Allocate host and device memory
    float *hostInput, *hostOutput, *deviceInput, *deviceOutput;
    
    hipHostMalloc(&hostInput, totalSize * sizeof(float));
    hipHostMalloc(&hostOutput, totalSize * sizeof(float));
    hipMalloc(&deviceInput, segmentSize * sizeof(float));
    hipMalloc(&deviceOutput, segmentSize * sizeof(float));
    
    // Initialize host data
    for (int i = 0; i < totalSize; i++) {
        hostInput[i] = static_cast<float>(i);
    }
    
    // Create streams
    hipStream_t streams[2];
    for (int i = 0; i < 2; i++) {
        hipStreamCreate(&streams[i]);
    }
    
    // Process data in segments, alternating between streams
    for (int i = 0; i < segments; i++) {
        const int offset = i * segmentSize;
        const int streamIdx = i % 2;
        
        // Copy segment to device asynchronously
        hipMemcpyAsync(
            deviceInput, 
            hostInput + offset, 
            segmentSize * sizeof(float), 
            hipMemcpyHostToDevice, 
            streams[streamIdx]
        );
        
        // Launch kernel in the same stream
        const int threadsPerBlock = 256;
        const int blocks = (segmentSize + threadsPerBlock - 1) / threadsPerBlock;
        
        hipLaunchKernelGGL(
            processData,  // Kernel function (not shown)
            dim3(blocks), 
            dim3(threadsPerBlock), 
            0, 
            streams[streamIdx], 
            deviceInput, 
            deviceOutput, 
            segmentSize
        );
        
        // Copy results back asynchronously
        hipMemcpyAsync(
            hostOutput + offset, 
            deviceOutput, 
            segmentSize * sizeof(float), 
            hipMemcpyDeviceToHost, 
            streams[streamIdx]
        );
    }
    
    // Synchronize all streams
    for (int i = 0; i < 2; i++) {
        hipStreamSynchronize(streams[i]);
        hipStreamDestroy(streams[i]);
    }
    
    // Clean up
    hipHostFree(hostInput);
    hipHostFree(hostOutput);
    hipFree(deviceInput);
    hipFree(deviceOutput);
    
    return 0;
}

This example processes data in segments, using two streams to overlap operations: while one stream is executing a kernel, another can be transferring data. This pattern keeps both the CPU-to-GPU data bus and the GPU compute units busy, potentially improving overall throughput dramatically.

Querying Device Properties for Optimization #

Understanding your specific hardware’s capabilities is essential for targeted optimizations. HIP provides device query functions that reveal important properties:

#include <hip/hip_runtime.h>
#include <iostream>

void printDeviceProperties() {
    int deviceCount;
    hipGetDeviceCount(&deviceCount);
    
    std::cout << "Found " << deviceCount << " HIP-compatible devices:" << std::endl;
    
    for (int i = 0; i < deviceCount; i++) {
        hipDeviceProp_t props;
        hipGetDeviceProperties(&props, i);
        
        std::cout << "\nDevice " << i << ": " << props.name << std::endl;
        std::cout << "  Compute capability: " << props.major << "." << props.minor << std::endl;
        std::cout << "  Multiprocessors: " << props.multiProcessorCount << std::endl;
        std::cout << "  Max threads per block: " << props.maxThreadsPerBlock << std::endl;
        std::cout << "  Max threads per multiprocessor: " << props.maxThreadsPerMultiProcessor << std::endl;
        std::cout << "  Warp size: " << props.warpSize << std::endl;
        std::cout << "  Shared memory per block: " << props.sharedMemPerBlock / 1024.0 << " KB" << std::endl;
        std::cout << "  Total global memory: " << props.totalGlobalMem / (1024.0 * 1024.0) << " MB" << std::endl;
        std::cout << "  Memory clock rate: " << props.memoryClockRate / 1000.0 << " GHz" << std::endl;
        std::cout << "  Memory bus width: " << props.memoryBusWidth << " bits" << std::endl;
    }
}

int main() {
    printDeviceProperties();
    return 0;
}

These properties can guide key decisions like block size, shared memory allocation, and algorithm selection. For example, knowing the warp size helps ensure thread counts align with hardware execution units.