Attachment 112474 Details for Bug 88587 – Test case

Test case

0001-Test-for-bug-reading-arrays-in-loop.patch (text/plain), 3.04 KB, created by Philip Taylor on 2015-01-19 14:39:06 UTC

(hide)

Description:

Filename:

MIME Type:

Creator: Philip Taylor

Created: 2015-01-19 14:39:06 UTC

Size: 3.04 KB

patch

obsolete

>From ddffbfb8bc5c2c388cbdf0782767d63b5787dc53 Mon Sep 17 00:00:00 2001
>From: Philip Taylor <pttaylo@amazon.co.uk>
>Date: Mon, 19 Jan 2015 13:55:51 +0000
>Subject: [PATCH] Test for bug reading arrays in loop.
>
>Expected output:
>
>  # 1240 1240
>  (...repeated 16 times)
>
>Actual output:
>
>  # 1240 1240
>  (...repeated 8 times)
>  # 560 1240
>      [FAILED]
>      Error: ((int64_t *)buf_data[2])[j] == sum
>---
> kernels/compiler_buggy_loop.cl |   14 ++++++++++++
> utests/CMakeLists.txt          |    1 +
> utests/compiler_buggy_loop.cpp |   47 ++++++++++++++++++++++++++++++++++++++++
> 3 files changed, 62 insertions(+)
> create mode 100644 kernels/compiler_buggy_loop.cl
> create mode 100644 utests/compiler_buggy_loop.cpp
>
>diff --git a/kernels/compiler_buggy_loop.cl b/kernels/compiler_buggy_loop.cl
>new file mode 100644
>index 0000000..f60a534
>--- /dev/null
>+++ b/kernels/compiler_buggy_loop.cl
>@@ -0,0 +1,14 @@
>+kernel void compiler_buggy_loop(global short *src0, global short *src1, global long *dst) {
>+  long sum = 0;
>+  int j = get_global_id(0);
>+  for (int i = 0; i < 16; i += 2) {
>+    short a0 = src0[i];
>+    short a1 = src1[i];
>+    short b0 = src0[i+1];
>+    short b1 = src1[i+1];
>+    sum += a0 * a1;
>+    sum += b0 * b1;
>+  }
>+  dst[j] = sum;
>+}
>+
>diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
>index 867c0cc..6173f39 100644
>--- a/utests/CMakeLists.txt
>+++ b/utests/CMakeLists.txt
>@@ -30,6 +30,7 @@ set (utests_basic_sources
> set (utests_binary_kernel_sources load_program_from_bin_file.cpp enqueue_built_in_kernels.cpp)
> 
> set (utests_sources
>+  compiler_buggy_loop.cpp
>   compiler_basic_arithmetic.cpp
>   compiler_displacement_map_element.cpp
>   compiler_mandelbrot.cpp
>diff --git a/utests/compiler_buggy_loop.cpp b/utests/compiler_buggy_loop.cpp
>new file mode 100644
>index 0000000..aae5b7a
>--- /dev/null
>+++ b/utests/compiler_buggy_loop.cpp
>@@ -0,0 +1,47 @@
>+#include <cstdint>
>+#include <cstring>
>+#include "utest_helper.hpp"
>+
>+void compiler_buggy_loop(void)
>+{
>+  const size_t s = 16;
>+  const size_t n = 16;
>+  int16_t src[n];
>+
>+  // Setup kernel and buffers
>+  OCL_CREATE_KERNEL("compiler_buggy_loop");
>+  OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(int16_t), NULL);
>+  OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(int16_t), NULL);
>+  OCL_CREATE_BUFFER(buf[2], 0, s * sizeof(int64_t), NULL);
>+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
>+  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
>+  OCL_SET_ARG(2, sizeof(cl_mem), &buf[2]);
>+  globals[0] = s;
>+  locals[0] = s;
>+
>+  for (size_t i = 0; i < n; ++i)
>+    src[i] = i;
>+
>+  OCL_MAP_BUFFER(0);
>+  OCL_MAP_BUFFER(1);
>+  memcpy(buf_data[0], src, sizeof(src));
>+  memcpy(buf_data[1], src, sizeof(src));
>+  OCL_UNMAP_BUFFER(0);
>+  OCL_UNMAP_BUFFER(1);
>+
>+  // Run the kernel on GPU
>+  OCL_NDRANGE(1);
>+
>+  // Compare
>+  OCL_MAP_BUFFER(2);
>+  for (size_t j = 0; j < s; ++j) {
>+    int64_t sum = 0;
>+    for (size_t i = 0; i < n; ++i)
>+      sum += i * i;
>+    printf("# %d %d\n", ((int64_t *)buf_data[2])[j], sum);
>+    OCL_ASSERT(((int64_t *)buf_data[2])[j] == sum);
>+  }
>+  OCL_UNMAP_BUFFER(2);
>+}
>+
>+MAKE_UTEST_FROM_FUNCTION(compiler_buggy_loop);
>-- 
>1.7.9.5
>

Actions: View

Attachments on bug 88587: 112474

Use of freedesktop.org services, including Bugzilla, is subject to our Code of Conduct. How we collect and use information is described in our Privacy Policy.