我的第一个CUDA程序

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 和 Bfor (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);}}// 输出结果矩阵 Cprintf("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) {// 获取线程的全局IDint 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;// 将结果写入矩阵 CC[index] = result;}
}void MatAdd(int height, int width) {// 在主机内存中分配 A、B 和 Cfloat* A = (float*)malloc(height * width * sizeof(float));float* B = (float*)malloc(height * width * sizeof(float));float* C = (float*)malloc(height * width * sizeof(float));// 初始化输入矩阵 A 和 Bfor (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); // 复制 AcudaMemcpy(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);// 输出结果矩阵 Cprintf("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) {// 获取线程的全局IDint 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;// 将结果写入矩阵 CC[index] = result;}
}void MatAdd(int height, int width) {// 在主机内存中分配 A、B 和 Cfloat* A = (float*)malloc(height * width * sizeof(float));float* B = (float*)malloc(height * width * sizeof(float));float* C = (float*)malloc(height * width * sizeof(float));// 初始化输入矩阵 A 和 Bfor (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); // 复制 AcudaMemcpy(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);// 输出结果矩阵 Cprintf("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) {// 获取线程的全局IDint 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;// 将结果写入矩阵 CC[index] = result;}
}void MatAdd(int height, int width) {// 在主机内存中分配 A、B 和 Cfloat* A = (float*)malloc(height * width * sizeof(float));float* B = (float*)malloc(height * width * sizeof(float));float* C = (float*)malloc(height * width * sizeof(float));// 初始化输入矩阵 A 和 Bfor (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); // 复制 AcudaMemcpy(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);// 输出结果矩阵 Cprintf("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)
{// 获取线程的全局IDint 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 和 Bfor (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);
// 输出结果矩阵 Cprintf("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)
{// 获取线程的全局IDint 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和Bfor (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 IDint 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 dataint 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 Aint 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 matricesfloat 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];// Computefloat result = A1 * B1 + A2 * B2 + A3 * B3 + A4 * B4;// Write resultC[index_C] = result;}
}int main()
{int M = 4; // Example dimensions for the matricesint 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 memoryfloat *h_A = (float*)malloc(size_A);float *h_B = (float*)malloc(size_B);float *h_C = (float*)malloc(size_C);// Initialize matrices A and Bfor (int i = 0; i < M * N; ++i){h_A[i] = i + 1; // Example initializationh_B[i] = i + 1;}// Allocate device memoryfloat *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 devicecudaMemcpy(d_A, h_A, size_A, cudaMemcpyHostToDevice);cudaMemcpy(d_B, h_B, size_B, cudaMemcpyHostToDevice);// Define block and grid sizedim3 blockDim(BLOCK_SIZE, BLOCK_SIZE);dim3 gridDim((N / 2 + BLOCK_SIZE - 1) / BLOCK_SIZE, (M / 2 + BLOCK_SIZE - 1) / BLOCK_SIZE);// Launch kernelMatAdd<<<gridDim, blockDim>>>(d_A, d_B, d_C, M, N);// Copy result from device to hostcudaMemcpy(h_C, d_C, size_C, cudaMemcpyDeviceToHost);// Print resultfor (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 memorycudaFree(d_A);cudaFree(d_B);cudaFree(d_C);// Free host memoryfree(h_A);free(h_B);free(h_C);return 0;
}

 输出：

66 138 
546 746 

• 异构计算整成为当前计算领域的重点方向
• GPGPU是异构计算的主要形式
• GPGPU是一款大规模细粒度并行处理器，并行思维是进行GPGPU编程的重要前提
• NVIDIA是当前GPGPU领域当之无愧的霸主
• GPGPU编程的重点是定义明确的线程和数据间的映射