Skip to content

Explicitly Use __ldg on Pointer Derefs#1800

Closed
fthaler wants to merge 1 commit into
GridTools:masterfrom
fthaler:ldg-on-derefs
Closed

Explicitly Use __ldg on Pointer Derefs#1800
fthaler wants to merge 1 commit into
GridTools:masterfrom
fthaler:ldg-on-derefs

Conversation

@fthaler

@fthaler fthaler commented Jul 24, 2024

Copy link
Copy Markdown
Contributor

Used the __ldg intrinsic on (scalar) pointer derefs in SID neighbor tables and derefs within the fn cartesian and unstructured backends.

@fthaler

fthaler commented Jul 24, 2024

Copy link
Copy Markdown
Contributor Author

launch perftest

@fthaler

fthaler commented Jul 24, 2024

Copy link
Copy Markdown
Contributor Author

launch jenkins

@gridtoolsjenkins

Copy link
Copy Markdown
Collaborator

Hi there, this is jenkins continuous integration...
Do you want me to verify this patch?

#include "./common_interface.hpp"
#include "./executor.hpp"

#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 350

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

We only support cuda >= 11.0, therefore anyway no support for < 350

Comment on lines +51 to +56
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 350
if constexpr (std::is_pointer_v<decltype(it.m_ptr)> &&
is_texture_type<std::decay_t<std::remove_pointer_t<decltype(it.m_ptr)>>>::value) {
return __ldg(it.m_ptr);
}
#endif

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

what about introducing a function like
deref_as_const?

@fthaler

fthaler commented Aug 14, 2024

Copy link
Copy Markdown
Contributor Author

Closed in favor of #1802

@fthaler fthaler closed this Aug 14, 2024
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants