Skip to content

Commit 7efbe03

Browse files
author
ptx
committed
First import
0 parents  commit 7efbe03

File tree

7 files changed

+252
-0
lines changed

7 files changed

+252
-0
lines changed

‎CMakeLists.txt

+21
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,21 @@
1+
2+
SET( EXTENSION_NAME "Extensions_LevelSetCUDA")
3+
4+
CUDA_ADD_LIBRARY("CUDA_KERNELS"
5+
LevelSet/LevelSet.cu
6+
# Renderers/CUDA/RayCaster.cu
7+
# Utils/CUDA/DoseCalc.cu
8+
# Utils/CUDA/DozeSetup.cu
9+
# # Utils/CUDA/Superposition.cu
10+
)
11+
12+
13+
# Create the extension library
14+
ADD_LIBRARY(${EXTENSION_NAME}
15+
LevelSet/CUDAStrategy.cpp
16+
LevelSet/CUDAStrategy.h
17+
)
18+
19+
TARGET_LINK_LIBRARIES( ${EXTENSION_NAME}
20+
"CUDA_KERNELS"
21+
)

‎LevelSet/CUDAStrategy.cpp

+36
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,36 @@
1+
//
2+
// -------------------------------------------------------------------
3+
// Copyright (C) 2007 OpenEngine.dk (See AUTHORS)
4+
//
5+
// This program is free software; It is covered by the GNU General
6+
// Public License version 2 or any later version.
7+
// See the GNU General Public License for more details (see LICENSE).
8+
//--------------------------------------------------------------------
9+
10+
#include "CUDAStrategy.h"
11+
#include <Meta/CUDA.h>
12+
#include "LevelSet.h"
13+
#include <Resources/Tex.h>
14+
#include <LevelSet/SDF.h>
15+
#include <Logging/Logger.h>
16+
17+
namespace OpenEngine {
18+
namespace LevelSet {
19+
20+
CUDAStrategy::CUDAStrategy() {
21+
INITIALIZE_CUDA();
22+
logger.info << PRINT_CUDA_DEVICE_INFO() << logger.end;
23+
//cu_Init();
24+
}
25+
26+
void CUDAStrategy::Reinitialize(SDF* sdf, unsigned int iterations) {
27+
28+
// copy phi to cuda...
29+
Tex<float> phi = sdf->GetPhi();
30+
cu_Reinit(phi.GetData(),phi.GetWidth(),phi.GetHeight(),iterations);
31+
sdf->SetPhi(phi);
32+
33+
}
34+
35+
} // NS LevelSet
36+
} // NS OpenEngine

‎LevelSet/CUDAStrategy.h

+27
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,27 @@
1+
//
2+
// -------------------------------------------------------------------
3+
// Copyright (C) 2007 OpenEngine.dk (See AUTHORS)
4+
//
5+
// This program is free software; It is covered by the GNU General
6+
// Public License version 2 or any later version.
7+
// See the GNU General Public License for more details (see LICENSE).
8+
//--------------------------------------------------------------------
9+
10+
#include <LevelSet/Strategy.h>
11+
12+
namespace OpenEngine {
13+
namespace LevelSet {
14+
15+
/**
16+
* Short description.
17+
*
18+
* @class CUDAStrategy CUDAStrategy.h ons/LevelSetCUDA/LevelSet/CUDAStrategy.h
19+
*/
20+
class CUDAStrategy : public Strategy {
21+
public:
22+
CUDAStrategy();
23+
void Reinitialize(SDF* sdf, unsigned int iterations);
24+
};
25+
26+
} // NS LevelSet
27+
} // NS OpenEngine

‎LevelSet/LevelSet.cu

