zara
2 years ago
commit
596bd34c58
5 changed files with 666 additions and 0 deletions
@ -0,0 +1,363 @@
@@ -0,0 +1,363 @@
|
||||
## Ignore Visual Studio temporary files, build results, and |
||||
## files generated by popular Visual Studio add-ons. |
||||
## |
||||
## Get latest from https://github.com/github/gitignore/blob/master/VisualStudio.gitignore |
||||
|
||||
# User-specific files |
||||
*.rsuser |
||||
*.suo |
||||
*.user |
||||
*.userosscache |
||||
*.sln.docstates |
||||
|
||||
# User-specific files (MonoDevelop/Xamarin Studio) |
||||
*.userprefs |
||||
|
||||
# Mono auto generated files |
||||
mono_crash.* |
||||
|
||||
# Build results |
||||
[Dd]ebug/ |
||||
[Dd]ebugPublic/ |
||||
[Rr]elease/ |
||||
[Rr]eleases/ |
||||
x64/ |
||||
x86/ |
||||
[Ww][Ii][Nn]32/ |
||||
[Aa][Rr][Mm]/ |
||||
[Aa][Rr][Mm]64/ |
||||
bld/ |
||||
[Bb]in/ |
||||
[Oo]bj/ |
||||
[Oo]ut/ |
||||
[Ll]og/ |
||||
[Ll]ogs/ |
||||
|
||||
# Visual Studio 2015/2017 cache/options directory |
||||
.vs/ |
||||
# Uncomment if you have tasks that create the project's static files in wwwroot |
||||
#wwwroot/ |
||||
|
||||
# Visual Studio 2017 auto generated files |
||||
Generated\ Files/ |
||||
|
||||
# MSTest test Results |
||||
[Tt]est[Rr]esult*/ |
||||
[Bb]uild[Ll]og.* |
||||
|
||||
# NUnit |
||||
*.VisualState.xml |
||||
TestResult.xml |
||||
nunit-*.xml |
||||
|
||||
# Build Results of an ATL Project |
||||
[Dd]ebugPS/ |
||||
[Rr]eleasePS/ |
||||
dlldata.c |
||||
|
||||
# Benchmark Results |
||||
BenchmarkDotNet.Artifacts/ |
||||
|
||||
# .NET Core |
||||
project.lock.json |
||||
project.fragment.lock.json |
||||
artifacts/ |
||||
|
||||
# ASP.NET Scaffolding |
||||
ScaffoldingReadMe.txt |
||||
|
||||
# StyleCop |
||||
StyleCopReport.xml |
||||
|
||||
# Files built by Visual Studio |
||||
*_i.c |
||||
*_p.c |
||||
*_h.h |
||||
*.ilk |
||||
*.meta |
||||
*.obj |
||||
*.iobj |
||||
*.pch |
||||
*.pdb |
||||
*.ipdb |
||||
*.pgc |
||||
*.pgd |
||||
*.rsp |
||||
*.sbr |
||||
*.tlb |
||||
*.tli |
||||
*.tlh |
||||
*.tmp |
||||
*.tmp_proj |
||||
*_wpftmp.csproj |
||||
*.log |
||||
*.vspscc |
||||
*.vssscc |
||||
.builds |
||||
*.pidb |
||||
*.svclog |
||||
*.scc |
||||
|
||||
# Chutzpah Test files |
||||
_Chutzpah* |
||||
|
||||
# Visual C++ cache files |
||||
ipch/ |
||||
*.aps |
||||
*.ncb |
||||
*.opendb |
||||
*.opensdf |
||||
*.sdf |
||||
*.cachefile |
||||
*.VC.db |
||||
*.VC.VC.opendb |
||||
|
||||
# Visual Studio profiler |
||||
*.psess |
||||
*.vsp |
||||
*.vspx |
||||
*.sap |
||||
|
||||
# Visual Studio Trace Files |
||||
*.e2e |
||||
|
||||
# TFS 2012 Local Workspace |
||||
$tf/ |
||||
|
||||
# Guidance Automation Toolkit |
||||
*.gpState |
||||
|
||||
# ReSharper is a .NET coding add-in |
||||
_ReSharper*/ |
||||
*.[Rr]e[Ss]harper |
||||
*.DotSettings.user |
||||
|
||||
# TeamCity is a build add-in |
||||
_TeamCity* |
||||
|
||||
# DotCover is a Code Coverage Tool |
||||
*.dotCover |
||||
|
||||
# AxoCover is a Code Coverage Tool |
||||
.axoCover/* |
||||
!.axoCover/settings.json |
||||
|
||||
# Coverlet is a free, cross platform Code Coverage Tool |
||||
coverage*.json |
||||
coverage*.xml |
||||
coverage*.info |
||||
|
||||
# Visual Studio code coverage results |
||||
*.coverage |
||||
*.coveragexml |
||||
|
||||
# NCrunch |
||||
_NCrunch_* |
||||
.*crunch*.local.xml |
||||
nCrunchTemp_* |
||||
|
||||
# MightyMoose |
||||
*.mm.* |
||||
AutoTest.Net/ |
||||
|
||||
# Web workbench (sass) |
||||
.sass-cache/ |
||||
|
||||
# Installshield output folder |
||||
[Ee]xpress/ |
||||
|
||||
# DocProject is a documentation generator add-in |
||||
DocProject/buildhelp/ |
||||
DocProject/Help/*.HxT |
||||
DocProject/Help/*.HxC |
||||
DocProject/Help/*.hhc |
||||
DocProject/Help/*.hhk |
||||
DocProject/Help/*.hhp |
||||
DocProject/Help/Html2 |
||||
DocProject/Help/html |
||||
|
||||
# Click-Once directory |
||||
publish/ |
||||
|
||||
# Publish Web Output |
||||
*.[Pp]ublish.xml |
||||
*.azurePubxml |
||||
# Note: Comment the next line if you want to checkin your web deploy settings, |
||||
# but database connection strings (with potential passwords) will be unencrypted |
||||
*.pubxml |
||||
*.publishproj |
||||
|
||||
# Microsoft Azure Web App publish settings. Comment the next line if you want to |
||||
# checkin your Azure Web App publish settings, but sensitive information contained |
||||
# in these scripts will be unencrypted |
||||
PublishScripts/ |
||||
|
||||
# NuGet Packages |
||||
*.nupkg |
||||
# NuGet Symbol Packages |
||||
*.snupkg |
||||
# The packages folder can be ignored because of Package Restore |
||||
**/[Pp]ackages/* |
||||
# except build/, which is used as an MSBuild target. |
||||
!**/[Pp]ackages/build/ |
||||
# Uncomment if necessary however generally it will be regenerated when needed |
||||
#!**/[Pp]ackages/repositories.config |
||||
# NuGet v3's project.json files produces more ignorable files |
||||
*.nuget.props |
||||
*.nuget.targets |
||||
|
||||
# Microsoft Azure Build Output |
||||
csx/ |
||||
*.build.csdef |
||||
|
||||
# Microsoft Azure Emulator |
||||
ecf/ |
||||
rcf/ |
||||
|
||||
# Windows Store app package directories and files |
||||
AppPackages/ |
||||
BundleArtifacts/ |
||||
Package.StoreAssociation.xml |
||||
_pkginfo.txt |
||||
*.appx |
||||
*.appxbundle |
||||
*.appxupload |
||||
|
||||
# Visual Studio cache files |
||||
# files ending in .cache can be ignored |
||||
*.[Cc]ache |
||||
# but keep track of directories ending in .cache |
||||
!?*.[Cc]ache/ |
||||
|
||||
# Others |
||||
ClientBin/ |
||||
~$* |
||||
*~ |
||||
*.dbmdl |
||||
*.dbproj.schemaview |
||||
*.jfm |
||||
*.pfx |
||||
*.publishsettings |
||||
orleans.codegen.cs |
||||
|
||||
# Including strong name files can present a security risk |
||||
# (https://github.com/github/gitignore/pull/2483#issue-259490424) |
||||
#*.snk |
||||
|
||||
# Since there are multiple workflows, uncomment next line to ignore bower_components |
||||
# (https://github.com/github/gitignore/pull/1529#issuecomment-104372622) |
||||
#bower_components/ |
||||
|
||||
# RIA/Silverlight projects |
||||
Generated_Code/ |
||||
|
||||
# Backup & report files from converting an old project file |
||||
# to a newer Visual Studio version. Backup files are not needed, |
||||
# because we have git ;-) |
||||
_UpgradeReport_Files/ |
||||
Backup*/ |
||||
UpgradeLog*.XML |
||||
UpgradeLog*.htm |
||||
ServiceFabricBackup/ |
||||
*.rptproj.bak |
||||
|
||||
# SQL Server files |
||||
*.mdf |
||||
*.ldf |
||||
*.ndf |
||||
|
||||
# Business Intelligence projects |
||||
*.rdl.data |
||||
*.bim.layout |
||||
*.bim_*.settings |
||||
*.rptproj.rsuser |
||||
*- [Bb]ackup.rdl |
||||
*- [Bb]ackup ([0-9]).rdl |
||||
*- [Bb]ackup ([0-9][0-9]).rdl |
||||
|
||||
# Microsoft Fakes |
||||
FakesAssemblies/ |
||||
|
||||
# GhostDoc plugin setting file |
||||
*.GhostDoc.xml |
||||
|
||||
# Node.js Tools for Visual Studio |
||||
.ntvs_analysis.dat |
||||
node_modules/ |
||||
|
||||
# Visual Studio 6 build log |
||||
*.plg |
||||
|
||||
# Visual Studio 6 workspace options file |
||||
*.opt |
||||
|
||||
# Visual Studio 6 auto-generated workspace file (contains which files were open etc.) |
||||
*.vbw |
||||
|
||||
# Visual Studio LightSwitch build output |
||||
**/*.HTMLClient/GeneratedArtifacts |
||||
**/*.DesktopClient/GeneratedArtifacts |
||||
**/*.DesktopClient/ModelManifest.xml |
||||
**/*.Server/GeneratedArtifacts |
||||
**/*.Server/ModelManifest.xml |
||||
_Pvt_Extensions |
||||
|
||||
# Paket dependency manager |
||||
.paket/paket.exe |
||||
paket-files/ |
||||
|
||||
# FAKE - F# Make |
||||
.fake/ |
||||
|
||||
# CodeRush personal settings |
||||
.cr/personal |
||||
|
||||
# Python Tools for Visual Studio (PTVS) |
||||
__pycache__/ |
||||
*.pyc |
||||
|
||||
# Cake - Uncomment if you are using it |
||||
# tools/** |
||||
# !tools/packages.config |
||||
|
||||
# Tabs Studio |
||||
*.tss |
||||
|
||||
# Telerik's JustMock configuration file |
||||
*.jmconfig |
||||
|
||||
# BizTalk build output |
||||
*.btp.cs |
||||
*.btm.cs |
||||
*.odx.cs |
||||
*.xsd.cs |
||||
|
||||
# OpenCover UI analysis results |
||||
OpenCover/ |
||||
|
||||
# Azure Stream Analytics local run output |
||||
ASALocalRun/ |
||||
|
||||
# MSBuild Binary and Structured Log |
||||
*.binlog |
||||
|
||||
# NVidia Nsight GPU debugger configuration file |
||||
*.nvuser |
||||
|
||||
# MFractors (Xamarin productivity tool) working folder |
||||
.mfractor/ |
||||
|
||||
# Local History for Visual Studio |
||||
.localhistory/ |
||||
|
||||
# BeatPulse healthcheck temp database |
||||
healthchecksdb |
||||
|
||||
# Backup folder for Package Reference Convert tool in Visual Studio 2017 |
||||
MigrationBackup/ |
||||
|
||||
# Ionide (cross platform F# VS Code tools) working folder |
||||
.ionide/ |
||||
|
||||
# Fody - auto-generated XML schema |
||||
FodyWeavers.xsd |
||||
@ -0,0 +1,182 @@
@@ -0,0 +1,182 @@
|
||||
#include "bila.cuh" |
||||
|
||||
|
||||
|
||||
#include "cuda_runtime.h" |
||||
#include "device_launch_parameters.h" |
||||
|
||||
#include <stdio.h> |
||||
|
||||
#define threadX 16 |
||||
#define threadY 16 |
||||
#define threadZ 1 |
||||
|
||||
__global__ void bilakernel(uchar3 * src, uchar3* dst, int width, int height, int r, double sigmaC, double sigmaS) |
||||
{ |
||||
int x = threadIdx.x + blockDim.x * blockIdx.x; |
||||
int y = threadIdx.y + blockDim.y * blockIdx.y; |
||||
|
||||
if (x >= width || y >= height)return; |
||||
double mask = 0; |
||||
double3 result; |
||||
|
||||
// mask.x = 0; |
||||
// mask.y = 0; |
||||
// mask.z = 0; |
||||
result.x = 0; |
||||
result.y = 0; |
||||
result.z = 0; |
||||
|
||||
uchar3 origin = src[x + y * width]; |
||||
|
||||
for(int i = -r; i<=r;i++ ) |
||||
{ |
||||
for(int j = -r; j<=r;j++) |
||||
{ |
||||
if(x+i>=0&&x+i<width&&y+j>=0&&y+j<height) |
||||
{ |
||||
uchar3 tmp = src[x + i + (y + j) * width]; |
||||
int3 sub; |
||||
sub.x = tmp.x - origin.x; |
||||
sub.y = tmp.y - origin.y; |
||||
sub.z = tmp.z - origin.z; |
||||
double masktmp = (exp(-0.5 / sigmaC / sigmaC * sub.x * sub.x) + |
||||
exp(-0.5 / sigmaC / sigmaC * sub.y * sub.y) + |
||||
exp(-0.5 / sigmaC / sigmaC * sub.z * sub.z)) * |
||||
exp(-0.5 / sigmaS / sigmaS * (i * i + j * j)); |
||||
mask += masktmp; |
||||
result.x += tmp.x * masktmp; |
||||
result.y += tmp.y * masktmp; |
||||
result.z += tmp.z * masktmp; |
||||
} |
||||
} |
||||
} |
||||
|
||||
if(mask>0) |
||||
{ |
||||
dst[x + y * width].x = result.x / mask; |
||||
dst[x + y * width].y = result.y / mask; |
||||
dst[x + y * width].z = result.z / mask; |
||||
} |
||||
else |
||||
{ |
||||
dst[x + y * width].x = 0; |
||||
dst[x + y * width].y = 0; |
||||
dst[x + y * width].z = 0; |
||||
} |
||||
// if(x == 0 && y == 0) |
||||
// { |
||||
// printf("\n\n\n\n\n\n\n\n/////////////// on gpu /////////////////// \n\n grid dim = (%d, %d) blockdim = (%d,%d)\n\n\n\n\n",gridDim.x,gridDim.y,blockDim.x,blockDim.y); |
||||
// } |
||||
} |
||||
|
||||
|
||||
__global__ void bilakernel_v2(uchar3 * src, uchar3* dst,int width,int height, int r, double sigmaC,double sigmaD ) |
||||
{ |
||||
|
||||
int x = threadIdx.x + blockDim.x * blockIdx.x; |
||||
int y = threadIdx.y + blockDim.y * blockIdx.y; |
||||
|
||||
|
||||
|
||||
if (x >= width || y >= height)return; |
||||
int id = x + y * width; |
||||
__shared__ uchar3 box[(2 * threadX) * (2 * threadY)]; |
||||
for(int i = 0; i< 999999;i++) |
||||
{ |
||||
int localIndex = threadIdx.x + threadIdx.y * blockIdx.x + i*blockIdx.x*blockIdx.y; |
||||
if(localIndex<(2*r+blockDim.x)*(2*r+blockDim.y)) |
||||
{ |
||||
int x_temp = localIndex % (blockDim.x + 2 * r); |
||||
int y_temp = localIndex / (blockDim.x + 2 * r); |
||||
// if(x_temp-r+blockDim.x*blockIdx.x<0 || y_temp-r+blockDim.y*blockIdx.y|| |
||||
// x_temp+r+blockDim.x*(blockIdx.x+1)-1>width) |
||||
int x_real = x_temp - r + blockDim.x*blockIdx.x; |
||||
int y_real = y_temp - r + blockDim.y * blockIdx.y; |
||||
if(x_real<0||x_real>=width||y_real<0||y>=height) |
||||
{ |
||||
box[localIndex].x = 0; |
||||
box[localIndex].y = 0; |
||||
box[localIndex].z = 0; |
||||
} |
||||
else |
||||
{ |
||||
box[localIndex] = src[x_temp + y_temp * width]; |
||||
} |
||||
} |
||||
else |
||||
{ |
||||
break; |
||||
} |
||||
} |
||||
// __threadsync(); |
||||
double coeff = 0; |
||||
double sum[3] = { 0 }; |
||||
for(int i = -r; i<=r;i++) |
||||
{ |
||||
for(int j = -r;j<=r;j++) |
||||
{ |
||||
if(x+i>=0&&x+i<width&&y+j>=0&&y+j<height) |
||||
{ |
||||
double tmp = ( |
||||
exp(-0.5 / sigmaC / sigmaC * (box[i + r + (j + r) * (2 * r + blockDim.x)].x - box[r + r * (2 * r + blockDim.x)].x)* (box[i + r + (j + r) * (2 * r + blockDim.x)].x - box[r + r * (2 * r + blockDim.x)].x)) + |
||||
exp(-0.5 / sigmaC / sigmaC * (box[i + r + (j + r) * (2 * r + blockDim.x)].x - box[r + r * (2 * r + blockDim.x)].y)* (box[i + r + (j + r) * (2 * r + blockDim.x)].x - box[r + r * (2 * r + blockDim.x)].y))+ |
||||
exp(-0.5 / sigmaC / sigmaC * (box[i + r + (j + r) * (2 * r + blockDim.x)].x - box[r + r * (2 * r + blockDim.x)].z)* (box[i + r + (j + r) * (2 * r + blockDim.x)].x - box[r + r * (2 * r + blockDim.x)].z)) |
||||
)* |
||||
exp(-0.5/sigmaD/sigmaD* |
||||
(i*i+j*j)); |
||||
coeff += tmp; |
||||
sum[0] += tmp * box[i * r + (j + r) * (2 * r + blockDim.x)].x; |
||||
sum[1] += tmp * box[i * r + (j + r) * (2 * r + blockDim.x)].y; |
||||
sum[2] += tmp * box[i * r + (j + r) * (2 * r + blockDim.x)].z; |
||||
} |
||||
|
||||
} |
||||
} |
||||
if (coeff > 0) { |
||||
dst[id].x = (uchar)(sum[0] / coeff); |
||||
dst[id].y = (uchar)(sum[1] / coeff); |
||||
dst[id].z = (uchar)(sum[2] / coeff); |
||||
} |
||||
else |
||||
{ |
||||
dst[id].x = 0; |
||||
dst[id].y = 0; |
||||
dst[id].z = 0; |
||||
} |
||||
|
||||
if (x == 0 && y == 0) |
||||
{ |
||||
printf("\n\n\n\n\n\n\n\n/////////////// on gpu /////////////////// \n\n grid dim = (%d, %d) blockdim = (%d,%d)\n\n\n\n\n", gridDim.x, gridDim.y, blockDim.x, blockDim.y); |
||||
} |
||||
|
||||
} |
||||
|
||||
|
||||
void mycu::bilateralFilter(cv::Mat src, cv::Mat dst, int r, double sigmaC, double sigmaS) |
||||
{ |
||||
// double sigmaCC = exp(-0.5 / sigmaC / sigmaC); |
||||
// double sigmaSS = exp(-0.5 / sigmaS / sigmaS); |
||||
|
||||
if(src.type()!=CV_8UC3) |
||||
{ |
||||
return; |
||||
} |
||||
int width = src.cols; |
||||
int height = src.rows; |
||||
|
||||
uchar3* srcGpu; |
||||
uchar3* dstGpu; |
||||
|
||||
cudaMalloc(&srcGpu, 3 * width * height); |
||||
cudaMalloc(&dstGpu, 3 * width * height); |
||||
|
||||
cudaMemcpy(srcGpu, src.data, 3 * width * height, cudaMemcpyHostToDevice); |
||||
|
||||
bilakernel << <dim3((width - 1) / threadX + 1, (height - 1) / threadY + 1), dim3(threadX, threadY) >> > (srcGpu, dstGpu, width, height, r, sigmaC, sigmaS); |
||||
cudaDeviceSynchronize(); |
||||
cudaMemcpy(dst.data, dstGpu, 3 * width * height, cudaMemcpyDeviceToHost); |
||||
|
||||
cudaFree(srcGpu); |
||||
cudaFree(dstGpu); |
||||
} |
||||
@ -0,0 +1,11 @@
@@ -0,0 +1,11 @@
|
||||
#pragma once |
||||
#ifndef __BILA__ |
||||
#define __BILA__ |
||||
#include <opencv2/opencv.hpp> |
||||
|
||||
namespace mycu |
||||
{ |
||||
void bilateralFilter(cv::Mat src, cv::Mat dst, int r, double sigmaC, double sigmaD); |
||||
} |
||||
|
||||
#endif |
||||
Loading…
Reference in new issue