[SYCLomatic #1534] Add test for cub::{BFE, BFI} (#584)

Signed-off-by: Wang, Yihan <yihan.wang@intel.com>

Author: Yihan Wang

features/feature_case/cub/cub_intrinsic.cu

@@ -109,6 +109,58 @@ bool test_ptx_version() {
   return true;
 }
 
+__global__ void bfe_kernel(int *res) {
+  if (cub::BFE((uint8_t)0xF0, 4, 8) != 15) {
+    *res = 1;
+    return;
+  }
+  if (cub::BFE((uint16_t)0x0FF0u, 4, 12) != 255) {
+    *res = 2;
+    return;
+  }
+  if (cub::BFE(0x00FFFF00u, 8, 16) != 65535u) {
+    *res = 3;
+    return;
+  }
+  if (cub::BFE(0x000000FFull, 0, 9) != 255) {
+    *res = 4;
+    return;
+  }
+  *res = 0;
+}
+
+__global__ void bfi_kernel(int *res) {
+  unsigned d = 0;
+  cub::BFI(d, 0x00FF0000u, 0x0000FFFFu, 0, 16);
+  if (d != 0x00FFFFFFu) {
+    *res = 1;
+    return;
+  }
+
+  cub::BFI(d, 0x00FF0000u, 0x000000FFu, 0, 8);
+  if (d != 0x00FF00FFu) {
+    *res = 2;
+    return;
+  }
+  *res = 0;
+}
+
+bool test_bfe() {
+  int *res;
+  cudaMallocManaged(&res, sizeof(int));
+  bfe_kernel<<<1, 1>>>(res);
+  cudaDeviceSynchronize();
+  return *res == 0;
+}
+
+bool test_bfi() {
+  int *res;
+  cudaMallocManaged(&res, sizeof(int));
+  bfi_kernel<<<1, 1>>>(res);
+  cudaDeviceSynchronize();
+  return *res == 0;
+}
+
 #define TEST(FUNC)                                                             \
   if (!FUNC()) {                                                               \
     printf(#FUNC " failed\n");                                                 \
@@ -122,5 +174,7 @@ int main() {
   TEST(test_device_count);
   TEST(test_sync_stream);
   TEST(test_ptx_version);
+  TEST(test_bfe);
+  TEST(test_bfi);
   return 0;
 }