[SYCLomatic] Migrate device ID and UUID (#732)

* [SYCLomatic] Migrate device ID and UUID

Signed-off-by: Michael Aziz <[email protected]>

* [SYCLomatic] Disable device info migration with option

Signed-off-by: Michael Aziz <[email protected]>

---------

Signed-off-by: Michael Aziz <[email protected]>

Author: 0x12CC

clang/lib/DPCT/ASTTraversal.cpp

@@ -2020,7 +2020,7 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) {
                   "cublasAtomicsMode_t", "CUmem_advise_enum", "CUmem_advise",
                   "thrust::tuple_element", "thrust::tuple_size", "cublasMath_t",
                   "cudaPointerAttributes", "thrust::zip_iterator",
-                  "cusolverEigRange_t")
+                  "cusolverEigRange_t", "cudaUUID_t")
               )))))
           .bind("cudaTypeDef"),
       this);
@@ -3524,8 +3524,7 @@ void DeviceInfoVarRule::runRule(const MatchFinder::MatchResult &Result) {
     emplaceTransformation(
         new ReplaceToken(ME->getBeginLoc(), ME->getEndLoc(), "false"));
     return;
-  } else if (MemberName == "pciDomainID" || MemberName == "pciBusID" ||
-             MemberName == "pciDeviceID") {
+  } else if (MemberName == "pciDomainID" || MemberName == "pciBusID") {
     report(ME->getBeginLoc(), Diagnostics::UNCOMPATIBLE_DEVICE_PROP, false,
            MemberName, "-1");
     emplaceTransformation(
@@ -3565,6 +3564,13 @@ void DeviceInfoVarRule::runRule(const MatchFinder::MatchResult &Result) {
     report(ME->getBeginLoc(), Diagnostics::MAX_GRID_SIZE, false);
   }
 
+  if (!DpctGlobalInfo::useDeviceInfo() &&
+      (MemberName == "pciDeviceID" || MemberName == "uuid")) {
+    report(ME->getBeginLoc(), Diagnostics::UNMIGRATED_DEVICE_PROP, false,
+           MemberName);
+    return;
+  }
+
   auto Search = PropNamesMap.find(MemberName);
   if (Search == PropNamesMap.end()) {
     return;
@@ -3589,6 +3595,15 @@ void DeviceInfoVarRule::runRule(const MatchFinder::MatchResult &Result) {
       emplaceTransformation(new ReplaceText(BO->getOperatorLoc(), 1, "("));
       emplaceTransformation(new InsertAfterStmt(BO, ")"));
     }
+  } else if (auto *OCE = Parents[0].get<clang::CXXOperatorCallExpr>()) {
+    // migrate to set_XXX() for types with an overloaded = operator
+    if (OCE->getOperator() == clang::OverloadedOperatorKind::OO_Equal) {
+      requestFeature(PropToSetFeatureMap.at(MemberName), ME);
+      emplaceTransformation(
+          new RenameFieldInMemberExpr(ME, "set_" + Search->second));
+      emplaceTransformation(new ReplaceText(OCE->getOperatorLoc(), 1, "("));
+      emplaceTransformation(new InsertAfterStmt(OCE, ")"));
+    }
   }
   if ((Search->second.compare(0, 13, "major_version") == 0) ||
       (Search->second.compare(0, 13, "minor_version") == 0)) {
@@ -15264,3 +15279,22 @@ void CudaExtentRule::runRule(
 }
 
 REGISTER_RULE(CudaExtentRule, PassKind::PK_Analysis)
+
+void CudaUuidRule::registerMatcher(ast_matchers::MatchFinder &MF) {
+  MF.addMatcher(memberExpr(hasObjectExpression(hasType(namedDecl(
+                               hasAnyName("CUuuid_st", "cudaUUID_t")))),
+                           member(hasName("bytes")))
+                    .bind("UUID_bytes"),
+                this);
+}
+
+void CudaUuidRule::runRule(
+    const ast_matchers::MatchFinder::MatchResult &Result) {
+  if (auto ME = Result.Nodes.getNodeAs<MemberExpr>("UUID_bytes")) {
+    const auto SM = Result.SourceManager;
+    const auto Begin = SM->getSpellingLoc(ME->getOperatorLoc());
+    return emplaceTransformation(new ReplaceText(Begin, 6, ""));
+  }
+}
+
+REGISTER_RULE(CudaUuidRule, PassKind::PK_Analysis)

clang/lib/DPCT/ASTTraversal.h

@@ -583,6 +583,12 @@ class CudaExtentRule : public NamedMigrationRule<CudaExtentRule> {
   }
 };
 
+class CudaUuidRule : public NamedMigrationRule<CudaUuidRule> {
+public:
+  void registerMatcher(ast_matchers::MatchFinder &MF) override;
+  void runRule(const ast_matchers::MatchFinder::MatchResult &Result);
+};
+
 /// Migration rule for return types replacements.
 class ReturnTypeRule : public NamedMigrationRule<ReturnTypeRule> {
 public:

clang/lib/DPCT/CustomHelperFiles.cpp

@@ -1128,6 +1128,10 @@ const std::unordered_map<std::string, HelperFeatureEnum> PropToGetFeatureMap = {
      HelperFeatureEnum::Device_device_info_get_memory_clock_rate},
     {"memoryBusWidth",
      HelperFeatureEnum::Device_device_info_get_memory_bus_width},
+    {"pciDeviceID",
+     HelperFeatureEnum::Device_device_info_get_device_id},
+    {"uuid",
+     HelperFeatureEnum::Device_device_info_get_uuid},
 };
 
 const std::unordered_map<std::string, HelperFeatureEnum> PropToSetFeatureMap = {
@@ -1156,6 +1160,10 @@ const std::unordered_map<std::string, HelperFeatureEnum> PropToSetFeatureMap = {
      HelperFeatureEnum::Device_device_info_set_memory_clock_rate},
     {"memoryBusWidth",
      HelperFeatureEnum::Device_device_info_set_memory_bus_width},
+    {"pciDeviceID",
+     HelperFeatureEnum::Device_device_info_set_device_id},
+    {"uuid",
+     HelperFeatureEnum::Device_device_info_set_uuid},
 };
 
 const std::unordered_map<std::string, HelperFeatureEnum>

clang/lib/DPCT/MapNames.cpp

@@ -431,6 +431,7 @@ void MapNames::setExplicitNamespaceMap() {
       {"cuda::std::complex", std::make_shared<TypeNameRule>("std::complex")},
       {"cuda::std::array", std::make_shared<TypeNameRule>("std::array")},
       {"cusolverEigRange_t", std::make_shared<TypeNameRule>("oneapi::mkl::rangev")},
+      {"cudaUUID_t", std::make_shared<TypeNameRule>("std::array<unsigned char, 16>")},
       // ...
   };
 
@@ -4228,6 +4229,8 @@ const MapNames::MapTy DeviceInfoVarRule::PropNamesMap{
     {"maxThreadsDim", "max_work_item_sizes"},
     {"memoryClockRate", "memory_clock_rate"},
     {"memoryBusWidth", "memory_bus_width"},
+    {"pciDeviceID", "device_id"},
+    {"uuid", "uuid"},
     // ...
 };
 

clang/runtime/dpct-rt/include/device.hpp.inc

@@ -40,6 +40,7 @@
 // DPCT_CODE
 #include <sycl/sycl.hpp>
 #include <algorithm>
+#include <array>
 #include <cstring>
 #include <iostream>
 #include <mutex>
@@ -303,6 +304,20 @@ public:
   /// Returns the maximum bus width between device and memory in bits. If
   /// compiler does not support this API then returns default value 64 bits.
   unsigned int get_memory_bus_width() const { return _memory_bus_width; }
+// DPCT_LABEL_END
+// DPCT_LABEL_BEGIN|device_info_get_device_id|dpct
+// DPCT_DEPENDENCY_BEGIN
+// Device|device_info
+// DPCT_DEPENDENCY_END
+// DPCT_CODE
+  uint32_t get_device_id() const { return _device_id; }
+// DPCT_LABEL_END
+// DPCT_LABEL_BEGIN|device_info_get_uuid|dpct
+// DPCT_DEPENDENCY_BEGIN
+// Device|device_info
+// DPCT_DEPENDENCY_END
+// DPCT_CODE
+  std::array<unsigned char, 16> get_uuid() const { return _uuid; }
 // DPCT_LABEL_END
   // set interface
 // DPCT_LABEL_BEGIN|device_info_set_name|dpct
@@ -478,6 +493,26 @@ public:
     _max_register_size_per_work_group = max_register_size_per_work_group;
   }
 // DPCT_LABEL_END
+// DPCT_LABEL_BEGIN|device_info_set_device_id|dpct
+// DPCT_PARENT_FEATURE|device_info
+// DPCT_DEPENDENCY_BEGIN
+// Device|device_info
+// DPCT_DEPENDENCY_END
+// DPCT_CODE
+  void set_device_id(uint32_t device_id) {
+    _device_id = device_id;
+  }
+// DPCT_LABEL_END
+// DPCT_LABEL_BEGIN|device_info_set_uuid|dpct
+// DPCT_PARENT_FEATURE|device_info
+// DPCT_DEPENDENCY_BEGIN
+// Device|device_info
+// DPCT_DEPENDENCY_END
+// DPCT_CODE
+  void set_uuid(std::array<unsigned char, 16> uuid) {
+    _uuid = std::move(uuid);
+  }
+// DPCT_LABEL_END
 // DPCT_LABEL_BEGIN|device_info_1|dpct
 // DPCT_PARENT_FEATURE|device_info
 // DPCT_DEPENDENCY_BEGIN
@@ -506,6 +541,8 @@ private:
   size_t _local_mem_size;
   size_t _max_nd_range_size[3];
   int _max_nd_range_size_i[3];
+  uint32_t _device_id;
+  std::array<unsigned char, 16> _uuid;
 };
 // DPCT_LABEL_END
 
