diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 99f3c5204dc74..d6f063e5fada6 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -1321,9 +1321,15 @@ RTDeviceBinaryImage *getBinImageFromMultiMap( reinterpret_cast(&CompileTargetByteArray[0]), CompileTargetByteArray.size()); // Note: there are no explicit targets for CPUs, so on x86_64, - // so we use a spir64_x86_64 compile target image. + // intel_cpu_spr, and intel_cpu_gnr, we use a spir64_x86_64 + // compile target image. + // TODO: When dedicated targets for CPU are added, (i.e. + // -fsycl-targets=intel_cpu_spr etc.) remove this special + // handling of CPU targets. if ((ArchName == CompileTarget) || - (ArchName == "x86_64" && CompileTarget == "spir64_x86_64")) { + (CompileTarget == "spir64_x86_64" && + (ArchName == "x86_64" || ArchName == "intel_cpu_spr" || + ArchName == "intel_cpu_gnr"))) { AddImg(); } } diff --git a/sycl/unittests/program_manager/CompileTarget.cpp b/sycl/unittests/program_manager/CompileTarget.cpp index 095a4af2a996c..281530a8e43e4 100644 --- a/sycl/unittests/program_manager/CompileTarget.cpp +++ b/sycl/unittests/program_manager/CompileTarget.cpp @@ -97,20 +97,25 @@ static sycl::unittest::UrImage Img[] = { static sycl::unittest::UrImageArray ImgArray{Img}; -ur_device_handle_t MockSklDeviceHandle = - reinterpret_cast(1); -ur_device_handle_t MockPvcDeviceHandle = - reinterpret_cast(2); -ur_device_handle_t MockX86DeviceHandle = - reinterpret_cast(3); -constexpr int SklIp = 0x02400009; -constexpr int PvcIp = 0x030f0000; -constexpr int X86Ip = 0; - -ur_device_handle_t MockDevices[] = { - MockSklDeviceHandle, - MockPvcDeviceHandle, - MockX86DeviceHandle, +struct MockDeviceData { + int Ip; + ur_device_type_t DeviceType; + ur_device_handle_t getHandle() { + return reinterpret_cast(this); + } + static MockDeviceData *fromHandle(ur_device_handle_t handle) { + return reinterpret_cast(handle); + } +}; + +// IP are from IntelGPUArchitectures/IntelCPUArchitectures in +// sycl/source/detail/device_info.hpp +MockDeviceData MockDevices[] = { + {0x02400009, UR_DEVICE_TYPE_GPU}, // Skl + {0x030f0000, UR_DEVICE_TYPE_GPU}, // Pvc + {0, UR_DEVICE_TYPE_CPU}, // X86 + {8, UR_DEVICE_TYPE_CPU}, // Spr + {9, UR_DEVICE_TYPE_CPU}, // Gnr }; static ur_result_t redefinedDeviceGet(void *pParams) { @@ -123,7 +128,7 @@ static ur_result_t redefinedDeviceGet(void *pParams) { if (*params.pphDevices) { assert(*params.pNumEntries <= std::size(MockDevices)); for (uint32_t i = 0; i < *params.pNumEntries; ++i) { - (*params.pphDevices)[i] = MockDevices[i]; + (*params.pphDevices)[i] = MockDevices[i].getHandle(); } } @@ -149,27 +154,22 @@ static ur_result_t redefinedDeviceGetInfo(void *pParams) { auto params = *static_cast(pParams); if (*params.ppropName == UR_DEVICE_INFO_IP_VERSION && *params.ppPropValue) { int &ret = *static_cast(*params.ppPropValue); - if (*params.phDevice == MockSklDeviceHandle) - ret = SklIp; - if (*params.phDevice == MockPvcDeviceHandle) - ret = PvcIp; - if (*params.phDevice == MockX86DeviceHandle) - ret = X86Ip; + ret = MockDeviceData::fromHandle(*params.phDevice)->Ip; } - if (*params.ppropName == UR_DEVICE_INFO_TYPE && - *params.phDevice == MockX86DeviceHandle) { + if (*params.ppropName == UR_DEVICE_INFO_TYPE) { if (*params.ppPropValue) *static_cast(*params.ppPropValue) = - UR_DEVICE_TYPE_CPU; + MockDeviceData::fromHandle(*params.phDevice)->DeviceType; if (*params.ppPropSizeRet) - **params.ppPropSizeRet = sizeof(UR_DEVICE_TYPE_CPU); + **params.ppPropSizeRet = sizeof(ur_device_type_t); } return UR_RESULT_SUCCESS; } static ur_result_t redefinedDeviceSelectBinary(void *pParams) { auto params = *static_cast(pParams); - auto target = *params.phDevice == MockX86DeviceHandle + auto target = MockDeviceData::fromHandle(*params.phDevice)->DeviceType == + UR_DEVICE_TYPE_CPU ? UR_DEVICE_BINARY_TARGET_SPIRV64_X86_64 : UR_DEVICE_BINARY_TARGET_SPIRV64_GEN; uint32_t fallback = *params.pNumBinaries; @@ -246,6 +246,16 @@ TEST_F(CompileTargetTest, SingleTask) { checkUsedImageWithCompileTarget("spir64_x86_64", [&]() { launchSingleTaskKernel(queue{archSelector(syclex::architecture::x86_64)}); }); + + checkUsedImageWithCompileTarget("spir64_x86_64", [&]() { + launchSingleTaskKernel( + queue{archSelector(syclex::architecture::intel_cpu_spr)}); + }); + + checkUsedImageWithCompileTarget("spir64_x86_64", [&]() { + launchSingleTaskKernel( + queue{archSelector(syclex::architecture::intel_cpu_gnr)}); + }); } void launchNDRangeKernel(queue q) { @@ -268,6 +278,16 @@ TEST_F(CompileTargetTest, NDRangeKernel) { checkUsedImageWithCompileTarget("spir64_x86_64", [&]() { launchNDRangeKernel(queue{archSelector(syclex::architecture::x86_64)}); }); + + checkUsedImageWithCompileTarget("spir64_x86_64", [&]() { + launchNDRangeKernel( + queue{archSelector(syclex::architecture::intel_cpu_spr)}); + }); + + checkUsedImageWithCompileTarget("spir64_x86_64", [&]() { + launchNDRangeKernel( + queue{archSelector(syclex::architecture::intel_cpu_gnr)}); + }); } void launchRangeKernel(queue q) { @@ -288,6 +308,14 @@ TEST_F(CompileTargetTest, RangeKernel) { checkUsedImageWithCompileTarget("spir64_x86_64", [&]() { launchRangeKernel(queue{archSelector(syclex::architecture::x86_64)}); }); + + checkUsedImageWithCompileTarget("spir64_x86_64", [&]() { + launchRangeKernel(queue{archSelector(syclex::architecture::intel_cpu_spr)}); + }); + + checkUsedImageWithCompileTarget("spir64_x86_64", [&]() { + launchRangeKernel(queue{archSelector(syclex::architecture::intel_cpu_gnr)}); + }); } TEST_F(CompileTargetTest, NoDeviceKernel) {