📄 convolutionseparable.cu
字号:
/*
这个例子利用高斯内核执行一个二维信号的可分离卷积滤波
*/
#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <cutil.h>
////////////////////////////////////////////////////////////////////////////////
// 共同的主机和设备函数
////////////////////////////////////////////////////////////////////////////////
//Round a / b to nearest higher integer value ？？a/b是一个高度接近整型的值？？？？？？
int iDivUp(int a, int b){
return (a % b != 0) ? (a / b + 1) : (a / b);
}
//Round a / b to nearest lower integer value
int iDivDown(int a, int b){
return a / b;
}
//Align a to nearest higher multiple of b
int iAlignUp(int a, int b){
return (a % b != 0) ? (a - a % b + b) : a;
}
//Align a to nearest lower multiple of b
int iAlignDown(int a, int b){
return a - a % b;
}
/////////////////////////////////////////////
//读图片
//////////////////////////////////////////////
extern "C" void LoadBMPFile(uchar4 **, int *, int *, const char *);
//////////////////////////////////////////////
// 涉及到的cpu卷积
/////////////////////////////////////////////
extern "C" void convolutionRowCPU(
float *h_Result,
float *h_Data,
float *h_Kernel,
int dataW,
int dataH,
int kernelR
);
extern "C" void convolutionColumnCPU(
float *h_Result,
float *h_Data,
float *h_Kernel,
int dataW,
int dataH,
int kernelR
);
//////////////////////////////////////////////
//GPU卷积
//////////////////////////////////////////////
//全局宏，打开卷积内部控制循环
#define UNROLL_INNER //打开内部
#include <convolutionSeparable_kernel.cu>
#define M 12
//在主函数循环计算开始时传出一个虚拟计算目的是唤醒硬件驱动
#define WARMUP
/////////////////////////////////////////////
// 主函数
/////////////////////////////////////////////
int main(int argc, char **argv){
float
*h_Kernel,
*h_DataA,
*h_DataB,
*h_ResultGPU;
float
*d_DataA,
*d_DataB,
*d_Temp;
double gpuTime;
int i;
uchar4 *h_Src;
cudaArray *a_Src;
int imageW, imageH;
float Gtime[M];
float Ctime[M];
int x[M];
int y[M];
char picname[M][10]={{"1.bmp"},{"2.bmp"},{"3.bmp"},{"4.bmp"},{"5.bmp"},{"6.bmp"},
{"7.bmp"},{"8.bmp"},{"9.bmp"},{"10.bmp"},{"11.bmp"},{"12.bmp"}};
for(int t=0;t<M;t++)
{
printf("导入图片%-8s ",picname[t]);
unsigned int hTimer;
CUT_DEVICE_INIT();
CUT_SAFE_CALL(cutCreateTimer(&hTimer));
const char *image_path = cutFindFilePath(picname[t], argv[0]);//设置读取文件
LoadBMPFile(&h_Src, &imageW, &imageH, image_path);
printf(" 分别在GPU和CPU上作卷积运算，并存储数据\n");
int DATA_W=imageW;
int DATA_H=imageH;
x[t]=imageW;
y[t]=imageH;
const int DATA_SIZE = DATA_W * DATA_H * sizeof(float);
const int KERNEL_SIZE = KERNEL_W * sizeof(float);
//printf("图片格式为：%i x %i\n", x[t], y[t]);
//printf("初始化数据\n");
h_Kernel = (float *)malloc(KERNEL_SIZE);
h_DataA = (float *)malloc(DATA_SIZE);
h_DataB = (float *)malloc(DATA_SIZE);
h_ResultGPU = (float *)malloc(DATA_SIZE);
CUDA_SAFE_CALL( cudaMalloc( (void **)&d_DataA, DATA_SIZE) );
CUDA_SAFE_CALL( cudaMalloc( (void **)&d_DataB, DATA_SIZE) );
CUDA_SAFE_CALL( cudaMalloc( (void **)&d_Temp , DATA_SIZE) );
//具体是什么算法？？高斯？？？？
float kernelSum = 0;
for(i = 0; i < KERNEL_W; i++)
{
float dist = (float)(i - KERNEL_RADIUS) / (float)KERNEL_RADIUS;
h_Kernel[i] = expf(- dist * dist / 2);
kernelSum += h_Kernel[i];
}
for(i = 0; i < KERNEL_W; i++)
h_Kernel[i] /= kernelSum;
CUDA_SAFE_CALL( cudaMemcpyToSymbol(d_Kernel, h_Kernel, KERNEL_SIZE) );
CUDA_SAFE_CALL( cudaMemcpy(d_DataA, h_Src, DATA_SIZE, cudaMemcpyHostToDevice) );//导入图片数据
dim3 blockGridRows(iDivUp(DATA_W, ROW_TILE_W), DATA_H);
dim3 blockGridColumns(iDivUp(DATA_W, COLUMN_TILE_W), iDivUp(DATA_H, COLUMN_TILE_H));
dim3 threadBlockRows(KERNEL_RADIUS_ALIGNED + ROW_TILE_W + KERNEL_RADIUS);
dim3 threadBlockColumns(COLUMN_TILE_W, 8);
#ifdef WARMUP
convolutionRowGPU<<<blockGridRows, threadBlockRows>>>(
d_Temp,
d_DataA,
DATA_W,
DATA_H
);
CUT_CHECK_ERROR("GPU行卷积失败！\n");
convolutionColumnGPU<<<blockGridColumns, threadBlockColumns>>>(
d_Temp,
d_DataA,
DATA_W,
DATA_H,
COLUMN_TILE_W * threadBlockColumns.y,
DATA_W * threadBlockColumns.y
);
CUT_CHECK_ERROR("GPU列卷积失败！\n");
CUDA_SAFE_CALL( cudaThreadSynchronize() );
#endif
//printf("\n在GPU上执行可分离卷积运算：\n");
CUDA_SAFE_CALL( cudaThreadSynchronize() );
CUT_SAFE_CALL( cutResetTimer(hTimer) );
CUT_SAFE_CALL( cutStartTimer(hTimer) );
convolutionRowGPU<<<blockGridRows, threadBlockRows>>>(
d_DataB,
d_DataA,
DATA_W,
DATA_H
);
CUT_CHECK_ERROR("GPU行卷积失败！\n");
convolutionColumnGPU<<<blockGridColumns, threadBlockColumns>>>(
d_DataA,
d_DataB,
DATA_W,
DATA_H,
COLUMN_TILE_W * threadBlockColumns.y,
DATA_W * threadBlockColumns.y
);
CUT_CHECK_ERROR("GPU列卷积失败！\n");
CUDA_SAFE_CALL( cudaThreadSynchronize() );
CUT_SAFE_CALL(cutStopTimer(hTimer));
gpuTime = cutGetTimerValue(hTimer);
//printf("GPU卷积时间: %f 毫秒 || %f像素/每秒\n", gpuTime, 1e-6 * DATA_W * DATA_H / (gpuTime * 0.001));
//float q=gpuTime;
Gtime[t]=gpuTime;
//printf("读回GPU结果\n");
CUDA_SAFE_CALL( cudaMemcpy(h_ResultGPU, d_DataA, DATA_SIZE, cudaMemcpyDeviceToHost) );
//printf("\n在CPU上执行普通卷积运算：\n");
CUT_SAFE_CALL( cutResetTimer(hTimer) );
CUT_SAFE_CALL( cutStartTimer(hTimer) );
convolutionRowCPU(
h_DataB,
h_DataA,
h_Kernel,
DATA_W,
DATA_H,
KERNEL_RADIUS
);
convolutionColumnCPU(
h_DataA,
h_DataB,
h_Kernel,
DATA_W,
DATA_H,
KERNEL_RADIUS
);
CUT_SAFE_CALL(cutStopTimer(hTimer));
gpuTime = cutGetTimerValue(hTimer);
//printf("CPU卷积时间: %f 毫秒 || %f像素/每秒\n", gpuTime, 1e-6 * DATA_W * DATA_H / (gpuTime * 0.001));
//float w=gpuTime;
Ctime[t]=gpuTime;
//printf("\n比较结果：\n");
//printf("q=%f\nw=%f",q,w);
//printf("在GPU上执行可分离卷积运算速度是在CPU上执行卷积速度的%f倍！",w/q);
//printf("\n\n结束\n");
CUDA_SAFE_CALL( cudaFree(d_Temp ) );
CUDA_SAFE_CALL( cudaFree(d_DataB) );
CUDA_SAFE_CALL( cudaFree(d_DataA) );
free(h_ResultGPU);
free(h_DataB);
free(h_DataA);
free(h_Kernel);
CUT_SAFE_CALL(cutDeleteTimer(hTimer));
// CUT_EXIT(argc, argv);
}
printf("\n图片名称 大小 GPU上卷积速率 CPU上卷积速率 GPU时间是CPU上时间的多少倍\n\n");
FILE *fp;
fp=fopen("date.txt","a+");
fprintf(fp,"\n图片名称 大小 GPU上卷积时间 GPU上卷积速率 CPU上卷积时间 CPU上卷积速率 GPU时间是CPU上时间的多少倍\n\n");
fprintf(fp,"\KERNEL_RADIUS=%d\n", KERNEL_RADIUS);
for(int t=0;t<M;t++)
{
//printf("图片 %s 大小为 %d × %d 作卷积运算数据如下：\n",picname[t],x[t],y[t]);
//printf("GPU卷积时间: %f 毫秒 || %f像素/每秒\n", Gtime[t], 1e-6 * x[t] * y[t] / (Gtime[t] * 0.001));
//printf("CPU卷积时间: %f 毫秒 || %f像素/每秒\n", Ctime[t], 1e-6 * x[t] * y[t] / (Ctime[t] * 0.001));
//printf("在GPU上执行可分离卷积运算速度是在CPU上执行卷积速度的%f倍！\n\n", Ctime[t]/Gtime[t]);
//printf("%-8s %-4d× %-4d %-10f毫秒 %-10f毫秒\t %-10f\n",picname[t],x[t],y[t],Gtime[t],Ctime[t],Ctime[t]/Gtime[t]);
printf("%-8s %-4d× %-4d %-10f像素/每秒 %-10f像素/每秒\t %-10f\n",picname[t],x[t],y[t],1e-6 * x[t] * y[t] / (Gtime[t] * 0.001), 1e-6 * x[t] * y[t] / (Ctime[t] * 0.001),Ctime[t]/Gtime[t]);
fprintf(fp,"%-8s %-4d× %-4d %-10f毫秒 %-10f像素/每秒 %f 毫秒 %-10f像素/每秒\t %-10f\n",picname[t],x[t],y[t],Gtime[t],1e-6 * x[t] * y[t] / (Gtime[t] * 0.001), Ctime[t],1e-6 * x[t] * y[t] / (Ctime[t] * 0.001),Ctime[t]/Gtime[t]);
if(t%4==3)
printf("\n");
}
CUT_EXIT(argc, argv);
}
⌨️ 快捷键说明
复制代码
Ctrl + C
搜索代码
Ctrl + F
全屏模式
F11
切换主题
Ctrl + Shift + D
显示快捷键
?
增大字号
Ctrl + =
减小字号
Ctrl + -