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;
}