FPGA: Update the gzip code sample to use LSU builtins (#2545)

The sample was relying on undocumented compiler flags.
This change updated the sample to instead use the documented LSU builtins to achieve the same goal.

Author: yuguen

DirectProgramming/C++SYCL_FPGA/ReferenceDesigns/gzip/CMakeLists.txt

@@ -135,7 +135,7 @@ message(STATUS "NUM_REORDER=${NUM_REORDER}")
 
 # Use cmake -DUSER_FPGA_FLAGS=<flags> to set extra flags for FPGA backend
 # compilation. 
-set(USER_FPGA_FLAGS ${USER_FPGA_FLAGS};${SEED};-Xsopt-arg=\"-nocaching\";${NUM_REORDER})
+set(USER_FPGA_FLAGS ${USER_FPGA_FLAGS};${SEED};${NUM_REORDER})
 
 # Use cmake -DUSER_FLAGS=<flags> to set extra flags for general compilation.
 set(USER_FLAGS ${USER_FLAGS};-DNUM_ENGINES=${NUM_ENGINES})

DirectProgramming/C++SYCL_FPGA/ReferenceDesigns/gzip/src/gzipkernel_ll.cpp

@@ -76,6 +76,10 @@ using acc_lz_to_crc_channel_array =
       32,
       NUM_ENGINES
     >;
+    
+using NonCachingLSU = sycl::ext::intel::lsu<sycl::ext::intel::burst_coalesce<true>,
+                                            sycl::ext::intel::cache<0>,
+                                            sycl::ext::intel::statically_coalesce<true>>;
 
 template <int Begin, int End>
 struct Unroller {
@@ -2164,7 +2168,8 @@ event SubmitLZReduction(queue &q, size_t block_size, bool last_block,
         // load in new data
         struct LzInput in;
         Unroller<0, kVec>::step([&](int i) {
-          in.data[i] = acc_pibuf[inpos++];  // Reads 16 bytes, just one time.
+          in.data[i] = NonCachingLSU::load(acc_pibuf + inpos); // Reads 16 bytes, just one time.
+          inpos++;
           input_data.arr[i] =
               in.data[i];  // Send a copy of the data to the CRC kernel, through
                            // a pipe. input_data is used to gang the data
@@ -2185,7 +2190,8 @@ event SubmitLZReduction(queue &q, size_t block_size, bool last_block,
 
           // load in new data
           Unroller<0, kVec>::step([&](int i) {
-            in.data[i] = acc_pibuf[inpos++];
+            in.data[i] = NonCachingLSU::load(acc_pibuf + inpos); // Reads 16 bytes, just one time.
+            inpos++;
             input_data.arr[16 * (int)crc_ch_load_upper + i] = in.data[i];
           });