+137
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,137 @@
1+
// hello world
2+
3+
#include <Meta/CUDA.h>
4+
5+
#define GetPhi(phi,x,y,w) phi[x+w*(y)]
6+
7+
void cu_Init() {
8+
9+
10+
}
11+
12+
__global__ void reinit(float *phi,float* phi0, float* phin,
13+
unsigned int width, unsigned int height) {
14+
uint x = __umul24(blockIdx.x, blockDim.x) + threadIdx.x;
15+
uint y = __umul24(blockIdx.y, blockDim.y) + threadIdx.y;
16+
17+
if (x > width || y > height)
18+
return;
19+
20+
float xy = GetPhi(phi,x,y,width);
21+
22+
float phiXPlus = 0.0f;
23+
float phiXMinus = 0.0f;
24+
float phiYPlus = 0.0f;
25+
float phiYMinus = 0.0f;
26+
if (x != width-1) phiXPlus = (GetPhi(phi,x+1, y,width) - xy);
27+
if (x != 0) phiXMinus = (xy - GetPhi(phi,x-1, y,width));
28+
29+
if (y !=height-1) phiYPlus = (GetPhi(phi,x, y+1,width) - xy);
30+
if (y != 0) phiYMinus = (xy - GetPhi(phi,x, y-1,width));
31+
32+
/* GetPhi(phin,x,y,width) = phiYPlus; */
33+
/* return; */
34+
35+
36+
float dXSquared = 0;
37+
float dYSquared = 0;
38+
float a = GetPhi(phi0,x,y,width);
39+
if (a > 0) {
40+
// formula 6.3 page 58
41+
float _max = max(phiXMinus, 0.0f);
42+
float _min = min(phiXPlus, 0.0f);
43+
dXSquared = max(_max*_max, _min*_min);
44+
45+
_max = max(phiYMinus, 0.0f);
46+
_min = min(phiYPlus, 0.0f);
47+
dYSquared = max(_max*_max, _min*_min);
48+
} else {
49+
// formula 6.4 page 58
50+
float _max = max(phiXPlus, 0.0f);
51+
float _min = min(phiXMinus, 0.0f);
52+
dXSquared = max(_max*_max, _min*_min);
53+
54+
_max = max(phiYPlus, 0.0f);
55+
_min = min(phiYMinus, 0.0f);
56+
dYSquared = max(_max*_max, _min*_min);
57+
}
58+
59+
float normSquared = dXSquared + dYSquared;
60+
float norm = sqrt(normSquared);
61+
62+
// Using the S(phi) sign formula 7.6 on page 67
63+
//float sign = phi(x,y) / sqrt(phi(x,y)*phi(x,y) + normSquared);
64+
float sign = GetPhi(phi0,x,y,width) /
65+
sqrt(GetPhi(phi0,x,y,width)*GetPhi(phi0,x,y,width) + 1);
66+
float t = 0.3; // A stabil CFL condition
67+
GetPhi(phin,x,y,width) = GetPhi(phi,x,y,width) - sign*(norm - 1)*t;
68+
69+
70+
}
71+
72+
void cu_Reinit(float* data,
73+
unsigned int w,
74+
unsigned int h,
75+
unsigned int iterations) {
76+
float* phiData;
77+
float* phi0Data;
78+
float* phinData;
79+
/* int phiPitch; */
80+
/* int phi0Pitch; */
81+
/* int phinPitch; */
82+
83+
/* cudaArray* phiData; */
84+
/* cudaArray* phi0Data; */
85+
/* cudaArray* phinData; */
86+
87+
/* cudaChannelFormatDesc channelDesc = */
88+
/* cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat); */
89+
90+
91+
/* cudaMallocArray(&phiData, &channelDesc, w, h); */
92+
/* cudaMallocArray(&phi0Data, &channelDesc, w, h); */
93+
/* cudaMallocArray(&phinData, &channelDesc, w, h); */
94+
95+
/* cudaMemcpyToArray(phiData, 0, 0, data, sizeof(float)*w*h, cudaMemcpyHostToDevice); */
96+
/* cudaMemcpyToArray(phi0Data, 0, 0, data, sizeof(float)*w*h, cudaMemcpyHostToDevice); */
97+
/* cudaMemcpyToArray(phinData, 0, 0, data, sizeof(float)*w*h, cudaMemcpyHostToDevice); */
98+
99+
100+
cudaMalloc((void**)&phiData, sizeof(float)*w*h);
101+
cudaMalloc((void**)&phi0Data, sizeof(float)*w*h);
102+
cudaMalloc((void**)&phinData, sizeof(float)*w*h);
103+
cudaMemcpy((void*)phiData,(void*)data, sizeof(float)*w*h,cudaMemcpyHostToDevice);
104+
cudaMemcpy((void*)phi0Data,(void*)data, sizeof(float)*w*h,cudaMemcpyHostToDevice);
105+
cudaMemcpy((void*)phinData,(void*)data, sizeof(float)*w*h,cudaMemcpyHostToDevice);
106+
107+
/* cudaMallocPitch((void**)&phiData, &phiPitch, sizeof(float)*w,h); */
108+
/* cudaMallocPitch((void**)&phi0Data, &phi0Pitch, sizeof(float)*w,h); */
109+
/* cudaMallocPitch((void**)&phinData, &phinPitch, sizeof(float)*w,h); */
110+
/* cudaMemcpy((void*)phiData,(void*)data, sizeof(float)*w*h,cudaMemcpyHostToDevice); */
111+
/* cudaMemcpy((void*)phi0Data,(void*)data, sizeof(float)*w*h,cudaMemcpyHostToDevice); */
112+
/* cudaMemcpy((void*)phinData,(void*)data, sizeof(float)*w*h,cudaMemcpyHostToDevice); */
113+
114+
115+
CHECK_FOR_CUDA_ERROR();
116+
117+
const dim3 blockSize(32,16,1);
118+
const dim3 gridSize(w/blockSize.x, h/blockSize.y);
119+
120+
//printf("%i,%i\n",w,h);
121+
122+
//iterations=1;
123+
for (unsigned int i=0;i<iterations;i++) {
124+
reinit<<<gridSize,blockSize>>>(phiData,phi0Data,phinData,w,h);
125+
cudaMemcpy((void*)phiData,(void*)phinData,sizeof(float)*w*h,cudaMemcpyDeviceToDevice);
126+
cudaThreadSynchronize();
127+
CHECK_FOR_CUDA_ERROR();
128+
}
129+
130+
cudaMemcpy((void*)data,(void*)phiData, sizeof(float)*w*h,cudaMemcpyDeviceToHost);
131+
CHECK_FOR_CUDA_ERROR();
132+
cudaFree(phiData);
133+
cudaFree(phi0Data);
134+
cudaFree(phinData);
135+
136+
137+
}