@@ -769,6 +806,14 @@ public:
       prop.set_memory_bus_width(
           this->get_info<sycl::ext::intel::info::device::memory_bus_width>());
     }
+    if (this->has(sycl::aspect::ext_intel_device_id)) {
+      prop.set_device_id(
+          this->get_info<sycl::ext::intel::info::device::device_id>());
+    }
+    if (this->has(sycl::aspect::ext_intel_device_info_uuid)) {
+      prop.set_uuid(
+          this->get_info<sycl::ext::intel::info::device::uuid>());
+    }
 #elif defined(_MSC_VER) && !defined(__clang__)
 #pragma message("get_device_info: querying memory_clock_rate and \
 memory_bus_width are not supported by the compiler used. \

clang/test/dpct/device001.cu

@@ -1,3 +1,5 @@
+// UNSUPPORTED: cuda-8.0, cuda-9.1, cuda-9.2
+// UNSUPPORTED: v8.0, v9.1, v9.2
 // RUN: dpct --format-range=none --usm-level=none -out-root %T/device001 %s --cuda-include-path="%cuda-path/include" --sycl-named-lambda -- -x cuda --cuda-host-only
 // RUN: FileCheck %s --match-full-lines --input-file %T/device001/device001.dp.cpp
 
@@ -8,6 +10,16 @@ int main(int argc, char **argv) {
   // CHECK: dpct::device_info deviceProp;
   cudaDeviceProp deviceProp;
 
+  // CHECK: std::array<unsigned char, 16> uuid = deviceProp.get_uuid();
+  cudaUUID_t uuid = deviceProp.uuid;
+  // CHECK: deviceProp.set_uuid(uuid);
+  deviceProp.uuid=uuid;
+
+  // CHECK: int device_id = deviceProp.get_device_id();
+  int device_id = deviceProp.pciDeviceID;
+  // CHECK: deviceProp.set_device_id(device_id);
+  deviceProp.pciDeviceID=device_id;
+
   // CHECK: /*
   // CHECK-NEXT: DPCT1035:{{[0-9]+}}: All SYCL devices can be used by the host to submit tasks. You may need to adjust this code.
   // CHECK-NEXT: */

clang/test/dpct/device004.cu

@@ -0,0 +1,33 @@
+// UNSUPPORTED: cuda-8.0, cuda-9.1, cuda-9.2
+// UNSUPPORTED: v8.0, v9.1, v9.2
+// RUN: dpct --no-dpcpp-extensions=device_info -out-root %T/device004 %s --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only
+// RUN: FileCheck %s --match-full-lines --input-file %T/device004/device004.dp.cpp
+
+#include <cuda.h>
+#include <iostream>
+
+int main() {
+  cudaDeviceProp properties;
+  cudaGetDeviceProperties(&properties, 0);
+
+  // CHECK: /*
+  // CHECK-NEXT: DPCT1090:{{[0-9]+}}: SYCL does not support the device property that would be functionally compatible with pciDeviceID. It was not migrated. You need to rewrite the code.
+  // CHECK-NEXT: */
+  // CHECK-NEXT: const int id = properties.pciDeviceID;
+  const int id = properties.pciDeviceID;
+  // CHECK: /*
+  // CHECK-NEXT: DPCT1090:{{[0-9]+}}: SYCL does not support the device property that would be functionally compatible with uuid. It was not migrated. You need to rewrite the code.
+  // CHECK-NEXT: */
+  // CHECK-NEXT: const std::array<unsigned char, 16> uuid = properties.uuid;
+  const cudaUUID_t uuid = properties.uuid;
+  // CHECK: /*
+  // CHECK-NEXT: DPCT1090:{{[0-9]+}}: SYCL does not support the device property that would be functionally compatible with pciDeviceID. It was not migrated. You need to rewrite the code.
+  // CHECK-NEXT: */
+  // CHECK-NEXT: properties.pciDeviceID = id;
+  properties.pciDeviceID = id;
+  // CHECK: /*
+  // CHECK-NEXT: DPCT1090:{{[0-9]+}}: SYCL does not support the device property that would be functionally compatible with uuid. It was not migrated. You need to rewrite the code.
+  // CHECK-NEXT: */
+  // CHECK-NEXT: properties.uuid = uuid;
+  properties.uuid = uuid;
+}

clang/test/dpct/helper_files_ref/include/device.hpp

@@ -11,6 +11,7 @@
 
 #include <sycl/sycl.hpp>
 #include <algorithm>
+#include <array>
 #include <cstring>
 #include <iostream>
 #include <mutex>
@@ -125,6 +126,8 @@ class device_info {
   /// Returns the maximum bus width between device and memory in bits. If
   /// compiler does not support this API then returns default value 64 bits.
   unsigned int get_memory_bus_width() const { return _memory_bus_width; }
+  uint32_t get_device_id() const { return _device_id; }
+  std::array<unsigned char, 16> get_uuid() const { return _uuid; }
   // set interface
   void set_name(const char* name) {
     size_t length = strlen(name);
@@ -182,6 +185,12 @@ class device_info {
   set_max_register_size_per_work_group(int max_register_size_per_work_group) {
     _max_register_size_per_work_group = max_register_size_per_work_group;
   }
+  void set_device_id(uint32_t device_id) {
+    _device_id = device_id;
+  }
+  void set_uuid(std::array<unsigned char, 16> uuid) {
+    _uuid = std::move(uuid);
+  }
 private:
   char _name[256];
   sycl::id<3> _max_work_item_sizes;
@@ -204,6 +213,8 @@ class device_info {
   size_t _local_mem_size;
   size_t _max_nd_range_size[3];
   int _max_nd_range_size_i[3];
+  uint32_t _device_id;
+  std::array<unsigned char, 16> _uuid;
 };
 
 /// dpct device extension
@@ -332,6 +343,14 @@ class device_ext : public sycl::device {
       prop.set_memory_bus_width(
           this->get_info<sycl::ext::intel::info::device::memory_bus_width>());
     }
+    if (this->has(sycl::aspect::ext_intel_device_id)) {
+      prop.set_device_id(
+          this->get_info<sycl::ext::intel::info::device::device_id>());
+    }
+    if (this->has(sycl::aspect::ext_intel_device_info_uuid)) {
+      prop.set_uuid(
+          this->get_info<sycl::ext::intel::info::device::uuid>());
+    }
 #elif defined(_MSC_VER) && !defined(__clang__)
 #pragma message("get_device_info: querying memory_clock_rate and \
 memory_bus_width are not supported by the compiler used. \

clang/test/dpct/test_api_level/Device/api_test6.cu

@@ -5,9 +5,15 @@
 
 // CHECK: 35
 // TEST_FEATURE: Device_device_ext_get_device_info_return_info
+// TEST_FEATURE: Device_device_info_get_device_id
+// TEST_FEATURE: Device_device_info_set_device_id
+// TEST_FEATURE: Device_device_info_get_uuid
+// TEST_FEATURE: Device_device_info_set_uuid
 
 int main() {
   cudaDeviceProp deviceProp;
   cudaGetDeviceProperties(&deviceProp, 0);
+  deviceProp.uuid;
+  deviceProp.pciDeviceID;
   return 0;
 }

clang/test/dpct/types001.cu

@@ -548,10 +548,7 @@ void fun3() {
   // CHECK-NEXT: /*
   // CHECK-NEXT: DPCT1051:{{[0-9]+}}: SYCL does not support a device property functionally compatible with pciBusID. It was migrated to -1. You may need to adjust the value of -1 for the specific device.
   // CHECK-NEXT: */
-  // CHECK-NEXT: /*
-  // CHECK-NEXT: DPCT1051:{{[0-9]+}}: SYCL does not support a device property functionally compatible with pciDeviceID. It was migrated to -1. You may need to adjust the value of -1 for the specific device.
-  // CHECK-NEXT: */
-  // CHECK-NEXT: sprintf(devstr, "pci %x:%x:%x", -1, -1, -1);
+  // CHECK-NEXT: sprintf(devstr, "pci %x:%x:%x", -1, -1, deviceProp.get_device_id());
   // CHECK-NEXT: /*
   // CHECK-NEXT: DPCT1051:{{[0-9]+}}: SYCL does not support a device property functionally compatible with concurrentKernels. It was migrated to true. You may need to adjust the value of true for the specific device.
   // CHECK-NEXT: */