Add OpenMP versions of axpby, bincount, dwconv, pointerchase, and sundials

src/bincount-omp/CMakeLists.txt

@@ -6,3 +6,15 @@ add_hecbench_benchmark(
     SOURCES main.cpp
     CATEGORIES algorithms
 )
+
+# Exact histogram validation requires the host reference and Intel GPU kernel
+# to use the same floating-point division semantics.
+if(CMAKE_CXX_COMPILER_ID STREQUAL "IntelLLVM")
+    target_compile_options(bincount-omp PRIVATE -fp-model=precise)
+    target_compile_definitions(bincount-omp PRIVATE
+        __STRICT_ANSI__
+    )
+    target_link_options(bincount-omp PRIVATE
+        "SHELL:-Xopenmp-target-backend -cl-fp32-correctly-rounded-divide-sqrt"
+    )
+endif()

src/bincount-omp/Makefile

@@ -43,7 +43,7 @@ ifeq ($(OPTIMIZE),yes)
 endif
 
 ifeq ($(DEVICE),gpu)
-  CFLAGS +=-fiopenmp -fopenmp-targets=spir64 -D__STRICT_ANSI__ -DHAVE_OMPX_DEVICE_INFO
+  CFLAGS +=-fiopenmp -fopenmp-targets=spir64 -D__STRICT_ANSI__
   LDFLAGS +=-Xopenmp-target-backend "-cl-fp32-correctly-rounded-divide-sqrt"
 else
   CFLAGS +=-qopenmp

src/dwconv-cuda/main.cu

@@ -227,7 +227,7 @@ void dwconv2d_forward (const int m,
   int errors = 0;
   for (int i = 0; i < output_size; i++) {
     const scalar_t tolerance = 1e-4f + 1e-4f * fabs(h_reference[i]);
-    if (fabs(h_output[i] - h_reference[i]) > tolerance)
+    if (!(fabs(h_output[i] - h_reference[i]) <= tolerance))
       errors++;
   }
   printf("%s\n", errors == 0 ? "PASS" : "FAIL");

src/dwconv-hip/main.cu

@@ -227,7 +227,7 @@ void dwconv2d_forward (const int m,
   int errors = 0;
   for (int i = 0; i < output_size; i++) {
     const scalar_t tolerance = 1e-4f + 1e-4f * fabs(h_reference[i]);
-    if (fabs(h_output[i] - h_reference[i]) > tolerance)
+    if (!(fabs(h_output[i] - h_reference[i]) <= tolerance))
       errors++;
   }
   printf("%s\n", errors == 0 ? "PASS" : "FAIL");

src/dwconv-omp/main.cpp

@@ -199,7 +199,7 @@ void dwconv2d_forward (const int m,
   int errors = 0;
   for (int i = 0; i < output_size; i++) {
     const scalar_t tolerance = 1e-4f + 1e-4f * fabs(h_reference[i]);
-    if (fabs(h_output[i] - h_reference[i]) > tolerance)
+    if (!(fabs(h_output[i] - h_reference[i]) <= tolerance))
       errors++;
   }
   printf("%s\n", errors == 0 ? "PASS" : "FAIL");

src/dwconv-sycl/main.cpp

@@ -246,7 +246,7 @@ void dwconv2d_forward (sycl::queue &q,
   int errors = 0;
   for (int i = 0; i < output_size; i++) {
     const scalar_t tolerance = 1e-4f + 1e-4f * fabs(h_reference[i]);
-    if (fabs(h_output[i] - h_reference[i]) > tolerance)
+    if (!(fabs(h_output[i] - h_reference[i]) <= tolerance))
       errors++;
   }
   printf("%s\n", errors == 0 ? "PASS" : "FAIL");