我的第一个CUDA程序

已于 2024-08-15 16:03:24 修改
阅读量1.1k 收藏 8
点赞数 21
分类专栏： GPU编程 文章标签： 算法 人工智能 c++ c语言 vscode
于 2024-08-14 19:56:53 首次发布
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/qq_62704693/article/details/141198039
版权

MatAdd算法

实现两个矩阵对应元素相加

 

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

// 矩阵加法函数
void MatAdd(int height, int width)
{
    // 在主机内存中为 A、B 和 C 分配内存
    float* A = (float*)malloc(height * width * sizeof(float));
    float* B = (float*)malloc(height * width * sizeof(float));
    float* C = (float*)malloc(height * width * sizeof(float));

    // 初始化输入矩阵 A 和 B
    for (int i = 0; i < height; i++)
    {
        for (int j = 0; j < width; j++)
        {
            // 例如，初始化为 i+j 的和
            *(A + i * width + j) = i + j;
            *(B + i * width + j) = i - j;
        }
    }

    // 执行矩阵加法运算，将结果存储在矩阵 C 中
    for (int i = 0; i < height; i++)
    {
        for (int j = 0; j < width; j++)
        {
            *(C + i * width + j) = *(A + i * width + j) + *(B + i * width + j);
        }
    }

    // 输出结果矩阵 C
    printf("Matrix C (result of A + B):\n");
    for (int i = 0; i < height; i++)
    {
        for (int j = 0; j < width; j++)
        {
            printf("%f ", *(C + i * width + j));
        }
        printf("\n");
    }

    // 释放分配的内存
    free(A);
    free(B);
    free(C);
}

int main()
{
    int height = 3; // 矩阵的高度
    int width = 3;  // 矩阵的宽度

    // 执行矩阵加法
    MatAdd(height, width);

    return 0;
}

 输出：

Matrix C (result of A + B):
0.000000 0.000000 0.000000 
2.000000 2.000000 2.000000 
4.000000 4.000000 4.000000

MatAdd算法的GPU实现

MatAdd算法的GPU实现

  • CPU端为输入矩阵A和B、输出矩阵C分配空间，并进行初始化
  • CPU端分配设备端内存，并将A和B传输到GPU上
  • 定义数据和线程的映射关系，并确定线程的开启数量和组织方式
  1. 每个线程负责输出矩阵C的一个元素的计算，全局ID为（i，j）的线程计算索引为（i，j）的矩阵元素
  2. 当矩阵规模为width*height时，共开启width * height个线程
  3. 每个block包含256个线程，采用（16,16）的组织形式
  • 编写计算kernel，完成计算任务
  • CPU端将计算结果从Device内存拷贝到Host内存

内存分配 数据传输

开启线程 启动kernel 结果返回

GPU kernel

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

// CUDA 核函数，用于矩阵加法
__global__ void MatAddKernel(float* A, float* B, float* C, int height, int width) {
    // 获取线程的全局ID
    int i = blockIdx.x * blockDim.x + threadIdx.x; // 计算全局行索引
    int j = blockIdx.y * blockDim.y + threadIdx.y; // 计算全局列索引

    // 确保索引在矩阵范围内
    if (i < width && j < height) {
        // 计算当前线程对应的元素索引
        int index = j * width + i;

        // 从矩阵 A 和 B 中读取数据
        float src_data_A = A[index];
        float src_data_B = B[index];

        // 执行加法运算
        float result = src_data_A + src_data_B;

        // 将结果写入矩阵 C
        C[index] = result;
    }
}

