Skip to content
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

[pre-commit.ci] pre-commit autoupdate #3248

Merged
merged 5 commits into from
Jan 9, 2025
Merged

[pre-commit.ci] pre-commit autoupdate #3248

Show file tree
Hide file tree
Changes from 2 commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
5cdb4ab
[pre-commit.ci] pre-commit autoupdate
pre-commit-ci[bot] Jan 7, 2025
7035bf6
[pre-commit.ci] auto code formatting
pre-commit-ci[bot] Jan 7, 2025
908c56a
[pre-commit.ci] auto code formatting
pre-commit-ci[bot] Jan 7, 2025
79b9191
Merge branch 'main' into pre-commit-ci-update-config
miscco Jan 8, 2025
f309d1d
Merge branch 'main' into pre-commit-ci-update-config
miscco Jan 9, 2025
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: 3 additions & 3 deletions .pre-commit-config.yaml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,7 +17,7 @@ repos:
- id: mixed-line-ending
- id: trailing-whitespace
- repo: https://github.com/pre-commit/mirrors-clang-format
rev: v18.1.8
rev: v19.1.6
hooks:
- id: clang-format
types_or: [file]
Expand All @@ -39,7 +39,7 @@ repos:
# TODO/REMINDER: add the Ruff vscode extension to the devcontainers
# Ruff, the Python auto-correcting linter/formatter written in Rust
- repo: https://github.com/astral-sh/ruff-pre-commit
rev: v0.8.3
rev: v0.8.6
hooks:
- id: ruff # linter
- id: ruff-format # formatter
Expand All @@ -57,7 +57,7 @@ repos:


- repo: https://github.com/pre-commit/mirrors-mypy
rev: 'v1.13.0'
rev: 'v1.14.1'
hooks:
- id: mypy
additional_dependencies: [types-cachetools, numpy]
Expand Down
2 changes: 1 addition & 1 deletion cub/cub/agent/agent_histogram.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -629,7 +629,7 @@ struct AgentHistogram

// Set valid flags
MarkValid<IS_FULL_TILE>(
is_valid, valid_samples, Int2Type<AgentHistogramPolicyT::LOAD_ALGORITHM == BLOCK_LOAD_STRIPED>{});
is_valid, valid_samples, Int2Type < AgentHistogramPolicyT::LOAD_ALGORITHM == BLOCK_LOAD_STRIPED > {});
Copy link
Collaborator

Choose a reason for hiding this comment

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

Those look bad :(


// Accumulate samples
if (prefer_smem)
Expand Down
8 changes: 4 additions & 4 deletions cub/cub/agent/agent_reduce.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -382,8 +382,8 @@ struct AgentReduce
even_share.template BlockInit<TILE_ITEMS>(block_offset, block_end);

return (IsAligned(d_in + block_offset, Int2Type<ATTEMPT_VECTORIZATION>()))
? ConsumeRange(even_share, Int2Type < true && ATTEMPT_VECTORIZATION > ())
: ConsumeRange(even_share, Int2Type < false && ATTEMPT_VECTORIZATION > ());
? ConsumeRange(even_share, Int2Type<true && ATTEMPT_VECTORIZATION>())
: ConsumeRange(even_share, Int2Type<false && ATTEMPT_VECTORIZATION>());
}

