Vermeiden Sie undefiniertes Verhalten bei der Verlagerung in CUDA ⇐ C++

[Anonymous]

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>

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