void MatAdd(int height, int width) {
    // 在主机内存中分配 A、B 和 C
    float* A = (float*)malloc(height * width * sizeof(float));
    float* B = (float*)malloc(height * width * sizeof(float));
    float* C = (float*)malloc(height * width * sizeof(float));

    // 初始化输入矩阵 A 和 B
    for (int i = 0; i < height; i++) {
        for (int j = 0; j < width; j++) {
            A[i * width + j] = i + j; // 简单初始化，A的元素为行索引+列索引
            B[i * width + j] = i - j; // 简单初始化，B的元素为行索引-列索引
        }
    }

    // 第一步：在设备内存中为矩阵 A、B 和 C 分配内存
    float* d_A;
    cudaMalloc(&d_A, height * width * sizeof(float)); // 分配矩阵 A 的设备内存
    float* d_B;
    cudaMalloc(&d_B, height * width * sizeof(float)); // 分配矩阵 B 的设备内存
    float* d_C;
    cudaMalloc(&d_C, height * width * sizeof(float)); // 分配矩阵 C 的设备内存

    // 第二步：将矩阵 A 和 B 从主机内存复制到设备内存
    cudaMemcpy(d_A, A, height * width * sizeof(float), cudaMemcpyHostToDevice); // 复制 A
    cudaMemcpy(d_B, B, height * width * sizeof(float), cudaMemcpyHostToDevice); // 复制 B

    // 第三步：调用 CUDA 核函数
    dim3 threadsPerBlock(16, 16); // 定义每个块中的线程数
    dim3 numBlocks((width + threadsPerBlock.x - 1) / threadsPerBlock.x, 
                   (height + threadsPerBlock.y - 1) / threadsPerBlock.y); // 计算网格中的块数
    MatAddKernel<<<numBlocks, threadsPerBlock>>>(d_A, d_B, d_C, height, width); // 启动 CUDA 核函数

    // 第四步：将结果从设备内存复制回主机内存
    cudaMemcpy(C, d_C, height * width * sizeof(float), cudaMemcpyDeviceToHost);

    // 输出结果矩阵 C
    printf("Matrix C (result of A + B):\n");
    for (int i = 0; i < height; i++) {
        for (int j = 0; j < width; j++) {
            printf("%f ", C[i * width + j]);
        }
        printf("\n");
    }

    // 释放设备内存
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    // 释放主机内存
    free(A);
    free(B);
    free(C);
}

int main() {
    int height = 3; // 矩阵的高度
    int width = 3;  // 矩阵的宽度

    // 调用矩阵加法函数
    MatAdd(height, width);

    return 0;
}

输出：

Matrix C (result of A + B):
0.000000 0.000000 0.000000 
2.000000 2.000000 2.000000 
4.000000 4.000000 4.000000

CUDA程序编译

CUDA程序性能测试

 使用 CUDA GPU Timers 实际要循环100次求平均值

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

// CUDA 核函数，用于矩阵加法
__global__ void MatAddKernel(float* A, float* B, float* C, int height, int width) {
    // 获取线程的全局ID
    int i = blockIdx.x * blockDim.x + threadIdx.x; // 计算全局行索引
    int j = blockIdx.y * blockDim.y + threadIdx.y; // 计算全局列索引

    // 确保索引在矩阵范围内
    if (i < width && j < height) {
        // 计算当前线程对应的元素索引
        int index = j * width + i;

        // 从矩阵 A 和 B 中读取数据
        float src_data_A = A[index];
        float src_data_B = B[index];

        // 执行加法运算
        float result = src_data_A + src_data_B;

        // 将结果写入矩阵 C
        C[index] = result;
    }
}

void MatAdd(int height, int width) {
    // 在主机内存中分配 A、B 和 C
    float* A = (float*)malloc(height * width * sizeof(float));
    float* B = (float*)malloc(height * width * sizeof(float));
    float* C = (float*)malloc(height * width * sizeof(float));

    // 初始化输入矩阵 A 和 B
    for (int i = 0; i < height; i++) {
        for (int j = 0; j < width; j++) {
            A[i * width + j] = i + j; // 简单初始化，A的元素为行索引+列索引
            B[i * width + j] = i - j; // 简单初始化，B的元素为行索引-列索引
        }
    }

    // 第一步：在设备内存中为矩阵 A、B 和 C 分配内存
    float* d_A;
    cudaMalloc(&d_A, height * width * sizeof(float)); // 分配矩阵 A 的设备内存
    float* d_B;
    cudaMalloc(&d_B, height * width * sizeof(float)); // 分配矩阵 B 的设备内存
    float* d_C;
    cudaMalloc(&d_C, height * width * sizeof(float)); // 分配矩阵 C 的设备内存

    // 第二步：将矩阵 A 和 B 从主机内存复制到设备内存
    cudaMemcpy(d_A, A, height * width * sizeof(float), cudaMemcpyHostToDevice); // 复制 A
    cudaMemcpy(d_B, B, height * width * sizeof(float), cudaMemcpyHostToDevice); // 复制 B

    // 第三步：定义事件变量，用于测量核函数执行时间
    cudaEvent_t start, stop;
    float time;
    cudaEventCreate(&start); // 创建开始事件
    cudaEventCreate(&stop);  // 创建停止事件

    // 第四步：调用 CUDA 核函数
    dim3 threadsPerBlock(16, 16); // 定义每个块中的线程数
    dim3 numBlocks((width + threadsPerBlock.x - 1) / threadsPerBlock.x, 
                   (height + threadsPerBlock.y - 1) / threadsPerBlock.y); // 计算网格中的块数

    cudaEventRecord(start, 0); // 记录开始时间
    MatAddKernel<<<numBlocks, threadsPerBlock>>>(d_A, d_B, d_C, height, width); // 启动 CUDA 核函数
    cudaEventRecord(stop, 0);  // 记录结束时间

    cudaEventSynchronize(stop); // 等待事件完成
    cudaEventElapsedTime(&time, start, stop); // 计算核函数执行时间
    printf("Kernel execution time: %f ms\n", time); // 输出核函数执行时间

    // 销毁事件
    cudaEventDestroy(start);
    cudaEventDestroy(stop);

    // 第五步：将结果从设备内存复制回主机内存
    cudaMemcpy(C, d_C, height * width * sizeof(float), cudaMemcpyDeviceToHost);

    // 输出结果矩阵 C
    printf("Matrix C (result of A + B):\n");
    for (int i = 0; i < height; i++) {
        for (int j = 0; j < width; j++) {
            printf("%f ", C[i * width + j]);
        }
        printf("\n");
    }

    // 释放设备内存
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    // 释放主机内存
    free(A);
    free(B);
    free(C);
}