/**
Expand All @@ -396,8 +396,8 @@ struct AgentReduce
even_share.template BlockInit<TILE_ITEMS, GRID_MAPPING_STRIP_MINE>();

return (IsAligned(d_in, Int2Type<ATTEMPT_VECTORIZATION>()))
? ConsumeRange(even_share, Int2Type < true && ATTEMPT_VECTORIZATION > ())
: ConsumeRange(even_share, Int2Type < false && ATTEMPT_VECTORIZATION > ());
? ConsumeRange(even_share, Int2Type<true && ATTEMPT_VECTORIZATION>())
: ConsumeRange(even_share, Int2Type<false && ATTEMPT_VECTORIZATION>());
}

private:
Expand Down
3 changes: 1 addition & 2 deletions cub/cub/block/block_radix_rank.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -606,8 +606,7 @@ private:
{
volatile DigitCounterT warp_digit_counters[RADIX_DIGITS][PADDED_WARPS];
DigitCounterT raking_grid[BLOCK_THREADS][PADDED_RAKING_SEGMENT];
}
aliasable;
} aliasable;
};
#endif // !_CCCL_DOXYGEN_INVOKED

Expand Down
153 changes: 46 additions & 107 deletions cub/cub/detail/strong_load.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -59,14 +59,14 @@ static _CCCL_DEVICE _CCCL_FORCEINLINE uint4 load_relaxed(uint4 const* ptr)
uint4 retval;
NV_IF_TARGET(
NV_PROVIDES_SM_70,
(asm volatile("ld.relaxed.gpu.v4.u32 {%0, %1, %2, %3}, [%4];"
: "=r"(retval.x), "=r"(retval.y), "=r"(retval.z), "=r"(retval.w)
: "l"(ptr)
: "memory");),
(asm volatile("ld.cg.v4.u32 {%0, %1, %2, %3}, [%4];"
: "=r"(retval.x), "=r"(retval.y), "=r"(retval.z), "=r"(retval.w)
: "l"(ptr)
: "memory");));
(asm volatile("ld.relaxed.gpu.v4.u32 {%0, %1, %2, %3}, [%4];" : "=r"(retval.x),
"=r"(retval.y),
"=r"(retval.z),
"=r"(retval.w) : "l"(ptr) : "memory");),
(asm volatile("ld.cg.v4.u32 {%0, %1, %2, %3}, [%4];" : "=r"(retval.x),
"=r"(retval.y),
"=r"(retval.z),
"=r"(retval.w) : "l"(ptr) : "memory");));
return retval;
}

Expand All @@ -75,14 +75,8 @@ static _CCCL_DEVICE _CCCL_FORCEINLINE ulonglong2 load_relaxed(ulonglong2 const*
ulonglong2 retval;
NV_IF_TARGET(
NV_PROVIDES_SM_70,
(asm volatile("ld.relaxed.gpu.v2.u64 {%0, %1}, [%2];"
: "=l"(retval.x), "=l"(retval.y)
: "l"(ptr)
: "memory");),
(asm volatile("ld.cg.v2.u64 {%0, %1}, [%2];"
: "=l"(retval.x), "=l"(retval.y)
: "l"(ptr)
: "memory");));
(asm volatile("ld.relaxed.gpu.v2.u64 {%0, %1}, [%2];" : "=l"(retval.x), "=l"(retval.y) : "l"(ptr) : "memory");),
(asm volatile("ld.cg.v2.u64 {%0, %1}, [%2];" : "=l"(retval.x), "=l"(retval.y) : "l"(ptr) : "memory");));
return retval;
}

Expand All @@ -91,14 +85,14 @@ static _CCCL_DEVICE _CCCL_FORCEINLINE ushort4 load_relaxed(ushort4 const* ptr)
ushort4 retval;
NV_IF_TARGET(
NV_PROVIDES_SM_70,
(asm volatile("ld.relaxed.gpu.v4.u16 {%0, %1, %2, %3}, [%4];"
: "=h"(retval.x), "=h"(retval.y), "=h"(retval.z), "=h"(retval.w)
: "l"(ptr)
: "memory");),
(asm volatile("ld.cg.v4.u16 {%0, %1, %2, %3}, [%4];"
: "=h"(retval.x), "=h"(retval.y), "=h"(retval.z), "=h"(retval.w)
: "l"(ptr)
: "memory");));
(asm volatile("ld.relaxed.gpu.v4.u16 {%0, %1, %2, %3}, [%4];" : "=h"(retval.x),
"=h"(retval.y),
"=h"(retval.z),
"=h"(retval.w) : "l"(ptr) : "memory");),
(asm volatile("ld.cg.v4.u16 {%0, %1, %2, %3}, [%4];" : "=h"(retval.x),
"=h"(retval.y),
"=h"(retval.z),
"=h"(retval.w) : "l"(ptr) : "memory");));
return retval;
}

Expand All @@ -107,63 +101,36 @@ static _CCCL_DEVICE _CCCL_FORCEINLINE uint2 load_relaxed(uint2 const* ptr)
uint2 retval;
NV_IF_TARGET(
NV_PROVIDES_SM_70,
(asm volatile("ld.relaxed.gpu.v2.u32 {%0, %1}, [%2];"
: "=r"(retval.x), "=r"(retval.y)
: "l"(ptr)
: "memory");),
(asm volatile("ld.cg.v2.u32 {%0, %1}, [%2];"
: "=r"(retval.x), "=r"(retval.y)
: "l"(ptr)
: "memory");));
(asm volatile("ld.relaxed.gpu.v2.u32 {%0, %1}, [%2];" : "=r"(retval.x), "=r"(retval.y) : "l"(ptr) : "memory");),
(asm volatile("ld.cg.v2.u32 {%0, %1}, [%2];" : "=r"(retval.x), "=r"(retval.y) : "l"(ptr) : "memory");));
return retval;
}

static _CCCL_DEVICE _CCCL_FORCEINLINE unsigned long long load_relaxed(unsigned long long const* ptr)
{
unsigned long long retval;
NV_IF_TARGET(
NV_PROVIDES_SM_70,
(asm volatile("ld.relaxed.gpu.u64 %0, [%1];"
: "=l"(retval)
: "l"(ptr)
: "memory");),
(asm volatile("ld.cg.u64 %0, [%1];"
: "=l"(retval)
: "l"(ptr)
: "memory");));
NV_IF_TARGET(NV_PROVIDES_SM_70,
(asm volatile("ld.relaxed.gpu.u64 %0, [%1];" : "=l"(retval) : "l"(ptr) : "memory");),
(asm volatile("ld.cg.u64 %0, [%1];" : "=l"(retval) : "l"(ptr) : "memory");));
return retval;
}

static _CCCL_DEVICE _CCCL_FORCEINLINE unsigned int load_relaxed(unsigned int const* ptr)
{
unsigned int retval;
NV_IF_TARGET(
NV_PROVIDES_SM_70,
(asm volatile("ld.relaxed.gpu.u32 %0, [%1];"
: "=r"(retval)
: "l"(ptr)
: "memory");),
(asm volatile("ld.cg.u32 %0, [%1];"
: "=r"(retval)
: "l"(ptr)
: "memory");));
NV_IF_TARGET(NV_PROVIDES_SM_70,
(asm volatile("ld.relaxed.gpu.u32 %0, [%1];" : "=r"(retval) : "l"(ptr) : "memory");),
(asm volatile("ld.cg.u32 %0, [%1];" : "=r"(retval) : "l"(ptr) : "memory");));

return retval;
}

static _CCCL_DEVICE _CCCL_FORCEINLINE unsigned short load_relaxed(unsigned short const* ptr)
{
unsigned short retval;
NV_IF_TARGET(
NV_PROVIDES_SM_70,
(asm volatile("ld.relaxed.gpu.u16 %0, [%1];"
: "=h"(retval)
: "l"(ptr)
: "memory");),
(asm volatile("ld.cg.u16 %0, [%1];"
: "=h"(retval)
: "l"(ptr)
: "memory");));
NV_IF_TARGET(NV_PROVIDES_SM_70,
(asm volatile("ld.relaxed.gpu.u16 %0, [%1];" : "=h"(retval) : "l"(ptr) : "memory");),
(asm volatile("ld.cg.u16 %0, [%1];" : "=h"(retval) : "l"(ptr) : "memory");));
return retval;
}

Expand All @@ -172,24 +139,16 @@ static _CCCL_DEVICE _CCCL_FORCEINLINE unsigned char load_relaxed(unsigned char c
unsigned short retval;
NV_IF_TARGET(
NV_PROVIDES_SM_70,
(asm volatile(
"{"
" .reg .u8 datum;"
" ld.relaxed.gpu.u8 datum, [%1];"
" cvt.u16.u8 %0, datum;"
"}"
: "=h"(retval)
: "l"(ptr)
: "memory");),
(asm volatile(
"{"
" .reg .u8 datum;"
" ld.cg.u8 datum, [%1];"
" cvt.u16.u8 %0, datum;"
"}"
: "=h"(retval)
: "l"(ptr)
: "memory");));
(asm volatile("{"
" .reg .u8 datum;"
" ld.relaxed.gpu.u8 datum, [%1];"
" cvt.u16.u8 %0, datum;"
"}" : "=h"(retval) : "l"(ptr) : "memory");),
(asm volatile("{"
" .reg .u8 datum;"
" ld.cg.u8 datum, [%1];"
" cvt.u16.u8 %0, datum;"
"}" : "=h"(retval) : "l"(ptr) : "memory");));
return (unsigned char) retval;
}

Expand All @@ -198,14 +157,8 @@ static _CCCL_DEVICE _CCCL_FORCEINLINE ulonglong2 load_acquire(ulonglong2 const*
ulonglong2 retval;
NV_IF_TARGET(
NV_PROVIDES_SM_70,
(asm volatile("ld.acquire.gpu.v2.u64 {%0, %1}, [%2];"
: "=l"(retval.x), "=l"(retval.y)
: "l"(ptr)
: "memory");),
(asm volatile("ld.cg.v2.u64 {%0, %1}, [%2];"
: "=l"(retval.x), "=l"(retval.y)
: "l"(ptr)
: "memory");
(asm volatile("ld.acquire.gpu.v2.u64 {%0, %1}, [%2];" : "=l"(retval.x), "=l"(retval.y) : "l"(ptr) : "memory");),
(asm volatile("ld.cg.v2.u64 {%0, %1}, [%2];" : "=l"(retval.x), "=l"(retval.y) : "l"(ptr) : "memory");
__threadfence();));
return retval;
}
Expand All @@ -215,32 +168,18 @@ static _CCCL_DEVICE _CCCL_FORCEINLINE uint2 load_acquire(uint2 const* ptr)
uint2 retval;
NV_IF_TARGET(
NV_PROVIDES_SM_70,
(asm volatile("ld.acquire.gpu.v2.u32 {%0, %1}, [%2];"
: "=r"(retval.x), "=r"(retval.y)
: "l"(ptr)
: "memory");),
(asm volatile("ld.cg.v2.u32 {%0, %1}, [%2];"
: "=r"(retval.x), "=r"(retval.y)
: "l"(ptr)
: "memory");
(asm volatile("ld.acquire.gpu.v2.u32 {%0, %1}, [%2];" : "=r"(retval.x), "=r"(retval.y) : "l"(ptr) : "memory");),
(asm volatile("ld.cg.v2.u32 {%0, %1}, [%2];" : "=r"(retval.x), "=r"(retval.y) : "l"(ptr) : "memory");
__threadfence();));
return retval;
}

static _CCCL_DEVICE _CCCL_FORCEINLINE unsigned int load_acquire(unsigned int const* ptr)
{
unsigned int retval;
NV_IF_TARGET(
NV_PROVIDES_SM_70,
(asm volatile("ld.acquire.gpu.u32 %0, [%1];"
: "=r"(retval)
: "l"(ptr)
: "memory");),
(asm volatile("ld.cg.u32 %0, [%1];"
: "=r"(retval)
: "l"(ptr)
: "memory");
__threadfence();));
NV_IF_TARGET(NV_PROVIDES_SM_70,
(asm volatile("ld.acquire.gpu.u32 %0, [%1];" : "=r"(retval) : "l"(ptr) : "memory");),
(asm volatile("ld.cg.u32 %0, [%1];" : "=r"(retval) : "l"(ptr) : "memory"); __threadfence();));

return retval;
}
Expand Down
Loading