Skip to content

Commit

Permalink
[cudadev] Removed cache access style support code in SoA.
Browse files Browse the repository at this point in the history 
The cache access function at the SoA level interfere with the __restrict__ support. It could still be re-introduced as a separate tool.
  • Loading branch information
@ericcano
ericcano committed Jan 14, 2022
1 parent 5a2d472 commit f89caec
Show file tree
Hide file tree
Showing 4 changed files with 6 additions and 52 deletions.
2 changes: 1 addition & 1 deletion src/cudadev/DataFormats/SoA.md
View file
Original file line number Diff line number Diff line change
Expand Up @@ -163,7 +163,7 @@ scenarios where only a subset of columns are used in a given GPU kernel.
- The layout and views support scalars and columns, alignment and alignment enforcement and hinting.
- Automatic `__restrict__` compiler hinting is supported.
- A shortcut alloCreate a mechanism to derive trivial views and const views from a single layout.
- Cache access style, which was explored will be removed as this not-yet-used feature interferes with `__restrict__` support (which is already in used in existing code)
- Cache access style, which was explored, was abandoned as this not-yet-used feature interferes with `__restrict__` support (which is already in used in existing code). It could be made available as a separate tool that can be used directly by the module developer, orthogonally from SoA.

### Planned additions
- Optional range checking will be added later. This implies adding support for size to views and will restrict views to columns of
Expand Down
25 changes: 0 additions & 25 deletions src/cudadev/DataFormats/SoACommon.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -23,23 +23,10 @@
#define SOA_DEVICE_RESTRICT
#endif

#if defined(__CUDACC__) && defined(__CUDA_ARCH__)
// Read a pointer content via read-only (non coherent) cache.
#define LOAD_NONCOHERENT(A) __ldg(A)
#define LOAD_STREAMING(A) __ldcs(A)
#define STORE_STREAMING(A, V) __stcs(A, V)
#else
#define LOAD_NONCOHERENT(A) *(A)
#define LOAD_STREAMING(A) *(A)
#define STORE_STREAMING(A, V) *(A) = (V)
#endif

// compile-time sized SoA

