?? 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("導入圖片%-10s ",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);
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");
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]);
if(t%4==3)
printf("\n");
}
CUT_EXIT(argc, argv);
}
?? 快捷鍵說明
復制代碼
Ctrl + C
搜索代碼
Ctrl + F
全屏模式
F11
切換主題
Ctrl + Shift + D
顯示快捷鍵
?
增大字號
Ctrl + =
減小字號
Ctrl + -