Skip to content

[SYCL] Make internal device_global pointer mutable #8780

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 3 commits into from
Mar 28, 2023
Merged

[SYCL] Make internal device_global pointer mutable #8780

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
7847203
[SYCL] Make internal device_global pointer mutable
steffenlarsen Mar 27, 2023
7bf6921
Merge remote-tracking branch 'intel/sycl' into steffen/mutable_dg_ptr
steffenlarsen Mar 28, 2023
887da7e
Add comment
steffenlarsen Mar 28, 2023
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
6 changes: 5 additions & 1 deletion sycl/include/sycl/ext/oneapi/device_global/device_global.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -44,7 +44,11 @@ template <typename T, typename PropertyListT, typename = void>
class device_global_base {
protected:
using pointer_t = typename decorated_global_ptr<T>::pointer;
pointer_t usmptr{};

// The pointer member is mutable to avoid the compiler optimizing it out when
// accessing const-qualified device_global variables.
mutable pointer_t usmptr{};
Copy link
Contributor

Choose a reason for hiding this comment

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

I think it would be useful to have a comment here that briefly mentions why this has to be mutable.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Good idea! It has been added.


pointer_t get_ptr() noexcept { return usmptr; }
const pointer_t get_ptr() const noexcept { return usmptr; }

Expand Down
20 changes: 20 additions & 0 deletions sycl/test/check_device_code/device_global_ptr_use.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,20 @@
// RUN: %clangxx -fsycl -c -fsycl-device-only -S -emit-llvm %s -o - | FileCheck %s

// Tests that the underlying pointer in a const-qualified shared device_global
// is not optimized out during access.

#include <sycl/sycl.hpp>

using namespace sycl;
using namespace sycl::ext::oneapi::experimental;

const device_global<int> DeviceGlobalVar;

int main() {
queue Q;
Q.single_task([]() {
// CHECK: load i32 {{.*}}({{.*}}* @_ZL15DeviceGlobalVar, i64 0, i32 0)
volatile int ReadVal = DeviceGlobalVar;
});
return 0;
}