int main() {
    int height = 3; // 矩阵的高度
    int width = 3;  // 矩阵的宽度

    // 调用矩阵加法函数
    MatAdd(height, width);

    return 0;
}

输出： 

Kernel execution time: 0.086016 ms
Matrix C (result of A + B):
0.000000 0.000000 0.000000 
2.000000 2.000000 2.000000 
4.000000 4.000000 4.000000

 使用 CPU Timers 实际要循环100次求平均值

#include <stdio.h>
#include <stdlib.h>
#include <sys/time.h> // 包含 gettimeofday 的头文件
#include <cuda_runtime.h>

// CUDA 核函数，用于矩阵加法
__global__ void MatAddKernel(float* A, float* B, float* C, int height, int width) {
    // 获取线程的全局ID
    int i = blockIdx.x * blockDim.x + threadIdx.x; // 计算全局行索引
    int j = blockIdx.y * blockDim.y + threadIdx.y; // 计算全局列索引

    // 确保索引在矩阵范围内
    if (i < width && j < height) {
        // 计算当前线程对应的元素索引
        int index = j * width + i;

        // 从矩阵 A 和 B 中读取数据
        float src_data_A = A[index];
        float src_data_B = B[index];

        // 执行加法运算
        float result = src_data_A + src_data_B;

        // 将结果写入矩阵 C
        C[index] = result;
    }
}

void MatAdd(int height, int width) {
    // 在主机内存中分配 A、B 和 C
    float* A = (float*)malloc(height * width * sizeof(float));
    float* B = (float*)malloc(height * width * sizeof(float));
    float* C = (float*)malloc(height * width * sizeof(float));

    // 初始化输入矩阵 A 和 B
    for (int i = 0; i < height; i++) {
        for (int j = 0; j < width; j++) {
            A[i * width + j] = i + j; // 简单初始化，A的元素为行索引+列索引
            B[i * width + j] = i - j; // 简单初始化，B的元素为行索引-列索引
        }
    }

    // 第一步：在设备内存中为矩阵 A、B 和 C 分配内存
    float* d_A;
    cudaMalloc(&d_A, height * width * sizeof(float)); // 分配矩阵 A 的设备内存
    float* d_B;
    cudaMalloc(&d_B, height * width * sizeof(float)); // 分配矩阵 B 的设备内存
    float* d_C;
    cudaMalloc(&d_C, height * width * sizeof(float)); // 分配矩阵 C 的设备内存

    // 第二步：将矩阵 A 和 B 从主机内存复制到设备内存
    cudaMemcpy(d_A, A, height * width * sizeof(float), cudaMemcpyHostToDevice); // 复制 A
    cudaMemcpy(d_B, B, height * width * sizeof(float), cudaMemcpyHostToDevice); // 复制 B

    // 第三步：定义时间测量变量
    struct timeval start, end;
    double elapsedTime;

    // 记录开始时间
    gettimeofday(&start, NULL);

    // 启动 CUDA 核函数
    dim3 threadsPerBlock(16, 16); // 定义每个块中的线程数
    dim3 numBlocks((width + threadsPerBlock.x - 1) / threadsPerBlock.x, 
                   (height + threadsPerBlock.y - 1) / threadsPerBlock.y); // 计算网格中的块数

    MatAddKernel<<<numBlocks, threadsPerBlock>>>(d_A, d_B, d_C, height, width); // 启动 CUDA 核函数
    cudaDeviceSynchronize(); // 等待核函数完成

    // 记录结束时间
    gettimeofday(&end, NULL);

    // 计算时间差
    elapsedTime = (end.tv_sec - start.tv_sec) * 1000.0; // 秒转毫秒
    elapsedTime += (end.tv_usec - start.tv_usec) / 1000.0; // 微秒转毫秒

    // 输出核函数执行时间
    printf("Kernel execution time: %f ms\n", elapsedTime);

    // 第四步：将结果从设备内存复制回主机内存
    cudaMemcpy(C, d_C, height * width * sizeof(float), cudaMemcpyDeviceToHost);

    // 输出结果矩阵 C
    printf("Matrix C (result of A + B):\n");
    for (int i = 0; i < height; i++) {
        for (int j = 0; j < width; j++) {
            printf("%f ", C[i * width + j]);
        }
        printf("\n");
    }

    // 释放设备内存
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    // 释放主机内存
    free(A);
    free(B);
    free(C);
}

