CUDA学习笔记5——CUDA程序错误检测

已于 2023-11-17 09:07:44 修改
阅读量1.3k 收藏 1
点赞数
分类专栏： CUDA 文章标签： 学习 笔记 CUDA
于 2023-10-19 21:19:08 首次发布
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/akadiao/article/details/133907441
版权
CUDA程序错误检测

所有CUDA的API函数都有一个类型为cudaError_t的返回值，代表了一种错误信息；只有返回cudaSuccess时，才是成功调用。

  • cudaGetLastError()用来检测核函数的执行是否出错
  • cudaGetErrorString()输出错误信息
#include <stdio.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include<math.h>
#include <malloc.h> 
#include <opencv2/opencv.hpp>
#include <stdlib.h>



#define BLOCK_SIZE 1


//图像卷积 GPU
__global__ void sobel_gpu(unsigned char* in, unsigned char* out, const int Height, const int Width)
{
	int x = blockDim.x * blockIdx.x + threadIdx.x;
	int y = blockDim.y + blockIdx.y + threadIdx.y;
	int index = y * Width + x;

	int Gx = 0;
	int Gy = 0;

	unsigned char x0, x1, x2, x3, x4, x5, x6, x7, x8;

	if (x>0 && x<(Width-1) && y>0 && y<(Height-1))
	{
		x0 = in[(y - 1)*Width + (x - 1)];
		x1 = in[(y - 1)*Width + (x)];
		x2 = in[(y - 1)*Width + (x + 1)];
		x3 = in[(y)*Width + (x - 1)];

		x5 = in[(y)*Width + (x + 1)];
		x6 = in[(y + 1)*Width + (x - 1)];
		x7 = in[(y + 1)*Width + (x)];
		x8 = in[(y + 1)*Width + (x + 1)];

		Gx = (x0 + 2 * x3 + x6) - (x2 + 2 * x5 + x8);
		Gy = (x0 + 2 * x1 + x2) - (x6 + 2 * x7 + x8);

		out[index] = (abs(Gx) + abs(Gy)) / 2;
	}
}



int main()
{
	cv::Mat src;
	src = cv::imread("complete004.jpg");

	cv::Mat grayImg,gaussImg;
	cv::cvtColor(src, grayImg, cv::COLOR_BGR2GRAY);
	cv::GaussianBlur(grayImg, gaussImg, cv::Size(3,3), 0, 0, cv::BORDER_DEFAULT);

	int height = src.rows;
	int width = src.cols;
	//输出图像
	cv::Mat dst_gpu(height, width, CV_8UC1, cv::Scalar(0));
	//GPU存储空间
	int memsize = height * width * sizeof(unsigned char);
	//输入 输出
	unsigned char* in_gpu;
	unsigned char* out_gpu;

	cudaMalloc((void**)&in_gpu, memsize);
	cudaMalloc((void**)&out_gpu, memsize);
	cudaError_t error_code;

	dim3 threadsPreBlock(BLOCK_SIZE, BLOCK_SIZE);
	dim3 blocksPreGrid((width + threadsPreBlock.x - 1)/threadsPreBlock.x, (height + threadsPreBlock.y - 1)/threadsPreBlock.y);
	
	cudaMemcpy(in_gpu, gaussImg.data, memsize, cudaMemcpyHostToDevice);

	sobel_gpu <<<blocksPreGrid, threadsPreBlock>>> (in_gpu, out_gpu, height, width);
	error_code = cudaGetLastError();
	
	printf("Error: %s\n", cudaGetErrorString(error_code));
	printf("FILE: %s\n", __FILE__);
	printf("LINE: %d\n", __LINE__);
	printf("Error code: %d\n", error_code);

	cudaMemcpy(dst_gpu.data, out_gpu, memsize, cudaMemcpyDeviceToHost);

	cv::imwrite("dst_gpu_save.png", dst_gpu);

	//cv::namedWindow("src", cv::WINDOW_NORMAL);
	cv::imshow("src", src);
	cv::imshow("dst_gpu", dst_gpu);
	cv::waitKey();

	cudaFree(in_gpu);
	cudaFree(out_gpu);

	return 0;
}

