Return to signed types

Signed-off-by: Larsen, Steffen <[email protected]>

Author: steffenlarsen

clang/lib/Sema/SPIRVBuiltins.td

@@ -910,7 +910,7 @@ foreach name = ["BitCount"] in {
 }
 
 def : SPVBuiltin<"BitwiseFunctionINTEL",
-                 [AUIGenTypeN, AUIGenTypeN, AUIGenTypeN, AUIGenTypeN, Int]>;
+                 [AIGenTypeN, AIGenTypeN, AIGenTypeN, AIGenTypeN, Int]>;
 
 // 3.32.20. Barrier Instructions
 

sycl/include/sycl/ext/oneapi/experimental/ternary_bitwise.hpp

@@ -555,15 +555,11 @@ sycl::detail::builtin_enable_integer_t<T> ternary_bitwise(T A, T B, T C) {
 #if defined(__SYCL_DEVICE_ONLY__) && !defined(__NVPTX__) && !defined(__AMDGCN__)
     // TODO: Implement __spirv_BitwiseFunctionINTEL for NVPTX and AMDGCN.
     using ret_type = sycl::detail::builtin_enable_integer_t<T>;
-    using UnsignedT = sycl::detail::make_unsigned_t<T>;
-    auto UnsignedA = sycl::bit_cast<UnsignedT>(A);
-    auto UnsignedB = sycl::bit_cast<UnsignedT>(B);
-    auto UnsignedC = sycl::bit_cast<UnsignedT>(C);
-    return sycl::bit_cast<ret_type>(__spirv_BitwiseFunctionINTEL(
-        sycl::detail::builtins::convert_arg(UnsignedA),
-        sycl::detail::builtins::convert_arg(UnsignedB),
-        sycl::detail::builtins::convert_arg(UnsignedC),
-        static_cast<uint32_t>(LUTIndex)));
+    return sycl::bit_cast<ret_type>(
+        __spirv_BitwiseFunctionINTEL(sycl::detail::builtins::convert_arg(A),
+                                     sycl::detail::builtins::convert_arg(B),
+                                     sycl::detail::builtins::convert_arg(C),
+                                     static_cast<uint32_t>(LUTIndex)));
 #else
     return sycl::detail::applyTernaryBitwise(
         LUTIndex, sycl::detail::simplify_if_swizzle_t<T>{A},

sycl/test-e2e/Experimental/ternary_bitwise.cpp

@@ -113,8 +113,6 @@ template <typename T> int checkResult(MemObj<T> &Mem, std::string_view TName) {
 int main() {
   sycl::queue Q;
 
-  auto *CharObj = createMem<char>(Q);
-  auto *UCharObj = createMem<unsigned char>(Q);
   auto *ShortObj = createMem<short>(Q);
   auto *UIntObj = createMem<unsigned int>(Q);
   auto *LongObj = createMem<long>(Q);
@@ -127,8 +125,6 @@ int main() {
 #define APPLY(MEM_OBJ)                                                         \
   if ((WorkCounter++) == Idx[0])                                               \
     MEM_OBJ->Out = apply(MEM_OBJ->A, MEM_OBJ->B, MEM_OBJ->C, IdxSeq);
-     APPLY(CharObj)
-     APPLY(UCharObj)
      APPLY(ShortObj)
      APPLY(UIntObj)
      APPLY(LongObj)
@@ -138,16 +134,12 @@ int main() {
 
   int Failed = 0;
 
-  Failed += checkResult(*CharObj, "char");
-  Failed += checkResult(*UCharObj, "unsigned char");
   Failed += checkResult(*ShortObj, "short");
   Failed += checkResult(*UIntObj, "unsigned int");
   Failed += checkResult(*LongObj, "long");
   Failed += checkResult(*UShort8Obj, "sycl::vec<uint16_t, 8>");
   Failed += checkResult(*IntMarrayObj, "sycl::marray<int, 3>");
 
-  sycl::free(CharObj, Q);
-  sycl::free(UCharObj, Q);
   sycl::free(ShortObj, Q);
   sycl::free(UIntObj, Q);
   sycl::free(LongObj, Q);