int main() {
    int height = 3; // 矩阵的高度
    int width = 3;  // 矩阵的宽度

    // 调用矩阵加法函数
    MatAdd(height, width);

    return 0;
}

 输出：

Kernel execution time: 0.101000 ms
Matrix C (result of A + B):
0.000000 0.000000 0.000000 
2.000000 2.000000 2.000000 
4.000000 4.000000 4.000000

 MatAdd程序之变一：每个线程处理4个元素

  • CPU端为输入矩阵A和B、输出矩阵C分配空间，并进行初始化
  • CPU端分配设备端内存，并将A和B传输到GPU上
  • 定义数据和线程的映射关系，并确定线程的开启数量和组织方式
  1. 每个线程负责输出矩阵C的四个元素的计算，全局ID为（i，j）的线程计算索引为（i，4*j~4*j+3）的矩阵元素
  2. 当矩阵规模为width*height时，共开启width/4 * height个线程
  3. 每个block包含256个线程，采用（16,16）的组织形式
  • 编写计算kernel，完成计算任务
  • CPU端将计算结果从Device内存拷贝到Host内存

#include <cuda_runtime.h>
#include <iostream>

// 自定义向上取整除法函数
int div_up(int a, int b) {
    return (a + b - 1) / b;
}
// CUDA核函数，执行矩阵加法
__global__ void MatAdd(float *A, float *B, float *C, int height, int width)
{
    // 获取线程的全局ID
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;

    // 检查线程是否在有效的计算范围内
    // 每个线程处理4个连续的元素，因此检查4*i是否在宽度范围内
    if ((4 * i) < width && j < height)
    {
        // 获取线程与数据之间的映射关系
        int index = j * width + 4 * i;

        // 从源矩阵读取数据
        float4 src_data_A = *((float4 *)(A + index));
        float4 src_data_B = *((float4 *)(B + index));

        // 执行加法运算
          // 执行加法运算，逐个对float4的分量进行加法
        float4 result;
        result.x = src_data_A.x + src_data_B.x;
        result.y = src_data_A.y + src_data_B.y;
        result.z = src_data_A.z + src_data_B.z;
        result.w = src_data_A.w + src_data_B.w;

        // 将结果写回结果矩阵
        *((float4 *)(C + index)) = result;
    }
}