在这里插入图片描述在这里插入图片描述

樊哲勇大牛的检测CUDA运行时错误的宏函数：

#pragma once
#include<stdio.h>

#define CHECK(call)																\
do                                                                              \
{																				\
	const cudaError_t error_code = call;										\
	if (error_code != cudaSuccess)												\
	{																			\
		printf("CUDA Error:\n");												\
		printf("	File:	%s\n", __FILE__);									\
		printf("	Line:	%d\n",__LINE__);									\
		printf("	Error code: %d\n",error_code);								\
		printf("	Error text: %s\n", cudaGetErrorString(error_code));			\
		exit(1);																\
	}																			\
} while (0)
采用检测CUDA运行时错误的宏函数，图像卷积：
#include <stdio.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include<math.h>
#include <malloc.h> 
#include <opencv2/opencv.hpp>
#include <stdlib.h>

#include "error.cuh"

#define BLOCK_SIZE 1


//图像卷积 GPU
__global__ void sobel_gpu(unsigned char* in, unsigned char* out, const int Height, const int Width)
{
	int x = blockDim.x * blockIdx.x + threadIdx.x;
	int y = blockDim.y + blockIdx.y + threadIdx.y;
	int index = y * Width + x;

	int Gx = 0;
	int Gy = 0;

	unsigned char x0, x1, x2, x3, x4, x5, x6, x7, x8;

	if (x>0 && x<(Width-1) && y>0 && y<(Height-1))
	{
		x0 = in[(y - 1)*Width + (x - 1)];
		x1 = in[(y - 1)*Width + (x)];
		x2 = in[(y - 1)*Width + (x + 1)];
		x3 = in[(y)*Width + (x - 1)];

		x5 = in[(y)*Width + (x + 1)];
		x6 = in[(y + 1)*Width + (x - 1)];
		x7 = in[(y + 1)*Width + (x)];
		x8 = in[(y + 1)*Width + (x + 1)];

		Gx = (x0 + 2 * x3 + x6) - (x2 + 2 * x5 + x8);
		Gy = (x0 + 2 * x1 + x2) - (x6 + 2 * x7 + x8);

		out[index] = (abs(Gx) + abs(Gy)) / 2;
	}
}



int main()
{
	cv::Mat src;
	src = cv::imread("complete004.jpg");

	cv::Mat grayImg,gaussImg;
	cv::cvtColor(src, grayImg, cv::COLOR_BGR2GRAY);
	cv::GaussianBlur(grayImg, gaussImg, cv::Size(3,3), 0, 0, cv::BORDER_DEFAULT);

	int height = src.rows;
	int width = src.cols;
	//输出图像
	cv::Mat dst_gpu(height, width, CV_8UC1, cv::Scalar(0));
	//GPU存储空间
	int memsize = height * width * sizeof(unsigned char);
	//输入 输出
	unsigned char* in_gpu;
	unsigned char* out_gpu;

	cudaMalloc((void**)&in_gpu, memsize);
	cudaMalloc((void**)&out_gpu, memsize);

	dim3 threadsPreBlock(BLOCK_SIZE, BLOCK_SIZE);
	dim3 blocksPreGrid((width + threadsPreBlock.x - 1)/threadsPreBlock.x, (height + threadsPreBlock.y - 1)/threadsPreBlock.y);
	
	cudaMemcpy(in_gpu, gaussImg.data, memsize, cudaMemcpyHostToDevice);

	sobel_gpu <<<blocksPreGrid, threadsPreBlock>>> (in_gpu, out_gpu, height, width);

	CHECK(cudaMemcpy(dst_gpu.data, out_gpu, memsize*10, cudaMemcpyDeviceToHost));//增大size值  引起报错

	cv::imwrite("dst_gpu_save.png", dst_gpu);

	//cv::namedWindow("src", cv::WINDOW_NORMAL);
	cv::imshow("src", src);
	cv::imshow("dst_gpu", dst_gpu);
	cv::waitKey();

	cudaFree(in_gpu);
	cudaFree(out_gpu);

	return 0;
}

