Skip to content

[SYCLomatic] Support migration of attribute "cudaDevAttrMaxSharedMemoryPerBlockOptin" referred in API "cudaDeviceGetAttribute" #2409

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 2 commits into from
Oct 22, 2024
Merged

[SYCLomatic] Support migration of attribute "cudaDevAttrMaxSharedMemoryPerBlockOptin" referred in API "cudaDeviceGetAttribute" #2409

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
8c35934
[SYCLomatic] Support migration of attribute "cudaDevAttrMaxSharedMemo…
tomflinda Oct 14, 2024
1a2d360
Fix CI test failure for cuda-8.0
tomflinda Oct 14, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
5 changes: 5 additions & 0 deletions clang/lib/DPCT/ASTTraversal.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 += ").";
Expand Down
3 changes: 3 additions & 0 deletions clang/lib/DPCT/MapNames.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1295,6 +1295,9 @@ void MapNames::setExplicitNamespaceMap(
{"cudaDevAttrComputeCapabilityMajor",
std::make_shared<EnumNameRule>("get_major_version",
HelperFeatureEnum::device_ext)},
{"cudaDevAttrMaxSharedMemoryPerBlockOptin",
std::make_shared<EnumNameRule>("get_local_mem_size",
HelperFeatureEnum::device_ext)},
{"cudaDevAttrComputeCapabilityMinor",
std::make_shared<EnumNameRule>("get_minor_version",
HelperFeatureEnum::device_ext)},
Expand Down
4 changes: 4 additions & 0 deletions clang/runtime/dpct-rt/include/dpct/device.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 {
Copy link
Contributor

@zhiweij1 zhiweij1 Oct 15, 2024

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Please also submit PR for SYCLcompat.

Copy link
Contributor Author

@tomflinda tomflinda Oct 15, 2024 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

intel/llvm#15695 is submitted to add this helper function to SYCLcompat.

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.
Expand Down
9 changes: 9 additions & 0 deletions clang/test/dpct/device004.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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) \
Expand Down