diff --git a/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td b/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td index 286546e0dbc9b..c0b324427464b 100644 --- a/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td +++ b/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td @@ -33,7 +33,6 @@ def AspectExt_intel_gpu_slices : Aspect<"ext_intel_gpu_slices">; def AspectExt_intel_gpu_subslices_per_slice : Aspect<"ext_intel_gpu_subslices_per_slice">; def AspectExt_intel_gpu_eu_count_per_subslice : Aspect<"ext_intel_gpu_eu_count_per_subslice">; def AspectExt_intel_max_mem_bandwidth : Aspect<"ext_intel_max_mem_bandwidth">; -def AspectExt_intel_mem_channel : Aspect<"ext_intel_mem_channel">; def AspectUsm_atomic_host_allocations : Aspect<"usm_atomic_host_allocations">; def AspectUsm_atomic_shared_allocations : Aspect<"usm_atomic_shared_allocations">; def AspectAtomic64 : Aspect<"atomic64">; @@ -82,7 +81,6 @@ def AspectExt_oneapi_private_alloca : Aspect<"ext_oneapi_private_alloca">; def AspectExt_oneapi_queue_profiling_tag : Aspect<"ext_oneapi_queue_profiling_tag">; def AspectExt_oneapi_virtual_mem : Aspect<"ext_oneapi_virtual_mem">; def AspectExt_oneapi_cuda_cluster_group : Aspect<"ext_oneapi_cuda_cluster_group">; -def AspectExt_intel_fpga_task_sequence : Aspect<"ext_intel_fpga_task_sequence">; def AspectExt_oneapi_atomic16 : Aspect<"ext_oneapi_atomic16">; def AspectExt_oneapi_virtual_functions : Aspect<"ext_oneapi_virtual_functions">; def AspectExt_intel_spill_memory_size : Aspect<"ext_intel_spill_memory_size">; @@ -139,7 +137,7 @@ def : TargetInfo<"__TestAspectList", AspectUsm_shared_allocations, AspectUsm_system_allocations, AspectExt_intel_pci_address, AspectExt_intel_gpu_eu_count, AspectExt_intel_gpu_eu_simd_width, AspectExt_intel_gpu_slices, AspectExt_intel_gpu_subslices_per_slice, AspectExt_intel_gpu_eu_count_per_subslice, - AspectExt_intel_max_mem_bandwidth, AspectExt_intel_mem_channel, AspectUsm_atomic_host_allocations, + AspectExt_intel_max_mem_bandwidth, AspectUsm_atomic_host_allocations, AspectUsm_atomic_shared_allocations, AspectAtomic64, AspectExt_intel_device_info_uuid, AspectExt_oneapi_srgb, AspectExt_oneapi_native_assert, AspectHost_debuggable, AspectExt_intel_gpu_hw_threads_per_eu, AspectExt_oneapi_cuda_async_barrier, AspectExt_intel_free_memory, @@ -163,7 +161,6 @@ def : TargetInfo<"__TestAspectList", AspectExt_oneapi_tangle, AspectExt_intel_matrix, AspectExt_oneapi_is_composite, AspectExt_oneapi_is_component, AspectExt_oneapi_graph, AspectExt_oneapi_limited_graph, AspectExt_oneapi_private_alloca, AspectExt_oneapi_queue_profiling_tag, AspectExt_oneapi_virtual_mem, AspectExt_oneapi_cuda_cluster_group, - AspectExt_intel_fpga_task_sequence, AspectExt_oneapi_atomic16, AspectExt_oneapi_virtual_functions, AspectExt_intel_spill_memory_size, diff --git a/llvm/include/llvm/SYCLLowerIR/HostPipes.h b/llvm/include/llvm/SYCLLowerIR/HostPipes.h deleted file mode 100644 index 9942a99187c3b..0000000000000 --- a/llvm/include/llvm/SYCLLowerIR/HostPipes.h +++ /dev/null @@ -1,59 +0,0 @@ -//===------- HostPipes.h - get required info about FPGA Host Pipes --------===// -// -// 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 -// -//===----------------------------------------------------------------------===// -// -// The file contains a number of functions to extract corresponding attributes -// of the host pipe global variables and save them as a property set for the -// runtime. -//===----------------------------------------------------------------------===// - -#pragma once - -#include "llvm/ADT/MapVector.h" - -#include -#include - -namespace llvm { - -class GlobalVariable; -class Module; -class StringRef; - -// Represents a host pipe variable - at SYCL RT level host pipe -// variables are being represented as a byte-array. -struct HostPipeProperty { - HostPipeProperty(uint32_t Size) : Size(Size) {} - - // Encodes size of the underlying type T of the host pipe variable. - uint32_t Size; -}; - -using HostPipePropertyMapTy = - MapVector>; - -/// Return \c true if the variable @GV is a host pipe variable. -/// -/// The function checks whether the variable has the LLVM IR attribute \c -/// sycl-host-pipe -/// @param GV [in] A variable to test. -/// -/// @return \c true if the variable is a host pipe variable, \c false -/// otherwise. -bool isHostPipeVariable(const GlobalVariable &GV); - -/// Searches given module for occurrences of host pipe variable-specific -/// metadata and builds "host pipe variable name" -> -/// vector<"variable properties"> map. -/// -/// @param M [in] LLVM Module. -/// -/// @returns the "host pipe variable name" -> vector<"variable properties"> -/// map. -HostPipePropertyMapTy collectHostPipeProperties(const Module &M); - -} // end namespace llvm diff --git a/llvm/lib/SYCLLowerIR/CMakeLists.txt b/llvm/lib/SYCLLowerIR/CMakeLists.txt index 4320eedd69bda..020e4c378b1c2 100644 --- a/llvm/lib/SYCLLowerIR/CMakeLists.txt +++ b/llvm/lib/SYCLLowerIR/CMakeLists.txt @@ -39,7 +39,6 @@ add_llvm_component_library(LLVMSYCLLowerIR DeviceGlobals.cpp ESIMD/LowerESIMDVLoadVStore.cpp ESIMD/LowerESIMDSlmReservation.cpp - HostPipes.cpp LowerInvokeSimd.cpp LowerWGLocalMemory.cpp LowerWGScope.cpp diff --git a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp index c049ad34434f2..7ee80814739cc 100644 --- a/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp +++ b/llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp @@ -11,7 +11,6 @@ #include "llvm/SYCLLowerIR/CompileTimePropertiesPass.h" #include "llvm/SYCLLowerIR/DeviceGlobals.h" #include "llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h" -#include "llvm/SYCLLowerIR/HostPipes.h" #include "llvm/SYCLLowerIR/TargetHelpers.h" #include "llvm/ADT/APInt.h" @@ -602,12 +601,6 @@ PreservedAnalyses CompileTimePropertiesPass::run(Module &M, HostAccessDecorValue, VarName)); } - if (isHostPipeVariable(GV)) { - auto VarName = getGlobalVariableUniqueId(GV); - MDOps.push_back(buildSpirvDecorMetadata( - Ctx, SpirvHostAccessDecor, SpirvHostAccessDefaultValue, VarName)); - } - // Add the generated metadata to the variable if (!MDOps.empty()) { GV.addMetadata(MDKindID, *MDNode::get(Ctx, MDOps)); diff --git a/llvm/lib/SYCLLowerIR/DeviceGlobals.cpp b/llvm/lib/SYCLLowerIR/DeviceGlobals.cpp index 2dd2bb0560552..87d2f7881a23b 100644 --- a/llvm/lib/SYCLLowerIR/DeviceGlobals.cpp +++ b/llvm/lib/SYCLLowerIR/DeviceGlobals.cpp @@ -72,19 +72,19 @@ bool hasDeviceImageScopeProperty(const GlobalVariable &GV) { return hasProperty(GV, SYCL_DEVICE_IMAGE_SCOPE_ATTR); } -/// Returns the unique id for the device global or host pipe variable. +/// Returns the unique id for the device global variable. /// /// The function gets this value from the LLVM IR attribute \c /// sycl-unique-id. /// -/// @param GV [in] Device Global or Hostpipe variable. +/// @param GV [in] Device Global variable. /// -/// @returns the unique id of the device global or hostpipe variable +/// @returns the unique id of the device global variable /// represented in the LLVM IR by \c GV. StringRef getGlobalVariableUniqueId(const GlobalVariable &GV) { assert(GV.hasAttribute(SYCL_UNIQUE_ID_ATTR) && "a 'sycl-unique-id' string must be associated with every device " - "global or hostpipe variable"); + "global variable"); return GV.getAttribute(SYCL_UNIQUE_ID_ATTR).getValueAsString(); } diff --git a/llvm/lib/SYCLLowerIR/HostPipes.cpp b/llvm/lib/SYCLLowerIR/HostPipes.cpp deleted file mode 100644 index ca116c702c1d4..0000000000000 --- a/llvm/lib/SYCLLowerIR/HostPipes.cpp +++ /dev/null @@ -1,80 +0,0 @@ -//===------------- HostPipes.cpp - SYCL Host Pipes Pass -------------------===// -// -// 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 -// -//===----------------------------------------------------------------------===// -// See comments in the header. -//===----------------------------------------------------------------------===// - -#include "llvm/SYCLLowerIR/HostPipes.h" -#include "llvm/SYCLLowerIR/CompileTimePropertiesPass.h" -#include "llvm/SYCLLowerIR/DeviceGlobals.h" - -#include "llvm/ADT/STLExtras.h" -#include "llvm/ADT/StringRef.h" -#include "llvm/IR/Module.h" - -#include - -using namespace llvm; - -namespace { - -constexpr StringRef SYCL_HOST_PIPE_ATTR = "sycl-host-pipe"; -constexpr StringRef SYCL_HOST_PIPE_SIZE_ATTR = "sycl-host-pipe-size"; - -/// Returns the size (in bytes) of the type \c T of the host -/// pipe variable. -/// -/// The function gets this value from the LLVM IR attribute \c -/// sycl-host-pipe-size. -/// -/// @param GV [in] Host Pipe variable. -/// -/// @returns the size (int bytes) of the underlying type \c T of the -/// host pipe variable represented in the LLVM IR by @GV. -uint32_t getHostPipeTypeSize(const GlobalVariable &GV) { - assert(GV.hasAttribute(SYCL_HOST_PIPE_SIZE_ATTR) && - "The host pipe variable must have the 'sycl-host-pipe-size' " - "attribute that must contain a number representing the size of the " - "underlying type T of the host pipe variable"); - return getAttributeAsInteger(GV, SYCL_HOST_PIPE_SIZE_ATTR); -} - -} // anonymous namespace - -namespace llvm { - -/// Return \c true if the variable @GV is a host pipe variable. -/// -/// The function checks whether the variable has the LLVM IR attribute \c -/// sycl-host-pipe. -/// @param GV [in] A variable to test. -/// -/// @return \c true if the variable is a host pipe variable, \c false -/// otherwise. -bool isHostPipeVariable(const GlobalVariable &GV) { - return GV.hasAttribute(SYCL_HOST_PIPE_ATTR); -} - -HostPipePropertyMapTy collectHostPipeProperties(const Module &M) { - HostPipePropertyMapTy HPM; - auto HostPipeNum = count_if(M.globals(), isHostPipeVariable); - if (HostPipeNum == 0) - return HPM; - - HPM.reserve(HostPipeNum); - - for (auto &GV : M.globals()) { - if (!isHostPipeVariable(GV)) - continue; - - HPM[getGlobalVariableUniqueId(GV)] = {getHostPipeTypeSize(GV)}; - } - - return HPM; -} - -} // namespace llvm diff --git a/llvm/lib/SYCLPostLink/ComputeModuleRuntimeInfo.cpp b/llvm/lib/SYCLPostLink/ComputeModuleRuntimeInfo.cpp index f3172adc7769e..c38f920973e5d 100644 --- a/llvm/lib/SYCLPostLink/ComputeModuleRuntimeInfo.cpp +++ b/llvm/lib/SYCLPostLink/ComputeModuleRuntimeInfo.cpp @@ -14,7 +14,6 @@ #include "llvm/IR/PassInstrumentation.h" #include "llvm/SYCLLowerIR/CompileTimePropertiesPass.h" #include "llvm/SYCLLowerIR/DeviceGlobals.h" -#include "llvm/SYCLLowerIR/HostPipes.h" #include "llvm/SYCLLowerIR/LowerWGLocalMemory.h" #include "llvm/SYCLLowerIR/SYCLKernelParamOptInfo.h" #include "llvm/SYCLLowerIR/SYCLUtils.h" @@ -388,10 +387,6 @@ PropSetRegTy computeModuleProperties(const Module &M, PropSet.add(PropSetRegTy::SYCL_DEVICE_GLOBALS, DevGlobalPropertyMap); } - auto HostPipePropertyMap = collectHostPipeProperties(M); - if (!HostPipePropertyMap.empty()) { - PropSet.add(PropSetRegTy::SYCL_HOST_PIPES, HostPipePropertyMap); - } bool IsSpecConstantDefault = M.getNamedMetadata( SpecConstantsPass::SPEC_CONST_DEFAULT_VAL_MODULE_MD_STRING) != diff --git a/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/fpga-single-task-property.ll b/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/fpga-single-task-property.ll deleted file mode 100644 index 85175e2b8efe0..0000000000000 --- a/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/fpga-single-task-property.ll +++ /dev/null @@ -1,15 +0,0 @@ -; RUN: opt -passes=compile-time-properties --mtriple=spir64_fpga-unknown-unknown %s -S | FileCheck %s --check-prefix CHECK-FPGA-IR -; RUN: opt -passes=compile-time-properties %s -S | FileCheck %s --check-prefix CHECK-DEFAULT-IR - -; CHECK-DEFAULT-IR-NOT: !max_global_work_dim - -; CHECK-FPGA-IR-DAG: @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE9TheKernel0"() #0 {{.*}}!max_global_work_dim ![[MaxGlobWorkDim:[0-9]+]] -; Function Attrs: convergent norecurse -define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE9TheKernel0"() #0 { -entry: - ret void -} - -attributes #0 = { convergent norecurse "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="fpga_single_task_property.cpp" "uniform-work-group-size"="true" "sycl-single-task" } - -; CHECK-FPGA-IR-DAG: ![[MaxGlobWorkDim]] = !{i32 0} diff --git a/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/host-pipes/basic.ll b/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/host-pipes/basic.ll deleted file mode 100644 index 8619156b79e7d..0000000000000 --- a/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/host-pipes/basic.ll +++ /dev/null @@ -1,21 +0,0 @@ -; RUN: opt -passes=compile-time-properties %s -S | FileCheck %s --check-prefix CHECK-IR - -; This test is intended to check that CompileTimePropertiesPass adds all the required -; metadata nodes to host pipe vars decorated with the "sycl-host-pipe" attribute - -source_filename = "basic.cpp" -target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64" -target triple = "spir64_fpga-unknown-unknown" - -%struct.BasicKernel = type { i8 } - -$_ZN4sycl3_V13ext5intel12experimental9host_pipeI9D2HPipeIDiNS1_6oneapi12experimental10propertiesISt5tupleIJEEEEE6__pipeE = comdat any - -@_ZN4sycl3_V13ext5intel12experimental9host_pipeI9D2HPipeIDiNS1_6oneapi12experimental10propertiesISt5tupleIJEEEEE6__pipeE = linkonce_odr dso_local addrspace(1) constant %struct.BasicKernel zeroinitializer, comdat, align 1 #0 -; CHECK-IR: @_ZN4sycl3_V13ext5intel12experimental9host_pipeI9D2HPipeIDiNS1_6oneapi12experimental10propertiesISt5tupleIJEEEEE6__pipeE = linkonce_odr dso_local addrspace(1) constant %struct.BasicKernel zeroinitializer, comdat, align 1, !spirv.Decorations ![[#MN0:]] - -attributes #0 = { "sycl-host-pipe" "sycl-host-pipe-size"="4" "sycl-unique-id"="_ZN4sycl3_V13ext5intel12experimental9host_pipeI9H2DPipeIDiNS1_6oneapi12experimental10propertiesISt5tupleIJEEEEE6__pipeE" } - -; Ensure that the generated metadata nodes are correct -; CHECK-IR-DAG: ![[#MN0]] = !{![[#MN1:]]} -; CHECK-IR-DAG: ![[#MN1]] = !{i32 6147, i32 2, !"_ZN4sycl3_V13ext5intel12experimental9host_pipeI9H2DPipeIDiNS1_6oneapi12experimental10propertiesISt5tupleIJEEEEE6__pipeE"} diff --git a/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/kernel-attributes/fpga-cluster.ll b/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/kernel-attributes/fpga-cluster.ll deleted file mode 100644 index cad7c4398e929..0000000000000 --- a/llvm/test/SYCLLowerIR/CompileTimePropertiesPass/kernel-attributes/fpga-cluster.ll +++ /dev/null @@ -1,33 +0,0 @@ -; Check conversion of sycl-fpga-cluster attribute -; RUN: opt -passes="compile-time-properties" %s -S -o - | FileCheck %s --check-prefix CHECK-IR - -; CHECK-IR-DAG: @stallFree() #0 {{.*}}!stall_free [[MD_TRUE:![0-9]+]] { -; Function Attrs: convergent norecurse -define weak_odr dso_local spir_kernel void @stallFree() #0 { -entry: - ret void -} - -; CHECK-IR-DAG: @stallEnable() #1 {{.*}}!stall_enable [[MD_TRUE:![0-9]+]] { -; Function Attrs: convergent norecurse -define weak_odr dso_local spir_kernel void @stallEnable() #1 { -entry: - ret void -} - -attributes #0 = { convergent norecurse "frame-pointer"="all" "sycl-fpga-cluster"="0" } -attributes #1 = { convergent norecurse "frame-pointer"="all" "sycl-fpga-cluster"="1" } - -!opencl.spir.version = !{!0, !0, !0, !0, !0, !0, !0, !0, !0, !0, !0} -!spirv.Source = !{!1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1} -!llvm.ident = !{!2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2} -!llvm.module.flags = !{!3, !4} - -!0 = !{i32 1, i32 2} -!1 = !{i32 4, i32 100000} -!2 = !{!"clang version 13.0.0 (https://github.com/intel/llvm)"} -!3 = !{i32 1, !"wchar_size", i32 4} -!4 = !{i32 7, !"frame-pointer", i32 2} - -; Confirm the decorations for the functions -; CHECK-IR-DAG: [[MD_TRUE]] = !{i32 1} diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp index 8f21122f9e996..c49c552c1d2db 100644 --- a/llvm/tools/sycl-post-link/sycl-post-link.cpp +++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp @@ -31,7 +31,6 @@ #include "llvm/SYCLLowerIR/CompileTimePropertiesPass.h" #include "llvm/SYCLLowerIR/DeviceConfigFile.hpp" #include "llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h" -#include "llvm/SYCLLowerIR/HostPipes.h" #include "llvm/SYCLLowerIR/LowerInvokeSimd.h" #include "llvm/SYCLLowerIR/SYCLDeviceLibBF16.h" #include "llvm/SYCLLowerIR/SYCLJointMatrixTransform.h" diff --git a/sycl/doc/GetStartedGuide.md b/sycl/doc/GetStartedGuide.md index 85d3db3256a1e..80d5b9513e503 100644 --- a/sycl/doc/GetStartedGuide.md +++ b/sycl/doc/GetStartedGuide.md @@ -1,7 +1,7 @@ # Getting Started with oneAPI DPC++ The DPC++ Compiler compiles C++ and SYCL\* source files with code for both CPU -and a wide range of compute accelerators such as GPU and FPGA. +and GPU. ## Table of contents @@ -23,7 +23,6 @@ and a wide range of compute accelerators such as GPU and FPGA. * [Obtain prerequisites for ahead of time (AOT) compilation](#obtain-prerequisites-for-ahead-of-time-aot-compilation) * [GPU](#gpu) * [CPU](#cpu) - * [Accelerator](#accelerator) * [Test DPC++ toolchain](#test-dpc-toolchain) * [Run in-tree LIT tests](#run-in-tree-lit-tests) * [Run DPC++ E2E tests](#run-dpc-e2e-tests) @@ -407,8 +406,8 @@ To run DPC++ applications on Level Zero devices, Level Zero implementation(s) must be present in the system. You can find the link to the Level Zero spec in the following section [Find More](#find-more). -The Level Zero RT for `GPU`, OpenCL RT for `GPU`, OpenCL RT for `CPU`, FPGA -emulation RT and TBB runtime which are needed to run DPC++ application +The Level Zero RT for `GPU`, OpenCL RT for `GPU`, OpenCL RT for `CPU` +and TBB runtime which are needed to run DPC++ application on Intel `GPU` or Intel `CPU` devices can be downloaded using links in [the dependency configuration file](../../devops/dependencies.json) and installed following the instructions below. The same versions are used in @@ -416,15 +415,10 @@ PR testing. **Linux**: -1) Extract the archive. For example, for the archives -`oclcpuexp_.tar.gz` and `fpgaemu_.tar.gz` you would -run the following commands +1) Extract the archive. For example, for the archive +`oclcpuexp_.tar.gz` you would run the following commands ```bash - # Extract OpenCL FPGA emulation RT - mkdir -p /opt/intel/oclfpgaemu_ - cd /opt/intel/oclfpgaemu_ - tar zxvf fpgaemu_.tar.gz # Extract OpenCL CPU RT mkdir -p /opt/intel/oclcpuexp_ cd /opt/intel/oclcpuexp_ @@ -434,9 +428,6 @@ run the following commands 2) Create ICD file pointing to the new runtime (requires sudo access) ```bash - # OpenCL FPGA emulation RT - echo /opt/intel/oclfpgaemu_/x64/libintelocl_emu.so | sudo tee - /etc/OpenCL/vendors/intel_fpgaemu.icd # OpenCL CPU RT echo /opt/intel/oclcpuexp_/x64/libintelocl.so | sudo tee /etc/OpenCL/vendors/intel_expcpu.icd @@ -456,15 +447,6 @@ for the archive oneapi-tbb--lin.tgz: folder: ```bash - # OpenCL FPGA emulation RT - ln -s /opt/intel/oneapi-tbb-/lib/intel64/gcc4.8/libtbb.so - /opt/intel/oclfpgaemu_/x64/libtbb.so - ln -s /opt/intel/oneapi-tbb-/lib/intel64/gcc4.8/libtbbmalloc.so - /opt/intel/oclfpgaemu_/x64/libtbbmalloc.so - ln -s /opt/intel/oneapi-tbb-/lib/intel64/gcc4.8/libtbb.so.12 - /opt/intel/oclfpgaemu_/x64/libtbb.so.12 - ln -s /opt/intel/oneapi-tbb-/lib/intel64/gcc4.8/libtbbmalloc.so.2 - /opt/intel/oclfpgaemu_/x64/libtbbmalloc.so.2 # OpenCL CPU RT ln -s /opt/intel/oneapi-tbb-/lib/intel64/gcc4.8/libtbb.so /opt/intel/oclcpuexp_/x64/libtbb.so @@ -479,8 +461,6 @@ folder: 5) Configure library paths (requires sudo access) ```bash - echo /opt/intel/oclfpgaemu_/x64 | sudo tee - /etc/ld.so.conf.d/libintelopenclexp.conf echo /opt/intel/oclcpuexp_/x64 | sudo tee -a /etc/ld.so.conf.d/libintelopenclexp.conf sudo ldconfig -f /etc/ld.so.conf.d/libintelopenclexp.conf @@ -494,8 +474,7 @@ OpenCL runtime for Intel `GPU` installer may re-write some important files or settings and make existing OpenCL runtime for Intel `CPU` runtime not working properly. -2) Extract the archive with OpenCL runtime for Intel `CPU` and/or for Intel -`FPGA` emulation using links in +2) Extract the archive with OpenCL runtime for Intel `CPU` using links in [the dependency configuration file](../../devops/dependencies.json). For example, to `c:\oclcpu_rt_`. @@ -513,9 +492,6 @@ extracted files are in `c:\oclcpu_rt_\` folder, then type the command: ```bash - # Install OpenCL FPGA emulation RT - # Answer Y to clean previous OCL_ICD_FILENAMES configuration and ICD records cleanup - c:\oclfpga_rt_\install.bat c:\oneapi-tbb-\redist\intel64\vc14 # Install OpenCL CPU RT # Answer N for ICD records cleanup c:\oclcpu_rt_\install.bat c:\oneapi-tbb-\redist\intel64\vc14 @@ -529,7 +505,6 @@ AOT compiler for each device type: * `GPU`, Level Zero and OpenCL runtimes are supported, * `CPU`, OpenCL runtime is supported, -* `Accelerator` (FPGA or FPGA emulation), OpenCL runtime is supported. #### GPU @@ -558,15 +533,6 @@ AOT compiler for each device type: * CPU AOT compiler `opencl-aot` is enabled by default. For more, see [opencl-aot documentation](https://github.com/intel/llvm/blob/sycl/opencl/opencl-aot/README.md). -#### Accelerator - -* Accelerator AOT compiler `aoc` is a part of -[Intel® oneAPI Base Toolkit](https://software.intel.com/content/www/us/en/develop/tools/oneapi/base-toolkit.html) -(Intel® oneAPI DPC++/C++ Compiler component). -Make sure that these binaries are available in `PATH` environment variable: - - * `aoc` from `/compiler///lib/oclfpga/bin` - * `aocl-ioc64` from `/compiler///bin` ### Test DPC++ toolchain @@ -720,8 +686,7 @@ The results are correct! ``` **NOTE**: oneAPI DPC++/SYCL developers can specify SYCL device for execution -using device selectors (e.g. `sycl::cpu_selector_v`, `sycl::gpu_selector_v`, -[Intel FPGA selector(s)](extensions/supported/sycl_ext_intel_fpga_device_selector.asciidoc)) +using device selectors (e.g. `sycl::cpu_selector_v`, `sycl::gpu_selector_v`) as explained in following section [Code the program for a specific GPU](#code-the-program-for-a-specific-gpu). @@ -760,14 +725,13 @@ simplify passing the specific architectures, for example [Users Manual](UsersManual.md#generic-options), for the `-fsycl-targets` option. -To build simple-sycl-app ahead of time for GPU, CPU or Accelerator devices, +To build simple-sycl-app ahead of time for GPU or CPU devices, specify the target architecture. The examples provided use a supported alias for the target, representing a full triple. Additional details can be found in the [Users Manual](UsersManual.md#generic-options). ```-fsycl-targets=spir64_gen``` for GPU, -```-fsycl-targets=spir64_x86_64``` for CPU, -```-fsycl-targets=spir64_fpga``` for Accelerator. +```-fsycl-targets=spir64_x86_64``` for CPU. Multiple target architectures are supported. @@ -785,7 +749,6 @@ more. To find available options, execute: ```ocloc compile --help``` for GPU, ```opencl-aot --help``` for CPU, -```aoc -help -sycl``` for Accelerator. The `simple-sycl-app.exe` application doesn't specify SYCL device for execution, so SYCL runtime will use `default_selector` logic to select one @@ -838,8 +801,8 @@ available, a "device selector" may be used. A "device selector" is a ranking function (C++ Callable) that will give an integer ranking value to all the devices on the system. It can be passed to `sycl::queue`, `sycl::device` and `sycl::platform` constructors. The highest ranking device is then selected. SYCL -has built-in device selectors for selecting a generic GPU, CPU, or accelerator -device, as well as one for a default device. Additionally, a user can define +has built-in device selectors for selecting a generic GPU, CPU, as well as one +for a default device. Additionally, a user can define their own as function, lambda, or functor class. Device selectors returning negative values will "reject" a device ensuring it is not selected, but values 0 or higher will be selected by the highest score with ties resolved by an diff --git a/sycl/doc/design/DeviceConfigFile.md b/sycl/doc/design/DeviceConfigFile.md index 35273bbddcc84..06da5e5fa73ba 100644 --- a/sycl/doc/design/DeviceConfigFile.md +++ b/sycl/doc/design/DeviceConfigFile.md @@ -166,7 +166,6 @@ def AspectExt_intel_gpu_slices : Aspect<"ext_intel_gpu_slices">; def AspectExt_intel_gpu_subslices_per_slice : Aspect<"ext_intel_gpu_subslices_per_slice">; def AspectExt_intel_gpu_eu_count_per_subslice : Aspect<"ext_intel_gpu_eu_count_per_subslice">; def AspectExt_intel_max_mem_bandwidth : Aspect<"ext_intel_max_mem_bandwidth">; -def AspectExt_intel_mem_channel : Aspect<"ext_intel_mem_channel">; def AspectUsm_atomic_host_allocations : Aspect<"usm_atomic_host_allocations">; def AspectUsm_atomic_shared_allocations : Aspect<"usm_atomic_shared_allocations">; def AspectAtomic64 : Aspect<"atomic64">; diff --git a/sycl/doc/design/DeviceGlobal.md b/sycl/doc/design/DeviceGlobal.md index 3bbcc55b9a289..36c73d851a776 100644 --- a/sycl/doc/design/DeviceGlobal.md +++ b/sycl/doc/design/DeviceGlobal.md @@ -396,15 +396,6 @@ IR for this decoration. The first SPIR-V operand is set according to the have that property). The second SPIR-V operation is set to the value of the device global's `sycl-unique-id`. -The `sycl-post-link` tool also generates idiomatic LLVM IR for the -**InitModeINTEL** decoration (if the device global has the `init_mode` -property) and for the **ImplementInCSRINTEL** decoration (if the device global -has the `implement_in_csr` property). See the -[SPV\_INTEL\_global\_variable\_decorations][6] SPIR-V extension for details -about all of these decorations. - -[6]: - The `sycl-post-link` tool also create a "SYCL/device globals" property set for each device code module that contains at least one device global variable. diff --git a/sycl/doc/design/HostPipes.md b/sycl/doc/design/HostPipes.md deleted file mode 100644 index 5a634a224f67e..0000000000000 --- a/sycl/doc/design/HostPipes.md +++ /dev/null @@ -1,191 +0,0 @@ -# Implementation design for "Host Pipes" - -This document describes the implementation design for the host pipes section -of the DPC++ extension [SYCL_INTEL_data_flow_pipes][1]. Pipes are a FIFO construct -that provide links between elements of a design that are accessed through read -and write application programming interfaces (APIs), without the notion of a -memory address/pointer to elements within the FIFO. A host pipe is a pipe that -links a device kernel with a host program. - -[1]: <../extensions/supported/sycl_ext_intel_dataflow_pipes.asciidoc> - -## Requirements - -The extension specification document referenced above contains the full set of -requirements for this feature, but a requirement that is particularly -relevant to the design, and similar in nature to one raised in the [device_global][2] -design is called out here. - -This issue relates to the mechanism for integrating host and device code. -Like device global variables, host pipes are referenced in both -host and device code, so they require some mechanism to correlate the pipe -instance in device code with the pipe instance in host code. We will use -a similar mechanism as the device global implementation that creates a map -database in the integration headers and footers. - -[2]: - -## Design - -### Changes to DPC++ headers - -#### Attributes attached to the class - -The `pipe` class uses a new C++ attribute `[[__sycl_detail__::host_pipe]]` on the -`pipe::__pipeType` type to identify the `static const __pipeType` member `__pipe` -as a host pipe. Similar to `[[__sycl_detail__::device_global]]`, this will inform -the front end to generate a `sycl-unique-id` for each `__pipe`. The `pipe` class -also introduces the global variable attribute `sycl-host-pipe` attribute to inform the sycl-post-link tool -to generate the SPIR-V decoration `HostAccessINTEL` for each `__pipe` using the -`sycl-unique-id` generated. - -As these attributes are only needed for the device compiler, the `#ifdef __SYCL_DEVICE_ONLY__` -allows the customer to use another host compiler, even if it does not recognize these attributes. -Also note that these attributes are all in the `__sycl_detail__` namespace, so -they are considered implementation details of DPC++. We do not intend to -support them as general attributes that customer code can use. - -``` -template -class pipe {/*...*/}; - -// Partial specialization to make propertiesT visible as a parameter pack -// of properties. -template -class pipe -{ - struct -#ifdef __SYCL_DEVICE_ONLY__ - [[__sycl_detail__::add_ir_attributes_global_variable( - "sycl-host-pipe", - Props::meta_name..., - nullptr, - Props::meta_value... - )]] - [[__sycl_detail__::host_pipe]] - [[__sycl_detail__::global_variable_allowed]] // may not be needed -#endif - __pipeType { const char __p; }; - - static constexpr __pipeType __pipe = {0}; - ... -}; -``` -The `[[__sycl_detail__::add_ir_attributes_global_variable()]]` attribute is -described more fully by the [compile-time properties][3] design -document. This attribute is also used for other classes that have properties, -so it is not specific to the `pipe` class. - -The address of `static const __pipeType` member `__pipe` will be used to identify the pipe -in host code, and provide one half of the host-to-device mapping of the pipe -(see the section on __New content in the integration header and footer__ below). - -[3]: - -### Changes to the DPC++ front-end - -There are several changes to the device compiler front-end: - -* The front-end adds a new LLVM IR attribute `sycl-unique-id` to the definition - of each `pipe` variable, which provides a unique string identifier - for each. - -* The front-end generates new content in both the integration header and the - integration footer, which is described in more detail below. - -#### New content in the integration header and footer - -New content in the integration header and footer provides a mapping from the -host address of each pipe variable to the unique string for that -variable. To illustrate, consider a translation unit that defines two -`pipe` classes: - -``` -#include - -class some_pipe; -namespace inner { - class some_other_pipe; -} // namespace inner -... -pipe::write(...); // a usage of pipe -... -pipe::read(...); // a usage of pipe -... - -``` - -The corresponding integration header defines a namespace scope variable of type -`__sycl_host_pipe_registration` (referred to below as the __host pipe registrar__) -whose sole purpose is to run its constructor before the application's main() function: - -``` -namespace sycl::detail { -namespace { - -class __sycl_host_pipe_registration { - public: - __sycl_host_pipe_registration() noexcept; -}; -__sycl_host_pipe_registration __sycl_host_pipe_registrar; - -} // namespace (unnamed) -} // namespace sycl::detail -``` - -The integration footer contains the definition of the constructor, which calls -a function in the DPC++ runtime with the following information for each host -pipe that is used in the translation unit: - -* The (host) address of the static member variable `__pipe`. -* The variable's string from the `sycl-unique-id` attribute. - -``` -namespace sycl::detail { -namespace { - -__sycl_host_pipe_registration::__sycl_host_pipe_registration() noexcept { - host_pipe_map::add(&pipe::__pipe, - /* same string returned from __builtin_sycl_unique_pipe_id(pipe::__pipe) */); - host_pipe_map::add(&inner::pipe::__pipe, - /* same string returned from __builtin_sycl_unique_pipe_id(pipe::__pipe) */); -} - -} // namespace (unnamed) -} // namespace sycl::detail -``` - -Further details on adherence to C++ rules for unconstructed objects can be found -in the [device_global][2] design. - -Unique pipe ids will be generated by the same method as [device_global][2] uses to generate `sycl-unique-id`s. - -### Changes to the DPC++ runtime - -Several changes are needed to the DPC++ runtime - -* As we noted above, the front-end generates new content in the integration - footer which calls the function `sycl::detail::host_pipe_map::add()`. - The runtime defines this function and maintains information about all the - host pipe variables in the application. This information includes: - - - The host address of the variable. - - The string which uniquely identifies the variable. - -* The runtime implements the `read` and `write` functions of the pipe - class. These will use this [host pipe API][4]. These functions will - need to retrieve the mapping added to the __host pipe registrar__ - for the pipe being read or written to, and pass it to the corresponding - underlying OpenCL API call - -[4]: https://github.com/intel-sandbox/ip-authoring-specs/blob/MJ_ChangeDocs4/Pipe/Spec/cl_intel_host_pipe_symbol.asciidoc - -### Changes to the sycl-post-link tool - -As mentioned in the __Attributes attached to the class__ section, the sycl-post-link tool -will generate the `HostAccessINTEL` decoration for each variable declared of a -type marked with the global variable attribute `sycl-host-pipe`. The name operand -should be filled with the id generated by the front end when the `host-pipe` attribute -is encountered. Since there is no current use for specific host access information, -the access field can be set to `1` (read/write). If a use for this information -is found, this can be changed in the future. diff --git a/sycl/doc/design/PropertySets.md b/sycl/doc/design/PropertySets.md index 9c51b7ea86915..1130f2aa35ea6 100644 --- a/sycl/doc/design/PropertySets.md +++ b/sycl/doc/design/PropertySets.md @@ -194,26 +194,6 @@ Set of device requirements for the entire module: See also [OptionalDeviceFeatures.md](OptionalDeviceFeatures.md). -### [SYCL/host pipes] - -__Key:__ Host pipe variable name. - -__Value type:__ Byte array. ("2") - -__Value:__ Information about the host pipe variable with the following -fields: - -```c++ - // Encodes size of the underlying type T of the host pipe variable. - uint32_t Size; -``` - -__Notes:__ - -1. If this property set is missing, the binary does not contain any host pipe -variables. - - ### [SYCL/virtual functions] Set of information about virtual function usage in the module. diff --git a/sycl/doc/design/SYCLPipesLoweringToSPIRV.rst b/sycl/doc/design/SYCLPipesLoweringToSPIRV.rst deleted file mode 100644 index fdeaf611295ae..0000000000000 --- a/sycl/doc/design/SYCLPipesLoweringToSPIRV.rst +++ /dev/null @@ -1,119 +0,0 @@ -SYCL INTEL spatial pipes -======================== - -Introduction -============ - -SPIR-V is first class target in which SYCL pipes should be representable, and -pipes are already exposed within SPIR-V. For this implementation API functions -call for SPIR-V friendly mangled functions instead of OpenCL built-ins. -This document describes how SYCL pipes are being lowered to SPIR-V. - -OpenCL 2.2 program pipe representation in SPIR-V -================================================ - -The SPIR-V program pipe representation is used to be an underlying -representation of intra-kernel and inter-kernel static pipe connectivity. -The SPIR-V pipe representation exists in a series of pieces: - - - OpTypePipeStorage: Type representing memory allocated for storage of data - within a pipe. Used for OpenCL 2.2 program pipes (program-scope pipes) that - the host program is not aware of, but that enables connectivity between - kernels. - - - OpConstantPipeStorage: Instruction that creates an OpTypePipeStorage object. - Requires packet size (number of bytes) and capacity (number of packets) to be - defined. - - - OpTypePipe: A pipe object that can act as a read/write endpoint of some pipe - storage, either allocated by the host and passed as a kernel argument, or - allocated at "program scope" through a pipe storage object. - - - OpCreatePipeFromPipeStorage: Creates a pipe object (that can be read/written) - from an OpTypePipeStorage instance. - - - OpReadPipe / OpWritePipe: Read packet from or write packet to a pipe object. - -Lowering of kernel to kernel pipes to SPIR-V (non-blocking) -=========================================================== - -This connectivity is achieved through OpTypePipeStorage which allows a SPIR-V -device consumer to leverage static connectivity. An OpConstantPipeStorage -instruction must create a single instance of OpPipeStorage for each kernel to -kernel pipe type used by any kernel within the application. - -OpTypePipe objects is created from OpPipeStorage using -OpCreatePipeFromPipeStorage. The number of OpTypePipe objects created from an -OpPipeStorage object is an implementation detail, as are the access qualifiers -applied to those types. For example, an implementation is free to create a -different OpTypePipe corresponding to each read and write, with unidirectional -access qualifiers annotated, or it can create fewer OpTypePipe objects, although -read and write pipes must be distinct according to OpReadPipe and OpWritePipe -rules. - -NOTE: The SPIR-V OpReadPipe and OpWritePipe instructions are non-blocking. - -Details SPIR-V representation in LLVM IR -======================================== - -Pipe built-ins are mangled in LLVM IR to make it SPIR-V friendly. -As an example: - - SPIR-V built-in | Mangled built-in in LLVM IR - ----------------------------+----------------------------------------------- - OpReadPipe | __spirv_ReadPipe - ----------------------------+----------------------------------------------- - OpWritePipe | __spirv_WritePipe - ----------------------------+----------------------------------------------- - OpCreatePipeFromPipeStorage | __spirv_CreatePipeFromPipeStorage_{read|write} - -More about SPIR-V representation in LLVM IR can be found under the link: -.. _SPIRVRepresentationInLLVM.rst: https://github.com/KhronosGroup/SPIRV-LLVM-Translator/blob/master/docs/SPIRVRepresentationInLLVM.rst/ - -In SYCL headers the built-ins are declared as external functions with the -appropriate mangling. The translator will transform calls of these built-ins -into calls of SPIR-V instructions. - -Example of SYCL -> LLVM IR -> SPIR-V -> LLVM-IR code transformations -==================================================================== -Consider following SYCL device code: -.. code:: cpp - pipe::write(42, SuccessCode); - -After compiling this code with clang we will be given following piece of IR for -the write pipe function call (NOTE: for this implementation clang-known -OpenCL 2.0 pipe types are reused): -.. code:: cpp - define internal spir_func void @_ZN2cl4sycl4pipeIZ4mainE9some_pipeiLi1EE5writeEiRb(i32, i8* dereferenceable(1)) #4 align 2 { - //... - %12 = call spir_func i32 @_Z17__spirv_WritePipeIiEi8ocl_pipePT_ii(%opencl.pipe_wo_t addrspace(1)* %10, i32 addrspace(4)* %11, i32 4, i32 4) #8 - //... - } - -with following declaration: -.. code:: cpp - %12 = call spir_func i32 @_Z17__spirv_WritePipeIiEi8ocl_pipePT_ii(%opencl.pipe_wo_t addrspace(1)* %10, i32 addrspace(4)* %11, i32 4, i32 4) #8 - -SPIR-V translator will drop all of these manglings, just making a call of SPIR-V -write pipe built-in: -.. code:: cpp - 7 WritePipe 51 158 156 157 52 52 - -Resulting code for translation back to LLVM IR from SPIR-V are calls of OpenCL -built-ins: -.. code:: cpp - define internal spir_func void @_ZN2cl4sycl4pipeIZ4mainE9some_pipeiLi1EE5writeEiRb(i32, i8*) #0 { - //... - %9 = call spir_func i32 @__write_pipe_2(%opencl.pipe_wo_t addrspace(1)* %6, i8 addrspace(4)* %8, i32 4, i32 4) //... - } - -again with write pipe declaration (but now it's built-in!): -.. code:: cpp - declare spir_func i32 @__write_pipe_2(%opencl.pipe_wo_t addrspace(1)*, i8 addrspace(4)*, i32, i32) #0 - -The first argument in a call of __write_pipe_2 OpenCL built-in is a pipe object, -which is created as a result of SPIR-V built-in call -__spirv_CreatePipeFromPipeStorage_{read|write} which has no OpenCL -representation and therefore stays in IR before and after SPIR-V tool-chain as: -.. code:: cpp - %9 = call spir_func %opencl.pipe_wo_t addrspace(1)* @_Z39__spirv_CreatePipeFromPipeStorage_writeIiE8ocl_pipe11PipeStorage(%struct._ZTS11PipeStorage.PipeStorage* byval align 4 %6) #8 diff --git a/sycl/doc/design/fpga_io_pipes_design.rst b/sycl/doc/design/fpga_io_pipes_design.rst deleted file mode 100644 index 5c67729bbf4ef..0000000000000 --- a/sycl/doc/design/fpga_io_pipes_design.rst +++ /dev/null @@ -1,81 +0,0 @@ -IO pipes design -=============== - -Requirements ------------- - - Device shall be able to distinguish kernel-to-kernel pipes and I/O pipes; - - No changes shall be in SYCL pipe specification; - - I/O pipe namings/IDs are provided by a vendor in a separated header, like: -.. code:: cpp - namespace intelfpga { - template - struct ethernet_pipe_id { - static constexpr unsigned id = ID; - }; - using ethernet_read_pipe = - sycl::intel::kernel_readable_io_pipe, int, 0>; - using ethernet_write_pipe = - sycl::intel::kernel_writeable_io_pipe, int, 0>; - } - - Thus, the user interacts only with vendor-defined pipe objects. - -Links ------ -.. _Spec: ../extensions/supported/sycl_ext_intel_dataflow_pipes.asciidoc -.. _Interesting comment from Ronan: https://github.com/intel/llvm/pull/635#discussion_r325851766 - -Summary -------- -This document describes a design of I/O pipes implementation in SYCL compiler. -It includes changes in SYCL library, clang and SPIR-V to LLVM IR translator. -It adds extra attribute '__attribute__((io_pipe_id(ID)))' that generates -metadata, attached to a pipe storage declaration. By this metadata a backend can -recognize I/O pipe and distinguish different I/O pipes from each other. - -There is another notable solution which was proposed by Ronan (see the link -above): don't make any compiler/library changes - just make backend recognizing -I/O pipe by demangling it's name. This proposal wasn't picked, because it will -make backend support of the feature more difficult. So far we already have -two devices' backends that support SYCL_INTEL_data_flow_pipes extension -(Intel FPGA HW and Intel FPGA emulator) and in the future this number may -increase. So efforts put in compiler implementation shall be payed off even -more. - -clang ------ -Need to implement additional attribute, that possible to attach to pipe storage -declaration. The attribute shall accept a compile-time known integer argument -(the pipe ID). With the attribute applied, clang generates a metadata attached -the to pipe storage declaration, that contains the I/O pipe ID (argument). - -llvm-spirv translator ---------------------- -Need to implement additional decoration, that saves the I/O pipe ID information, -that can be collected from a metadata attached to the pipe storage object. - -SYCL implementation in headers ------------------------------- -Following the spec, we need to add two more classes for pipes: - - 'kernel_readable_io_pipe' - - 'kernel_writeable_io_pipe' - - with the same member functions and fields as it is already done for pipe class. - -The attribute should be attached to a pipe storage declaration in the headers -and it would be looking like: -.. code:: cpp - static constexpr int32_t ID = name::id; - static constexpr struct ConstantPipeStorage - m_Storage __attribute__((io_pipe_id(ID))) = {m_Size, m_Alignment, m_Capacity}; - - - where 'name' is some class used in pipe type construction, like: -.. code:: cpp - using pipe_type = pipe; - -When specific io_pipe_def.h is included in a user's code and 'name' is mapped to -the pipe name defined in this header (for example 'ethernet_pipe_id' structure -defined above) 'name::id' returns the actual I/O pipe ID (compile-time known -integer constant) that is passed as the attribute's argument and used to -identify the I/O pipe in some backend. diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc index c700285da6a89..57dc70a9f1445 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_global.asciidoc @@ -646,24 +646,6 @@ struct host_access_key { std::integral_constant>; }; -enum class init_mode_enum : /* unspecified */ { - reprogram, - reset -}; - -struct init_mode_key { - template - using value_t = - property_value>; -}; - -struct implement_in_csr_key { - template - using value_t = - property_value>; -}; - inline constexpr device_image_scope_key::value_t device_image_scope; template @@ -677,29 +659,12 @@ inline constexpr host_access_key::value_t inline constexpr host_access_key::value_t host_access_none; -template -inline constexpr init_mode_key::value_t init_mode; -inline constexpr init_mode_key::value_t - init_mode_reprogram; -inline constexpr init_mode_key::value_t init_mode_reset; - -template -inline constexpr implement_in_csr_key::value_t implement_in_csr; -inline constexpr implement_in_csr_key::value_t implement_in_csr_on; -inline constexpr implement_in_csr_key::value_t implement_in_csr_off; - template struct is_property_key_of> : std::true_type {}; template struct is_property_key_of> : std::true_type {}; -template -struct is_property_key_of> - : std::true_type {}; -template -struct is_property_key_of> - : std::true_type {}; } // namespace sycl::ext::oneapi::experimental ---- @@ -714,8 +679,7 @@ a| device_image_scope ---- a| -This property is most useful for kernels that are submitted to an FPGA device, -but it may be used with any kernel. Normally, a single instance of a device +Normally, a single instance of a device global variable is allocated for each device, and that instance is shared by all kernels that belong to the same context and are submitted to the same device, regardless of which _device image_ contains the kernel. @@ -723,7 +687,7 @@ When this property is specified, it is an assertion by the user that on a given device a given device_global decorated with this property is only ever accessed in a single _device_image_. An implementation may be able to optimize accesses to the device global when this -property is specified (especially on an FPGA device), but the user must be aware +property is specified, but the user must be aware of which _device image_ contains the kernels that use the variable. A device global that is decorated with this property may not be accessed from @@ -759,58 +723,18 @@ host_access a| This property provides an assertion by the user telling the implementation whether the host code copies to or from the device global. As a result, the -implementation may be able to perform certain optimizations. Although this -property may be used with any device, it is generally only beneficial when used -on FPGA devices. +implementation may be able to perform certain optimizations. The following values are supported: * `read`: The user asserts that the host code may copy from (read) the - variable, but it will never copy to (write) it. For an FPGA device, only a - read port is exposed. + variable, but it will never copy to (write) it. * `write`: The user asserts that the host code may copy to (write) the - variable, but it never copy from (read) it. For an FPGA device, only a write - port is exposed. + variable, but it never copy from (read) it. * `none`: The user asserts that the host code will never copy to or copy - from the variable. For an FPGA device, no external ports are exposed. + from the variable. * `read_write`: The user provides no assertions, and the host code may either - copy to or copy from the variable. This is the default. For an FPGA device, - a read/write port is exposed. - -a| -[source,c++] ----- -init_mode ----- -a| -This property is only meaningful when used with an FPGA device. It is ignored -for other devices. The following values are supported: - -* `reprogram`: Initialization is performed by reprogramming the device. This - may require more frequent reprogramming but may reduce area. -* `reset`: Initialization is performed by sending a reset signal to the device. - This may increase area but may reduce reprogramming frequency. - -If the `init_mode` property is not specified, the default behavior is -equivalent to one of the values listed above, but the choice is implementation -defined. - -a| -[source,c++] ----- -implement_in_csr ----- -a| -This property is only meaningful when used with an FPGA device. It is ignored -for other devices. The following values are supported: - -* `true`: Access to this memory is done through a CSR interface shared with - kernel arguments. -* `false`: Access to this memory is done through a dedicated interface. - -If the `implement_in_csr` property is not specified, the default behavior is -equivalent to one of the values listed above, but the choice is implementation -defined. + copy to or copy from the variable. This is the default. |=== @@ -848,7 +772,7 @@ _device images_: * A kernel that uses specialization constants may have a new instance in a new _device image_ each time the application sets a new value for the specialization constant. However, this happens only if the device supports - native specialization constants, which is not the case for FPGA devices. + native specialization constants. ==== === Relax language restrictions for SYCL device functions diff --git a/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_kernel_arg_properties.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_kernel_arg_properties.asciidoc deleted file mode 100644 index 6139647e9f9cb..0000000000000 --- a/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_kernel_arg_properties.asciidoc +++ /dev/null @@ -1,526 +0,0 @@ -= sycl_ext_intel_fpga_kernel_arg_properties - -:source-highlighter: coderay -:coderay-linenums-mode: table - -// This section needs to be after the document title. -:doctype: book -:toc2: -:toc: left -:encoding: utf-8 -:lang: en - -:blank: pass:[ +] - -// Set the default source code type in this document to C++, -// for syntax highlighting purposes. This is needed because -// docbook uses c++ and html5 uses cpp. -:language: {basebackend@docbook:c++:cpp} - -// This is necessary for asciidoc, but not for asciidoctor -:cpp: C++ -:dpcpp: DPC++ - -== Notice - -[%hardbreaks] -Copyright (C) 2022 Intel Corporation. All rights reserved. - -Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks -of The Khronos Group Inc. OpenCL(TM) is a trademark of Apple Inc. used by -permission by Khronos. - -== Contact - -To report problems with this extension, please open a new issue at: - -https://github.com/intel/llvm/issues - -== Contributors - -Abhishek Tiwari, Intel + -Joseph Garvey, Intel - - -== Dependencies - -This extension is written against the SYCL 2020 specification, revision 5. - -It depends on the following extensions: - - - link:../experimental/sycl_ext_oneapi_properties.asciidoc[sycl_ext_oneapi_properties] - - link:sycl_ext_oneapi_annotated_ptr.asciidoc[sycl_ext_oneapi_annotated_ptr] - - link:sycl_ext_oneapi_annotated_arg.asciidoc[sycl_ext_oneapi_annotated_arg] - -== Status - -This is a proposed extension specification, intended to gather community -feedback. Interfaces defined in this specification may not be implemented yet -or may be in a preliminary state. The specification itself may also change in -incompatible ways before it is finalized. *Shipping software products should -not rely on APIs defined in this specification.* - -== Overview - -This extension introduces properties for the classes -`sycl::ext::oneapi::experimental::annotated_ptr` and -`sycl::ext::oneapi::experimental::annotated_arg`. The properties will influence -the kernel argument interfaces for FPGA kernels and can be ignored for other -devices. - -Some examples of the syntax are shown below. - -[source,c++] ----- -auto ptr = ... -annotated_ptr arg_a{ - ptr, properties{buffer_location<1>, awidth<32>, dwidth<64>}}; - -... - -auto ptr2 = ... -annotated_arg arg_b{ - ptr2, properties{buffer_location<2>, awidth<32>, dwidth<64>}}; - -... - -int val = 5; -annotated_arg arg_c{val, properties{register_map, stable}}; - -... - -annotated_ptr, awidth<32>, - dwidth<64>> arg_d; - -annotated_arg> arg_e; - -annotated_arg, awidth<32>, - dwidth<64>> arg_f; - - ----- - - -== Specification - -=== Feature test macro - -This extension provides a feature-test macro as described in the core SYCL -specification. An implementation supporting this extension must predefine the -macro `SYCL_EXT_INTEL_FPGA_KERNEL_ARG_PROPERTIES` to one of the values -defined in the table below. Applications can test for the existence of this -macro to determine if the implementation supports this feature, or applications -can test the macro's value to determine which of the extension's features the -implementation supports. - -[%header,cols="1,5"] -|=== -|Value -|Description - -|1 -|Initial version of this extension. -|=== - -=== `annotated_ptr` and `annotated_arg` Properties - -Below is a list of compile-time constant properties supported by -the `annotated_ptr` and `annotated_arg` classes. These properties control the -kernel argument interface on FPGA devices. The properties are allowed even on -kernels that are submitted to other devices, but they are silently ignored when -the kernel is submitted to a non-FPGA device. - -All of the properties defined in this extension are meaningful only on the -kernel argument and are not meaningful within the kernel body. - -The properties that are defined for pointers will be meaningful for -`annotated_arg` only when T is a pointer type. - -```c++ -namespace sycl::ext::intel::experimental { -struct conduit_key { - using value_t = sycl::ext::oneapi::experimental::properties::property_value< - conduit_key>; -}; - -struct register_map_key { - using value_t = sycl::ext::oneapi::experimental::properties::property_value< - register_map_key>; -}; - -struct stable_key { - using value_t = sycl::ext::oneapi::experimental::properties::property_value< - stable_key>; -}; - -struct buffer_location_key { - template - using value_t = sycl::ext::oneapi::experimental::properties::property_value< - buffer_location_key, std::integral_constant>; -}; - -struct awidth_key { - template - using value_t = sycl::ext::oneapi::experimental::properties::property_value< - awidth_key, std::integral_constant>; -}; - -struct dwidth_key { - template - using value_t = sycl::ext::oneapi::experimental::properties::property_value< - dwidth_key, std::integral_constant>; -}; - -enum class read_write_mode_enum { - read, - write, - read_write -}; - -struct read_write_mode_key { - template - using value_t = sycl::ext::oneapi::experimental::properties::property_value< - read_write_mode_key, std::integral_constant>; -}; - -struct latency_key { - template - using value_t = sycl::ext::oneapi::experimental::properties::property_value< - latency_key, std::integral_constant>; -}; - -struct maxburst_key { - template - using value_t = sycl::ext::oneapi::experimental::properties::property_value< - maxburst_key, std::integral_constant>; -}; - -struct wait_request_key { - template - using value_t = sycl::ext::oneapi::experimental::properties::property_value< - wait_request_key, std::integral_constant>; -}; - -inline constexpr conduit_key::value_t - conduit; -inline constexpr register_map_key::value_t - register_map; -inline constexpr stable_key::value_t stable; -template inline constexpr buffer_location_key::value_t - buffer_location; -template inline constexpr awidth_key::value_t - awidth; -template inline constexpr dwidth_key::value_t - dwidth; -template -inline constexpr read_write_mode_key::value_t - read_write_mode; -inline constexpr read_write_mode_key::value_t< - read_write_mode_enum::read> read_write_mode_read; -inline constexpr read_write_mode_key::value_t< - read_write_mode_enum::write> read_write_mode_write; -inline constexpr read_write_mode_key::value_t< - read_write_mode_enum::read_write> - read_write_mode_readwrite; -template inline constexpr latency_key::value_t - latency; -template inline constexpr maxburst_key::value_t - maxburst; -template inline constexpr wait_request_key::value_t - wait_request; -inline constexpr wait_request_key::value_t - wait_request_requested; -inline constexpr wait_request_key::value_t - wait_request_not_requested; -} // namespace sycl::ext::intel::experimental - -// Type trait specializations -namespace sycl::ext::oneapi::experimental { -template -struct is_property_key_of< - sycl::ext::intel::experimental::conduit_key, - annotated_ptr> : std::true_type {}; -template -struct is_property_key_of< - sycl::ext::intel::experimental::register_map_key, - annotated_ptr> : std::true_type {}; -template -struct is_property_key_of< - sycl::ext::intel::experimental::stable_key, - annotated_ptr> : std::true_type {}; -template -struct is_property_key_of< - sycl::ext::intel::experimental::buffer_location_key, - annotated_ptr> : std::true_type {}; -template -struct is_property_key_of< - sycl::ext::intel::experimental::awidth_key, - annotated_ptr> : std::true_type {}; -template -struct is_property_key_of< - sycl::ext::intel::experimental::dwidth_key, - annotated_ptr> : std::true_type {}; -template -struct is_property_key_of< - sycl::ext::intel::experimental::read_write_mode_key, - annotated_ptr> : std::true_type {}; -template -struct is_property_key_of< - sycl::ext::intel::experimental::latency_key, - annotated_ptr> : std::true_type {}; -template -struct is_property_key_of< - sycl::ext::intel::experimental::maxburst_key, - annotated_ptr> : std::true_type {}; -template -struct is_property_key_of< - sycl::ext::intel::experimental::wait_request_key, - annotated_ptr> : std::true_type {}; - -template -struct is_property_key_of< - sycl::ext::intel::experimental::conduit_key, - annotated_arg> : std::true_type {}; -template -struct is_property_key_of< - sycl::ext::intel::experimental::register_map_key, - annotated_arg> : std::true_type {}; -template -struct is_property_key_of< - sycl::ext::intel::experimental::stable_key, - annotated_arg> : std::true_type {}; -template -struct is_property_key_of< - sycl::ext::intel::experimental::buffer_location_key, - annotated_arg> : std::true_type {}; -template -struct is_property_key_of< - sycl::ext::intel::experimental::awidth_key, - annotated_arg> : std::true_type {}; -template -struct is_property_key_of< - sycl::ext::intel::experimental::dwidth_key, - annotated_arg> : std::true_type {}; -template -struct is_property_key_of< - sycl::ext::intel::experimental::read_write_mode_key, - annotated_arg> : std::true_type {}; -template -struct is_property_key_of< - sycl::ext::intel::experimental::latency_key, - annotated_arg> : std::true_type {}; -template -struct is_property_key_of< - sycl::ext::intel::experimental::maxburst_key, - annotated_arg> : std::true_type {}; -template -struct is_property_key_of< - sycl::ext::intel::experimental::wait_request_key, - annotated_arg> : std::true_type {}; -} // namespace sycl::ext::oneapi::experimental -``` --- - -[frame="topbot",options="header"] -|=== -|Property |Description - -a| -[source,c++] ----- -conduit ----- -a| -Directs the compiler to create a dedicated input port on the kernel for the -input. - -a| -[source,c++] ----- -register_map ----- -a| -Directs the compiler to create a register to store the input as opposed to -creating a dedicated input port on the kernel. - -a| -[source,c++] ----- -stable ----- -a| -While the SYCL software model makes kernel arguments read-only, the IP which is -output by the FPGA device compiler can be plugged into external systems where -kernel arguments can change while the kernel executes. - -This property specifies that the input to the kernel will not change between -pipelined invocations of the kernel. The input can still change after all active -kernel invocations have finished. - -If the input is changed while the pipelined kernel invocations are executing, -the behavior is undefined. - -a| -[source,c++] ----- -buffer_location ----- -a| -Specifies a global memory identifier for the pointer interface. - -This property is only meaningful on pointer kernel arguments. - -a| -[source,c++] ----- -awidth ----- -a| -Specifies the width of the memory-mapped address bus in bits. The default is -determined by the implementation. - -This property is only meaningful for pointer kernel arguments and only -when the `buffer_location` property is specified. - -a| -[source,c++] ----- -dwidth ----- -a| -Specifies the width of the memory-mapped data bus in bits. The default is set -to 64. - -This property is only meaningful for pointer kernel arguments and only -when the `buffer_location` property is specified. - -a| -[source,c++] ----- -read_write_mode ----- -a| -Specifies the port direction of the memory interface associated with the input -pointer. `mode` can be one of: - -`read_write` - Interface can be used for read and write operations. - -`read` - Interface can only be used for read operations. - -`write` - Interface can only be used for write operations. - -The default is set to `read_write`. - -For convenience, the following are provided: - - - read_write_mode_read - - read_write_mode_write - - read_write_mode_readwrite - -This property is only meaningful for pointer kernel arguments and only -when the `buffer_location` property is specified. - -a| -[source,c++] ----- -latency ----- -a| -Specifies the guaranteed latency in cycles, from when a read command exits -the kernel to when the external memory returns valid read data. The default -is set to 1. - -A value of 0 specifies a variable latency and a positive value specifies a -fixed latency. - -This property is only meaningful for pointer kernel arguments and only -when the `buffer_location` property is specified. - -a| -[source,c++] ----- -maxburst ----- -a| -Specifies the maximum number of data transfers that can be associated with a -read or write transaction. The default is set to 1. - -This property is only meaningful for pointer kernel arguments and only -when the `buffer_location` property is specified. - -a| -[source,c++] ----- -wait_request ----- -a| -Specifies whether the 'wait request' signal is generated or not. This signal is -asserted by the memory system when it is unable to respond to a read or write -request. The default is set to `false`. - -For convenience, the following are provided: - - - wait_request_requested - - wait_request_not_requested - -This property is only meaningful for pointer kernel arguments and only -when the `buffer_location` property is specified. -|=== --- - -=== Usage Examples - -The example below shows a simple kernel with one `annotated_ptr` kernel -argument and one `annotated_arg` kernel argument. - -.Usage Example -```c++ -using sycl::ext::intel::experimental; -{ - sycl::queue q{...}; - - // Allocate memory - auto ptr_a = ... - constexpr int kN = 10; - - // Add properties - auto arg_a = annotated_ptr(ptr_a, properties{ - register_map, buffer_location<1>, awidth<18>, dwidth<64>}); - auto arg_n = annotated_arg(kN, properties{register_map, stable}); - - q.single_task([=] { - for (int i=0; i A, int32_t MA, int32_t RoundingMode = 0, int32_t RoundingAccuracy = 0) noexcept; -template -extern __DPCPP_SYCL_EXTERNAL int32_t -__spirv_ReadPipe(__ocl_RPipeTy Pipe, dataT *Data, int32_t Size, - int32_t Alignment) noexcept; -template -extern __DPCPP_SYCL_EXTERNAL int32_t -__spirv_WritePipe(__ocl_WPipeTy Pipe, const dataT *Data, int32_t Size, - int32_t Alignment) noexcept; -template -extern __DPCPP_SYCL_EXTERNAL void -__spirv_ReadPipeBlockingINTEL(__ocl_RPipeTy Pipe, dataT *Data, - int32_t Size, int32_t Alignment) noexcept; -template -extern __DPCPP_SYCL_EXTERNAL void -__spirv_WritePipeBlockingINTEL(__ocl_WPipeTy Pipe, const dataT *Data, - int32_t Size, int32_t Alignment) noexcept; -template -extern __DPCPP_SYCL_EXTERNAL __ocl_RPipeTy -__spirv_CreatePipeFromPipeStorage_read( - const ConstantPipeStorage *Storage) noexcept; -template -extern __DPCPP_SYCL_EXTERNAL __ocl_WPipeTy -__spirv_CreatePipeFromPipeStorage_write( - const ConstantPipeStorage *Storage) noexcept; - extern __DPCPP_SYCL_EXTERNAL float __spirv_ConvertBF16ToFINTEL(uint16_t) noexcept; extern __DPCPP_SYCL_EXTERNAL uint16_t diff --git a/sycl/include/sycl/__spirv/spirv_types.hpp b/sycl/include/sycl/__spirv/spirv_types.hpp index 1aeb4cdd88c50..23a905a0b293d 100644 --- a/sycl/include/sycl/__spirv/spirv_types.hpp +++ b/sycl/include/sycl/__spirv/spirv_types.hpp @@ -142,24 +142,10 @@ struct __spirv_TaskSequenceINTEL; } // namespace __spv #ifdef __SYCL_DEVICE_ONLY__ -// OpenCL pipe types -template -using __ocl_RPipeTy = __attribute__((pipe("read_only"))) const dataT; -template -using __ocl_WPipeTy = __attribute__((pipe("write_only"))) const dataT; - // OpenCL vector types template using __ocl_vec_t = dataT __attribute__((ext_vector_type(dims))); -// Struct representing layout of pipe storage -// TODO: rename to __spirv_ConstantPipeStorage -struct ConstantPipeStorage { - int32_t _PacketSize; - int32_t _PacketAlignment; - int32_t _Capacity; -}; - namespace sycl { inline namespace _V1 { namespace detail { diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp index ad938b2685e07..0069b4801a1fb 100644 --- a/sycl/include/sycl/accessor.hpp +++ b/sycl/include/sycl/accessor.hpp @@ -1163,7 +1163,6 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : const property_list &PropertyList = {}, const detail::code_location CodeLoc = detail::code_location::current()) : accessor(BufferRef, PropertyList, CodeLoc) { - adjustAccPropsInBuf(BufferRef); } template AccessOffset, TagT, const property_list &PropertyList = {}, const detail::code_location CodeLoc = detail::code_location::current()) : accessor(BufferRef, AccessRange, AccessOffset, PropertyList, CodeLoc) { - adjustAccPropsInBuf(BufferRef); } template @@ -1935,26 +1923,6 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : "accessor cannot be both read_only and no_init"); } } - - template - void adjustAccPropsInBuf(BufT &Buffer) { - if constexpr (PropertyListT::template has_property< - sycl::ext::intel::property::buffer_location>()) { - auto location = (PropertyListT::template get_property< - sycl::ext::intel::property::buffer_location>()) - .get_location(); - property_list PropList{ - sycl::property::buffer::detail::buffer_location(location)}; - Buffer.addOrReplaceAccessorProperties(PropList); - } else { - deleteAccPropsFromBuf(Buffer); - } - } - - template void deleteAccPropsFromBuf(BufT &Buffer) { - Buffer.deleteAccProps( - sycl::detail::PropWithDataKind::AccPropBufferLocation); - } }; template diff --git a/sycl/include/sycl/buffer.hpp b/sycl/include/sycl/buffer.hpp index da1814996c807..c94f6091ae95c 100644 --- a/sycl/include/sycl/buffer.hpp +++ b/sycl/include/sycl/buffer.hpp @@ -152,10 +152,6 @@ class __SYCL_EXPORT buffer_plain { const std::unique_ptr &get_allocator_internal() const; - void deleteAccProps(const sycl::detail::PropWithDataKind &Kind); - - void addOrReplaceAccessorProperties(const property_list &PropertyList); - size_t getSize() const; void handleRelease() const; @@ -769,14 +765,6 @@ class buffer : public detail::buffer_plain, dimensions, sizeof(T), detail::rangeToArray(Range).data()); } - void addOrReplaceAccessorProperties(const property_list &PropertyList) { - buffer_plain::addOrReplaceAccessorProperties(PropertyList); - } - - void deleteAccProps(const sycl::detail::PropWithDataKind &Kind) { - buffer_plain::deleteAccProps(Kind); - } - // Reinterpret contructor buffer(const std::shared_ptr &Impl, range reinterpretRange, size_t reinterpretOffset, diff --git a/sycl/include/sycl/detail/cg_types.hpp b/sycl/include/sycl/detail/cg_types.hpp index 6125cf327e03c..2dbcf86aa7690 100644 --- a/sycl/include/sycl/detail/cg_types.hpp +++ b/sycl/include/sycl/detail/cg_types.hpp @@ -58,15 +58,14 @@ enum class CGType : unsigned int { Memset2DUSM = 18, CopyToDeviceGlobal = 19, CopyFromDeviceGlobal = 20, - ReadWriteHostPipe = 21, - ExecCommandBuffer = 22, - CopyImage = 23, - SemaphoreWait = 24, - SemaphoreSignal = 25, - ProfilingTag = 26, - EnqueueNativeCommand = 27, - AsyncAlloc = 28, - AsyncFree = 29, + ExecCommandBuffer = 21, + CopyImage = 22, + SemaphoreWait = 23, + SemaphoreSignal = 24, + ProfilingTag = 25, + EnqueueNativeCommand = 26, + AsyncAlloc = 27, + AsyncFree = 28, }; template struct check_fn_signature { diff --git a/sycl/include/sycl/detail/host_pipe_map.hpp b/sycl/include/sycl/detail/host_pipe_map.hpp deleted file mode 100644 index cfbb75f2c004a..0000000000000 --- a/sycl/include/sycl/detail/host_pipe_map.hpp +++ /dev/null @@ -1,23 +0,0 @@ -//==-------------------- host_pipe_map.hpp -----------------------------==// -// -// 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 - -namespace sycl { -inline namespace _V1 { -namespace detail { -namespace host_pipe_map { - -__SYCL_EXPORT void add(const void *HostPipePtr, const char *UniqueId); - -} // namespace host_pipe_map -} // namespace detail -} // namespace _V1 -} // namespace sycl diff --git a/sycl/include/sycl/detail/property_list_base.hpp b/sycl/include/sycl/detail/property_list_base.hpp index 59bc13b978fcb..9be47ffbec843 100644 --- a/sycl/include/sycl/detail/property_list_base.hpp +++ b/sycl/include/sycl/detail/property_list_base.hpp @@ -103,30 +103,6 @@ class PropertyListBase { "The property is not found"); } - void add_or_replace_accessor_properties_helper( - const std::vector> &PropsWithData) { - for (auto &Prop : PropsWithData) { - if (Prop->isSame(sycl::detail::PropWithDataKind::AccPropBufferLocation)) { - delete_accessor_property_helper( - sycl::detail::PropWithDataKind::AccPropBufferLocation); - MPropsWithData.push_back(Prop); - break; - } - } - } - - void delete_accessor_property_helper(const PropWithDataKind &Kind) { - auto It = MPropsWithData.begin(); - for (; It != MPropsWithData.end(); ++It) { - if ((*It)->isSame(Kind)) - break; - } - if (It != MPropsWithData.end()) { - std::iter_swap(It, MPropsWithData.end() - 1); - MPropsWithData.pop_back(); - } - } - void checkPropsAndThrow(std::function FunctionForDataless, std::function FunctionForData) const { static const auto ErrorCode = sycl::make_error_code(errc::invalid); diff --git a/sycl/include/sycl/ext/intel/esimd/memory_properties.hpp b/sycl/include/sycl/ext/intel/esimd/memory_properties.hpp index da9ca3d5506d3..f54a46b6fa922 100644 --- a/sycl/include/sycl/ext/intel/esimd/memory_properties.hpp +++ b/sycl/include/sycl/ext/intel/esimd/memory_properties.hpp @@ -9,7 +9,6 @@ #pragma once #include -#include #include #include #include diff --git a/sycl/include/sycl/ext/intel/experimental/fpga_annotated_properties.hpp b/sycl/include/sycl/ext/intel/experimental/fpga_annotated_properties.hpp deleted file mode 100644 index bbf55d469809a..0000000000000 --- a/sycl/include/sycl/ext/intel/experimental/fpga_annotated_properties.hpp +++ /dev/null @@ -1,349 +0,0 @@ -//==-- fpga_annotated_properties.hpp - SYCL properties associated with -// annotated_arg/ptr --==// -// -// 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 -#include -#include - -#include -#include -#include -#include - -namespace sycl { -inline namespace _V1 { -namespace ext { -namespace intel { -namespace experimental { - -template -using property_value = - sycl::ext::oneapi::experimental::property_value; -//===----------------------------------------------------------------------===// -// FPGA properties of annotated_arg/annotated_ptr -//===----------------------------------------------------------------------===// -struct register_map_key - : oneapi::experimental::detail::compile_time_property_key< - oneapi::experimental::detail::PropKind::RegisterMap> { - using value_t = property_value; -}; - -struct conduit_key : oneapi::experimental::detail::compile_time_property_key< - oneapi::experimental::detail::PropKind::Conduit> { - using value_t = property_value; -}; - -struct stable_key : oneapi::experimental::detail::compile_time_property_key< - oneapi::experimental::detail::PropKind::Stable> { - using value_t = property_value; -}; - -struct buffer_location_key - : oneapi::experimental::detail::compile_time_property_key< - oneapi::experimental::detail::PropKind::BufferLocation> { - template - using value_t = - property_value>; -}; - -struct awidth_key : oneapi::experimental::detail::compile_time_property_key< - oneapi::experimental::detail::PropKind::AddrWidth> { - template - using value_t = property_value>; -}; - -struct dwidth_key : oneapi::experimental::detail::compile_time_property_key< - oneapi::experimental::detail::PropKind::DataWidth> { - template - using value_t = property_value>; -}; - -struct latency_key : oneapi::experimental::detail::compile_time_property_key< - oneapi::experimental::detail::PropKind::Latency> { - template - using value_t = property_value>; -}; - -enum class read_write_mode_enum : std::uint16_t { read, write, read_write }; - -struct read_write_mode_key - : oneapi::experimental::detail::compile_time_property_key< - oneapi::experimental::detail::PropKind::RWMode> { - template - using value_t = - property_value>; -}; - -struct maxburst_key : oneapi::experimental::detail::compile_time_property_key< - oneapi::experimental::detail::PropKind::MaxBurst> { - template - using value_t = property_value>; -}; - -struct wait_request_key - : oneapi::experimental::detail::compile_time_property_key< - oneapi::experimental::detail::PropKind::WaitRequest> { - template - using value_t = - property_value>; -}; - -// non-mmhost properties -inline constexpr register_map_key::value_t register_map; -inline constexpr conduit_key::value_t conduit; -inline constexpr stable_key::value_t stable; - -// mmhost properties -template -inline constexpr buffer_location_key::value_t buffer_location; -template inline constexpr awidth_key::value_t awidth; -template inline constexpr dwidth_key::value_t dwidth; -template inline constexpr latency_key::value_t latency; -template inline constexpr maxburst_key::value_t maxburst; -template -inline constexpr wait_request_key::value_t wait_request; -inline constexpr wait_request_key::value_t<1> wait_request_requested; -inline constexpr wait_request_key::value_t<0> wait_request_not_requested; - -template -inline constexpr read_write_mode_key::value_t read_write_mode; -inline constexpr read_write_mode_key::value_t - read_write_mode_read; -inline constexpr read_write_mode_key::value_t - read_write_mode_write; -inline constexpr read_write_mode_key::value_t - read_write_mode_readwrite; - -} // namespace experimental -} // namespace intel - -namespace oneapi { -namespace experimental { -template class annotated_arg; -template class annotated_ptr; - -struct alignment_key; -using register_map_key = intel::experimental::register_map_key; -using conduit_key = intel::experimental::conduit_key; -using stable_key = intel::experimental::stable_key; -using buffer_location_key = intel::experimental::buffer_location_key; -using awidth_key = intel::experimental::awidth_key; -using dwidth_key = intel::experimental::dwidth_key; -using latency_key = intel::experimental::latency_key; -using read_write_mode_key = intel::experimental::read_write_mode_key; -using maxburst_key = intel::experimental::maxburst_key; -using wait_request_key = intel::experimental::wait_request_key; -using read_write_mode_enum = intel::experimental::read_write_mode_enum; - -template -struct is_property_key_of> - : std::true_type {}; - -template -struct is_property_key_of> - : std::true_type {}; - -template -struct is_property_key_of> - : std::true_type {}; - -template -struct is_property_key_of> - : std::true_type {}; - -template -struct is_property_key_of> - : std::true_type {}; - -template -struct is_property_key_of> - : std::true_type {}; - -template -struct is_property_key_of> - : std::true_type {}; - -template -struct is_property_key_of> - : std::true_type {}; - -template -struct is_property_key_of> - : std::true_type {}; - -template -struct is_property_key_of> - : std::true_type {}; - -template -struct is_property_key_of> - : std::true_type {}; - -template -struct is_property_key_of> - : std::true_type {}; - -template -struct is_property_key_of> - : std::true_type {}; - -template -struct is_property_key_of> - : std::true_type {}; - -template -struct is_property_key_of> - : std::true_type {}; - -template -struct is_property_key_of> - : std::true_type {}; - -template -struct is_property_key_of> - : std::true_type {}; - -template -struct is_property_key_of> - : std::true_type {}; - -template -struct is_property_key_of> - : std::true_type {}; - -template -struct is_property_key_of> - : std::true_type {}; - -namespace detail { -template <> struct PropertyMetaInfo { - static constexpr const char *name = "sycl-register-map"; - static constexpr std::nullptr_t value = nullptr; -}; -template <> struct PropertyMetaInfo { - static constexpr const char *name = "sycl-conduit"; - static constexpr std::nullptr_t value = nullptr; -}; -template <> struct PropertyMetaInfo { - static constexpr const char *name = "sycl-stable"; - static constexpr std::nullptr_t value = nullptr; -}; - -template struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-buffer-location"; - static constexpr int value = N; -}; -template struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-awidth"; - static constexpr int value = W; -}; -template struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-dwidth"; - static constexpr int value = W; -}; -template struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-latency"; - static constexpr int value = N; -}; -template struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-maxburst"; - static constexpr int value = N; -}; -template -struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-wait-request"; - static constexpr int value = Enable; -}; -template -struct PropertyMetaInfo> { - static constexpr const char *name = "sycl-read-write-mode"; - static constexpr read_write_mode_enum value = Mode; -}; - -} // namespace detail - -// 'buffer_location' and mmhost properties are pointers-only -template -struct is_valid_property> - : std::bool_constant> {}; - -template -struct is_valid_property> - : std::bool_constant> {}; - -template -struct is_valid_property> - : std::bool_constant> {}; - -template -struct is_valid_property> - : std::bool_constant> {}; - -template -struct is_valid_property> - : std::bool_constant> {}; - -template -struct is_valid_property> - : std::bool_constant> {}; - -template -struct is_valid_property> - : std::bool_constant> {}; - -// 'register_map', 'conduit', 'stable' are common properties for pointers -// and non pointers; -template -struct is_valid_property : std::true_type {}; -template -struct is_valid_property : std::true_type {}; -template -struct is_valid_property : std::true_type {}; - -// buffer_location is applied on PtrAnnotation -template <> -struct propagateToPtrAnnotation : std::true_type {}; - -//===----------------------------------------------------------------------===// -// Utility for FPGA properties -//===----------------------------------------------------------------------===// -// -namespace detail { -template struct checkValidFPGAPropertySet { - template - static constexpr bool has_one_of = - ((property_list_t::template has_property() || ...)); - - static constexpr bool has_BufferLocation = has_one_of; - - static constexpr bool has_InterfaceConfig = - has_one_of; - - static constexpr bool value = !(!has_BufferLocation && has_InterfaceConfig); -}; - -template struct checkHasConduitAndRegisterMap { - static constexpr bool has_Conduit = - property_list_t::template has_property(); - static constexpr bool has_RegisterMap = - property_list_t::template has_property(); - static constexpr bool value = !(has_Conduit && has_RegisterMap); -}; -} // namespace detail - -} // namespace experimental -} // namespace oneapi -} // namespace ext -} // namespace _V1 -} // namespace sycl diff --git a/sycl/include/sycl/ext/intel/experimental/fpga_kernel_properties.hpp b/sycl/include/sycl/ext/intel/experimental/fpga_kernel_properties.hpp deleted file mode 100644 index e60989acdb835..0000000000000 --- a/sycl/include/sycl/ext/intel/experimental/fpga_kernel_properties.hpp +++ /dev/null @@ -1,193 +0,0 @@ -//===--------------------- fpga_kernel_properties.hpp ---------------------===// -// SYCL properties associated with FPGA kernel properties -// -// 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 - -#include -#include - -namespace sycl { -inline namespace _V1 { -namespace ext::intel::experimental { - -template class fpga_kernel_attribute; -template class task_sequence; - -enum class streaming_interface_options_enum : uint16_t { - accept_downstream_stall, - remove_downstream_stall -}; - -enum class register_map_interface_options_enum : uint16_t { - do_not_wait_for_done_write, - wait_for_done_write, -}; - -enum class fpga_cluster_options_enum : std::uint16_t { - stall_free, - stall_enable -}; - -struct streaming_interface_key - : oneapi::experimental::detail::compile_time_property_key< - oneapi::experimental::detail::PropKind::StreamingInterface> { - template - using value_t = ext::oneapi::experimental::property_value< - streaming_interface_key, - std::integral_constant>; -}; - -struct register_map_interface_key - : oneapi::experimental::detail::compile_time_property_key< - oneapi::experimental::detail::PropKind::RegisterMapInterface> { - template - using value_t = ext::oneapi::experimental::property_value< - register_map_interface_key, - std::integral_constant>; -}; - -struct pipelined_key : oneapi::experimental::detail::compile_time_property_key< - oneapi::experimental::detail::PropKind::Pipelined> { - template - using value_t = ext::oneapi::experimental::property_value< - pipelined_key, - std::integral_constant>; -}; - -struct fpga_cluster_key - : oneapi::experimental::detail::compile_time_property_key< - oneapi::experimental::detail::PropKind::FPGACluster> { - template - using value_t = ext::oneapi::experimental::property_value< - fpga_cluster_key, - std::integral_constant>; -}; - -template -inline constexpr streaming_interface_key::value_t