Skip to content

Commit 5f750e1

Browse files
author
Konstantin Zverev
committed
Fix CornellBoxShadow tutorial.
1 parent 2a8765b commit 5f750e1

File tree

2 files changed

+24
-117
lines changed

2 files changed

+24
-117
lines changed

‎Tutorials/CornellBoxShadow/kernel.cl

Lines changed: 15 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -22,6 +22,8 @@
2222
#ifndef KERNEL_CL
2323
#define KERNEL_CL
2424

25+
#define EPSILON 0.001f
26+
2527
typedef struct _Ray
2628
{
2729
/// xyz - origin, w - max range
@@ -78,7 +80,6 @@ float4 ConvertFromBarycentric(__global const float* vec,
7880
return a * (1 - uvwt->x - uvwt->y) + b * uvwt->x + c * uvwt->y;
7981
}
8082

81-
8283
__kernel void GeneratePerspectiveRays(__global Ray* rays,
8384
__global const Camera* cam,
8485
int width,
@@ -124,7 +125,7 @@ __kernel void GenerateShadowRays(__global Ray* rays,
124125
int width,
125126
int height)
126127
{
127-
int2 globalid;
128+
int2 globalid;
128129
globalid.x = get_global_id(0);
129130
globalid.y = get_global_id(1);
130131

@@ -141,12 +142,15 @@ __kernel void GenerateShadowRays(__global Ray* rays,
141142
return;
142143
}
143144

144-
// Calculate position of the intersection point
145+
// Calculate position and normal of the intersection point
145146
int ind = indents[shape_id];
146-
float4 pos = ConvertFromBarycentric(positions + ind*3, ids + ind*3, prim_id, &isect[k].uvwt);
147+
float4 pos = ConvertFromBarycentric(positions + ind*3, ids + ind, prim_id, &isect[k].uvwt);
148+
float4 norm = ConvertFromBarycentric(normals + ind*3, ids + ind, prim_id, &isect[k].uvwt);
149+
norm = normalize(norm);
150+
147151
float4 dir = light - pos;
148152
rays[k].d = normalize(dir);
149-
rays[k].o = pos + rays[k].d * FLT_EPSILON;
153+
rays[k].o = pos + norm * EPSILON;
150154
rays[k].o.w = length(dir);
151155

152156
rays[k].extra.x = 0xFFFFFFFF;
@@ -181,23 +185,21 @@ __kernel void Shading(//scene
181185
int shape_id = isect[k].shapeid;
182186
int prim_id = isect[k].primid;
183187

184-
if (shape_id != -1 && prim_id != -1/* && occl[k] != -1*/)
188+
if (shape_id != -1 && prim_id != -1 && occl[k] == -1)
185189
{
186190
// Calculate position and normal of the intersection point
187191
int ind = indents[shape_id];
188-
float4 pos = ConvertFromBarycentric(positions + ind*3, ids + ind*3, prim_id, &isect[k].uvwt);
189-
float4 norm = ConvertFromBarycentric(normals + ind*3, ids + ind*3, prim_id, &isect[k].uvwt);
192+
193+
float4 pos = ConvertFromBarycentric(positions + ind*3, ids + ind, prim_id, &isect[k].uvwt);
194+
float4 norm = ConvertFromBarycentric(normals + ind*3, ids + ind, prim_id, &isect[k].uvwt);
190195
norm = normalize(norm);
191196

192-
// float4 diff_col = (float4)( 0.5f,
193-
// 0.5f,
194-
// 0.5f, 1.f);
195-
int color_id = (ind + prim_id)*3;
197+
//triangle diffuse color
198+
int color_id = ind + prim_id*3;
196199
float4 diff_col = (float4)( colors[color_id],
197200
colors[color_id + 1],
198201
colors[color_id + 2], 1.f);
199202

200-
201203
// Calculate lighting
202204
float4 col = (float4)( 0.f, 0.f, 0.f, 0.f );
203205
float4 light_dir = normalize(light - pos);

‎Tutorials/CornellBoxShadow/main.cpp

Lines changed: 9 additions & 104 deletions
Original file line numberDiff line numberDiff line change
@@ -100,7 +100,7 @@ void InitData()
100100
colors.push_back(mat.diffuse[2]);
101101
}
102102

103-
//fill empty space to simplify indentation in arrays
103+
// add additional emty data to simplify indentation in arrays
104104
if (mesh.positions.size() / 3 < mesh.indices.size())
105105
{
106106
int count = mesh.indices.size() - mesh.positions.size() / 3;
@@ -113,7 +113,7 @@ void InitData()
113113
}
114114

115115
indents.push_back(indent);
116-
indent += mesh.indices.size() / 3;
116+
indent += mesh.indices.size();
117117
}
118118
g_positions = CLWBuffer<float>::Create(g_context, CL_MEM_READ_ONLY, verts.size(), verts.data());
119119
g_normals = CLWBuffer<float>::Create(g_context, CL_MEM_READ_ONLY, normals.size(), normals.data());
@@ -122,22 +122,6 @@ void InitData()
122122
g_indent = CLWBuffer<int>::Create(g_context, CL_MEM_READ_ONLY, indents.size(), indents.data());
123123
}
124124

125-
float3 ConvertFromBarycentric(const float* vec, const int* ind, int prim_id, const float4& uvwt)
126-
{
127-
float3 a = { vec[ind[prim_id * 3] * 3],
128-
vec[ind[prim_id * 3] * 3 + 1],
129-
vec[ind[prim_id * 3] * 3 + 2], 0.f};
130-
131-
float3 b = { vec[ind[prim_id * 3 + 1] * 3],
132-
vec[ind[prim_id * 3 + 1] * 3 + 1],
133-
vec[ind[prim_id * 3 + 1] * 3 + 2], 0.f };
134-
135-
float3 c = { vec[ind[prim_id * 3 + 2] * 3],
136-
vec[ind[prim_id * 3 + 2] * 3 + 1],
137-
vec[ind[prim_id * 3 + 2] * 3 + 2], 0.f };
138-
return a * (1 - uvwt.x - uvwt.y) + b * uvwt.x + c * uvwt.y;
139-
}
140-
141125
Buffer* GeneratePrimaryRays()
142126
{
143127
//prepare camera buf
@@ -228,7 +212,6 @@ Buffer* Shading(const CLWBuffer<Intersection> &isect, const CLWBuffer<int> &occl
228212
return CreateFromOpenClBuffer(g_api, out_buff);
229213
}
230214

231-
232215
void DrawScene()
233216
{
234217

@@ -397,112 +380,34 @@ int main(int argc, char* argv[])
397380
// Prepare rays. One for each texture pixel.
398381
Buffer* ray_buffer = GeneratePrimaryRays();
399382
// Intersection data
400-
std::vector<Intersection> isect(k_raypack_size);
401383
CLWBuffer<Intersection> isect_buffer_cl = CLWBuffer<Intersection>::Create(g_context, CL_MEM_READ_WRITE, g_window_width*g_window_height);
402384
Buffer* isect_buffer = CreateFromOpenClBuffer(g_api, isect_buffer_cl);
403385

404386
// Intersection
405387
g_api->QueryIntersection(ray_buffer, k_raypack_size, isect_buffer, nullptr, nullptr);
406388

407-
// Get results
408-
Event* e = nullptr;
409-
Intersection* tmp = nullptr;
410-
g_api->MapBuffer(isect_buffer, kMapRead, 0, isect.size() * sizeof(Intersection), (void**)&tmp, &e);
411-
// RadeonRays calls are asynchronous, so need to wait for calculation to complete.
412-
e->Wait();
413-
g_api->DeleteEvent(e);
414-
e = nullptr;
415-
416-
// Copy results
417-
for (int i = 0; i < k_raypack_size; ++i)
418-
{
419-
isect[i] = tmp[i];
420-
}
421-
g_api->UnmapBuffer(isect_buffer, tmp, nullptr);
422-
tmp = nullptr;
423-
424389
// Point light position
425390
float3 light = { -0.01f, 1.85f, 0.1f };
426-
Buffer* shadow_rays_buffer = GenerateShadowRays(isect_buffer_cl, light);
391+
427392
// Shadow rays
428-
//std::vector<ray> shadow_rays(k_raypack_size);
429-
//for (int i = 0; i < g_window_height; ++i)
430-
// for (int j = 0; j < g_window_width; ++j)
431-
// {
432-
// int k = i * g_window_width + j;
433-
// int shape_id = isect[k].shapeid;
434-
// int prim_id = isect[k].primid;
435-
//
436-
// // Need shadow rays only for intersections
437-
// if (shape_id == kNullId || prim_id == kNullId)
438-
// {
439-
// shadow_rays[k].SetActive(false);
440-
// continue;
441-
// }
442-
// mesh_t& mesh = g_objshapes[shape_id].mesh;
443-
// float3 pos = ConvertFromBarycentric(mesh.positions.data(), mesh.indices.data(), prim_id, isect[k].uvwt);
444-
// float3 dir = light - pos;
445-
// shadow_rays[k].d = dir;
446-
// shadow_rays[k].d.normalize();
447-
// shadow_rays[k].o = pos + shadow_rays[k].d * std::numeric_limits<float>::epsilon();
448-
// shadow_rays[k].SetMaxT(std::sqrt(dir.sqnorm()));
449-
450-
// }
451-
//Buffer* shadow_rays_buffer = g_api->CreateBuffer(shadow_rays.size() * sizeof(ray), shadow_rays.data());
393+
Buffer* shadow_rays_buffer = GenerateShadowRays(isect_buffer_cl, light);
452394
CLWBuffer<int> occl_buffer_cl = CLWBuffer<int>::Create(g_context, CL_MEM_READ_WRITE, g_window_width*g_window_height);
453395
Buffer* occl_buffer = CreateFromOpenClBuffer(g_api, occl_buffer_cl);
454396

455397
// Occlusion
456398
g_api->QueryOcclusion(shadow_rays_buffer, k_raypack_size, occl_buffer, nullptr, nullptr);
457399

400+
// Shading
458401
Buffer* tex_buf = Shading(isect_buffer_cl, occl_buffer_cl, light);
459-
// Draw
402+
403+
// Get image data
460404
std::vector<unsigned char> tex_data(k_raypack_size * 4);
461-
//for (int i = 0; i < k_raypack_size ; ++i)
462-
//{
463-
// int shape_id = isect[i].shapeid;
464-
// int prim_id = isect[i].primid;
465-
466-
// if (shape_id != kNullId && prim_id != kNullId)
467-
// {
468-
// mesh_t& mesh = g_objshapes[shape_id].mesh;
469-
// int mat_id = mesh.material_ids[prim_id];
470-
// material_t& mat = g_objmaterials[mat_id];
471-
//
472-
// // theck there is no shapes between intersection point and light source
473-
// if (tmp_occl[i] != kNullId)
474-
// {
475-
// continue;
476-
// }
477-
478-
// float3 diff_col = { mat.diffuse[0],
479-
// mat.diffuse[1],
480-
// mat.diffuse[2] };
481-
482-
// // Calculate position and normal of the intersection point
483-
// float3 pos = ConvertFromBarycentric(mesh.positions.data(), mesh.indices.data(), prim_id, isect[i].uvwt);
484-
// float3 norm = ConvertFromBarycentric(mesh.normals.data(), mesh.indices.data(), prim_id, isect[i].uvwt);
485-
// norm.normalize();
486-
487-
// // Calculate lighting
488-
// float3 col = { 0.f, 0.f, 0.f };
489-
// float3 light_dir = light - pos;
490-
// light_dir.normalize();
491-
// float dot_prod = dot(norm, light_dir);
492-
// if (dot_prod > 0)
493-
// col += dot_prod * diff_col;
494-
495-
// tex_data[i * 4] = col[0] * 255;
496-
// tex_data[i * 4 + 1] = col[1] * 255;
497-
// tex_data[i * 4 + 2] = col[2] * 255;
498-
// tex_data[i * 4 + 3] = 255;
499-
// }
500-
//}
501-
502405
unsigned char* pixels = nullptr;
406+
Event* e = nullptr;
503407
g_api->MapBuffer(tex_buf, kMapRead, 0, 4 * k_raypack_size * sizeof(unsigned char), (void**)&pixels, &e);
504408
e->Wait();
505409
memcpy(tex_data.data(), pixels, 4 * k_raypack_size * sizeof(unsigned char));
410+
506411
// Update texture data
507412
glBindTexture(GL_TEXTURE_2D, g_texture);
508413
glTexSubImage2D(GL_TEXTURE_2D, 0, 0, 0, g_window_width, g_window_height, GL_RGBA, GL_UNSIGNED_BYTE, tex_data.data());

0 commit comments

Comments
 (0)