[SYCLomatic] Fix an issue in migration of __ldcs (#2736)

Signed-off-by: Ziran Zhang <[email protected]>

Author: ziranzha

clang/lib/DPCT/RulesLang/Math/RewriterHalfPrecisionConversionAndDataMovement.cpp

@@ -816,7 +816,7 @@ RewriterMap dpct::createHalfPrecisionConversionAndDataMovementRewriterMap() {
                         EMPTY_FACTORY_ENTRY("__ldcs"),
                         EMPTY_FACTORY_ENTRY("__ldcs"),
                         WARNING_FACTORY_ENTRY(
-                            "__ldcs", DEREF_FACTORY_ENTRY("__ldcs", ARG_WC(0)),
+                            "__ldcs", DEREF_FACTORY_ENTRY("__ldcs", DEREF(0)),
                             Diagnostics::MATH_EMULATION_EXPRESSION,
                             std::string("__ldcs"), std::string("'*'"))))
       // __ldcv

clang/test/dpct/cpp_lang_extensions.cu

@@ -4,6 +4,20 @@
 
 #include "cpp_lang_extensions.cuh"
 
+template<class ElementType>
+struct alignas(16) Packed128 {
+    static constexpr const size_t size = sizeof(int4) / sizeof(ElementType);
+    ElementType payload[size];
+};
+
+
+// load a Packed128 from an aligned memory address with streaming cache hint
+template<class ElementType>
+__device__ Packed128<ElementType> load128cs(const ElementType* address) {
+    // CHECK: return Packed128<ElementType>{*(reinterpret_cast<const sycl::int4*>(address))};
+    return Packed128<ElementType>{__ldcs(reinterpret_cast<const int4*>(address))};
+}
+
 __device__ float df(float f) {
   float a[23];
   // CHECK: /*