namespace cms::soa {

enum class CacheAccessStyle : char { Default, NonCoherent, Streaming };

enum class RestrictQualify : bool { Enabled, Disabled, Default = Disabled };

template <typename T, RestrictQualify RESTRICT_QUALIFY>
Expand All @@ -65,21 +52,10 @@ struct add_restrict<T, RestrictQualify::Disabled> {
typedef const T & ReferenceToConst;
};

template <typename T, CacheAccessStyle CACHE_ACCESS_STYLE>
SOA_HOST_DEVICE_INLINE T readWithCacheStyle (const T * addr) {
if constexpr (CACHE_ACCESS_STYLE == CacheAccessStyle::NonCoherent) {
return LOAD_INCOHERENT(addr);
} else if constexpr (CACHE_ACCESS_STYLE == CacheAccessStyle::Streaming) {
return LOAD_STREAMING(addr);
}
return *addr;
}

// Helper template managing the value within it column
// The optional compile time alignment parameter enables informing the
// compiler of alignment (enforced by caller).
template <typename T, size_t ALIGNMENT,
CacheAccessStyle CACHE_STYLE = CacheAccessStyle::Default,
RestrictQualify RESTRICT_QUALIFY = RestrictQualify::Disabled>
class SoAValue {
public:
Expand Down Expand Up @@ -123,7 +99,6 @@ class SoAValue {

// Helper template managing the value within it column
template <typename T, size_t ALIGNMENT,
CacheAccessStyle CACHE_STYLE = CacheAccessStyle::Default,
RestrictQualify RESTRICT_QUALIFY = RestrictQualify::Disabled>
class SoAConstValue {
public:
Expand Down
14 changes: 5 additions & 9 deletions src/cudadev/DataFormats/SoAView.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -291,7 +291,6 @@ struct ConstValueTraits<C, SoAColumnType::eigen> {
#define GENERATE_SOA_VIEW(CLASS, LAYOUTS_LIST, VALUE_LIST) \
template <size_t ALIGNMENT = cms::soa::CacheLineSize::defaultSize, \
cms::soa::AlignmentEnforcement ALIGNMENT_ENFORCEMENT = cms::soa::AlignmentEnforcement::Relaxed, \
cms::soa::CacheAccessStyle CACHE_ACCESS_STYLE = cms::soa::CacheAccessStyle::Default, \
cms::soa::RestrictQualify RESTRICT_QUALIFY = cms::soa::RestrictQualify::Disabled> \
struct CLASS { \
/* these could be moved to an external type trait to free up the symbol names */ \
Expand All @@ -307,14 +306,13 @@ struct ConstValueTraits<C, SoAColumnType::eigen> {
constexpr static AlignmentEnforcement alignmentEnforcement = ALIGNMENT_ENFORCEMENT; \
constexpr static size_t conditionalAlignment = \
alignmentEnforcement == AlignmentEnforcement::Enforced ? byteAlignment : 0; \
constexpr static cms::soa::CacheAccessStyle cacheAccessStyle = CACHE_ACCESS_STYLE; \
constexpr static cms::soa::RestrictQualify restrictQualify = RESTRICT_QUALIFY; \
/* Those typedefs avoid having commas in macros (which is problematic) */ \
/* Those typedefs avoid having commas in macros (which is problematic) */ \
template <class C> \
using SoAValueWithConf = cms::soa::SoAValue<C, conditionalAlignment, cacheAccessStyle, restrictQualify>; \
using SoAValueWithConf = cms::soa::SoAValue<C, conditionalAlignment, restrictQualify>; \
\
template <class C> \
using SoAConstValueWithConf = cms::soa::SoAConstValue<C, conditionalAlignment, cacheAccessStyle, restrictQualify>; \
using SoAConstValueWithConf = cms::soa::SoAConstValue<C, conditionalAlignment, restrictQualify>; \
\
template <class C> \
using SoAEigenValueWithConf = cms::soa::SoAEigenValue<C, conditionalAlignment>; \
Expand Down Expand Up @@ -406,7 +404,6 @@ struct ConstValueTraits<C, SoAColumnType::eigen> {
#define GENERATE_SOA_CONST_VIEW(CLASS, LAYOUTS_LIST, VALUE_LIST) \
template <size_t ALIGNMENT = cms::soa::CacheLineSize::defaultSize, \
cms::soa::AlignmentEnforcement ALIGNMENT_ENFORCEMENT = cms::soa::AlignmentEnforcement::Relaxed, \
cms::soa::CacheAccessStyle CACHE_ACCESS_STYLE = cms::soa::CacheAccessStyle::NonCoherent, \
cms::soa::RestrictQualify RESTRICT_QUALIFY = cms::soa::RestrictQualify::Enabled> \
struct CLASS { \
/* these could be moved to an external type trait to free up the symbol names */ \
Expand All @@ -422,14 +419,13 @@ struct ConstValueTraits<C, SoAColumnType::eigen> {
constexpr static AlignmentEnforcement alignmentEnforcement = ALIGNMENT_ENFORCEMENT; \
constexpr static size_t conditionalAlignment = \
alignmentEnforcement == AlignmentEnforcement::Enforced ? byteAlignment : 0; \
constexpr static cms::soa::CacheAccessStyle cacheAccessStyle = CACHE_ACCESS_STYLE; \
constexpr static cms::soa::RestrictQualify restrictQualify = RESTRICT_QUALIFY; \
/* Those typedefs avoid having commas in macros (which is problematic) */ \
template <class C> \
using SoAValueWithConf = cms::soa::SoAValue<C, conditionalAlignment, cacheAccessStyle, restrictQualify>; \
using SoAValueWithConf = cms::soa::SoAValue<C, conditionalAlignment, restrictQualify>; \
\
template <class C> \
using SoAConstValueWithConf = cms::soa::SoAConstValue<C, conditionalAlignment, cacheAccessStyle, restrictQualify>; \
using SoAConstValueWithConf = cms::soa::SoAConstValue<C, conditionalAlignment, restrictQualify>; \
\
template <class C> \
using SoAEigenValueWithConf = cms::soa::SoAEigenValue<C, conditionalAlignment>; \
Expand Down
17 changes: 0 additions & 17 deletions src/cudadev/test/SoAStoreAndView_t.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -98,33 +98,16 @@ __device__ void addAndMulTemplate (

__global__ void aAMDef(SoA1ViewTemplate<cms::soa::CacheLineSize::defaultSize,
cms::soa::AlignmentEnforcement::Relaxed,
cms::soa::CacheAccessStyle::Default,
cms::soa::RestrictQualify::Disabled> soa, size_t size) {
addAndMulTemplate(soa, size);
}

__global__ void aAMRestrict(SoA1ViewTemplate<cms::soa::CacheLineSize::defaultSize,
cms::soa::AlignmentEnforcement::Relaxed,
cms::soa::CacheAccessStyle::Default,
cms::soa::RestrictQualify::Enabled> soa, size_t size) {
addAndMulTemplate(soa, size);
}

__global__ void aAMNC(SoA1ViewTemplate<cms::soa::CacheLineSize::defaultSize,
cms::soa::AlignmentEnforcement::Relaxed,
cms::soa::CacheAccessStyle::NonCoherent,
cms::soa::RestrictQualify::Disabled> soa, size_t size) {
addAndMulTemplate(soa, size);
}

__global__ void aAMRestrict(SoA1ViewTemplate<cms::soa::CacheLineSize::defaultSize,
cms::soa::AlignmentEnforcement::Relaxed,
cms::soa::CacheAccessStyle::NonCoherent,
cms::soa::RestrictQualify::Enabled> soa, size_t size) {
addAndMulTemplate(soa, size);
}


const size_t size=10000;

int main() {
Expand Down

0 comments on commit f89caec

Please sign in to comment.