// 矩阵加法函数，处理矩阵的加法运算
void MatAdd(int height, int width)
{
    float *A, *B, *C;         // 主机内存指针
    float *d_A, *d_B, *d_C;   // 设备内存指针

    // 分配主机内存
    A = (float*)malloc(height * width * sizeof(float));
    B = (float*)malloc(height * width * sizeof(float));
    C = (float*)malloc(height * width * sizeof(float));

    // 初始化矩阵A和B（这里省略具体的初始化代码）
    // 初始化输入矩阵 A 和 B
    for (int i = 0; i < height; i++) {
        for (int j = 0; j < width; j++) {
            A[i * width + j] = i + j; // 简单初始化，A的元素为行索引+列索引
            B[i * width + j] = i - j; // 简单初始化，B的元素为行索引-列索引
        }
    }
    // 分配设备内存
    cudaMalloc((void**)&d_A, height * width * sizeof(float));
    cudaMalloc((void**)&d_B, height * width * sizeof(float));
    cudaMalloc((void**)&d_C, height * width * sizeof(float));

    // 将数据从主机内存复制到设备内存
    cudaMemcpy(d_A, A, height * width * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, B, height * width * sizeof(float), cudaMemcpyHostToDevice);

    // 设置线程块的大小，每个块中有16x16个线程
    dim3 threadsPerBlock(16, 16);

    // 计算需要多少个线程块来覆盖整个矩阵区域
    // 宽度除以4是因为每个线程处理4个元素
    dim3 numBlocks(div_up(width / 4, threadsPerBlock.x), div_up(height, threadsPerBlock.y));

    // 调用核函数，执行矩阵加法运算
    MatAdd<<<numBlocks, threadsPerBlock>>>(d_A, d_B, d_C, height, width);

    // 将计算结果从设备内存复制到主机内存
    cudaMemcpy(C, d_C, height * width * sizeof(float), cudaMemcpyDeviceToHost);
// 输出结果矩阵 C
    printf("Matrix C (result of A + B):\n");
    for (int i = 0; i < height; i++) {
        for (int j = 0; j < width; j++) {
            printf("%f ", C[i * width + j]);
        }
        printf("\n");
    }
    // 释放设备内存
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    // 释放主机内存
    free(A);
    free(B);
    free(C);
}


int main() {
    int height = 8; // 示例矩阵的高度
    int width = 8;  // 示例矩阵的宽度

    // 调用矩阵加法函数
    MatAdd(height, width);

    return 0;
}

 输出：

Matrix C (result of A + B):
0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 
2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 
4.000000 4.000000 4.000000 4.000000 4.000000 4.000000 4.000000 4.000000 
6.000000 6.000000 6.000000 6.000000 6.000000 6.000000 6.000000 6.000000 
8.000000 8.000000 8.000000 8.000000 8.000000 8.000000 8.000000 8.000000 
10.000000 10.000000 10.000000 10.000000 10.000000 10.000000 10.000000 10.000000 
12.000000 12.000000 12.000000 12.000000 12.000000 12.000000 12.000000 12.000000 
14.000000 14.000000 14.000000 14.000000 14.000000 14.000000 14.000000 14.000000

MatAdd程序之变二

矩阵A、B、C都为NxN的方阵，A和B为已知矩阵，C[i][j] = A[i][j] + B[j][i]。

 

  • CPU端为输入矩阵A和B、输出矩阵C分配空间，并进行初始化
  • CPU端分配设备端内存，并将A和B传输到GPU上
  • 定义数据和线程的映射关系，并确定线程的开启数量和组织方式
  1. 每个线程负责输出矩阵C的一个元素的计算，全局ID为（i，j）的线程计算索引为（i，j）的矩阵元素
  2. 当矩阵规模为width*height时，共开启width * height个线程
  3. 每个block包含256个线程，采用（16,16）的组织形式
  • 编写计算kernel，完成计算任务
  • CPU端将计算结果从Device内存拷贝到Host内存

#include <cuda_runtime.h>
#include <iostream>

#define N 1024 // 矩阵的大小（N x N）

// CUDA核函数，执行矩阵加法运算
__global__ void MatAdd(float *A, float *B, float *C, int height, int width)
{
    // 获取线程的全局ID
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;

    // 确保线程在有效的计算范围内
    if(i < width && j < height)
    {
        // 获取矩阵元素的索引
        int index_A = j * width + i;
        int index_B = i * width + j;
        int index_C = index_A;

        // 从矩阵A和B读取数据并计算结果
        float src_data_A = *(A + index_A);
        float src_data_B = *(B + index_B);
        float result = src_data_A + src_data_B;

        // 将结果写入矩阵C
        *(C + index_C) = result;
    }
}

