Minor cmake adjustments

alpers-git · Dec 12, 2024 · 7563270 · 7563270
1 parent 1d4a2a8
commit 7563270
Show file tree

Hide file tree

Showing 5 changed files with 20 additions and 10 deletions.
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -20,6 +20,9 @@ if(ZFP_WITH_SYCL)
   set (CMAKE_C_FLAGS_RELEASE "${CMAKE_C_FLAGS_RELEASE} -O3")
 
   add_definitions(-DZFP_WITH_SYCL_DEVICE=${ZFP_WITH_SYCL_DEVICE})
+  if(NOT("${ZFP_WITH_SYCL_DEVICE}" STREQUAL "PVC" OR "${ZFP_WITH_SYCL_DEVICE}" STREQUAL "pvc"))
+    add_definitions(-DDISABLE_BLOCK_8)
+  endif()
 endif()
 
 # Enable MACOSX_RPATH by default

diff --git a/src/sycl/constants.h b/src/sycl/constants.h
@@ -9,14 +9,14 @@ namespace internal {
 
 using namespace ::sycl::ext::oneapi::experimental;
 
-// inline constexpr unsigned char perm_1[]={0, 1, 2, 3};
-inline static constexpr device_global<unsigned char[4]> perm_1{0, 1, 2, 3};
+inline constexpr unsigned char perm_1[]={0, 1, 2, 3};
+// inline static constexpr device_global<unsigned char[4]> perm_1{0, 1, 2, 3};
 
 #define index(i, j) ((i) + 4 * (j))
 
 // order coefficients (i, j) by i + j, then i^2 + j^2
-// inline constexpr unsigned char perm_2[]=
-inline static constexpr device_global<unsigned char[16]>perm_2
+inline constexpr unsigned char perm_2[]=
+// inline static constexpr device_global<unsigned char[16]>perm_2
 {
                                    index(0, 0), //  0 : 0
 
@@ -49,8 +49,8 @@ inline static constexpr device_global<unsigned char[16]>perm_2
 #define index(x, y, z) ((x) + 4 * ((y) + 4 * (z)))
 
 // order coefficients (i, j, k) by i + j + k, then i^2 + j^2 + k^2
-// inline constexpr unsigned char perm_3[]=
-inline static constexpr device_global<unsigned char[64]> perm_3
+inline constexpr unsigned char perm_3[]=
+// inline static constexpr device_global<unsigned char[64]> perm_3
 {
                                    index(0, 0, 0), //  0 : 0
 

diff --git a/src/sycl/decode.h b/src/sycl/decode.h
@@ -1059,9 +1059,11 @@ decode(
       break;
     case 2:
       switch (params->min_sub_group_size) {
+#ifdef DISABLE_BLOCK_8
         case 8:
           bits_read = internal::decode2<T, 8>(d_data, size, stride, params, d_stream, minbits, maxbits, maxprec, minexp, d_index, index_type, granularity);
         break;
+#endif
         case 16:
           bits_read = internal::decode2<T, 16>(d_data, size, stride, params, d_stream, minbits, maxbits, maxprec, minexp, d_index, index_type, granularity);
         break;
@@ -1071,10 +1073,11 @@ decode(
       break;
     case 3:
       switch (params->min_sub_group_size) {
+#ifdef DISABLE_BLOCK_8
         case 8:
           bits_read = internal::decode3<T, 8>(d_data, size, stride, params, d_stream, minbits, maxbits, maxprec, minexp, d_index, index_type, granularity);
           break;
-        case 16:
+#endif
           bits_read = internal::decode3<T, 16>(d_data, size, stride, params, d_stream, minbits, maxbits, maxprec, minexp, d_index, index_type, granularity);
           break;
         default:

diff --git a/src/sycl/encode.h b/src/sycl/encode.h
@@ -736,9 +736,11 @@ encode(
     case 2:
       switch (params->min_sub_group_size)
       {
+#ifdef DISABLE_BLOCK_8
       case 8:
         bits_written = internal::encode2<T, 8>(d_data, size, stride, params, d_stream, d_index, minbits, maxbits, maxprec, minexp);
         break;
+#endif
       case 16:
         bits_written = internal::encode2<T, 16>(d_data, size, stride, params, d_stream, d_index, minbits, maxbits, maxprec, minexp);
         break;
@@ -750,9 +752,11 @@ encode(
     case 3:
       switch (params->min_sub_group_size)
       {
+#ifdef DISABLE_BLOCK_8
       case 8:
         bits_written = internal::encode3<T, 8>(d_data, size, stride, params, d_stream, d_index, minbits, maxbits, maxprec, minexp);
         break;
+#endif
       case 16:
         bits_written = internal::encode3<T, 16>(d_data, size, stride, params, d_stream, d_index, minbits, maxbits, maxprec, minexp);
         break;

diff --git a/src/sycl/shared.h b/src/sycl/shared.h
@@ -172,21 +172,21 @@ template <>
 inline 
 const unsigned char* get_perm<4>()
 {
-  return perm_1.get();
+  return perm_1;
 }
 
 template <>
 inline 
 const unsigned char* get_perm<16>()
 {
-  return perm_2.get();
+  return perm_2;
 }
 
 template <>
 inline 
 const unsigned char* get_perm<64>()
 {
-  return perm_3.get();
+  return perm_3;
 }
 
 // maximum number of bit planes to encode/decode