remove incorrect subgroup private func names

Signed-off-by: jinge90 <[email protected]>

Author: jinge90

libdevice/fallback-gsort.cpp

@@ -6986,250 +6986,4 @@ void WG_PS_SD(p1i64_p1f32_u32_p3i8)(int64_t *keys, float *vals, uint32_t n,
                                       std::greater_equal<int64_t>{});
 }
 
-// 1-element version of subgroup private sort.
-DEVICE_EXTERN_C_INLINE
-void SG_PS_A(p1u8_p1u8_u32_p1i8)(uint8_t *key, uint8_t *val, uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, val, 1, scratch,
-                                     std::less_equal<uint8_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_D(p1u8_p1u8_u32_p1i8)(uint8_t *key, uint8_t *val, uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, val, 1, scratch,
-                                     std::greater_equal<uint8_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_A(p1u8_p1i8_u32_p1i8)(uint8_t *key, int8_t *val, uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, reinterpret_cast<uint8_t *>(val), 1,
-                                     scratch, std::less_equal<uint8_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_D(p1u8_p1i8_u32_p1i8)(uint8_t *key, int8_t *val, uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, reinterpret_cast<uint8_t *>(val), 1,
-                                     scratch, std::greater_equal<uint8_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_A(p1u8_p1u16_u32_p1i8)(uint8_t *key, uint16_t *val,
-                                  uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, val, 1, scratch,
-                                     std::less_equal<uint8_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_D(p1u8_p1u16_u32_p1i8)(uint8_t *key, uint16_t *val,
-                                  uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, val, 1, scratch,
-                                     std::greater_equal<uint8_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_A(p1u8_p1i16_u32_p1i8)(uint8_t *key, int16_t *val,
-                                  uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, reinterpret_cast<uint16_t *>(val), 1,
-                                     scratch, std::less_equal<uint8_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_D(p1u8_p1i16_u32_p1i8)(uint8_t *key, int16_t *val,
-                                  uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, reinterpret_cast<uint16_t *>(val), 1,
-                                     scratch, std::greater_equal<uint8_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_A(p1u8_p1u32_u32_p1i8)(uint8_t *key, uint32_t *val,
-                                  uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, val, 1, scratch,
-                                     std::less_equal<uint8_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_D(p1u8_p1u32_u32_p1i8)(uint8_t *key, uint32_t *val,
-                                  uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, val, 1, scratch,
-                                     std::greater_equal<uint8_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_A(p1u8_p1i32_u32_p1i8)(uint8_t *key, int32_t *val,
-                                  uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, reinterpret_cast<uint32_t *>(val), 1,
-                                     scratch, std::less_equal<uint8_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_D(p1u8_p1i32_u32_p1i8)(uint8_t *key, int32_t *val,
-                                  uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, reinterpret_cast<uint32_t *>(val), 1,
-                                     scratch, std::greater_equal<uint8_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_A(p1u8_p1u64_u32_p1i8)(uint8_t *key, uint64_t *val,
-                                  uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, val, 1, scratch,
-                                     std::less_equal<uint8_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_D(p1u8_p1u64_u32_p1i8)(uint8_t *key, uint64_t *val,
-                                  uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, val, 1, scratch,
-                                     std::greater_equal<uint8_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_A(p1u8_p1i64_u32_p1i8)(uint8_t *key, int64_t *val,
-                                  uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, reinterpret_cast<uint64_t *>(val), 1,
-                                     scratch, std::less_equal<uint8_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_D(p1u8_p1i64_u32_p1i8)(uint8_t *key, int64_t *val,
-                                  uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, reinterpret_cast<uint64_t *>(val), 1,
-                                     scratch, std::greater_equal<uint8_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_A(p1u8_p1f32_u32_p1i8)(uint8_t *key, float *val, uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, reinterpret_cast<uint32_t *>(val), 1,
-                                     scratch, std::less_equal<uint8_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_D(p1u8_p1f32_u32_p1i8)(uint8_t *key, float *val, uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, reinterpret_cast<uint32_t *>(val), 1,
-                                     scratch, std::greater_equal<uint8_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_A(p1u16_p1u8_u32_p1i8)(uint16_t *key, uint8_t *val,
-                                  uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, val, 1, scratch,
-                                     std::less_equal<uint16_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_D(p1u16_p1u8_u32_p1i8)(uint16_t *key, uint8_t *val,
-                                  uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, val, 1, scratch,
-                                     std::greater_equal<uint16_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_A(p1u16_p1i8_u32_p1i8)(uint16_t *key, int8_t *val,
-                                  uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, reinterpret_cast<uint8_t *>(val), 1,
-                                     scratch, std::less_equal<uint16_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_D(p1u16_p1i8_u32_p1i8)(uint16_t *key, int8_t *val,
-                                  uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, reinterpret_cast<uint8_t *>(val), 1,
-                                     scratch, std::greater_equal<uint16_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_A(p1u16_p1u16_u32_p1i8)(uint16_t *key, uint16_t *val,
-                                   uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, val, 1, scratch,
-                                     std::less_equal<uint16_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_D(p1u16_p1u16_u32_p1i8)(uint16_t *key, uint16_t *val,
-                                   uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, val, 1, scratch,
-                                     std::greater_equal<uint16_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_A(p1u16_p1i16_u32_p1i8)(uint16_t *key, int16_t *val,
-                                   uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, reinterpret_cast<uint16_t *>(val), 1,
-                                     scratch, std::less_equal<uint16_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_D(p1u16_p1i16_u32_p1i8)(uint16_t *key, int16_t *val,
-                                   uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, reinterpret_cast<uint16_t *>(val), 1,
-                                     scratch, std::greater_equal<uint16_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_A(p1u16_p1u32_u32_p1i8)(uint16_t *key, uint32_t *val,
-                                   uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, val, 1, scratch,
-                                     std::less_equal<uint16_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_D(p1u16_p1u32_u32_p1i8)(uint16_t *key, uint32_t *val,
-                                   uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, val, 1, scratch,
-                                     std::greater_equal<uint16_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_A(p1u16_p1i32_u32_p1i8)(uint16_t *key, int32_t *val,
-                                   uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, reinterpret_cast<uint32_t *>(val), 1,
-                                     scratch, std::less_equal<uint16_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_D(p1u16_p1i32_u32_p1i8)(uint16_t *key, int32_t *val,
-                                   uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, reinterpret_cast<uint32_t *>(val), 1,
-                                     scratch, std::greater_equal<uint16_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_A(p1u16_p1u64_u32_p1i8)(uint16_t *key, uint64_t *val,
-                                   uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, val, 1, scratch,
-                                     std::less_equal<uint16_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_D(p1u16_p1u64_u32_p1i8)(uint16_t *key, uint64_t *val,
-                                   uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, val, 1, scratch,
-                                     std::greater_equal<uint16_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_A(p1u16_p1i64_u32_p1i8)(uint16_t *key, int64_t *val,
-                                   uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, reinterpret_cast<uint64_t *>(val), 1,
-                                     scratch, std::less_equal<uint16_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_D(p1u16_p1i64_u32_p1i8)(uint16_t *key, int64_t *val,
-                                   uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, reinterpret_cast<uint64_t *>(val), 1,
-                                     scratch, std::greater_equal<uint16_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_A(p1u16_p1f32_u32_p1i8)(uint16_t *key, float *val,
-                                   uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, reinterpret_cast<uint32_t *>(val), 1,
-                                     scratch, std::less_equal<uint16_t>{});
-}
-
-DEVICE_EXTERN_C_INLINE
-void SG_PS_D(p1u16_p1f32_u32_p1i8)(uint16_t *key, float *val,
-                                   uint8_t *scratch) {
-  private_merge_sort_key_value_close(key, reinterpret_cast<uint32_t *>(val), 1,
-                                     scratch, std::greater_equal<uint16_t>{});
-}
 #endif // __SPIR__ || __SPIRV__

sycl/test-e2e/DeviceLib/group_sort/group_private_KV_sort_p1p1_p1.hpp

@@ -1132,7 +1132,6 @@ void test_work_group_KV_private_sort(sycl::queue &q, const KeyT keys[NUM],
        accessor ivals_acc{ivals_buf, h};
        accessor okeys_acc{okeys_buf, h};
        accessor ovals_acc{ovals_buf, h};
-       sycl::stream os(1024, 128, h);
        h.parallel_for(num_items, [=](nd_item<1> i) {
          KeyT pkeys[num_per_work_item];
          ValT pvals[num_per_work_item];

sycl/test-e2e/DeviceLib/group_sort/group_private_KV_sort_p1p1_p3.hpp

@@ -1134,7 +1134,6 @@ void test_work_group_KV_private_sort(sycl::queue &q, const KeyT keys[NUM],
        accessor okeys_acc{okeys_buf, h};
        accessor ovals_acc{ovals_buf, h};
        local_accessor<uint64_t, 1> scratch_acc(scratch_size >> 3, h);
-       sycl::stream os(1024, 128, h);
        h.parallel_for(num_items, [=](nd_item<1> i) {
          KeyT pkeys[num_per_work_item];
          ValT pvals[num_per_work_item];