// 主函数，处理矩阵的初始化、调用CUDA核函数以及结果输出
void MatAdd(int height, int width)
{
    // 在主机内存中为A、B和C分配空间
    float* A = (float*)malloc(height * width * sizeof(float));
    float* B = (float*)malloc(height * width * sizeof(float));
    float* C = (float*)malloc(height * width * sizeof(float));

    // 初始化矩阵A和B
    for (int i = 0; i < height; i++) {
        for (int j = 0; j < width; j++) {
            *(A + i * width + j) = static_cast<float>(i + j); // 示例初始化
            *(B + i * width + j) = static_cast<float>(i - j); // 示例初始化
        }
    }

    // 分配设备内存
    float *d_A, *d_B, *d_C;
    cudaMalloc((void**)&d_A, height * width * sizeof(float));
    cudaMalloc((void**)&d_B, height * width * sizeof(float));
    cudaMalloc((void**)&d_C, height * width * sizeof(float));

    // 将数据从主机内存复制到设备内存
    cudaMemcpy(d_A, A, height * width * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, B, height * width * sizeof(float), cudaMemcpyHostToDevice);

    // 配置CUDA网格和线程块大小
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks((width + threadsPerBlock.x - 1) / threadsPerBlock.x,
                   (height + threadsPerBlock.y - 1) / threadsPerBlock.y);

    // 调用CUDA核函数，执行矩阵加法运算
    MatAdd<<<numBlocks, threadsPerBlock>>>(d_A, d_B, d_C, height, width);

    // 将计算结果从设备内存复制回主机内存
    cudaMemcpy(C, d_C, height * width * sizeof(float), cudaMemcpyDeviceToHost);

    // 输出部分计算结果用于验证
    std::cout << "C[0][0] = " << C[0] << std::endl;
    std::cout << "C[N-1][N-1] = " << C[(N-1) * width + (N-1)] << std::endl;

    // 释放设备内存
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    // 释放主机内存
    free(A);
    free(B);
    free(C);
}

int main() {
    // 调用矩阵加法函数
    MatAdd(N, N);

    return 0;
}

输出：

C[0][0] = 0
C[N-1][N-1] = 2046

MatAdd程序之变三

  •  矩阵A、B都为MxN的矩阵，矩阵C为(M/2)*(N/2)的矩阵， A和B为已知矩阵，C[i][j] = A[2*i][2*j] *B[2*i][2*j] +  A[2*i][2*j+1] *B[2*i][2*j+1]   + A[2*i+1][2*j] *B[2*i+1][2*j]  + A[2*i+1][2*j +1] *B[2*i+1][2*j+1] 。 CPU端为输入矩阵A和B、输出矩阵C分配空间，并进行初始化
  • CPU端分配设备端内存，并将A和B传输到GPU上
  • 定义数据和线程的映射关系，并确定线程的开启数量和组织方式
  1. 每个线程负责输出矩阵C的一个元素的计算，全局ID为（i，j）的线程计算索引为（i，j）的矩阵元素
  2. 当矩阵规模为(M/2)*(N/2)时，共开启(M/2)*(N/2)个线程,每个线程对应A和B的四个元素
  3. 每个block包含256个线程，采用（16,16）的组织形式
  • 编写计算kernel，完成计算任务
  • CPU端将计算结果从Device内存拷贝到Host内存

#include <iostream>
#include <cuda_runtime.h>

#define BLOCK_SIZE 16

__global__ void MatAdd(float *A, float *B, float *C, int height, int width)
{
    // Get thread ID
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;

    if (i < width / 2 && j < height / 2)
    {
        // Get the mapping between thread and data
        int index_A_1 = 2 * j * width + 2 * i;
        int index_A_2 = 2 * j * width + 2 * i + 1;
        int index_A_3 = (2 * j + 1) * width + 2 * i;
        int index_A_4 = (2 * j + 1) * width + 2 * i + 1;

        int index_B_1 = index_A_1; // Indexes for matrix B are the same as A
        int index_B_2 = index_A_2;
        int index_B_3 = index_A_3;
        int index_B_4 = index_A_4;

        int index_C = j * (width / 2) + i;

        // Read data from source matrices
        float A1 = A[index_A_1];
        float A2 = A[index_A_2];
        float A3 = A[index_A_3];
        float A4 = A[index_A_4];

        float B1 = B[index_B_1];
        float B2 = B[index_B_2];
        float B3 = B[index_B_3];
        float B4 = B[index_B_4];

        // Compute
        float result = A1 * B1 + A2 * B2 + A3 * B3 + A4 * B4;

        // Write result
        C[index_C] = result;
    }
}

