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

3 participants