Skip to content

Commit 5326dc6

Browse files
author
Dmitry Kolzov
committed
Merge remote-tracking branch 'origin/tile-support' into baikal-next
# Conflicts: # App/Core/renderer.h # App/Renderers/PT/ptrenderer.cpp # App/Renderers/PT/ptrenderer.h
2 parents 950c693 + 21d1c84 commit 5326dc6

File tree

15 files changed

+446
-843
lines changed

15 files changed

+446
-843
lines changed

‎App/CL/camera.cl

Lines changed: 71 additions & 58 deletions
Original file line numberDiff line numberDiff line change
@@ -35,48 +35,54 @@ void PerspectiveCamera_GeneratePaths(
3535
// Camera
3636
GLOBAL Camera const* restrict camera,
3737
// Image resolution
38-
int img_width,
39-
int img_height,
38+
int output_width,
39+
int output_height,
40+
// Pixel domain buffer
41+
GLOBAL int const* restrict pixel_idx,
42+
// Size of pixel domain buffer
43+
GLOBAL int const* restrict num_pixels,
4044
// RNG seed value
4145
uint rng_seed,
4246
// Current frame
4347
uint frame,
4448
// Rays to generate
45-
GLOBAL ray* rays,
49+
GLOBAL ray* restrict rays,
4650
// RNG data
47-
GLOBAL uint* random,
51+
GLOBAL uint* restrict random,
4852
GLOBAL uint const* restrict sobolmat,
4953
// Path buffer
50-
GLOBAL Path* paths
54+
GLOBAL Path* restrict paths
5155
)
5256
{
53-
int2 global_id;
54-
global_id.x = get_global_id(0);
55-
global_id.y = get_global_id(1);
57+
int global_id = get_global_id(0);
5658

5759
// Check borders
58-
if (global_id.x < img_width && global_id.y < img_height)
60+
if (global_id < *num_pixels)
5961
{
62+
int idx = pixel_idx[global_id];
63+
int y = idx / output_width;
64+
int x = idx % output_width;
65+
6066
// Get pointer to ray & path handles
61-
GLOBAL ray* my_ray = rays + global_id.y * img_width + global_id.x;
62-
GLOBAL Path* my_path = paths + global_id.y * img_width + global_id.x;
67+
GLOBAL ray* my_ray = rays + global_id;
68+
GLOBAL Path* my_path = paths + y * output_width + x;
6369

6470
// Initialize sampler
6571
Sampler sampler;
6672
#if SAMPLER == SOBOL
67-
uint scramble = random[global_id.x + img_width * global_id.y] * 0x1fe3434f;
73+
uint scramble = random[x + output_width * y] * 0x1fe3434f;
6874

6975
if (frame & 0xF)
7076
{
71-
random[global_id.x + img_width * global_id.y] = WangHash(scramble);
77+
random[x + output_width * y] = WangHash(scramble);
7278
}
7379

7480
Sampler_Init(&sampler, frame, SAMPLE_DIM_CAMERA_OFFSET, scramble);
7581
#elif SAMPLER == RANDOM
76-
uint scramble = global_id.x + img_width * global_id.y * rng_seed;
82+
uint scramble = x + output_width * y * rng_seed;
7783
Sampler_Init(&sampler, scramble);
7884
#elif SAMPLER == CMJ
79-
uint rnd = random[global_id.x + img_width * global_id.y];
85+
uint rnd = random[x + output_width * y];
8086
uint scramble = rnd * 0x1fe3434f * ((frame + 133 * rnd) / (CMJ_DIM * CMJ_DIM));
8187
Sampler_Init(&sampler, frame % (CMJ_DIM * CMJ_DIM), SAMPLE_DIM_CAMERA_OFFSET, scramble);
8288
#endif
@@ -86,8 +92,8 @@ void PerspectiveCamera_GeneratePaths(
8692

8793
// Calculate [0..1] image plane sample
8894
float2 img_sample;
89-
img_sample.x = (float)global_id.x / img_width + sample0.x / img_width;
90-
img_sample.y = (float)global_id.y / img_height + sample0.y / img_height;
95+
img_sample.x = (float)x / output_width + sample0.x / output_width;
96+
img_sample.y = (float)y / output_height + sample0.y / output_height;
9197

9298
// Transform into [-0.5, 0.5]
9399
float2 h_sample = img_sample - make_float2(0.5f, 0.5f);
@@ -118,62 +124,69 @@ void PerspectiveCamera_GeneratePaths(
118124
// Physical camera implemenation.
119125
// This kernel is being used if aperture > 0.
120126
KERNEL void PerspectiveCameraDof_GeneratePaths(
121-
// Camera
122-
GLOBAL Camera const* camera,
123-
// Image resolution
124-
int img_width,
125-
int img_height,
126-
// RNG seed value
127-
uint rng_seed,
128-
// Current frame
129-
uint frame,
130-
// Rays to generate
131-
GLOBAL ray* rays,
132-
// RNG data
133-
GLOBAL uint* random,
134-
GLOBAL uint const* restrict sobolmat,
135-
GLOBAL Path* paths
127+
// Camera
128+
GLOBAL Camera const* restrict camera,
129+
// Image resolution
130+
int output_width,
131+
int output_height,
132+
// Pixel domain buffer
133+
GLOBAL int const* restrict pixel_idx,
134+
// Size of pixel domain buffer
135+
GLOBAL int const* restrict num_pixels,
136+
// RNG seed value
137+
uint rng_seed,
138+
// Current frame
139+
uint frame,
140+
// Rays to generate
141+
GLOBAL ray* restrict rays,
142+
// RNG data
143+
GLOBAL uint* restrict random,
144+
GLOBAL uint const* restrict sobolmat,
145+
// Path buffer
146+
GLOBAL Path* restrict paths
136147
)
137148
{
138-
int2 global_id;
139-
global_id.x = get_global_id(0);
140-
global_id.y = get_global_id(1);
149+
int global_id = get_global_id(0);
141150

142-
// Check borders
143-
if (global_id.x < img_width && global_id.y < img_height)
144-
{
145-
// Get pointer to ray & path handles
146-
GLOBAL ray* my_ray = rays + global_id.y * img_width + global_id.x;
147-
GLOBAL Path* my_path = paths + global_id.y * img_width + global_id.x;
151+
// Check borders
152+
if (global_id < *num_pixels)
153+
{
154+
int idx = pixel_idx[global_id];
155+
int y = idx / output_width;
156+
int x = idx % output_width;
148157

149-
// Initialize sampler
150-
Sampler sampler;
158+
// Get pointer to ray & path handles
159+
GLOBAL ray* my_ray = rays + global_id;
160+
GLOBAL Path* my_path = paths + y * output_width + x;
161+
162+
// Initialize sampler
163+
Sampler sampler;
151164
#if SAMPLER == SOBOL
152-
uint scramble = random[global_id.x + img_width * global_id.y] * 0x1fe3434f;
165+
uint scramble = random[x + output_width * y] * 0x1fe3434f;
153166

154-
if (frame & 0xF)
155-
{
156-
random[global_id.x + img_width * global_id.y] = WangHash(scramble);
157-
}
167+
if (frame & 0xF)
168+
{
169+
random[x + output_width * y] = WangHash(scramble);
170+
}
158171

159-
Sampler_Init(&sampler, frame, SAMPLE_DIM_CAMERA_OFFSET, scramble);
172+
Sampler_Init(&sampler, frame, SAMPLE_DIM_CAMERA_OFFSET, scramble);
160173
#elif SAMPLER == RANDOM
161-
uint scramble = global_id.x + img_width * global_id.y * rng_seed;
162-
Sampler_Init(&sampler, scramble);
174+
uint scramble = x + output_width * y * rng_seed;
175+
Sampler_Init(&sampler, scramble);
163176
#elif SAMPLER == CMJ
164-
uint rnd = random[global_id.x + img_width * global_id.y];
165-
uint scramble = rnd * 0x1fe3434f * ((frame + 133 * rnd) / (CMJ_DIM * CMJ_DIM));
166-
Sampler_Init(&sampler, frame % (CMJ_DIM * CMJ_DIM), SAMPLE_DIM_CAMERA_OFFSET, scramble);
177+
uint rnd = random[x + output_width * y];
178+
uint scramble = rnd * 0x1fe3434f * ((frame + 133 * rnd) / (CMJ_DIM * CMJ_DIM));
179+
Sampler_Init(&sampler, frame % (CMJ_DIM * CMJ_DIM), SAMPLE_DIM_CAMERA_OFFSET, scramble);
167180
#endif
168181

169182
// Generate pixel and lens samples
170183
float2 sample0 = Sampler_Sample2D(&sampler, SAMPLER_ARGS);
171184
float2 sample1 = Sampler_Sample2D(&sampler, SAMPLER_ARGS);
172185

173-
// Calculate [0..1] image plane sample
174-
float2 img_sample;
175-
img_sample.x = (float)global_id.x / img_width + sample0.x / img_width;
176-
img_sample.y = (float)global_id.y / img_height + sample0.y / img_height;
186+
// Calculate [0..1] image plane sample
187+
float2 img_sample;
188+
img_sample.x = (float)x / output_width + sample0.x / output_width;
189+
img_sample.y = (float)y / output_height + sample0.y / output_height;
177190

178191
// Transform into [-0.5, 0.5]
179192
float2 h_sample = img_sample - make_float2(0.5f, 0.5f);

‎App/CL/integrator_pt.cl

Lines changed: 49 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -772,8 +772,10 @@ __kernel void FillAOVs(
772772
__global ray const* rays,
773773
// Intersection data
774774
__global Intersection const* isects,
775+
// Pixel indices
776+
__global int const* pixel_idx,
775777
// Number of pixels
776-
int num_items,
778+
__global int const* num_items,
777779
// Vertices
778780
__global float3 const* vertices,
779781
// Normals
@@ -847,9 +849,10 @@ __kernel void FillAOVs(
847849
};
848850

849851
// Only applied to active rays after compaction
850-
if (globalid < num_items)
852+
if (globalid < *num_items)
851853
{
852854
Intersection isect = isects[globalid];
855+
int idx = pixel_idx[globalid];
853856

854857
if (isect.shapeid > -1)
855858
{
@@ -875,8 +878,8 @@ __kernel void FillAOVs(
875878

876879
if (world_position_enabled)
877880
{
878-
aov_world_position[globalid].xyz += diffgeo.p;
879-
aov_world_position[globalid].w += 1.f;
881+
aov_world_position[idx].xyz += diffgeo.p;
882+
aov_world_position[idx].w += 1.f;
880883
}
881884

882885
if (world_shading_normal_enabled)
@@ -902,28 +905,28 @@ __kernel void FillAOVs(
902905
DifferentialGeometry_ApplyBumpNormalMap(&diffgeo, TEXTURE_ARGS);
903906
DifferentialGeometry_CalculateTangentTransforms(&diffgeo);
904907

905-
aov_world_shading_normal[globalid].xyz += diffgeo.n;
906-
aov_world_shading_normal[globalid].w += 1.f;
908+
aov_world_shading_normal[idx].xyz += diffgeo.n;
909+
aov_world_shading_normal[idx].w += 1.f;
907910
}
908911

909912
if (world_geometric_normal_enabled)
910913
{
911-
aov_world_geometric_normal[globalid].xyz += diffgeo.ng;
912-
aov_world_geometric_normal[globalid].w += 1.f;
914+
aov_world_geometric_normal[idx].xyz += diffgeo.ng;
915+
aov_world_geometric_normal[idx].w += 1.f;
913916
}
914917

915918
if (wireframe_enabled)
916919
{
917920
bool hit = (isect.uvwt.x < 1e-2) || (isect.uvwt.y < 1e-2) || (1.f - isect.uvwt.x - isect.uvwt.y < 1e-2);
918921
float3 value = hit ? make_float3(1.f, 1.f, 1.f) : make_float3(0.f, 0.f, 0.f);
919-
aov_wireframe[globalid].xyz += value;
920-
aov_wireframe[globalid].w += 1.f;
922+
aov_wireframe[idx].xyz += value;
923+
aov_wireframe[idx].w += 1.f;
921924
}
922925

923926
if (uv_enabled)
924927
{
925-
aov_uv[globalid].xyz += diffgeo.uv.xyy;
926-
aov_uv[globalid].w += 1.f;
928+
aov_uv[idx].xyz += diffgeo.uv.xyy;
929+
aov_uv[idx].w += 1.f;
927930
}
928931

929932
if (albedo_enabled)
@@ -936,9 +939,41 @@ __kernel void FillAOVs(
936939

937940
const float3 kd = Texture_GetValue3f(diffgeo.mat.kx.xyz, diffgeo.uv, TEXTURE_ARGS_IDX(diffgeo.mat.kxmapidx));
938941

939-
aov_albedo[globalid].xyz += kd;
940-
aov_albedo[globalid].w += 1.f;
942+
aov_albedo[idx].xyz += kd;
943+
aov_albedo[idx].w += 1.f;
941944
}
942945
}
943946
}
944947
}
948+
949+
__kernel void GenerateTileDomain(
950+
int output_width,
951+
int output_height,
952+
int tile_offset_x,
953+
int tile_offset_y,
954+
int tile_width,
955+
int tile_height,
956+
int subtile_width,
957+
int subtile_height,
958+
__global int* pixelidx0,
959+
__global int* pixelidx1,
960+
__global int* count
961+
)
962+
{
963+
int tile_x = get_global_id(0);
964+
int tile_y = get_global_id(1);
965+
int tile_start_idx = output_width * tile_offset_y + tile_offset_x;
966+
967+
if (tile_x < tile_width && tile_y < tile_height)
968+
{
969+
// TODO: implement subtile support
970+
int idx = tile_start_idx + tile_y * output_width + tile_x;
971+
pixelidx0[tile_y * tile_width + tile_x] = idx;
972+
pixelidx1[tile_y * tile_width + tile_x] = idx;
973+
}
974+
975+
if (tile_x == 0 && tile_y == 0)
976+
{
977+
*count = tile_width * tile_height;
978+
}
979+
}

‎App/CLW/clw_scene_controller.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -29,9 +29,9 @@ namespace Baikal
2929
// Create intersection API
3030
m_api = CreateFromOpenClContext(m_context, id, queue);
3131

32-
m_api->SetOption("acc.type", "bvh");
32+
m_api->SetOption("acc.type", "fatbvh");
3333
m_api->SetOption("bvh.builder", "sah");
34-
m_api->SetOption("bvh.sah.num_bins", 64.f);
34+
m_api->SetOption("bvh.sah.num_bins", 16.f);
3535
}
3636

3737
Material const* ClwSceneController::GetDefaultMaterial() const

‎App/Core/renderer.h

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -92,6 +92,13 @@ namespace Baikal
9292
*/
9393
virtual void Render(Scene1 const& scene) = 0;
9494

95+
/**
96+
\brief Render single iteration.
97+
98+
\param scene Scene to render
99+
*/
100+
virtual void RenderTile(Scene1 const& scene, RadeonRays::int2 const& tile_origin, RadeonRays::int2 const& tile_size) = 0;
101+
95102
/**
96103
\brief Set the output for rendering.
97104

0 commit comments

Comments
 (0)