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批量转换
- 创业公司在开发软件app之前需要回答的 10 个问题
- 自然语言推断：注意力之注意（Attending）
- golang实现加密解密文档
- 【SV的虚拟口_2023.1.11】
- NFS 共享存储服务
- Redis双写一致性
- 21 UVM printer
- 芋道--如何自定义业务表单，配置对应的工作流程(详细步骤)
- 每日一练 | 华为认证真题练习Day156
- MyTinySTL 简单分析（四）--algobase.h uninitialized.h