1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
|
//-*-c-*-
texture<float4, 1, cudaReadModeElementType> mesh;
__device__ bool intersect_triangle(const float3 &x, const float3 &p, const float3 &v0, const float3 &v1, const float3 &v2, float3 &intersection)
{
Matrix m = make_matrix(v1-v0, v2-v0, -p);
float determinant = det(m);
if (determinant == 0.0)
return false;
float3 b = x-v0;
float u1 = ((m.a11*m.a22 - m.a12*m.a21)*b.x + (m.a02*m.a21 - m.a01*m.a22)*b.y + (m.a01*m.a12 - m.a02*m.a11)*b.z)/determinant;
if (u1 < 0.0)
return false;
float u2 = ((m.a12*m.a20 - m.a10*m.a22)*b.x + (m.a00*m.a22 - m.a02*m.a20)*b.y + (m.a02*m.a10 - m.a00*m.a12)*b.z)/determinant;
if (u2 < 0.0)
return false;
float u3 = ((m.a10*m.a21 - m.a11*m.a20)*b.x + (m.a01*m.a20 - m.a00*m.a21)*b.y + (m.a00*m.a11 - m.a01*m.a10)*b.z)/determinant;
if (u3 < 0.0 || (1-u1-u2) < 0.0)
return false;
intersection = x + p*u3;
return true;
}
__device__ int get_color(const float3 &p, const float3 &v0, const float3& v1, const float3 &v2)
{
float3 normal = cross(v1-v0,v2-v0);
float scale;
scale = dot(normal,-p)/(norm(normal)*norm(p));
if (scale < 0.0)
scale = dot(-normal,-p)/(norm(normal)*norm(p));
int rgb = floorf(255*scale);
return rgb*65536 + rgb*256 + rgb;
}
__device__ float3 make_float3(const float4 &a)
{
return make_float3(a.x, a.y, a.z);
}
extern "C"
{
__global__ void translate(int max_idx, float3 *x, float3 v)
{
int idx = blockIdx.x*blockDim.x + threadIdx.x;
if (idx > max_idx)
return;
x[idx] += v;
}
__global__ void rotate(int max_idx, float3 *x, float phi, float3 axis)
{
int idx = blockIdx.x*blockDim.x + threadIdx.x;
if (idx > max_idx)
return;
x[idx] = rotate(x[idx], phi, axis);
}
__global__ void intersect_triangle_mesh(int max_idx, float3 *xarr, float3 *parr, int n, int *pixelarr)
{
int idx = blockIdx.x*blockDim.x + threadIdx.x;
if (idx > max_idx)
return;
float3 x = xarr[idx];
float3 p = parr[idx];
int *pixel = pixelarr+idx;
bool hit = false;
float distance, min_distance;
float3 intersection, min_intersection;
int i;
for (i=0; i < n; i++)
{
float3 v0 = make_float3(tex1Dfetch(mesh, 3*i));
float3 v1 = make_float3(tex1Dfetch(mesh, 3*i+1));
float3 v2 = make_float3(tex1Dfetch(mesh, 3*i+2));
if (intersect_triangle(x, p, v0, v1, v2, intersection))
{
if (!hit)
{
*pixel = get_color(p, v0, v1, v2);
min_distance = norm(intersection-x);
min_intersection = intersection;
hit = true;
continue;
}
distance = norm(intersection-x);
if (distance < min_distance)
{
*pixel = get_color(p, v0, v1, v2);
min_distance = distance;
min_intersection = intersection;
}
}
}
if (!hit)
*pixel = 0;
}
} // extern "c"
|