在这里插入图片描述

采用检测CUDA运行时错误的宏函数，矩阵相加：

二维网络和二维线程块对二维矩阵进行索引，每个线程可负责一个矩阵元素的计算任务；
ix=threadIdx.x+blockIdx.xblockDim.x
iy=threadIdx.y+blockIdx.yblockDim.y
在这里插入图片描述

#include <stdio.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "error.cuh"
#include <stdlib.h>
#include <iostream>


__global__ void addMatrix(int *input_1, int *input_2, int *output, const int nx, const int ny)
{
	int ix = threadIdx.x + blockIdx.x * blockDim.x;
	int iy = threadIdx.y + blockIdx.y * blockDim.y;
	unsigned int idx = iy * nx + ix;

	if (ix < nx && iy < ny)
	{
		output[idx] = input_1[idx] + input_2[idx];
	}

}


int main(void)
{

	int nx = 16;
	int ny = 8;
	int nxy = nx * ny;
	size_t stBytesCount = nxy * sizeof(int);

	int *ipHost_A, *ipHost_B, *ipHost_C;
	ipHost_A = (int *)malloc(stBytesCount);
	ipHost_B = (int *)malloc(stBytesCount);
	ipHost_C = (int *)malloc(stBytesCount);
	if (ipHost_A != NULL && ipHost_B != NULL && ipHost_C != NULL)
	{
		for (int i = 0; i < nxy; i++)
		{
			ipHost_A[i] = i;
			ipHost_B[i] = i + 1;
		}
		memset(ipHost_C, 0, stBytesCount);
	}
	else
	{
		printf("Fail to allocate host memory! \n");
	}
	// 分配内存 初始化
	int *ipDevice_A, *ipDevice_B, *ipDevice_C;
	CHECK(cudaMalloc((int**)&ipDevice_A, stBytesCount));
	CHECK(cudaMalloc((int**)&ipDevice_B, stBytesCount));
	CHECK(cudaMalloc((int**)&ipDevice_C, stBytesCount));
	if (ipDevice_A != NULL && ipDevice_B != NULL && ipDevice_C != NULL)
	{
		CHECK(cudaMemcpy(ipDevice_A, ipHost_A, stBytesCount, cudaMemcpyHostToDevice));
		CHECK(cudaMemcpy(ipDevice_B, ipHost_B, stBytesCount, cudaMemcpyHostToDevice));
		CHECK(cudaMemcpy(ipDevice_C, ipHost_C, stBytesCount, cudaMemcpyHostToDevice));
	}
	else
	{
		printf("Fail to allocate memory \n");
		free(ipHost_A);
		free(ipHost_B);
		free(ipHost_C);
		exit(1);
	}
	//
	dim3 block(4,4);//线程块大小 4*4
	dim3 grid((nx + block.x-1)/block.x, (ny+block.y - 1)/block.y);
	printf("Thread config:grid:<%d,%d>, block<%d,%d>\n", grid.x, grid.y, block.x, block.y);
	// 调用核函数
	addMatrix <<< grid, block>>> (ipDevice_A, ipDevice_B, ipDevice_C, nx, ny);
	//拷贝出结果
	CHECK(cudaMemcpy(ipHost_C, ipDevice_C, stBytesCount,cudaMemcpyDeviceToHost));

	for (int i = 0; i < 10; i++)
	{
		printf("id:%d, martix_A: %d, martix_B: %d \n", i+1, ipHost_A[i], ipHost_B[i], ipHost_C[i]);
	}
	if (ipHost_C != NULL)
	{
		for (int i = 0; i < nxy; i++)
		{
			printf("C[%d]: %d ", i, ipHost_C[i]);
			if ((i+1)%16 == 0)
			{
				printf("\n");
			}
		}
	}
	//释放主机内存
	free(ipHost_A);
	free(ipHost_B);
	free(ipHost_C);

	CHECK(cudaFree(ipDevice_A));
	CHECK(cudaFree(ipDevice_B));
	CHECK(cudaFree(ipDevice_C));
	
	CHECK(cudaDeviceReset());//清空申请的当前关联gpu设备资源

	return 0;


}