‎LevelSet/LevelSet.h

+6
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,6 @@
1+
#ifndef _CU_LEVELSET_H_
2+
#define _CU_LEVELSET_H_
3+
4+
void cu_Reinit(float* data, unsigned int w, unsigned int h, unsigned int iterations);
5+
6+
#endif

‎LevelSetCUDA.dist

+13
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,13 @@
1+
# Hello World
2+
3+
dist /extensions/LevelSet http://openengine.dk/code/extensions/LevelSet/LevelSet.dist
4+
5+
darcs /extensions/LevelSetCUDA http://openengine.dk/code/extensions/LevelSetCUDA
6+
darcs /extensions/CUDA http://openengine.dk/code/extensions/CUDA
7+
8+
9+
darcs-dev /extensions/LevelSet fh.daimi.au.dk:/users/cgd/code/extensions/LevelSetCUDA
10+
darcs-dev /extensions/CUDA fh.daimi.au.dk:/users/cgd/code/extensions/CUDA
11+
12+
13+

‎README.txt

+12
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,12 @@
1+
Author: OpenEngine Team
2+
3+
Homepage: http://www.openengine.dk/wiki/Extensions/LevelSetCUDA
4+
5+
Get the latest version with:
6+
darcs get http://daimi.au.dk/~cgd/extensions/LevelSetCUDA
7+
8+
This is an example extension demonstrating how extensions work.
9+
For more infomation on how to create your own extensions checkout
10+
http://www.openengine.dk/wiki/CreatingExtension
11+
12+

0 commit comments

Comments
 (0)