[SYCL][ESIMD] Fix an issue causing a build break when building ESIMD …

…on Windows platform (#6789)
intel · Sep 27, 2022 · 3aa48db · 3aa48db
1 parent 3916d3b
commit 3aa48db
Show file tree

Hide file tree

Showing 2 changed files with 49 additions and 5 deletions.
diff --git a/sycl/include/sycl/ext/intel/experimental/esimd/detail/memory_intrin.hpp b/sycl/include/sycl/ext/intel/experimental/esimd/detail/memory_intrin.hpp
@@ -929,7 +929,7 @@ __esimd_lsc_load_stateless(__ESIMD_DNS::simd_mask_storage_t<N> pred,
       continue;
     }
 
-    constexpr uint MASK = loadstoreAlignMask<Ty, VS, DS, N>();
+    constexpr unsigned MASK = loadstoreAlignMask<Ty, VS, DS, N>();
     constexpr int ChanlCount = __ESIMD_EDNS::to_int<VS>();
 
     int ByteDistance = 0;
@@ -1159,7 +1159,7 @@ __ESIMD_INTRIN void __esimd_lsc_store_stateless(
       continue;
     }
 
-    constexpr uint MASK = loadstoreAlignMask<Ty, VS, DS, N>();
+    constexpr unsigned MASK = loadstoreAlignMask<Ty, VS, DS, N>();
     constexpr int ChanlCount = __ESIMD_EDNS::to_int<VS>();
 
     int ByteDistance = 0;
@@ -1636,7 +1636,7 @@ __esimd_lsc_xatomic_stateless_0(__ESIMD_DNS::simd_mask_storage_t<N> pred,
       continue;
     }
 
-    constexpr uint MASK = loadstoreAlignMask<Ty, VS, DS, N>();
+    constexpr unsigned MASK = loadstoreAlignMask<Ty, VS, DS, N>();
     constexpr int ChanlCount = __ESIMD_EDNS::to_int<VS>();
 
     int ByteDistance = 0;
@@ -1709,7 +1709,7 @@ __esimd_lsc_xatomic_stateless_1(
       continue;
     }
 
-    constexpr uint MASK = loadstoreAlignMask<Ty, VS, DS, N>();
+    constexpr unsigned MASK = loadstoreAlignMask<Ty, VS, DS, N>();
     constexpr int ChanlCount = __ESIMD_EDNS::to_int<VS>();
 
     int ByteDistance = 0;
@@ -1834,7 +1834,7 @@ __esimd_lsc_xatomic_stateless_2(
       continue;
     }
 
-    constexpr uint MASK = loadstoreAlignMask<Ty, VS, DS, N>();
+    constexpr unsigned MASK = loadstoreAlignMask<Ty, VS, DS, N>();
     constexpr int ChanlCount = __ESIMD_EDNS::to_int<VS>();
 
     int ByteDistance = 0;

diff --git a/sycl/test/esimd/regression/windows_build_test.cpp b/sycl/test/esimd/regression/windows_build_test.cpp
@@ -0,0 +1,44 @@
+//==------- Windows_build_test.cpp - DPC++ ESIMD build test ---------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+// REQUIRES: windows
+// RUN: %clangxx -fsycl -fsycl-device-only -fsyntax-only -Xclang -verify %s -I %sycl_include
+// expected-no-diagnostics
+
+// The tests validates an ability to build ESIMD code on windows platform
+
+#include <iostream>
+#include <CL/sycl.hpp>
+#include <sycl/ext/intel/esimd.hpp>
+#include <sycl/ext/intel/experimental/esimd/memory.hpp>
+
+class Kernel;
+
+int main()
+{
+  sycl::queue q;
+  sycl::device dev = q.get_device();
+  sycl::context ctx = q.get_context();
+  std::cout << "Device: " << dev.get_info<sycl::info::device::name>() << std::endl;
+
+  int* buffer = (int*)sycl::aligned_alloc_device(128, 1024, q);
+
+  q.parallel_for<Kernel>(
+    1,
+    [=](sycl::item<1> it) SYCL_ESIMD_KERNEL {
+      using namespace sycl::ext::intel::esimd;
+      using namespace sycl::ext::intel::experimental::esimd;
+
+      simd<int, 32> blk;
+      lsc_block_store<int, 32>(buffer, blk);
+    });
+
+  q.wait();
+  sycl::free(buffer, q);
+  std::cout << "Done" << std::endl;
+  return 0;
+}