Vermeiden Sie undefiniertes Verhalten bei der Verlagerung in CUDA
Posted: 22 May 2025, 11:38
In Cuda nutze ich regelmäßig die Tatsache, dass die Hardware die Breite der Verschiebungen nicht einschränkt. Verhaltensauslöser im folgenden Code, wobei der Druck (oder das unter ihm) immer auftritt, obwohl die Segmentmask immer 0 ist. < /p>
Wie kann ich den Kompiler nicht ausführen. https://cuda.godbolt.org/z/yextm7eeq
Code: Select all
__device__ int Deduplicate_literal_4_28_fast(int lit, uint32_t lane8) {
const auto shift_CDDD = uint32_t((24 - int(lane8)) * 3);
const auto segmentmask_CDDD = 0xffff'ff00u >> shift_CDDD;
//always prints due to undefined behavior,
//which eliminates the `&& (segmentmask_CDDD != 0)`
//because I told the compiler that lane8 < 9.
if ((lane8
Die PTX -Handbuch steht: < /p>
9.7.8.9. Logik- und Verschiebungsanweisungen: SHR
Schaltbits rechts, Zeichen oder Null-Fill links. Modulo (n)
Dies bedeutet, dass pro Dokumentation SHR dest, nicht signiert (a), 32 oder größer: dest = 0 Wie kann ich den Kompiler nicht ausführen. https://cuda.godbolt.org/z/yextm7eeq
Code: Select all
#include
#include
#include
#include
__device__ int Deduplicate_literal_4_28_fast(int lit, uint32_t lane8) { //4x faster than match_any_sync version
const auto shift_CDDD = uint32_t((24 - int(lane8)) * 3);
const auto segmentmask_CDDD = 0xffff'ff00u >> shift_CDDD;
if ((lane8