在这里插入图片描述

[9] CUDA性能测量与错误处理
yohnyang的博客
05-28 673
[9] CUDA性能测量与错误处理
英伟达DeepStream学习笔记53——RTSP流输入
耶律大石的博客
04-08 95
【代码】英伟达DeepStream学习笔记53——RTSP流输入。
nVIDIA显卡CUDA性能测试工具，可测试nVIDIA显卡的浮点运算性能
12-20
nVIDIA显卡CUDA性能测试工具，可测试nVIDIA显卡的参数，nVIDIA显卡的CUDA浮点运算性能
CUDA on Platform 学习笔记5--错误检测与事件
lansebingxuan的专栏
07-08 892
本文介绍CUDA运行时的错误检测函数，如何使用错误检测函数，以及利用CUDA的event进行计时
cmake 如何检测 CUDA 环境
最新发布
eloudy的专栏
03-19 211
这样就把具体的 nvcc的路径赋值为 变量 CMAKE_CUDA_COMPILER 的值，然后，这个result.cmake 会被 include进 check_language 脚本，于是 赋值后的CMAKE_CUDA_COMPILER就成为 check_languange 脚本中的一个携带了 nvcc 信息的目录。来找到变量 CMAKE_CUDA_COMPILER 的值，例如是 /usr/local/cuda/bin/nvcc。其中，cmProjectCommand 是一个cpp函数，
CUDA编程 4 错误检测，事件及存储单元 【参加CUDA线上训练营】
m0_60269818的博客
02-14 671
流中的任意点都可以通过API插入事件以及查询事件完成的函数，只有事件所在流中其之前的操作都完成后才能触发事件完成。在Cuda编程中，通常会使用尽量小的存储单元，提高程序的性能。利用共享存储单元加速矩阵乘法的思路是，将一个矩阵分块，并将每一块的部分数据加载到共享存储单元中，以便在单个 CUDA 核心中重复访问多次。Cuda编程模型中的事件。加载数据到共享存储单元中：每个 CUDA 核心加载其对应的小矩阵中的数据到共享存储单元中。通过使用共享存储单元，可以减少对全局存储单元的访问，从而提高矩阵乘法的性能。
CUDA错误检查
weixin_43915090的博客
01-31 521
CUDA api执行结束回返回一个返回值，用于标记api的执行状态。
基于《CUDA 编程-基础与实践》的 CUDA 学习笔记.zip
11-19
它包含了 GPU 硬件与 CUDA 程序开发工具、CUDA 中的线程组织、CUDA 程序的基本框架和错误检测、GPU 加速的关键、CUDA 内存组织、CUDA 标准库等内容。 主要功能点 提供 CUDA 编程的基础知识和实践案例 包含 CUDA ...
CUDA学习笔记（四）——CUDA性能
05-19
CUDA是一个可以大大提高GPU计算性能的编程平台，但是为了实现最高的性能，需要遵循一些最佳实践和优化技巧。以下是一些常用的CUDA性能优化方法： 1. 矢量化：使用SIMD指令和数据结构（如向量和矩阵）来并行计算。 ...
深度学习（基于Tensorflow2.0）学习笔记——Day2
01-20
本篇学习笔记主要关注如何在Tensorflow2.0环境下开始深度学习的学习和实践。 【深度学习基础知识回顾】 在开始实战前，对深度学习的基础知识进行回顾是必要的。这包括神经网络的基本概念、反向传播算法、损失函数...
英伟达DeepStream学习笔记52——在X86上运行deepstream 6.0.1（非docker）
耶律大石的博客
02-23 128
安装 librdkafka（为消息代理启用 Kafka 协议适配器）然后再次 sudo apt-get update ，成功！我安装了470.161.03，cuda版本11.4。安装完成之后，加入到环境变量，可用gedit。在.bashrc末尾添加两行环境变量。最后，使环境变量生效。
cuda检测工具 devicequery.exe
08-03
deviceQuery.exe Starting... CUDA Device Query (Runtime API) version (CUDART static linking) Detected 1 CUDA Capable device(s) Device 0: "GeForce GTX 650" CUDA Driver Version / Runtime Version 9.1 / 8.0 CUDA Capability Major/Minor version number: 3.0 Total amount of global memory: 2048 MBytes (2147483648 bytes) ( 2) Multiprocessors, (192) CUDA Cores/MP: 384 CUDA Cores GPU Max Clock rate: 1072 MHz (1.07 GHz) Memory Clock rate: 2500 Mhz Memory Bus Width: 128-bit L2 Cache Size: 262144 bytes Maximum Texture Dimension Size (x,y,z) 1D=(65536), 2D=(65536, 65536), 3D=(4096, 4096, 4096) Maximum Layered 1D Texture Size, (num) layers 1D=(16384), 2048 layers Maximum Layered 2D Texture Size, (num) layers 2D=(16384, 16384), 2048 layers Total amount of constant memory: 65536 bytes Total amount of shared memory per block: 49152 bytes Total number of registers available per block: 65536 Warp size: 32 Maximum number of threads per multiprocessor: 2048 Maximum number of threads per block: 1024 Max dimension size of a thread block (x,y,z): (1024, 1024, 64) Max dimension size of a grid size (x,y,z): (2147483647, 65535, 65535) Maximum memory pitch: 2147483647 bytes Texture alignment: 512 bytes Concurrent copy and kernel execution: Yes with 1 copy engine(s) Run time limit on kernels: Yes Integrated GPU sharing Host Memory: No Support host page-locked memory mapping: Yes Alignment requirement for Surfaces: Yes Device has ECC support: Disabled CUDA Device Driver Mode (TCC or WDDM): WDDM (Windows Display Driver Model) Device supports Unified Addressing (UVA): Yes Device PCI Domain ID / Bus ID / location ID: 0 / 1 / 0 Compute Mode: deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 9.1, CUDA Runtime Version = 8.0, NumDevs = 1, Device0 = GeForce GTX 650 Result = PASS
Cuda性能测试
04-28
基于VS2013得威cuda的性能测试，比较了不同规模运算时的cpu和gpu的运算时间，gpu在大规模运算更具优势
CUDA SDK测试程序全解析
04-13
CUDA SDK是CUDA加速常常使用的优化方法的例子，包括已经实现的一些好的并行算法，这里对CUDA SDK程序进行了详细解析。
《CUDA编程》4.CUDA程序的错误检测
立志成为炼丹师？
10-05 1286
在《CUDA编程》3.简单CUDA程序的基本框架 中列举的函数，返回值是，只有在返回cudaSuccess时，才表示调用成功，否则返回一个错误代码，下面新建一个CUDA头文件并编写一个错误检查的宏函数：①新建CUDA头文件 新建的文件是，注意后缀变化。 ②编写错误检查代码 在定义宏时，如果一行写不下，需要在行末写 \，表示续行*，错误检查代码如下： 该段代码会检查返回值是否为，如果不是，则返回错误代码的位置这里以《CUDA编程》3.简单CUDA程序的基本框架中的代码例子为例，注意在头文件中添加，并为代码中
CUDA错误检测与事件【参加CUDA线上训练营】
Mr_8j的博客
02-10 435
CUDA错误检测与事件【参加CUDA线上训练营】
cuda测试示例
hscoder的博客
01-04 4237
最近闲来无事，下载一个cuda，感受一下gpu强大的运算能力，选择cuda7.5，和vs2012的环境，参考网上的一些参考教程，费了一番周折终于配置成功。用一下代码测试/************************************************************************/ /* cuda测试
CUDA9.0 C++测试程序
qq_45298704的博客
08-19 648
【代码】CUDA9.0 C++测试程序。
Cuda API error detected: cudaLaunch returned
sinat_28750977的博客
04-04 1399
questions from stack overflow: cudaGetLastError returned (0xb) What is the canonical way to check for errors using the CUDA runtime API?
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值