Skip to content

Commit 0839c18

Browse files
committed
aobench: define grid size in ao.h
1 parent 6521692 commit 0839c18

File tree

5 files changed

+6
-5
lines changed

5 files changed

+6
-5
lines changed

samples/aobench/ao.c

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -58,7 +58,7 @@ FUNCTION static void vnormalize(vec *c)
5858
FUNCTION static void
5959
ray_sphere_intersect(Isect *isect, const Ray *ray, const Sphere *sphere)
6060
{
61-
vec rs;
61+
vec rs = { 0 };
6262

6363
rs.x = ray->org.x - sphere->center.x;
6464
rs.y = ray->org.y - sphere->center.y;

samples/aobench/ao.cl

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -17,7 +17,7 @@ void debug_printf_i32_i32(const __constant char*, int, int) __asm__("__shady::pr
1717

1818
}
1919

20-
__attribute__((reqd_work_group_size(16, 16, 1)))
20+
__attribute__((reqd_work_group_size(BLOCK_SIZE, BLOCK_SIZE, 1)))
2121
kernel void aobench_kernel(global unsigned char* out) {
2222
int x = get_global_id(0);
2323
int y = get_global_id(1);

samples/aobench/ao.cu

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -10,7 +10,7 @@ __global__ void aobench_kernel(unsigned char* out) {
1010
long int ptr = (long int) out;
1111
Ctx ctx = get_init_context();
1212
init_scene(&ctx);
13-
// render_pixel(&ctx, x, y, WIDTH, HEIGHT, NSUBSAMPLES, out);
13+
render_pixel(&ctx, x, y, WIDTH, HEIGHT, NSUBSAMPLES, out);
1414
}
1515

1616
}

samples/aobench/ao.h

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2,6 +2,7 @@
22
#define HEIGHT 2048
33
#define NSUBSAMPLES 1
44
#define NAO_SAMPLES 8
5+
#define BLOCK_SIZE 16
56

67
typedef float Scalar;
78

samples/aobench/ao_main.c

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -119,11 +119,11 @@ void render_device(Args* args, unsigned char *img, int w, int h, int nsubsamples
119119
Program* program = load_program_from_disk(runtime, &args->compiler_config, path);
120120

121121
// run it twice to compile everything and benefit from caches
122-
wait_completion(launch_kernel(program, device, "aobench_kernel", WIDTH / 16, HEIGHT / 16, 1, 1, (void*[]) { &buf_addr }));
122+
wait_completion(launch_kernel(program, device, "aobench_kernel", WIDTH / BLOCK_SIZE, HEIGHT / BLOCK_SIZE, 1, 1, (void*[]) { &buf_addr }));
123123
struct timespec ts;
124124
timespec_get(&ts, TIME_UTC);
125125
uint64_t tsn = timespec_to_nano(ts);
126-
wait_completion(launch_kernel(program, device, "aobench_kernel", WIDTH / 16, HEIGHT / 16, 1, 1, (void*[]) { &buf_addr }));
126+
wait_completion(launch_kernel(program, device, "aobench_kernel", WIDTH / BLOCK_SIZE, HEIGHT / BLOCK_SIZE, 1, 1, (void*[]) { &buf_addr }));
127127
struct timespec tp;
128128
timespec_get(&tp, TIME_UTC);
129129
uint64_t tpn = timespec_to_nano(tp);

0 commit comments

Comments
 (0)