[SYCL] Fix image selection for AOT on intel_cpu_{spr, gnr} (#15208)

When AOT compiling for cpu, the generic `spir64_x86_64` target is used
with `-fsycl-targets`. #14909,
functionality was added to select device images based on their
`compile_target` property in the image. The selection mechanism had to
consider CPU as a special case due to not having explicit targets.
However, the mechanism only considered `x86_64` and not `intel_cpu_spr`
or `intel_cpu_gnr`; therefore on a `intel_cpu_spr` or `intel_cpu_gnr`
device, trying to launch a program compiled with
`-fsycl-targets=spir64_x86_64`, device image selection would fail to
find an image (and thus fail to launch any kernels).

This PR updates the logic to include `intel_cpu_spr` and
`intel_cpu_gnr`. Note: for tests, this functionality is checked by any
test that AOT compiled for CPU and launches a kernel (includes
[AOT/cpu.cpp](https://github.com/intel/llvm/blob/sycl/sycl/test-e2e/AOT/cpu.cpp),
[AOT/double.cpp](https://github.com/intel/llvm/blob/sycl/sycl/test-e2e/AOT/double.cpp),
[AOT/half.cpp](https://github.com/intel/llvm/blob/sycl/sycl/test-e2e/AOT/half.cpp)).

sycl/source/detail/program_manager/program_manager.cpp

@@ -1321,9 +1321,15 @@ RTDeviceBinaryImage *getBinImageFromMultiMap(
  reinterpret_cast<const char *>(&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();
  }
  }

sycl/unittests/program_manager/CompileTarget.cpp

@@ -97,20 +97,25 @@ static sycl::unittest::UrImage Img[] = {
 
 static sycl::unittest::UrImageArray<std::size(Img)> ImgArray{Img};
 
-ur_device_handle_t MockSklDeviceHandle =
- reinterpret_cast<ur_device_handle_t>(1);
-ur_device_handle_t MockPvcDeviceHandle =
- reinterpret_cast<ur_device_handle_t>(2);
-ur_device_handle_t MockX86DeviceHandle =
- reinterpret_cast<ur_device_handle_t>(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<ur_device_handle_t>(this);
+ }
+ static MockDeviceData *fromHandle(ur_device_handle_t handle) {
+ return reinterpret_cast<MockDeviceData *>(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<ur_device_get_info_params_t *>(pParams);
  if (*params.ppropName == UR_DEVICE_INFO_IP_VERSION && *params.ppPropValue) {
  int &ret = *static_cast<int *>(*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<ur_device_type_t *>(*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<ur_device_select_binary_params_t *>(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) {