|
| 1 | +// RUN: %{build} -o %t.out |
| 2 | +// RUN: %if hip %{ env SYCL_JIT_AMDGCN_PTX_TARGET_CPU=%{amd_arch} %} %{run} %t.out |
| 3 | + |
| 4 | +// UNSUPPORTED: target-native_cpu |
| 5 | +// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/20142 |
| 6 | + |
| 7 | +// Verify that having #define before includes works as expected. |
| 8 | + |
| 9 | +#include <sycl/detail/core.hpp> |
| 10 | +#include <sycl/kernel_bundle.hpp> |
| 11 | + |
| 12 | +#include <chrono> |
| 13 | +#include <iostream> |
| 14 | +#include <sstream> |
| 15 | +#include <string_view> |
| 16 | + |
| 17 | +using namespace std::string_view_literals; |
| 18 | +namespace syclexp = sycl::ext::oneapi::experimental; |
| 19 | + |
| 20 | +int main(int argc, char **argv) { |
| 21 | + auto Test = [](std::string src) { |
| 22 | + sycl::queue q; |
| 23 | + // Two iterations to test pch creation/use: |
| 24 | + for (int i = 0; i < 2; ++i) { |
| 25 | + sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source> kb_src = |
| 26 | + syclexp::create_kernel_bundle_from_source( |
| 27 | + q.get_context(), syclexp::source_language::sycl, src); |
| 28 | + sycl::kernel_bundle<sycl::bundle_state::executable> kb_exe = |
| 29 | + syclexp::build(kb_src, syclexp::properties{syclexp::build_options{ |
| 30 | + std::vector<std::string>{"--auto-pch"}}} |
| 31 | + |
| 32 | + ); |
| 33 | + } |
| 34 | + }; |
| 35 | + |
| 36 | + Test(R"""( |
| 37 | +#define SYCL_SIMPLE_SWIZZLES |
| 38 | +#include <sycl/ext/oneapi/free_function_queries.hpp> |
| 39 | +#include <sycl/ext/oneapi/kernel_properties/properties.hpp> |
| 40 | +#include <sycl/vector.hpp> |
| 41 | +
|
| 42 | +namespace syclext = sycl::ext::oneapi; |
| 43 | +namespace syclexp = sycl::ext::oneapi::experimental; |
| 44 | +
|
| 45 | +auto foo(sycl::vec<int, 4> v) { |
| 46 | + // Wouldn't work without SYCL_SIMPLE_SWIZZLES |
| 47 | + return v.xx(); |
| 48 | +} |
| 49 | +
|
| 50 | +extern "C" |
| 51 | +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) |
| 52 | +void iota(float start, float *ptr) { |
| 53 | + size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); |
| 54 | + ptr[id] = start + static_cast<float>(id); |
| 55 | +} |
| 56 | +)"""); |
| 57 | +} |
0 commit comments