0% found this document useful (0 votes)
14 views2 pages

CUDA Matrix Multiplication Example

The document contains a CUDA program for performing matrix multiplication. It allocates memory for matrices on both the host and device, initializes them, and executes the multiplication using a kernel. The program also measures and prints the elapsed time for the operation.

Uploaded by

yaseeniqbal365
Copyright
© All Rights Reserved
We take content rights seriously. If you suspect this is your content, claim it here.
Available Formats
Download as PDF, TXT or read online on Scribd
0% found this document useful (0 votes)
14 views2 pages

CUDA Matrix Multiplication Example

The document contains a CUDA program for performing matrix multiplication. It allocates memory for matrices on both the host and device, initializes them, and executes the multiplication using a kernel. The program also measures and prints the elapsed time for the operation.

Uploaded by

yaseeniqbal365
Copyright
© All Rights Reserved
We take content rights seriously. If you suspect this is your content, claim it here.
Available Formats
Download as PDF, TXT or read online on Scribd

CUDA Program for Matrix Multiplication

#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>

#define BLOCK_SIZE 16

__global__ void matrix_multiply(float *a, float *b, float *c, int n) {


int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
float sum = 0.0f;

if (row < n && col < n) {


for (int i = 0; i < n; ++i) {
sum += a[row * n + i] * b[i * n + col];
}
c[row * n + col] = sum;
}
}

int main() {
int n = 1024;
size_t size = n * n * sizeof(float);

float *a, *b, *c;


float *d_a, *d_b, *d_c;
cudaEvent_t start, stop;
float elapsed_time;

// Allocate host memory


a = (float*)malloc(size);
b = (float*)malloc(size);
c = (float*)malloc(size);

// Initialize matrices
for (int i = 0; i < n * n; ++i) {
a[i] = i % n;
b[i] = i % n;
}

// Allocate device memory


cudaMalloc(&d_a, size);
cudaMalloc(&d_b, size);
cudaMalloc(&d_c, size);

// Copy data to device


cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice);

// Configure kernel launch parameters


dim3 threads(BLOCK_SIZE, BLOCK_SIZE);
dim3 blocks((n + threads.x - 1) / threads.x, (n + threads.y - 1) /
threads.y);

// Launch and time the kernel


cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start);

matrix_multiply<<<blocks, threads>>>(d_a, d_b, d_c, n);

cudaEventRecord(stop);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&elapsed_time, start, stop);

// Copy result to host


cudaMemcpy(c, d_c, size, cudaMemcpyDeviceToHost);

printf("Elapsed time for matrix multiplication: %.2f ms\n", elapsed_time);

// Free memory
cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
free(a); free(b); free(c);

return 0;
}

Output:
Elapsed time for matrix multiplication: 58.74 ms

Common questions

Powered by AI

The CUDA program relies on the natural synchronization provided within each block. Since the matrix multiplication kernel deals with independent computations per thread (each thread calculating a specific matrix element), there is no explicit need for thread synchronization commands within a block beyond ensuring no dependencies exist within the calculation. This independence allows all threads in a block to perform computations simultaneously without requiring additional synchronization .

Using cudaEvent_t is crucial for accurate performance measurement in CUDA programming because it records events that mark the start and end of GPU computation tasks. This timing mechanism considers only the time spent executing the kernel on the GPU, providing an accurate measure of actual computation time, excluding host-side execution time or other overheads. Such precise measurement is essential for optimizing and tuning applications for better performance .

Using float data types in the CUDA matrix multiplication program has implications for both precision and performance. While float provides a balance between range and computational speed, it introduces potential precision errors due to its limited significant digits compared to double. However, using float generally yields better performance on GPUs due to faster hardware support and lower memory bandwidth requirements. Developers must consider precision requirements of their application to determine if this trade-off is acceptable .

Choosing grid and block dimensions involves trade-offs between parallel efficiency and resource utilization. Smaller blocks might lead to higher overhead and underutilized resources due to increased number of blocks, while larger blocks can lead to increased register pressure and block dimension limits on certain devices. The optimum balance must consider the device's architecture, such as the number of multiprocessors and available memory, to maximize parallel occupancy and minimize execution time .

In the CUDA program, the cudaMemcpy function is responsible for transferring data between the host and the device (GPU). It copies matrices from host memory to device memory before the kernel execution and retrieves the result from the device back to the host afterward. Meanwhile, the cudaMalloc function allocates memory on the device for matrices, enabling the GPU to store the matrices temporarily during computation .

The use of dim3 data structure for specifying thread and block dimensions allows for a clear and structured definition of three-dimensional grids and blocks, even though the current program uses two dimensions. This abstraction simplifies the configuration of grid and block sizes, making it extensible for more complex problems that might require three-dimensional computations, thereby enhancing both clarity and flexibility in setting up problem scales .

Potential optimizations include using shared memory to reduce redundant global memory access by caching submatrices within each block, fine-tuning BLOCK_SIZE to achieve better occupancy by balancing between computation and memory latency, and leveraging Memory Coalescing to optimize memory bandwidth utilization. Additionally, employing asynchronous data transfers and overlapping computation with data transfer or using fast math intrinsics could significantly reduce the elapsed time .

The CUDA program utilizes parallel processing capabilities of GPUs by dividing the task into small blocks and threads that correspond to the dimensions of a matrix. It defines a BLOCK_SIZE (16), which determines the number of threads per block, and calculates grid size by determining how many blocks are necessary to cover the entire matrix. This allows each CUDA core to handle multiple threads, thereby improving efficiency through parallel execution. The matrix multiplication kernel computes each element of the result matrix concurrently by calculating the dot product of a row from the first matrix and a column of the second matrix in parallel .

Initializing matrices with a modulo operation (% n) ensures that the elements of the matrices have a predictable pattern. This pattern can sometimes help in debugging or verifying the correctness of the matrix multiplication since the expected results for simple indices are easier to deduce. However, it also means that the matrix has repeating elements, which may not reflect real-world datasets and could potentially skew performance measurement if the memory access patterns do not sufficiently mimic more complex matrices .

BLOCK_SIZE in the CUDA program defines the number of threads per block, which directly influences the parallel execution level. A sensible BLOCK_SIZE ensures adequate utilization of GPU cores by distributing workload evenly among them. It affects the granularity of performance optimization: too small might underutilize resources, while too large might exceed the GPU's maximum capacity, impacting parallel efficiency .

You might also like