CUDA：具有共享内存和矩阵大小（不是块大小的倍数）的平铺矩阵-矩阵乘法

CUDA：具有共享内存和矩阵大小（不是块大小的倍数）的平铺矩阵-矩阵乘法

cpp
#include 
// 定义矩阵维度
#define N 1024
// 定义块大小
#define BLOCK_SIZE 16
// 定义共享内存大小
#define SHARED_SIZE (BLOCK_SIZE * BLOCK_SIZE)
// CUDA核函数，实现矩阵乘法
__global__ void matrixMul(float* A, float* B, float* C, int n)
{
    // 计算当前线程的全局索引
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    // 定义共享内存
    __shared__ float shared_A[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ float shared_B[BLOCK_SIZE][BLOCK_SIZE];
    // 初始化结果矩阵
    float sum = 0.0;
    // 循环遍历每个子块
    for (int k = 0; k < n / BLOCK_SIZE; ++k)
    {
        // 从全局内存加载数据到共享内存
        shared_A[threadIdx.y][threadIdx.x] = A[row * n + (k * BLOCK_SIZE + threadIdx.x)];
        shared_B[threadIdx.y][threadIdx.x] = B[(k * BLOCK_SIZE + threadIdx.y) * n + col];
        // 等待所有线程加载完毕
        __syncthreads();
        // 计算乘积并累加到结果矩阵
        for (int i = 0; i < BLOCK_SIZE; ++i)
        {
            sum += shared_A[threadIdx.y][i] * shared_B[i][threadIdx.x];
        }
        // 等待所有线程计算完毕
        __syncthreads();
    }
    // 将结果写回全局内存
    C[row * n + col] = sum;
}
int main()
{
    // 定义输入矩阵和输出矩阵
    float* A, * B, * C;
    // 在主机上分配内存
    A = (float*)malloc(N * N * sizeof(float));
    B = (float*)malloc(N * N * sizeof(float));
    C = (float*)malloc(N * N * sizeof(float));
    // 在设备上分配内存
    float* d_A, * d_B, * d_C;
    cudaMalloc((void**)&d_A, N * N * sizeof(float));
    cudaMalloc((void**)&d_B, N * N * sizeof(float));
    cudaMalloc((void**)&d_C, N * N * sizeof(float));
    // 将输入矩阵从主机内存复制到设备内存
    cudaMemcpy(d_A, A, N * N * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, B, N * N * sizeof(float), cudaMemcpyHostToDevice);
    // 定义块和网格大小
    dim3 blockSize(BLOCK_SIZE, BLOCK_SIZE);
    dim3 gridSize(N / BLOCK_SIZE, N / BLOCK_SIZE);
    // 调用CUDA核函数
    matrixMul<<>>(d_A, d_B, d_C, N);
    // 将结果矩阵从设备内存复制到主机内存
    cudaMemcpy(C, d_C, N * N * sizeof(float), cudaMemcpyDeviceToHost);
    // 释放设备内存
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
    // 释放主机内存
    free(A);
    free(B);
    free(C);
    return 0;
}