int main()
{
    int M = 4; // Example dimensions for the matrices
    int N = 4;

    int size_A = M * N * sizeof(float);
    int size_B = M * N * sizeof(float);
    int size_C = (M / 2) * (N / 2) * sizeof(float);

    // Allocate host memory
    float *h_A = (float*)malloc(size_A);
    float *h_B = (float*)malloc(size_B);
    float *h_C = (float*)malloc(size_C);

    // Initialize matrices A and B
    for (int i = 0; i < M * N; ++i)
    {
        h_A[i] = i + 1; // Example initialization
        h_B[i] = i + 1;
    }

    // Allocate device memory
    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, size_A);
    cudaMalloc(&d_B, size_B);
    cudaMalloc(&d_C, size_C);

    // Copy data from host to device
    cudaMemcpy(d_A, h_A, size_A, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size_B, cudaMemcpyHostToDevice);

    // Define block and grid size
    dim3 blockDim(BLOCK_SIZE, BLOCK_SIZE);
    dim3 gridDim((N / 2 + BLOCK_SIZE - 1) / BLOCK_SIZE, (M / 2 + BLOCK_SIZE - 1) / BLOCK_SIZE);

    // Launch kernel
    MatAdd<<<gridDim, blockDim>>>(d_A, d_B, d_C, M, N);

    // Copy result from device to host
    cudaMemcpy(h_C, d_C, size_C, cudaMemcpyDeviceToHost);

    // Print result
    for (int i = 0; i < M / 2; ++i)
    {
        for (int j = 0; j < N / 2; ++j)
        {
            std::cout << h_C[i * (N / 2) + j] << " ";
        }
        std::cout << std::endl;
    }

    // Free device memory
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    // Free host memory
    free(h_A);
    free(h_B);
    free(h_C);

    return 0;
}

 输出：

66 138 
546 746
  • 异构计算整成为当前计算领域的重点方向
  • GPGPU是异构计算的主要形式
  • GPGPU是一款大规模细粒度并行处理器，并行思维是进行GPGPU编程的重要前提
  • NVIDIA是当前GPGPU领域当之无愧的霸主
  • GPGPU编程的重点是定义明确的线程和数据间的映射
