(7) cuda异常处理

文章目录

- 上节概要
- 异常处理
- 代码

上节概要

上一节 block_width = 64的时候，64×64=4096 > 1024（一个block里面最多只能有1024个线程，所以这里计算会有问题）
在这里插入图片描述

异常处理

__FILE__: 编译器内部定义的一个宏。表示的是当前文件的文件名
__LINE__: 编译器内部定义的一个宏。表示的是当前文件的行

在这里插入图片描述

cuda这些函数的返回类型都是cudaError_t

//用CUDA_CHECK把cudadriver的这些套一层
CUDA_CHECK(cudaMalloc(&M_device, size));
CUDA_CHECK(cudaMalloc(&N_device, size));

LAST_KERNEL_CHECK()检查我们写的核函数相关的报错
在这里插入图片描述

代码

//util.h
#pragma once#include <cuda_runtime.h>
#include <system_error>#define CUDA_CHECK(call)             __cudaCheck(call, __FILE__, __LINE__)
#define LAST_KERNEL_CHECK()          __kernelCheck(__FILE__, __LINE__)static void __cudaCheck(cudaError_t err, const char* file, const int line) {if (err != cudaSuccess) {printf("cuda ERROR: %s:%d, ", file, line);printf("CODE:%s, DETAIL:%s\n", cudaGetErrorName(err), cudaGetErrorString(err));exit(1);}
}static void __kernelCheck(const char* file, const int line) {/* * 在编写CUDA是，错误排查非常重要，默认的cuda runtime API中的函数都会返回cudaError_t类型的结果，* 但是在写kernel函数的时候，需要通过cudaPeekAtLastError或者cudaGetLastError来获取错误*/cudaError_t err = cudaGetLastError();if (err != cudaSuccess) {printf("kernel ERROR: %s:%d, ", file, line);printf("CODE:%s, DETAIL:%s\n", cudaGetErrorName(err), cudaGetErrorString(err));exit(1);}
}void initMatrix(float* data, int size, int low, int high, int seed);void printMat(float* data, int size);void compareMat(float* h_data, float* d_data, int size);

// matmul_gpu_basic.cu
#include "matmul_basic.h"
#include "cuda_runtime.h"
#include "cuda.h"
#include "stdio.h"
#include "util.h"/* matmul的函数实现*/
__global__ void MatmulKernel(float *M_device, float *N_device, float *P_device, int width){/* 我们设定每一个thread负责P中的一个坐标的matmul所以一共有width * width个thread并行处理P的计算*/int y = blockIdx.y * blockDim.y + threadIdx.y;int x = blockIdx.x * blockDim.x + threadIdx.x;float P_element = 0;/* 对于每一个P的元素，我们只需要循环遍历width次M和N中的元素就可以了*/for (int k = 0; k < width; k ++){float M_element = M_device[y * width + k];float N_element = N_device[k * width + x];P_element += M_element * N_element;}P_device[y * width + x] = P_element;
}void MatmulOnDevice(float *M_host, float *N_host, float* P_host, int width, int blockSize){/* 设置矩阵大小 */int size = width * width * sizeof(float);/* 分配M, N在GPU上的空间*/float *M_device;float *N_device;CUDA_CHECK(cudaMalloc(&M_device, size));CUDA_CHECK(cudaMalloc(&N_device, size));/* 分配M, N拷贝到GPU上*/CUDA_CHECK(cudaMemcpy(M_device, M_host, size, cudaMemcpyHostToDevice));CUDA_CHECK(cudaMemcpy(N_device, N_host, size, cudaMemcpyHostToDevice));/* 分配P在GPU上的空间*/float *P_device;CUDA_CHECK(cudaMalloc(&P_device, size));/* 调用kernel来进行matmul计算：将一个矩阵切分成多个blockSize * blockSize的大小 */dim3 dimBlock(blockSize, blockSize);dim3 dimGrid(width / blockSize, width / blockSize);  //有width*width个blockMatmulKernel <<<dimGrid, dimBlock>>> (M_device, N_device, P_device, width);/* 将结果从device拷贝回host*/CUDA_CHECK(cudaMemcpy(P_host, P_device, size, cudaMemcpyDeviceToHost));CUDA_CHECK(cudaDeviceSynchronize());LAST_KERNEL_CHECK();/* Free */cudaFree(P_device);cudaFree(N_device);cudaFree(M_device);
}