-
Notifications
You must be signed in to change notification settings - Fork 0
/
main.cl
123 lines (102 loc) · 3.54 KB
/
main.cl
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
#include "rt_math.cl"
#include "bsdf.cl"
#include "scene.cl"
//float3 li(float2 ndc,
// ray_t ray,
// unsigned long *rng_state,
// __global const scene_t *scene,
// __global const shape_t *shapes) {
// float3 scale = (float3)(1.f, 1.f, 1.f), total_radiance = (float3)(0.f, 0.f, 0.f);
// float3 wo = (float3)(0.f, 0.f, 0.f);
// int depth = 4;
// for (int i = 0; i <= depth; i++) {
// wo = -ray.dir;
// intersection_t current_isect = scene_intersect(ray, scene, shapes);
// if (current_isect.t < 0.f) {
// break;
// }
// //float3 isect_p = ray.p + current_isect.t * ray.dir;
// total_radiance += scale * shapes[current_isect.shape_index].le;
// if (shapes[current_isect.shape_index].bsdf.bxdfs_count == 0) {
// break;
// }
// float3 wi;
// float pdf = 1.f;
// int flags;
// float2 samp = get_xi(rng_state);
// float3 bsdf_f = bsdf_sample_f(
// &shapes[current_isect.shape_index].bsdf,
// current_isect.normal,
// wo,
// &wi,
// samp,
// &pdf);
// if (pdf == 0.f) {
// break;
// }
// //wi = normalize((float3)(0,7.45,0) - current_isect.point)l
// scale *= bsdf_f;
// scale *= fabs(dot(wi, current_isect.normal));
// scale /= pdf;
// //ray = spawn_ray(current_isect.point, wi);
// ray.p = current_isect.point + wi * 0.00001;
// ray.dir = wi;
// }
// return total_radiance;
//}
__kernel void pathtrace(
__global const scene_t *scene,
__global const shape_t *shapes,
__global const float2 *ndcs,
__global float3 *colors,
__global unsigned long *rngs,
const unsigned int iterations) {
int id = get_global_id(0);
unsigned long rng_state = rngs[id];
rng(&rng_state);
// float2 ndc = ndcs[id];
// float3 accum_color = (float3)(0.f, 0.f, 0.f);
// for (unsigned int i = 0; i < iterations; i++) {
// ray_t ray = raycast(scene->camera, ndcs[id], &rng_state);
// accum_color += li(ndc, ray, &rng_state, scene, shapes);
// }
// colors[id] = accum_color / iterations;
int local_id = get_local_id(0);
__local scene_t local_scene;
__local shape_t local_shapes[10];
if (local_id == 0) {
local_scene = *scene;
for (int i = 0; i < scene->shape_count; i++) {
local_shapes[i] = shapes[i];
}
}
float2 ndc = ndcs[id];
barrier(CLK_LOCAL_MEM_FENCE);
float3 accum_color = (float3)(0.f, 0.f, 0.f);
for (unsigned int i = 0; i < iterations; i++) {
ray_t ray = raycast(scene->camera, ndc, &rng_state);
intersection_t isect = scene_intersect(ray, scene->shape_count, local_shapes);
if (isect.t >= 0) {
accum_color += (isect.normal + 1) * 0.5;
}
}
colors[id] = accum_color / iterations;
//colors[id] = ray.dir;
}
// //colors[id] = ray.dir;
// intersection_t isect = scene_intersect(ray, scene, shapes);
// if (isect.t >= 0) {
// colors[id] = (isect.normal + 1) * 0.5;
// } else {
// colors[id] = (float3)(0, 1, 1);
// }
//unsigned int iterations = *iterations_ptr;
//float2 ndc = ndcs[id];
//ray_t ray = raycast(scene->camera, ndcs[id], &rng_state);
//intersection_t isect = scene_intersect(ray, scene, shapes);
//if (isect.t >= 0) {
// colors[id] = (isect.normal + 1) * 0.5;
//} else {
// colors[id] = (float3)(0, 1, 1);
//}
//colors[id] = ray.dir;