Polaris北极星少女
关注
专栏目录
ubuntu22.04+cuda11.5+gcc11.4第一个cuda程序示例
qq_43594926的博客
08-13 2112
VisualStudio 2019是微软的集成开发环境(IDE)，通常在Windows操作系统上使用。然而，并不直接支持在Linux上安装。如果想在Ubuntu上进行开发，可以考虑以下几个选项:使用替代的IDE或文本编辑器: Ubuntu上有许多适用于C++等编程语言的开发工具，比如Visual Studio Code、 Code.:Blocks、 Eclipse等。
CUDA学习笔记（三）：第一个cuda程序
Never-Giveup的博客
08-06 541
。。。。
cuda笔记-第一个cuda程序
IT1995的博客
12-09 1124
这里先说明下一些基本概念： 释放GPU中的内存cudaFree() CUDA函数的定义： __global__：定义在GPU上，可以在CPU上调用的函数； __device__：定义在GPU上，由GPU调用函数； __host__：在CPU上定义的函数，一般与__device__一起用 在GPU上开辟空间：cudaMalloc(**devPtr, byte_size) 如： int *gpu_int; cudaMalloc((void**)&gpu_int, sizeo..
编写第1个CUDA程序
小怪物爱吹牛的博客
02-13 581
最简CUDA程序
cuda编程学习——第一个cuda程序！干货向（一）
出门吃三碗饭的博客
05-18 880
最近在做三维重建，尤其是Nerf方面多视角合成工作的时候，意识到了cuda的编程计算可以大大提高其中渲染的计算，最明显的例子是Instant-ngp，Plenoxels等文章，因此后面会学Cuda一段时间，同时也就开了这个新坑。因为笔者也是cuda新手，所以大家有问题的话可以评论区指出，一起学习进步！
【CUDA编程1】年轻人的第一个CUDA程序
木盏
09-07 1265
说CUDA是人工智能算法行业的重要基石一点也不为过。老黄在图灵架构发布的时候说，图灵架构是英伟达历史上自2006年以来最伟大的飞跃。而2006年，正是CUDA发布的时间。
tx2+cuda+cmake:第一个cuda程序
01-09
在本项目中，我们将探讨如何在...通过这个项目，你可以学习如何在TX2上设置CUDA环境，编写混合C++/CUDA代码，使用CMake管理项目，并最终运行和测试CUDA程序。这将有助于提升你在嵌入式系统和高性能计算领域的技能。
fortran-cuda+fortran环境搭建+第一个fortran程序.zip
02-08
通过上述步骤，你将能够成功地在你的计算机上建立起一个支持Fortran+CUDA的开发环境，并编写和运行第一个Fortran+CUDA程序。在实践中，不断学习和理解CUDA的内存模型、线程结构以及如何有效利用GPU并行性是非常重要...
cuda程序示例代码.rar
05-24
在学习和研究这些代码时，可以结合《并行计算教程(第1版)》深入理解每个示例背后的原理，逐步掌握CUDA编程技巧，最终能够自如地运用CUDA解决各种计算密集型问题。同时，配合NVIDIA的CUDA C Programming Guide和CUDA ...
算法刷题：300. 最长递增子序列、674. 最长连续递增序列、718. 最长重复子数组、1143. 最长公共子序列
qq_61936886的博客
09-12 498
如果text1[i - 1] 与 text2[j - 1]不相同，那就看看text1[0, i - 2]与text2[0, j - 1]的最长公共子序列 和 text1[0, i - 1]与text2[0, j - 2]的最长公共子序列，取最大的。但dp[i][0] 和dp[0][j]要初始值，因为 为了方便递归公式dp[i][j] = dp[i - 1][j - 1] + 1;2.递推公式：当A[i - 1] 和B[j - 1]相等的时候，dp[i][j] = dp[i - 1][j - 1] + 1;
算法知识点——常用输入输出数据的方式
5233的博客
09-12 294
除了换行之外，还会清空缓存区，而清空缓存区这件事情会耗时很久，因此为了提高速度，一般直接：cout << “\n”;2、n组输入输出(n不确定）1、n组输入输出(n确定）3、 复杂的输入输出。
Java数据结构（十）——冒泡排序、快速排序
xiaokuer_的博客
09-11 1183
文章介绍了冒泡排序和快速排序，快速排序的内容包括：划分方式、优化策略以及非递归实现
算法基础-约数
qq_65186476的博客
09-11 320
试除法求约数、约数个数、约数之和、最大公约数
【Hot100】LeetCode—64. 最小路径和
weixin_44382896的博客
09-12 476
64. 最小路径和。
算法练习题24——查找杨辉三角中的组合数
hello77的blogggg 祝你得偿所愿
09-12 585
逐项递推法适合处理较小的数据量，计算较为直观，但当杨辉三角行数较大时，效率较低。 二分查找法利用组合数的单调性，显著提高查找效率，适合处理较大的数据范围。这两种解法在不同场景下都可以使用，二分查找法尤其适合大规模数据下的查找问题。
二分系列(二分查找)9/12
2301_78191305的博客
09-12 967
排序数组->使用二分。第一个位置:>=；第二个位置:=(target+1)-1；所以直接写一个>=的函数。
18068 选择排序
huang1xiao1sheng的博客
09-11 434
输入10个整数，编写一个实现对数组进行选择排序的函数 #include int sort(int a[], int n) { int i,j,k,tmp; for(i=0;i
选择排序
shylyly_的博客
09-13 1060
每一次从待排序的数据元素中选出最小（或最大）的一个元素，存放在序列的起始位置，直到全部待排序的数据元素排完
二分系列(二分答案最大值)9/15.16
最新发布
2301_78191305的博客
09-16 598
给你一批机器，及其它制作一块合金所需要的不同种类金属的个数。第一台机器制作一块合金，需要一块A金属、一块B金属、一块C金属。我们对每一类机器所需要的各种种类的金属进行判断，如果仓库里面有，就不需要花钱；如果mid-index<1的，要计算有多少个1。前n项和:(mid-1+1)*(mid-1)/2+leftCount-(mid-1);根据贪心:以index为中心，向两边依次-1，直到减到1为止。如果mid-rightCount<1,先计算多少个1，xxx。
GPU计算：CUDA程序生命周期与优化
GPU的发展可以分为三个阶段：第一代GPU主要用于硬件加速3D图像处理，不具备编程特性；第二代GPU增加了有限的编程性，如顶点级和像素级的可编程性；第三代GPU引入了像CUDA这样的编程环境，大大提高了GPU的易用性和...
评论 1
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包

打赏作者

Polaris北极星少女

你的鼓励将是我创作的最大动力

¥1 ¥2 ¥4 ¥6 ¥10 ¥20
扫码支付：¥1
获取中
扫码支付

您的余额不足，请更换扫码支付或充值

打赏作者

实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值