DAY3
一、错误检测(cudaError)与事件(cudaEvent)
1、错误检测(cudaError)
在CUDA代码中有着一个特殊类型:cudaError_t,该类型可以帮助我们检查CUDA函数在执行时发生的错误。通过观察《CUDA Programing Guide》(如下图),第一句中写到，所有的CUDA运行时函数都会返回一个cudaError_t 类型，有些异步的函数需要通过cudaDeviceSynchronize()函数的返回值来判断。
代码示例如下:
int* d_a;
cudaError_t cudaStatus = cudaMalloc((void**)&d_a, sizeof(int));
通过上述代码,我们将函数的返回值存放在了cudaStatus中，可以通过判断cudaStatus来判断我们当前的CUDA函数是否执行成功，若不成功我们可以使用相应的函数来获取相关参数，具体代码如下
if (cudaStatus != cudaSuccess) {
printf("Error: %s \n%s",cudaGetErrorName(cudaStatus), cudaGetErrorString(cudaStatus));
}
具体示例如下:
#include <cuda_runtime.h>
#include <stdio.h>
#include <device_launch_parameters.h>
#define INF 0x7fffffff
int main() {
int* d_a;
cudaError_t cudaStatus = cudaMalloc((void**)&d_a, sizeof(int)*INF);
if (cudaStatus != cudaSuccess) {
printf("Error: %s \n%s",cudaGetErrorName(cudaStatus), cudaGetErrorString(cudaStatus));
}
return 0;
}
2、事件(cudaEvent)
我们都知道通过GPU可以加速我们的程序，但是我们如何知道我们的程序到底快了多少呢，此时我们可以创建cudaevent来统计我们GPU代码的运行时间，再通过比较即可得出当前获得的性能加速，相关函数如下图所示:
首先声明两个cudaEvent_t类型的指针，再通过cudaEventCreate函数对其进行初始化
cudaEvent_t kernel_start;
cudaEvent_t kernel_end;
cudaEventCreate(&kernel_start);
cudaEventCreate(&kernel_end);
cudaEventRecord函数可以帮助我们记录当前的参数,记录过程中还需涉及到两个函数 cudaEventQuery()和cudaEventSynchronize(),前者是非阻塞式的，只要执行到就进行，后者是阻塞式的，只有到前边的任务运行完毕时，才会进行，具体代码实现如下:
cudaEventRecord(kernel_start);
cudaEventQuery(kernel_start);
cuda_transpose << < gird , block >> > (d_matrix, dtr_matrix, m, n);
cudaEventRecord(kernel_end);
cudaEventSynchronize(kernel_end);
通过上述的操作我们可以获得记载了信息的kernel_start和kernel_end，执行cudaEventElapsedTime()函数可以获取两个事件之间的时间差值（即为我们核函数的执行时间）
float ms;
cudaEventElapsedTime(&ms, kernel_start, kernel_end);
具体示例如下:
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <stdio.h>
#include <stdlib.h>
#include <string>
#include <iostream>
#define BLOCK_SIZE 32
using namespace std;
__global__ void cuda_transpose(int *matrix,int *tr_matrix,int m,int n) {
int row = blockDim.y * blockIdx.y + threadIdx.y;
int col = blockDim.x * blockIdx.x + threadIdx.x;
__shared__ int smem_matrix[BLOCK_SIZE][BLOCK_SIZE];
smem_matrix[threadIdx.y][threadIdx.x] = row < m&& col < n ? matrix[row*n+col] : 0;
__syncthreads();
if(blockIdx.x * blockDim.x + threadIdx.y < n && threadIdx.x + blockIdx.y * blockDim.x < m)
tr_matrix[threadIdx.x+blockIdx.y*blockDim.x+m*(blockIdx.x*blockDim.x+threadIdx.y)] = smem_matrix[threadIdx.x][threadIdx.y];
return;
}
__host__ void cpu_transpose(int *matrix,int *tr_matrix,int m,int n) {
for (int i = 0; i < n; i++) {
for (int j = 0; j < m; j++) {
tr_matrix[i * m + j] = matrix[j * n + i];
}
}
return;
}
__host__ void init_matrix(int* matrix,int m,int n) {
for (int i = 0; i < m; i++) {
for (int j = 0; j < n; j++) {
matrix[i*n+j] = rand();
}
}
}
void print(int*, string,int,int);
bool check(int*, int*, int, int);
int main() {
int m = 1111;
int n = 113;
int *matrix;
cudaMallocHost((void**)&matrix, sizeof(int) * m * n);
init_matrix(matrix,m,n);
//print(matrix, "init matrix", m, n);
int* htr_matrix;
cudaMallocHost((void**)&htr_matrix, sizeof(int) * m * n);
cpu_transpose(matrix, htr_matrix, m, n);
//print(htr_matrix, "CPU", n, m);
//将CPU端执行的结果存放在htr_matrix中
int* d_matrix, *dtr_matrix;
cudaMalloc((void**)&d_matrix, sizeof(int) * m * n);
cudaMalloc((void**)&dtr_matrix, sizeof(int) * m * n);
cudaMemcpy(d_matrix, matrix, sizeof(int) * m * n, cudaMemcpyHostToDevice);
dim3 gird = { (unsigned int)(n - 1 + BLOCK_SIZE) / BLOCK_SIZE, (unsigned int)(m - 1 + BLOCK_SIZE) / BLOCK_SIZE,1 };
dim3 block = { BLOCK_SIZE,BLOCK_SIZE,1 };
cudaEvent_t kernel_start;
cudaEvent_t kernel_end;
cudaEventCreate(&kernel_start);
cudaEventCreate(&kernel_end);
cudaEventRecord(kernel_start);
cudaEventQuery(kernel_start);
cuda_transpose << < gird , block >> > (d_matrix, dtr_matrix, m, n);
cudaEventRecord(kernel_end);
cudaEventSynchronize(kernel_end);
float ms;
cudaEventElapsedTime(&ms, kernel_start, kernel_end);
int* hdtr_matrix;
cudaMallocHost((void**)&hdtr_matrix, sizeof(int) * m * n);
cudaMemcpy(hdtr_matrix, dtr_matrix, sizeof(int) * m * n, cudaMemcpyDeviceToDevice);
//print(hdtr_matrix, "GPU", n, m);
if (check(hdtr_matrix, htr_matrix, n, m)) {
cout << "pass\n";
}
else {
cout << "error\n";
}
printf("GPU time is : %f \n", ms);
cudaFree(hdtr_matrix);
cudaFree(dtr_matrix);
cudaFree(matrix);
cudaFree(htr_matrix);
cudaFree(d_matrix);
return 0;
}
void print(int* a, string name,int m,int n) {
cout << "NAME : " << name << endl;
for (int i = 0; i < m; i++) {
for (int j = 0; j < n; j++) {
printf("%6d ", a[i * n + j]);
}
printf("\n");
}
}
bool check(int* a, int* b, int m, int n) {
bool check_flag = true;
for (int i = 0; i < m; i++) {
for (int j = 0; j < n; j++) {
if (a[i * n + j] != b[i * n + j]) {
return false;
}
}
}
return check_flag;
}
二、多种CUDA存储单元
| 存储单元 | 位置 | cache | 访问速度（时钟周期） | 权限 | 作用域 |
|---|---|---|---|---|---|
| Register | on chip | N/A | 0.19 | R/W | thread |
| Local Memory | off chip | 无 | 203 | R/W | thread |
| Shared memory | on chip | N/A | 47 | R/W | block |
| Constant memory | off chip | 有 | 110 | R | grid |
| Global memory | off chip | 有 | 218 | R/W | grid |
| Texture Memory | off chip | 有 | 115 | R | grid |
我的另一篇文章CUDA On Arm Platfrom —Day02简要摘录开头有详细的解释，想要了解的同学，可以点进去看看。
1、常量内存(constant memory)
常量内存只有读取的权限，是global Memory虚拟出来的，没有独立的存储单元，常用于需要大量读取并且不发生改变的数据，和CPU端的const差不多，GPU端的常量内存通过__constant__声明
具体的应用如光线跟踪这个例子，光线从无限远处传来，需要通过判断当前像素点穿过的小球，来得出最终的颜色，此时我们可以申请一个常量内存，用来存储小球的数据，这样就可以有效避免了，各个线程在运算时，对小球的数据发生意外的修改，减少了代码发生错误的概率。
2、纹理内存(Texture Memory)
Texture Memory实际上也是global Memory的一部分，但是有自己专用的cache，并且该内存是GPU的特性之一，是GPU编程优化的关键。该内存是专门为那些在内存访问模式中存在大量空间局部性的图形应用程序而设计的。这意味着一个Thread读取的位置可能与邻近的Thread读取的位置“非常接近”`，如下图所示:
具体的例子如热传导模型：
需要根据相邻的块来算出当前块的值，很好的符合了纹理内存的特性。
3、全局内存(Global Memory)
Global Memory是GPU中空间最大，延迟最高的内存，也是GPU中最基础的内存，有着独立的现存颗粒。由于它的延迟最大，我们需要考虑在使用时的顺序(row-major和col-major)和次数,相邻线程之间的row-major的效率高于col-major,如下图所示:
具体的例子可以访问CUDA:矩阵转置的GPU实现(Share Memory)来查看,该例子通过share memory的row-major和col-major效率几乎相同的特性，抵消了矩阵转置带来的global Memory上的两难问题。前几天所写的矩阵相乘的代码，访问global Memory的次数巨大，我们也可以采用相同的办法来解决这一问题。
利用共享内存来消除Global Memory的高latency是我们优化CUDA程序的一种常用手法
4、共享内存(Share Memory)
Shared Memory位于GPU芯片上，访问延迟仅次于寄存器。Shared Memory是可以被一个Block中的所有Thread来进行访问的，可以实现Block内的线程间的低开销通信。
Share Memory有着独立的储存芯片，采用多bank的架构，正式因此，可以做到row-major和col-major有着几乎相同的性能。但是仍然会出现bank conflict 即访问冲突
总的来说当一个warp中的线程访问同一个bank中的不同地址时，会发生bank conflict 。我们可以在定义share memory的时候行指针多加一位的方法有效错开的我们数据排布
三、Q&A 与 体会
1、与CPU端编程的不同点
GPU编程与CPU编程相比，最大的不同即为CPU端为串行的，GPU端为并行的，这个并行的思维不仅是在一个Grid中有着许多的Block，一个Block当中有着许多的thread体现，后续学习到的CUDA stream也体现了并行的思维，所以在进行GPU端编程的时候，需要忘掉CPU端的编程思维，要获得一种全新的思维。通过学习GPU端的编程可以有效地梳理编程的模式，具体体现在，内存的传输->数据的处理->内存的传输，该经典的GPU编程模式能有效的锻炼编程的思维，让我们的编程步骤更加清晰
2、整体思维的体现
thread在实际执行过程中，往往会扮演多种角色，如使用Share Memory优化的矩阵乘法中，各个线程先参加集体活动，将Global Memory中的数据拷贝到Share Memory当中，通过__syncthreads()函数切换各自的角色再参加到运算当中去，不参加集体活动的线程不一定不参加各自的运算，所以我们要有着整体思维，不能通过某一环节的执行条件而直接屏蔽掉整个线程。
3、GPU编程中的"局限性"
该局限性体现在内存方面，在GPU中Global Memory是最大的存储单元，其它存储单元的大小都是有限， 所以我们要掌握GPU编程中的“局限性”,具体体现在通过share memory优化的矩阵乘法中，tile需要在每个sub中进行移动，可以有效的避免我们的share memory 不够大，数据最够大的情况。而在CPU端的编程中没有相关的体现
再次感谢伟大的
Nvida开发者社区！！！
3909
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
