[SYCLomatic] Migrate device attribute cudaDevAttrTextureAlignment (#1752

) Signed-off-by: Jiang, Zhiwei <[email protected]>
oneapi-src · Mar 5, 2024 · 81a8a41 · 81a8a41
1 parent 498230e
commit 81a8a41
Show file tree

Hide file tree

Showing 5 changed files with 22 additions and 3 deletions.
diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp
@@ -6479,6 +6479,11 @@ void FunctionCallRule::runRule(const MatchFinder::MatchResult &Result) {
       ReplStr += Search->second->NewName;
       ReplStr += "()";
       requestFeature(HelperFeatureEnum::device_ext);
+
+      if (AttributeName == "cudaDevAttrTextureAlignment") {
+        report(CE->getBeginLoc(), Diagnostics::UNCOMPATIBLE_DEVICE_PROP, false,
+               AttributeName, Search->second->NewName);
+      }
     }
     if (IsAssigned)
       ReplStr = "DPCT_CHECK_ERROR(" + ReplStr + ")";

diff --git a/clang/lib/DPCT/MapNames.cpp b/clang/lib/DPCT/MapNames.cpp
@@ -900,6 +900,9 @@ void MapNames::setExplicitNamespaceMap() {
       {"cudaDevAttrConcurrentManagedAccess",
        std::make_shared<EnumNameRule>(
            "get_info<sycl::info::device::usm_shared_allocations>")},
+      {"cudaDevAttrTextureAlignment",
+       std::make_shared<EnumNameRule>("get_mem_base_addr_align_in_bytes",
+                                      HelperFeatureEnum::device_ext)},
       // enum Memcpy Kind
       {"cudaMemcpyHostToHost",
        std::make_shared<EnumNameRule>(getDpctNamespace() + "host_to_host")},
@@ -1008,7 +1011,7 @@ void MapNames::setExplicitNamespaceMap() {
        std::make_shared<EnumNameRule>("get_max_work_group_size",
                                       HelperFeatureEnum::device_ext)},
       {"CU_DEVICE_ATTRIBUTE_TEXTURE_ALIGNMENT",
-       std::make_shared<EnumNameRule>("get_mem_base_addr_align",
+       std::make_shared<EnumNameRule>("get_mem_base_addr_align_in_bytes",
                                       HelperFeatureEnum::device_ext)},
       {"CU_DEVICE_ATTRIBUTE_TOTAL_CONSTANT_MEMORY",
        std::make_shared<EnumNameRule>("get_global_mem_size",

diff --git a/clang/runtime/dpct-rt/include/dpct/device.hpp b/clang/runtime/dpct-rt/include/dpct/device.hpp
@@ -410,6 +410,10 @@ class device_ext : public sycl::device {
     return get_info<sycl::info::device::mem_base_addr_align>();
   }
 
+  int get_mem_base_addr_align_in_bytes() const {
+    return get_info<sycl::info::device::mem_base_addr_align>() / 8;
+  }
+
   size_t get_global_mem_size() const {
     return get_device_info().get_global_mem_size();
   }

diff --git a/clang/test/dpct/device002.cu b/clang/test/dpct/device002.cu
@@ -129,6 +129,13 @@ int integrated = -1;
 // CHECK: checkError(DPCT_CHECK_ERROR(integrated = dpct::dev_mgr::instance().get_device(dev_id).get_integrated()));
 checkError(cudaDeviceGetAttribute(&integrated, cudaDevAttrIntegrated, dev_id));
 
+int alignment;
+// CHECK: /*
+// CHECK-NEXT: DPCT1051:{{[0-9]+}}: SYCL does not support a device property functionally compatible with cudaDevAttrTextureAlignment. It was migrated to get_mem_base_addr_align_in_bytes. You may need to adjust the value of get_mem_base_addr_align_in_bytes for the specific device.
+// CHECK-NEXT: */
+// CHECK-NEXT: checkError(DPCT_CHECK_ERROR(alignment = dpct::dev_mgr::instance().get_device(dev_id).get_mem_base_addr_align_in_bytes()));
+checkError(cudaDeviceGetAttribute(&alignment, cudaDevAttrTextureAlignment, dev_id));
+
 int device1 = 0;
 int device2 = 1;
 int perfRank = 0;

diff --git a/clang/test/dpct/driver_device.cu b/clang/test/dpct/driver_device.cu
@@ -33,9 +33,9 @@ void test() {
   cuDeviceGetAttribute(&result2,CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_BLOCK, device);
   std::cout << " result2 " << result2 << std::endl;
   // CHECK: /*
-  // CHECK-NEXT: DPCT1051:{{[0-9]+}}: SYCL does not support a device property functionally compatible with CU_DEVICE_ATTRIBUTE_TEXTURE_ALIGNMENT. It was migrated to get_mem_base_addr_align. You may need to adjust the value of get_mem_base_addr_align for the specific device.
+  // CHECK-NEXT: DPCT1051:{{[0-9]+}}: SYCL does not support a device property functionally compatible with CU_DEVICE_ATTRIBUTE_TEXTURE_ALIGNMENT. It was migrated to get_mem_base_addr_align_in_bytes. You may need to adjust the value of get_mem_base_addr_align_in_bytes for the specific device.
   // CHECK-NEXT: */
-  // CHECK: result3 = dpct::dev_mgr::instance().get_device(device).get_mem_base_addr_align();
+  // CHECK: result3 = dpct::dev_mgr::instance().get_device(device).get_mem_base_addr_align_in_bytes();
   cuDeviceGetAttribute(&result3,CU_DEVICE_ATTRIBUTE_TEXTURE_ALIGNMENT, device);
   std::cout << " result3 " << result3 << std::endl;
   // CHECK: /*