diff options
| author | Zhigang Gong <zhigang.gong@intel.com> | 2014-08-27 10:33:42 +0800 |
|---|---|---|
| committer | Zhigang Gong <zhigang.gong@intel.com> | 2014-09-03 12:29:01 +0800 |
| commit | 44929d4847b7b5d0cf5929a3155b6193091c26b7 (patch) | |
| tree | 66f660a0bfaf670889e3ffe65ff1e8aa3e7ac426 /kernels | |
| parent | bfa6135c9e7a53d2aab21c8c3257e85eb61b1212 (diff) | |
Add new vload benchmark/test case.
v2:
refine the benchmark case and don't mix it with normal
unit test cases.
Signed-off-by: Zhigang Gong <zhigang.gong@intel.com>
Reviewed-by: "Song, Ruiling" <ruiling.song@intel.com>
Diffstat (limited to 'kernels')
| -rw-r--r-- | kernels/vload_bench.cl | 33 |
1 files changed, 33 insertions, 0 deletions
diff --git a/kernels/vload_bench.cl b/kernels/vload_bench.cl new file mode 100644 index 00000000..c906c752 --- /dev/null +++ b/kernels/vload_bench.cl @@ -0,0 +1,33 @@ +#define VLOAD_BENCH(T, N, M) \ +__kernel void \ +vload_bench_##M ##T ##N(__global T* src, __global uint* dst, uint offset) \ +{ \ + int id = (int)get_global_id(0); \ + uint ##N srcV = 0; \ + for(int i = 0; i < M; i++) \ + { \ + srcV += convert_uint ##N(vload ##N(id + (i & 0xFFFF), src + offset)); \ + } \ + vstore ##N(srcV, id, dst);\ + /*if (id < 16)*/ \ + /*printf("id %d %d %d\n", id, srcV.s0, srcV.s1);*/ \ +} + +#define VLOAD_BENCH_ALL_VECTOR(T, N_ITERATIONS) \ + VLOAD_BENCH(T, 2, N_ITERATIONS) \ + VLOAD_BENCH(T, 3, N_ITERATIONS) \ + VLOAD_BENCH(T, 4, N_ITERATIONS) \ + VLOAD_BENCH(T, 8, N_ITERATIONS) \ + VLOAD_BENCH(T, 16, N_ITERATIONS) + +#define VLOAD_BENCH_ALL_TYPES(N_ITERATIONS) \ + VLOAD_BENCH_ALL_VECTOR(uchar, N_ITERATIONS) \ + VLOAD_BENCH_ALL_VECTOR(char, N_ITERATIONS) \ + VLOAD_BENCH_ALL_VECTOR(ushort, N_ITERATIONS) \ + VLOAD_BENCH_ALL_VECTOR(short, N_ITERATIONS) \ + VLOAD_BENCH_ALL_VECTOR(uint, N_ITERATIONS) \ + VLOAD_BENCH_ALL_VECTOR(int, N_ITERATIONS) \ + VLOAD_BENCH_ALL_VECTOR(float, N_ITERATIONS) + +VLOAD_BENCH_ALL_TYPES(1) +VLOAD_BENCH_ALL_TYPES(10000) |