在CUDA中,为什么列加法比行加法快的原因是因为列优先存储的方式与内存访问模式的局部性有关。在很多情况下,列优先存储方式可以提供更好的内存访问局部性,从而提高数据读取速度。
以下是一个使用CUDA进行矩阵列加法的示例代码:
#include
__global__ void columnAdd(float* matrix, float* result, int numRows, int numCols) {
int col = blockIdx.x * blockDim.x + threadIdx.x;
if (col < numCols) {
float sum = 0.0f;
for (int row = 0; row < numRows; ++row) {
sum += matrix[row * numCols + col];
}
result[col] = sum;
}
}
int main() {
int numRows = 3;
int numCols = 4;
int size = numRows * numCols * sizeof(float);
// Allocate memory on the host
float* h_matrix = (float*)malloc(size);
float* h_result = (float*)malloc(numCols * sizeof(float));
// Initialize the matrix
for (int i = 0; i < numRows * numCols; ++i) {
h_matrix[i] = i + 1; // Example values
}
// Allocate memory on the device
float* d_matrix;
float* d_result;
cudaMalloc((void**)&d_matrix, size);
cudaMalloc((void**)&d_result, numCols * sizeof(float));
// Copy the matrix from host to device
cudaMemcpy(d_matrix, h_matrix, size, cudaMemcpyHostToDevice);
// Launch the kernel
int blockSize = 256;
int gridSize = (numCols + blockSize - 1) / blockSize;
columnAdd<<>>(d_matrix, d_result, numRows, numCols);
// Copy the result from device to host
cudaMemcpy(h_result, d_result, numCols * sizeof(float), cudaMemcpyDeviceToHost);
// Print the result
for (int i = 0; i < numCols; ++i) {
printf("Sum of column %d: %f\n", i, h_result[i]);
}
// Free memory
free(h_matrix);
free(h_result);
cudaFree(d_matrix);
cudaFree(d_result);
return 0;
}
在这个示例代码中,columnAdd
函数是在GPU上执行的CUDA核函数。每个线程负责计算一个列的和。为了实现并行计算,我们使用了CUDA的线程和块结构。
在主函数中,我们首先在主机上分配内存,然后将矩阵数据初始化,并在设备上分配内存。接下来,我们将矩阵数据从主机复制到设备,并在设备上启动核函数。最后,我们将结果从设备复制回主机,并打印出每列的和。
要编译和运行这个示例代码,您需要安装CUDA开发环境,并使用nvcc
编译器将代码编译为可执行文件。