* It seems that __cplusplus can be inconsistent with _MSVC_LANG when discerning C++17 version. See https://github.com/NVIDIA/cutlass/issues/1474. Added switch to check _MSVC_LANG in addition to __cplusplus
* Fixed typo.
* Oops, another typo.
* Changed incorrect logic, ifndef to ifdef
* Define CUTLAS_CPLUSPLUS for language version testing
Co-authored-by: Mark Hoemmen <mhoemmen@users.noreply.github.com>
---------
Co-authored-by: Mark Hoemmen <mhoemmen@users.noreply.github.com>
* add missing header for size_t in `numeric_types.h`
* make nvrtc happy
* add missing header for int types in `cutlass/arch/memory.h`
---------
Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
* fix uint128 operator add for 64-bit hilo implemenation
* add uint128 test for operator add
* make clang happy
---------
Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
* Allow per-column bias in EpilogueTensorBroadcast
EpilogueTensorBroadcast only supports per-row vector broadcast, because
the bias stride is hardcoded.
It can easily support both if the bias stride is made conditional, and
the original behavior is maintained by defaulting to per-row.
* Add unit test for EpilogueTensorBroadcast with per-col bias
---------
Co-authored-by: Ali Hassani <ahassanijr@gmail.com>
Co-authored-by: Ali Hassani <ali@hippoml.com>
* Fix inline ptx escaping for predicates.
Prevents `error: invalid % escape in inline assembly string` when compiling with clang.
* More double-quoting.
Currently, the default constructor of
`PredicatedTileAccessIteratorParams` will invoke undefined behavior in
its invocation of the `initialize` function. Specifically, it will
attempt to read from the uninitialized variables
`desc.element_size_bits` and `desc.advance_rank`. This commit changes
the default constructors of both `*Params` and `*Desc` to
zero-initialize all uninitialized members.
* Remove unused variables
* Qualify calls to make_fragment_? from templated base class.
Fixes clang build error.
* Add missing `#include <cstdio>`
* Various changes to fix clang compile errors.
* More changes to fix clang build.
Remaining issues:
- `params` initializer of `CollectiveEpilogue`.
- `ops` initializer of `Sm90VisitorImplBase`.
- `__usAtomicCAS` needs to be added to clang upstream.
* Fix remaining clang build issues.
* Qualify `cute::rank()` calls.
* Qualify some more calls that are otherwise ambiguous between `cute` and `std` namespace.
* Double-escape special registers in inline asm.
* small change
---------
Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
Followup to #1224.
A change in the stream-k threadblock swizzle ctor since 3.3 breaks
single source GEMM with fused epilogue and stream-k. Multi-source was
already corrected.
Co-authored-by: Ali Hassani <ahassanijr@gmail.com>
