diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp index 6224f1183290..4c9c98a26acc 100644 --- a/clang/lib/DPCT/ASTTraversal.cpp +++ b/clang/lib/DPCT/ASTTraversal.cpp @@ -6018,6 +6018,11 @@ void FunctionCallRule::runRule(const MatchFinder::MatchResult &Result) { } requestHelperFeatureForEnumNames(AttributeName); + if (AttributeName == "cudaDevAttrMaxSharedMemoryPerBlockOptin") { + report(CE->getBeginLoc(), Diagnostics::LOCAL_MEM_SIZE, false, + AttributeName); + } + ReplStr += " = " + MapNames::getDpctNamespace() + "get_device("; ReplStr += StmtStrArg2; ReplStr += ")."; diff --git a/clang/lib/DPCT/MapNames.cpp b/clang/lib/DPCT/MapNames.cpp index dd307b7d2f36..b8c0aeedc1e9 100644 --- a/clang/lib/DPCT/MapNames.cpp +++ b/clang/lib/DPCT/MapNames.cpp @@ -1295,6 +1295,9 @@ void MapNames::setExplicitNamespaceMap( {"cudaDevAttrComputeCapabilityMajor", std::make_shared("get_major_version", HelperFeatureEnum::device_ext)}, + {"cudaDevAttrMaxSharedMemoryPerBlockOptin", + std::make_shared("get_local_mem_size", + HelperFeatureEnum::device_ext)}, {"cudaDevAttrComputeCapabilityMinor", std::make_shared("get_minor_version", HelperFeatureEnum::device_ext)}, diff --git a/clang/runtime/dpct-rt/include/dpct/device.hpp b/clang/runtime/dpct-rt/include/dpct/device.hpp index d55ed31f74b0..b4fae4e4baff 100644 --- a/clang/runtime/dpct-rt/include/dpct/device.hpp +++ b/clang/runtime/dpct-rt/include/dpct/device.hpp @@ -512,6 +512,10 @@ class device_ext : public sycl::device { return get_device_info().get_global_mem_size(); } + size_t get_local_mem_size() const { + return get_device_info().get_local_mem_size(); + } + /// Get the number of bytes of free and total memory on the SYCL device. /// \param [out] free_memory The number of bytes of free memory on the SYCL device. /// \param [out] total_memory The number of bytes of total memory on the SYCL device. diff --git a/clang/test/dpct/device004.cu b/clang/test/dpct/device004.cu index 5c198049adad..af5a48aa700f 100644 --- a/clang/test/dpct/device004.cu +++ b/clang/test/dpct/device004.cu @@ -30,6 +30,15 @@ int main() { // CHECK-NEXT: */ // CHECK-NEXT: properties.uuid = uuid; properties.uuid = uuid; + + // Figure out how much shared memory is available on the device. + int maxSharedBytes = 0; + int dev_id; + // CHECK:/* + // CHECK-NEXT:DPCT1019:{{[0-9]+}}: local_mem_size in SYCL is not a complete equivalent of cudaDevAttrMaxSharedMemoryPerBlockOptin in CUDA. You may need to adjust the code. + // CHECK-NEXT:*/ + // CHECK-NEXT:maxSharedBytes = dpct::get_device(dev_id).get_local_mem_size(); + cudaDeviceGetAttribute(&maxSharedBytes, cudaDevAttrMaxSharedMemoryPerBlockOptin, dev_id); } #define CHECK_INTERNAL(err) \