-
Notifications
You must be signed in to change notification settings - Fork 407
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.
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
SYCL: preserve address space by using sycl::global_ptr where sensible #4396
SYCL: preserve address space by using sycl::global_ptr where sensible #4396
Conversation
e6729e6
to
bf053a0
Compare
d321bfb
to
758548c
Compare
93f25a5
to
96e00d4
Compare
96e00d4
to
699b444
Compare
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Looks good once you clarified removal of the first pointer to scratch memory in team policy specializations of the parallel patterns
char* const scratch_ptr[2] = {m_scratch_ptr[0], m_scratch_ptr[1]}; | ||
const auto shmem_begin = m_shmem_begin; | ||
const int scratch_size[2] = {m_scratch_size[0], m_scratch_size[1]}; | ||
sycl::global_ptr<char> const global_scratch_ptr = m_global_scratch_ptr; |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Can you comment about that change? We were not using scratch_ptr[0]
?
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
No, we are only using scratch_ptr[1]
which gets global memory assigned in the constructor. In Cuda
and HIP
, we set both pointers in the constructor but in SYCL
we can only interact with local memory close to kernel execution which is why we always used a different variable for it.
Try to minimize generic space operations