CUDA笔记3
发布时间：2024年01月23日
CUDA编程-入门、加法、乘法
例子1：入门
#include <stdio.h>
__global__ void hello_from_gpu()
{
const int bid = blockIdx.x;
const int tid = threadIdx.x;
printf("Hello world from block %d and thread %d!\n",bid,tid);
}
int main(void)
{
hello_from_gpu<<<2,3>>>();
cudaDeviceSynchronize();
return 0;
}
Hello world from block 1 and thread 0!
Hello world from block 1 and thread 1!
Hello world from block 1 and thread 2!
Hello world from block 0 and thread 0!
Hello world from block 0 and thread 1!
Hello world from block 0 and thread 2!
每个线程执行一次“Hello world”
例子2：数组相加
#include <math.h>
#include <stdio.h>
void __global__ add(const double *x, const double *y, double *z, int count)
{
const int n = blockDim.x * blockIdx.x + threadIdx.x;
if( n < count)
{
printf("Block ID: %d, blockDim.x: %d,Thread ID: %d, n: %d\n",
blockIdx.x, blockDim.x, threadIdx.x, n);
z[n] = x[n] + y[n];
printf("z[%d]: %f\n", n, z[n]);
}
// 第二种xy输出
// Add a barrier to ensure all threads have printed before exiting the kernel
// __syncthreads();
// // Only have one thread (threadIdx.x == 0) print the entire arrays
// if (threadIdx.x == 0 && blockIdx.x == 0)
// {
// for (int i = 0; i < count; ++i)
// {
// printf("Final Values - x[%d]: %f, y[%d]: %f, z[%d]: %f\n", i, x[i], i, y[i], i, z[i]);
// }
// }
}
// void check(const double *z, const int N)
// {
// bool error = false;
// for (int n = 0; n < N; ++n)
// {
// if (fabs(z[n] - 3) > (1.0e-10))
// {
// error = true;
// }
// }
// printf("%s\n", error ? "Errors" : "Pass");
// }
void check(const double *z, const double *h_x, const double *h_y, const int N)
{
bool error = false;
for (int n = 0; n < N; ++n)
{
// 修改此处的期望值计算方式
// 不会执行，即error=false，返回Pass
if (fabs(z[n] - (h_x[n] + h_y[n])) > (1.0e-10))
{
error = true;
}
}
printf("%s\n", error ? "Errors" : "Pass");
}
int main(void)
{
const int N = 6;
const int M = sizeof(double) * N;
double *h_x = (double*) malloc(M);
double *h_y = (double*) malloc(M);
double *h_z = (double*) malloc(M);
for (int n = 0; n < N; ++n)
{
h_x[n] = 1;
h_y[n] = n;
//第一种xy输出
printf("h_x[%d]: %f, h_y[%d]: %f\n", n, h_x[n], n, h_y[n]);
}
double *d_x, *d_y, *d_z;
cudaMalloc((void **)&d_x, M);
cudaMalloc((void **)&d_y, M);
cudaMalloc((void **)&d_z, M);
cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);
cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice);
const int block_size = 2;
/*N是数组大小
如果N能够被block_size整除，那么grid_size就等于N/block_size,
那么每个线程块处理一个数据。
如果N不能被block_size整除，那么(N+block_size-1)/block_size
的计算会确保生成的grid_size足够大，
保所有的输入数据都能够被覆盖到，并且没有数据被漏掉
如果 N = 10，block_size = 3，则 (N + block_size - 1) / block_size = 4。
这确保了有足够的线程块来处理所有的10个数据元素。*/
const int grid_size = (N + block_size - 1) / block_size;
add<<<grid_size, block_size>>>(d_x, d_y, d_z,N);
cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost);
// check(h_z, N);
check(h_z, h_x, h_y, N);
free(h_x);
free(h_y);
free(h_z);
cudaFree(d_x);
cudaFree(d_y);
cudaFree(d_z);
return 0;
}
// xy值
h_x[0]: 1.000000, h_y[0]: 0.000000
h_x[1]: 1.000000, h_y[1]: 1.000000
h_x[2]: 1.000000, h_y[2]: 2.000000
h_x[3]: 1.000000, h_y[3]: 3.000000
h_x[4]: 1.000000, h_y[4]: 4.000000
h_x[5]: 1.000000, h_y[5]: 5.000000
//线程块，线程，n的值
Block ID: 1, blockDim.x: 2, Thread ID: 0, n: 2
Block ID: 1, blockDim.x: 2, Thread ID: 1, n: 3
Block ID: 2, blockDim.x: 2, Thread ID: 0, n: 4
Block ID: 2, blockDim.x: 2, Thread ID: 1, n: 5
Block ID: 0, blockDim.x: 2, Thread ID: 0, n: 0
Block ID: 0, blockDim.x: 2, Thread ID: 1, n: 1
//输出的z值
z[2]: 3.000000
z[3]: 4.000000
z[4]: 5.000000
z[5]: 6.000000
z[0]: 1.000000
z[1]: 2.000000
Block ID: 0, blockDim.x: 2, Thread ID: 0, n: 0，即第1线程块，第1个线程，根据以下公式：
const int n = blockDim.x * blockIdx.x + threadIdx.x;
z[n] = x[n] + y[n];
算出n=0，进一步，z[0]=x[0]+y[0]=1+0=1
即第一步算出索引，第二步带入索引值求出z，其余同理。
请注意，由于线程和块的索引计算，每个线程处理不同的数据元素，从而实现了并行计算。在实际的大规模并行计算中，这种并行性能够带来显著的加速。
如Cuda笔记1所讲的，所有线程执行相同的核函数，且并行执行。
每个线程负责处理一个计算任务。在GPU中，大量的线程可以同时执行相同的指令，每个线程独立地处理不同的数据或任务，这被称为SIMD（Single Instruction, Multiple Data）架构。（分治思想，将大的任务分成多个一样的小任务，每个thread处理一个计算任务）
例如，假设有一个大数组需要排序，可以将数组分割成若干段，每个线程负责对其中一段进行排序。这样，每个线程独立执行排序任务，最后再将各段有序的数组合并成整体有序的数组。
通过规约(Reduction)操作理解GPU内存体系（源于知乎-紫气东来）
reduce 算法可以描述为
其中 ×○可表示为求 sum，min，max，avg 等操作，最后获得的输出相比于输入一般维度上会递减。
在GPU中，reduce采用了一种树形的计算方式，例如下面的求和问题：
由于GPU没有针对global数据的同步操作，只能针对block的数据进行同步。所以，一般而言将reduce分为两个阶段，其示意图如下：
假设给定一个长度为N的数组，需要计算该数组的所有元素之和。首先需要将数组分为m个小份。而后，在第一阶段中，开启m个block计算出m个小份的reduce值。最后，在第二阶段中，使用一个block将m个小份再次进行reduce，得到最终的结果。
chatgpt的回答：
在GPU编程中，Reduce（归约）是一种并行计算模式，用于将一个大型数据集的元素合并为一个单一的值。Reduce操作通常用于计算总和、平均值、最大值、最小值等。在GPU中，由于存在大量的并行处理单元（CUDA核心），可以通过归约操作充分利用并行性，加速数据处理。
典型的GPU中的Reduce流程如下：
数据分割： 将大型数据集划分成小块，每个块由一个线程块（Thread Block）处理。
并行计算： 每个线程块中的线程并行地对其分配的数据块执行归约操作。这可以是对数据进行相加、取最大值、取最小值等。
层级归约： 由于线程块的数量通常大于硬件所支持的线程数，需要进行多层级的归约。每一层级的归约将当前层级的中间结果传递到下一层级，直到最终得到一个单一的结果。
最终结果： 最后一层的归约结果即为整个数据集的最终结果。
CUDA编程中，可以使用**共享内存（shared memory）**来加速归约操作，因为共享内存是线程块内的线程共享的。这减少了对全局内存的访问次数，提高了性能。
归约操作的实现可能因问题的不同而有所差异，但其核心思想是通过并行计算和层级归约来充分利用GPU的并行计算能力。
例子3：矩阵乘法（全局内存vs共享内存）
1）全局内存
#include <stdio.h>
#include <math.h>
#include "error.cuh"
#define BLOCK_SIZE 16
__global__ void gpu_matrix_mult(int *a,int *b, int *c, int m, int n, int k)
{
//不是thread的一维索引，而是复合线程索引值，更符合矩阵运算等
//索引就是变值，不要想复杂了，就相当于一维数组，二维数组中的索引一回事
//一个索引对应一个线程（一个线程的位置，线程是最小单位），然后用这个线程计算，GPU只是很多线程同时计算而已。
//计算当前线程在矩阵中的行索引。
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
int sum = 0;
//确保当前线程的索引在矩阵的有效范围内，以避免越界访问。
if( col < k && row < m)
{
//对于每一列（或每一行，具体取决于矩阵乘法的顺序），
//将对应位置的元素相乘并累加到 sum 中。
//读取m的一行，n的一列
for(int i = 0; i < n; i++)
{
/*在实际的 GPU 编程中，选择线程块和网格的配置通常取决于算法和数据的特性，
以及对硬件架构的优化。
这儿矩阵相乘，因此选择二维的线程块和网格，线程块的 y 索引和 x 索引
可以自然地映射到结果矩阵中的行和列。
*/
sum += a[row * n + i] * b[i * k + col];
//即每个(row,col)位置的线程处理一个计算任务
/*
与CPU的对比，
void cpu_matrix_mult(int *h_a, int *h_b, int *h_result, int m, int n, int k) {
for (int i = 0; i < m; ++i)
{
for (int j = 0; j < k; ++j)
{
int tmp = 0.0;
for (int h = 0; h < n; ++h)
{
tmp += h_a[i * n + h] * h_b[h * k + j];
}
h_result[i * k + j] = tmp;
}
}
}
事实上，外面的两个for循环用线程的动态索引取代了，而线程不是按顺序执行的
而是多个线程同时执行的
cpu中的i和j分别相当于GPU中的row，col
*/
}
//row，col确定当前线程在输出矩阵 C 中的二维索引
//用的一维表示二维a[i*M+j] i是行下标，M是一行的元素个数，j是列下标
c[row * k + col] = sum;
}
}
__global__ void gpu_matrix_mult_shared(int *d_a, int *d_b, int *d_result, int n)
{
__shared__ int tile_a[BLOCK_SIZE][BLOCK_SIZE];
__shared__ int tile_b[BLOCK_SIZE][BLOCK_SIZE];
int row = blockIdx.y * BLOCK_SIZE + threadIdx.y;
int col = blockIdx.x * BLOCK_SIZE + threadIdx.x;
int tmp = 0;
int idx;
for (int sub = 0; sub < gridDim.x; ++sub)
{
idx = row * n + sub * BLOCK_SIZE + threadIdx.x;
tile_a[threadIdx.y][threadIdx.x] = row<n && (sub * BLOCK_SIZE + threadIdx.x)<n? d_a[idx]:0;
idx = (sub * BLOCK_SIZE + threadIdx.y) * n + col;
tile_b[threadIdx.y][threadIdx.x] = col<n && (sub * BLOCK_SIZE + threadIdx.y)<n? d_b[idx]:0;
__syncthreads();
for (int k = 0; k < BLOCK_SIZE; ++k)
{
tmp += tile_a[threadIdx.y][k] * tile_b[k][threadIdx.x];
}
__syncthreads();
}
if(row < n && col < n)
{
d_result[row * n + col] = tmp;
}
}
void cpu_matrix_mult(int *h_a, int *h_b, int *h_result, int m, int n, int k) {
for (int i = 0; i < m; ++i)
{
for (int j = 0; j < k; ++j)
{
int tmp = 0.0;
for (int h = 0; h < n; ++h)
{
tmp += h_a[i * n + h] * h_b[h * k + j];
}
h_result[i * k + j] = tmp;
}
}
}
int main(int argc, char const *argv[])
{
int m=1000;
int n=1000;
int k=1000;
int *h_a, *h_b, *h_c, *h_cc, *h_cs;
//void **) &h_a将分配的内存地址赋值给指针h_a
//sizeof(int)*m*n)分配一个 int 类型的数组，其大小为 m*n 个元素
//即主机端分配一个 int 类型的二维数组，大小为 m 行、n 列。
CHECK(cudaMallocHost((void **) &h_a, sizeof(int)*m*n));
CHECK(cudaMallocHost((void **) &h_b, sizeof(int)*n*k));
CHECK(cudaMallocHost((void **) &h_c, sizeof(int)*m*k));
CHECK(cudaMallocHost((void **) &h_cc, sizeof(int)*m*k));
CHECK(cudaMallocHost((void **) &h_cs, sizeof(int)*m*k));
/*CUDA事件
cudaEvent_t: 这是CUDA库中表示事件的数据类型。
CUDA事件通常用于测量 GPU 执行时间，用于确定 GPU 执行某个任务所花费的时间。
cudaEventCreate表示用于创建 CUDA 事件
原型“cudaError_t cudaEventCreate(cudaEvent_t* event);”
*/
cudaEvent_t start, stop,stop_share;
CHECK(cudaEventCreate(&start));
CHECK(cudaEventCreate(&stop));
CHECK(cudaEventCreate(&stop_share));
//将整个二维矩阵h_a赋值为1
//数组是在内存按顺序存储，因此二维其实也是一维
for (int i = 0; i < m; ++i) {
for (int j = 0; j < n; ++j) {
h_a[i * n + j] = 1;
}
}
for (int i = 0; i < n; ++i) {
for (int j = 0; j < k; ++j) {
h_b[i * k + j] = 0;
}
}
int *d_a, *d_b, *d_c, *d_c_share;
CHECK(cudaMalloc((void **) &d_a, sizeof(int)*m*n));
CHECK(cudaMalloc((void **) &d_b, sizeof(int)*n*k));
CHECK(cudaMalloc((void **) &d_c, sizeof(int)*m*k));
CHECK(cudaMalloc((void **) &d_c_share, sizeof(int)*m*k));
/* cudaDeviceSynchronize()目的是确保设备上的所有操作都已完成。
这段程序片段的目的是确保CUDA分配的设备内存已经成功完成。*/
cudaDeviceSynchronize();
// cudaEventRecord是CUDA中记录事件的函数，用于将当前时间戳记录到由start表示的事件中。
CHECK(cudaEventRecord(start));
// copy matrix A and B from host to device memory(cudaMemcpy全局内存)
CHECK(cudaMemcpy(d_a, h_a, sizeof(int)*m*n, cudaMemcpyHostToDevice));
CHECK(cudaMemcpy(d_b, h_b, sizeof(int)*n*k, cudaMemcpyHostToDevice));
unsigned int grid_rows = (m + BLOCK_SIZE - 1) / BLOCK_SIZE;
unsigned int grid_cols = (k + BLOCK_SIZE - 1) / BLOCK_SIZE;
/*
dim3三维网格和线程块，z=1，blockidx.z=0,threadidx.z=0
创建了一个二维的线程块与线程网格
*/
dim3 dimGrid(grid_cols, grid_rows);
dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
gpu_matrix_mult<<<dimGrid, dimBlock>>>(d_a, d_b, d_c, m,n,k);
CHECK(cudaMemcpy(h_c, d_c, (sizeof(int)*m*k), cudaMemcpyDeviceToHost));
//cudaThreadSynchronize();
CHECK(cudaEventRecord(stop));
CHECK(cudaEventSynchronize(stop));
gpu_matrix_mult_shared<<<dimGrid, dimBlock>>>(d_a, d_b, d_c_share, n);
CHECK(cudaMemcpy(h_cs, d_c_share, (sizeof(int)*m*k), cudaMemcpyDeviceToHost));
CHECK(cudaEventRecord(stop_share));
CHECK(cudaEventSynchronize(stop_share));
float elapsed_time, elapsed_time_share;
CHECK(cudaEventElapsedTime(&elapsed_time, start, stop));
CHECK(cudaEventElapsedTime(&elapsed_time_share, stop, stop_share));
printf("Time_global = %g ms.\n", elapsed_time);
printf("Time_share = %g ms.\n", elapsed_time_share);
CHECK(cudaEventDestroy(start));
CHECK(cudaEventDestroy(stop));
//cpu_matrix_mult(h_a, h_b, h_c, m, n, k);
int ok = 1;
/*for (int i = 0; i < m; ++i)
{
for (int j = 0; j < k; ++j)
{
if(fabs(h_cs[i*k + j] - 0)>(1.0e-10))
{
printf("hcs: %d hc: %d ",h_cs[i*k + j], h_c[i*k + j]);
ok = 0;
}
}
}
if(ok)
{
printf("Pass!!!\n");
}
else
{
printf("Error!!!\n");
}
*/
// free memory
CHECK(cudaFree(d_a));
CHECK(cudaFree(d_b));
CHECK(cudaFree(d_c));
CHECK(cudaFreeHost(h_a));
CHECK(cudaFreeHost(h_b));
CHECK(cudaFreeHost(h_c));
return 0;
}
重点：
①在实际的 GPU 编程中，选择线程块和网格的配置通常取决于算法和数据的特性，以及对硬件架构的优化。这儿矩阵相乘，因此选择二维的线程块和网格，线程块的 y 索引和 x 索引可以自然地映射到结果矩阵中的行和列。
②即每个(row,col)位置的线程处理一个计算任务。事实上，CPU外面的两个for循环用GPU线程的动态索引取代了，而线程不是按顺序执行的。而是多个线程同时执行的。CPU中的i和j分别相当于GPU中的row，col
③定义MK个thread，刚好动态映射_a，d_b的mn、nk，以及输出结果d_c中的mk个值
④copy matrix A and B from host to device memory(cudaMemcpy全局内存)
⑤逻辑上是二维的，实际内存物理上是一维的，因此要用一维索引形式表示
2）共享内存
理论：
整体思路：
申请很多线程，每个线程求取pd中的每个值，就需要从M和N读取一行和一列的值。那么要求出pdsut中每个值，就要读取m和n中by=1行，bx=1中的所有值。即每一行需要多很多次，放到共享内存。但是共享内存小，因此分块。做子矩阵（蓝色）线程。
1、在共享内存申请空间
2、pd相当于一个block，这里面的每个thread只读取一个元素，放入share memory中，所有线程同时行动
3、计算时，每个线程读取share memory读取行的元素（蓝色块中的黄色行）,结果存在temp中。然后黄色矩阵的行，结果放在temp中。计算完后，放到pd中。
#include <stdio.h>
#include <math.h>
#include "error.cuh"
#define BLOCK_SIZE 16
__global__ void gpu_matrix_mult(int *a,int *b, int *c, int m, int n, int k)
{
//不是thread的一维索引，而是复合线程索引值，更符合矩阵运算等
//索引就是变值，不要想复杂了，就相当于一维数组，二维数组中的索引一回事
//一个索引对应一个线程（一个线程的位置，线程是最小单位），然后用这个线程计算，GPU只是很多线程同时计算而已。
//计算当前线程在矩阵中的行索引。
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
int sum = 0;
//确保当前线程的索引在矩阵的有效范围内，以避免越界访问。
if( col < k && row < m)
{
//对于每一列（或每一行，具体取决于矩阵乘法的顺序），
//将对应位置的元素相乘并累加到 sum 中。
//读取m的一行，n的一列
for(int i = 0; i < n; i++)
{
/*在实际的 GPU 编程中，选择线程块和网格的配置通常取决于算法和数据的特性，
以及对硬件架构的优化。
这儿矩阵相乘，因此选择二维的线程块和网格，线程块的 y 索引和 x 索引
可以自然地映射到结果矩阵中的行和列。
*/
sum += a[row * n + i] * b[i * k + col];
//即每个(row,col)位置的线程处理一个计算任务
/*
与CPU的对比，
void cpu_matrix_mult(int *h_a, int *h_b, int *h_result, int m, int n, int k) {
for (int i = 0; i < m; ++i)
{
for (int j = 0; j < k; ++j)
{
int tmp = 0.0;
for (int h = 0; h < n; ++h)
{
tmp += h_a[i * n + h] * h_b[h * k + j];
}
h_result[i * k + j] = tmp;
}
}
}
事实上，外面的两个for循环用线程的动态索引取代了，而线程不是按顺序执行的
而是多个线程同时执行的
cpu中的i和j分别相当于GPU中的row，col
*/
}
//row，col确定当前线程在输出矩阵 C 中的二维索引
//用的一维表示二维a[i*M+j] i是行下标，M是一行的元素个数，j是列下标
c[row * k + col] = sum;
}
}
__global__ void gpu_matrix_mult_shared(int *d_a, int *d_b, int *d_result, int M, int N, int K)
{
/*
__shared__: 这是一个CUDA关键字，
表示声明的变量是共享内存中的变量。共享内存是一块对于线程块中的所有线程可见的内存。
int tile_a[BLOCK_SIZE][BLOCK_SIZE];这是一个二维整数数组，
表示在共享内存中分配的二维数组。tile_a 被用于存储一个矩阵的子矩阵。
怎么理解被用于存储一个矩阵的子矩阵？
矩阵乘法通常涉及到将大矩阵分解为子矩阵，
然后每个线程块负责计算一个或多个子矩阵的乘法。
在矩阵乘法的上下文中，tile_a 存储了一个矩阵的子矩阵。当每个线程块负责计算某个子矩阵的乘法时，
将这个子矩阵加载到 tile_a 中，以便线程块中的每个线程可以快速访问和共享这部分数据。
*/
__shared__ int tile_a[BLOCK_SIZE][BLOCK_SIZE];
__shared__ int tile_b[BLOCK_SIZE][BLOCK_SIZE];
int row = blockIdx.y * BLOCK_SIZE + threadIdx.y;
int col = blockIdx.x * BLOCK_SIZE + threadIdx.x;
int tmp = 0;
int idx;
//读取每个矩阵小块
// for (int sub = 0; sub < gridDim.x; ++sub)
for (int sub = 0; sub <= N/BLOCK_SIZE; ++sub)
{
// idx = row * n + sub * BLOCK_SIZE + threadIdx.x;
// tile_a[threadIdx.y][threadIdx.x] = row<n && (sub * BLOCK_SIZE + threadIdx.x)<n? d_a[idx]:0;
// idx = (sub * BLOCK_SIZE + threadIdx.y) * n + col;
// tile_b[threadIdx.y][threadIdx.x] = col<n && (sub * BLOCK_SIZE + threadIdx.y)<n? d_b[idx]:0;
//详细版
int r = row;
int c = sub * BLOCK_SIZE + threadIdx.x;
idx = r * N + c;
if (r >= M || c >= N)
{
tile_a[threadIdx.y][threadIdx.x]=0;
}
else
{
tile_a[threadIdx.y][threadIdx.x]=d_a[idx];
}
r = sub * BLOCK_SIZE + threadIdx.y;
c = col;
idx = r * K + c;
if (c >= K || c >= N)
{
tile_b[threadIdx.y][threadIdx.x]=0;
}
else
{
tile_b[threadIdx.y][threadIdx.x]=d_b[idx];
}
__syncthreads();
//相乘累加到一个temp中
for (int k = 0; k < BLOCK_SIZE; ++k)
{
tmp += tile_a[threadIdx.y][k] * tile_b[k][threadIdx.x];
}
__syncthreads();
}
if(row < M && col < K)
{
d_result[row * K + col] = tmp;
}
}
void cpu_matrix_mult(int *h_a, int *h_b, int *h_result, int m, int n, int k) {
for (int i = 0; i < m; ++i)
{
for (int j = 0; j < k; ++j)
{
int tmp = 0.0;
for (int h = 0; h < n; ++h)
{
tmp += h_a[i * n + h] * h_b[h * k + j];
}
h_result[i * k + j] = tmp;
}
}
}
int main(int argc, char const *argv[])
{
int m=1000;
int n=1000;
int k=1000;
int *h_a, *h_b, *h_c, *h_cc, *h_cs;
//void **) &h_a将分配的内存地址赋值给指针h_a
//sizeof(int)*m*n)分配一个 int 类型的数组，其大小为 m*n 个元素
//即主机端分配一个 int 类型的二维数组，大小为 m 行、n 列。
CHECK(cudaMallocHost((void **) &h_a, sizeof(int)*m*n));
CHECK(cudaMallocHost((void **) &h_b, sizeof(int)*n*k));
CHECK(cudaMallocHost((void **) &h_c, sizeof(int)*m*k));
CHECK(cudaMallocHost((void **) &h_cc, sizeof(int)*m*k));
CHECK(cudaMallocHost((void **) &h_cs, sizeof(int)*m*k));
/*CUDA事件
cudaEvent_t: 这是CUDA库中表示事件的数据类型。
CUDA事件通常用于测量 GPU 执行时间，用于确定 GPU 执行某个任务所花费的时间。
cudaEventCreate表示用于创建 CUDA 事件
原型“cudaError_t cudaEventCreate(cudaEvent_t* event);”
*/
cudaEvent_t start, stop,stop_share;
CHECK(cudaEventCreate(&start));
CHECK(cudaEventCreate(&stop));
CHECK(cudaEventCreate(&stop_share));
//将整个二维矩阵h_a赋值为1
//数组是在内存按顺序存储，因此二维其实也是一维
for (int i = 0; i < m; ++i) {
for (int j = 0; j < n; ++j) {
h_a[i * n + j] = 1;
}
}
for (int i = 0; i < n; ++i) {
for (int j = 0; j < k; ++j) {
h_b[i * k + j] = 0;
}
}
int *d_a, *d_b, *d_c, *d_c_share;
CHECK(cudaMalloc((void **) &d_a, sizeof(int)*m*n));
CHECK(cudaMalloc((void **) &d_b, sizeof(int)*n*k));
CHECK(cudaMalloc((void **) &d_c, sizeof(int)*m*k));
CHECK(cudaMalloc((void **) &d_c_share, sizeof(int)*m*k));
/* cudaDeviceSynchronize()目的是确保设备上的所有操作都已完成。
这段程序片段的目的是确保CUDA分配的设备内存已经成功完成。*/
cudaDeviceSynchronize();
// cudaEventRecord是CUDA中记录事件的函数，用于将当前时间戳记录到由start表示的事件中。
CHECK(cudaEventRecord(start));
// copy matrix A and B from host to device memory(cudaMemcpy全局内存)
CHECK(cudaMemcpy(d_a, h_a, sizeof(int)*m*n, cudaMemcpyHostToDevice));
CHECK(cudaMemcpy(d_b, h_b, sizeof(int)*n*k, cudaMemcpyHostToDevice));
unsigned int grid_rows = (m + BLOCK_SIZE - 1) / BLOCK_SIZE;
unsigned int grid_cols = (k + BLOCK_SIZE - 1) / BLOCK_SIZE;
/*
dim3三维网格和线程块，z=1，blockidx.z=0,threadidx.z=0
创建了一个二维的线程块与线程网格
*/
dim3 dimGrid(grid_cols, grid_rows);
dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
gpu_matrix_mult<<<dimGrid, dimBlock>>>(d_a, d_b, d_c, m,n,k);
CHECK(cudaMemcpy(h_c, d_c, (sizeof(int)*m*k), cudaMemcpyDeviceToHost));
/*使用这两个函数可以计算CUDA代码的执行时间，
方法是在cudaEventRecord(start)和cudaEventRecord(stop)之间执行CUDA代码
cudaEventSynchronize(stop)：等待stop事件完成，确保GPU执行到此处时
已完成前面的所有操作。
这是一种同步操作，保证我们能够准确地测量CUDA代码的执行时间。
*/
CHECK(cudaEventRecord(stop));
CHECK(cudaEventSynchronize(stop));
gpu_matrix_mult_shared<<<dimGrid, dimBlock>>> (d_a, d_b, d_c_share, m,n,k);
CHECK(cudaMemcpy(h_cs, d_c_share, (sizeof(int)*m*k), cudaMemcpyDeviceToHost));
CHECK(cudaEventRecord(stop_share));
CHECK(cudaEventSynchronize(stop_share));
//通过cudaEventElapsedTime计算两个事件之间的时间差。
float elapsed_time, elapsed_time_share;
CHECK(cudaEventElapsedTime(&elapsed_time, start, stop));
CHECK(cudaEventElapsedTime(&elapsed_time_share, stop, stop_share));
printf("Time_global = %g ms.\n", elapsed_time);
printf("Time_share = %g ms.\n", elapsed_time_share);
/*
cudaEventDestroy函数用于销毁CUDA事件。在CUDA程序中，为了避免内存泄漏，
应该在不再需要使用CUDA事件的地方调用cudaEventDestroy来释放相应的资源。*/
CHECK(cudaEventDestroy(start));
CHECK(cudaEventDestroy(stop));
//cpu_matrix_mult(h_a, h_b, h_c, m, n, k);
// int ok = 1;
/*for (int i = 0; i < m; ++i)
{
for (int j = 0; j < k; ++j)
{
if(fabs(h_cs[i*k + j] - 0)>(1.0e-10))
{
printf("hcs: %d hc: %d ",h_cs[i*k + j], h_c[i*k + j]);
ok = 0;
}
}
}
if(ok)
{
printf("Pass!!!\n");
}
else
{
printf("Error!!!\n");
}
*/
// free memory
CHECK(cudaFree(d_a));
CHECK(cudaFree(d_b));
CHECK(cudaFree(d_c));
CHECK(cudaFreeHost(h_a));
CHECK(cudaFreeHost(h_b));
CHECK(cudaFreeHost(h_c));
return 0;
}
解释部分代码：
参考链接：
https://zhuanlan.zhihu.com/p/654027980
https://www.bilibili.com/video/BV17T4y117vK?p=5
文章来源:https://blog.csdn.net/qq_44576434/article/details/135730581
本文来自互联网用户投稿，该文观点仅代表作者本人，不代表本站立场。本站仅提供信息存储空间服务，不拥有所有权，不承担相关法律责任。 如若内容造成侵权/违法违规/事实不符，请联系我的编程经验分享网邮箱：chenni525@qq.com进行投诉反馈，一经查实，立即删除！
本文来自互联网用户投稿，该文观点仅代表作者本人，不代表本站立场。本站仅提供信息存储空间服务，不拥有所有权，不承担相关法律责任。 如若内容造成侵权/违法违规/事实不符，请联系我的编程经验分享网邮箱：chenni525@qq.com进行投诉反馈，一经查实，立即删除！
最新文章
- Python教程
- 深入理解 MySQL 中的 HAVING 关键字和聚合函数
- Qt之QChar编码（1）
- MyBatis入门基础篇
- 用Python脚本实现FFmpeg批量转换
- GitHub上最适合初学者的10个最佳开源人工智能项目
- 【uniapp-scrollView滚动到底部+滚动到顶部加载更多】
- java - uniapp集成uni-push1.0
- 几种常见的UML图（类图 对象图 活动图 状态图）
- 怎么提升私域整体运营效率?
- 【通用】Linux，VSCode，IDEA，Eclipse等资源相对位置
- 内网穿透的应用-如何结合Cpolar内网穿透工具实现在IDEA中远程访问家里或者公司的数据库
- ORB-SLAM2局部建图LcalMapping函数说明
- GEE错误——使用PCA计算超时Layer error: Computation timed out.
- 聚道云软件连接器助力企业实现有赞商城与金蝶云星空系统无缝对接