diff --git a/sycl/test-e2e/LLVMIntrinsicLowering/bitreverse.cpp b/sycl/test-e2e/LLVMIntrinsicLowering/bitreverse.cpp new file mode 100644 index 0000000000000..239b018a9f30f --- /dev/null +++ b/sycl/test-e2e/LLVMIntrinsicLowering/bitreverse.cpp @@ -0,0 +1,236 @@ +// Test that llvm.bitreverse is lowered correctly by llvm-spirv. + +// UNSUPPORTED: hip || cuda + +// Make dump directory. +// RUN: rm -rf %t.spvdir && mkdir %t.spvdir + +// Ensure that SPV_KHR_bit_instructions is disabled so that translator +// will lower llvm.bitreverse.* intrinsics instead of relying on SPIRV +// BitReverse instruction. +// Also build executable with SPV dump. +// RUN: %{build} -o %t.out -O2 -Xspirv-translator --spirv-ext=-SPV_KHR_bit_instructions -fsycl-dump-device-code=%t.spvdir + +// Rename SPV file to explictly known filename. +// RUN: mv %t.spvdir/*.spv %t.spvdir/dump.spv + +// Convert to text. +// RUN: llvm-spirv -to-text %t.spvdir/dump.spv + +// Check that all lowerings are done by llvm-spirv. +// RUN: cat %t.spvdir/dump.spt | FileCheck %s --check-prefix CHECK-SPV --implicit-check-not=BitReverse + +// Execute to ensure lowering has correct functionality. +// RUN: %{run} %t.out + +///////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// + +// TODO FIXME Change NOT_READY to RUN when llvm.bitreverse.* is supported. + +// Build without lowering explicitly disabled. +// NOT_READY: %{build} -o %t.bitinstructions.out + +// Execution should still be correct. +// NOT_READY: %{run} %t.bitinstructions.out + +// CHECK-SPV: Name {{[0-9]+}} "llvm_bitreverse_i8" +// CHECK-SPV: Name {{[0-9]+}} "llvm_bitreverse_i16" +// CHECK-SPV: Name {{[0-9]+}} "llvm_bitreverse_i32" +// CHECK-SPV: Name {{[0-9]+}} "llvm_bitreverse_i64" + +// CHECK-SPV: Name {{[0-9]+}} "llvm_bitreverse_v2i8" +// CHECK-SPV: Name {{[0-9]+}} "llvm_bitreverse_v2i16" +// CHECK-SPV: Name {{[0-9]+}} "llvm_bitreverse_v2i32" + +// CHECK-SPV: Name {{[0-9]+}} "llvm_bitreverse_v3i8" +// CHECK-SPV: Name {{[0-9]+}} "llvm_bitreverse_v3i16" +// CHECK-SPV: Name {{[0-9]+}} "llvm_bitreverse_v3i32" + +// CHECK-SPV: Name {{[0-9]+}} "llvm_bitreverse_v4i8" +// CHECK-SPV: Name {{[0-9]+}} "llvm_bitreverse_v4i16" +// CHECK-SPV: Name {{[0-9]+}} "llvm_bitreverse_v4i32" + +// CHECK-SPV: Name {{[0-9]+}} "llvm_bitreverse_v8i8" +// CHECK-SPV: Name {{[0-9]+}} "llvm_bitreverse_v8i16" +// CHECK-SPV: Name {{[0-9]+}} "llvm_bitreverse_v8i32" + +// CHECK-SPV: Name {{[0-9]+}} "llvm_bitreverse_v16i8" +// CHECK-SPV: Name {{[0-9]+}} "llvm_bitreverse_v16i16" +// CHECK-SPV: Name {{[0-9]+}} "llvm_bitreverse_v16i32" + +// CHECK-SPV: LinkageAttributes "llvm_bitreverse_i8" Export +// CHECK-SPV: LinkageAttributes "llvm_bitreverse_i16" Export +// CHECK-SPV: LinkageAttributes "llvm_bitreverse_i32" Export +// CHECK-SPV: LinkageAttributes "llvm_bitreverse_i64" Export + +// CHECK-SPV: LinkageAttributes "llvm_bitreverse_v2i8" Export +// CHECK-SPV: LinkageAttributes "llvm_bitreverse_v2i16" Export +// CHECK-SPV: LinkageAttributes "llvm_bitreverse_v2i32" Export + +// CHECK-SPV: LinkageAttributes "llvm_bitreverse_v3i8" Export +// CHECK-SPV: LinkageAttributes "llvm_bitreverse_v3i16" Export +// CHECK-SPV: LinkageAttributes "llvm_bitreverse_v3i32" Export + +// CHECK-SPV: LinkageAttributes "llvm_bitreverse_v4i8" Export +// CHECK-SPV: LinkageAttributes "llvm_bitreverse_v4i16" Export +// CHECK-SPV: LinkageAttributes "llvm_bitreverse_v4i32" Export + +// CHECK-SPV: LinkageAttributes "llvm_bitreverse_v8i8" Export +// CHECK-SPV: LinkageAttributes "llvm_bitreverse_v8i16" Export +// CHECK-SPV: LinkageAttributes "llvm_bitreverse_v8i32" Export + +// CHECK-SPV: LinkageAttributes "llvm_bitreverse_v16i8" Export +// CHECK-SPV: LinkageAttributes "llvm_bitreverse_v16i16" Export +// CHECK-SPV: LinkageAttributes "llvm_bitreverse_v16i32" Export + +#include +#include +#include +#include "common.hpp" + +using namespace sycl; + +template +__attribute__((optnone, noinline)) TYPE reference_reverse(TYPE a, const int bitlength) { + TYPE ret = 0; + for (auto i = 0; i>=1; + } + return ret; +} + +template +__attribute__((noinline)) TYPE reverse(TYPE a, int bitlength) { + if (bitlength==8) { + // Avoid bug with __builtin_elementwise_bitreverse(a) on scalar 8-bit types. + a = ((0x55 & a) << 1) | (0x55 & (a >> 1)); + a = ((0x33 & a) << 2) | (0x33 & (a >> 2)); + return (a << 4) | (a >> 4); + } else if (bitlength==16) { + // Avoid bug with __builtin_elementwise_bitreverse(a) on scalar 16-bit types. + a = ((0x5555 & a) << 1) | (0x5555 & (a >> 1)); + a = ((0x3333 & a) << 2) | (0x3333 & (a >> 2)); + a = ((0x0F0F & a) << 4) | (0x0F0F & (a >> 4)); + return (a << 8) | (a >> 8); + } else + return __builtin_elementwise_bitreverse(a); +} + +template class BitreverseTest; + +#define NUM_TESTS 1024 + +template +void do_scalar_bitreverse_test() { + queue q; + + TYPE *Input = (TYPE *) malloc_shared(sizeof(TYPE) * NUM_TESTS, q.get_device(), q.get_context()); + TYPE *Output = (TYPE *) malloc_shared(sizeof(TYPE) * NUM_TESTS, q.get_device(), q.get_context()); + + for (unsigned i=0; i(); + q.submit([=](handler &cgh) { + cgh.single_task> ([=]() { + for (unsigned i=0; i +void do_vector_bitreverse_test() { + queue q; + + VTYPE *Input = (VTYPE *) malloc_shared(sizeof(VTYPE) * NUM_TESTS, q.get_device(), q.get_context()); + VTYPE *Output = (VTYPE *) malloc_shared(sizeof(VTYPE) * NUM_TESTS, q.get_device(), q.get_context()); + + for (unsigned i=0; i::type>(); + + q.submit([=](handler &cgh) { + cgh.single_task> ([=]() { + for (unsigned i=0; i(); + do_scalar_bitreverse_test(); + do_scalar_bitreverse_test(); + do_scalar_bitreverse_test(); + + do_vector_bitreverse_test(); + do_vector_bitreverse_test(); + do_vector_bitreverse_test(); + + do_vector_bitreverse_test(); + do_vector_bitreverse_test(); + do_vector_bitreverse_test(); + + do_vector_bitreverse_test(); + do_vector_bitreverse_test(); + do_vector_bitreverse_test(); + + do_vector_bitreverse_test(); + do_vector_bitreverse_test(); + do_vector_bitreverse_test(); + + do_vector_bitreverse_test(); + do_vector_bitreverse_test(); + do_vector_bitreverse_test(); + + return 0; +} + diff --git a/sycl/test-e2e/LLVMIntrinsicLowering/common.hpp b/sycl/test-e2e/LLVMIntrinsicLowering/common.hpp new file mode 100644 index 0000000000000..45c0a99840d93 --- /dev/null +++ b/sycl/test-e2e/LLVMIntrinsicLowering/common.hpp @@ -0,0 +1,27 @@ +//==------- common.hpp - DPC++ ESIMD on-device 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 +// +//===----------------------------------------------------------------------===// + +#pragma once + +#include +#include + +template class KernelID; + +template T get_rand() { + using Tuint = std::conditional_t< + sizeof(T) == 1, uint8_t, + std::conditional_t< + sizeof(T) == 2, uint16_t, + std::conditional_t>>>; + Tuint v = rand(); + if constexpr (sizeof(Tuint) > 4) + v = (v << 32) | rand(); + return sycl::bit_cast(v); +}