SHA512 salted hash from mkpasswd doesn't match an online version

Question

How nvcc handle const pointers in kernels?

According to nvidia, adding const and restrict for pointers during parameter-passing enable NVCC for aggressive optimizations, is this strictly following the C/C++ way?

Assuming A is pointer pointed to a buffer of data which will be freqenently updated by maybe other threads/streams, but the contents will not get modified during this test kernel call:

test<<<blocks, threads>>>(const int *__restrict__ A, int *__restrict__ B);

Then can NVCC maintain correctness of this: load the updated data in A at each kernel call, instead of loading some pre-cached out-dated data?

Solution

On Ubuntu/Debian mkpasswd is part of the package whois and implemented in mkpasswd.c which as actually just a sophisticated wrapper around the crypt() function in glibc declared in unistd.h. crypt() takes two arguments password and salt. Password is "test" in this case, salt is prepended by "$6$" for the SHA-512 hash (see SHA-crypt) so "$6$Zem197T4" is passed to crypt().

Maybe you noticed the -R option of mkpasswd which determines the number of rounds. In the document you'll find a default of 5000 rounds. This is the first hint why the result would never be equal to the simple concatenation of salt and password, it's not hashed only once. Actually if you pass -R 5000 you get the same result. In this case "$6$rounds=5000$Zem197T4" is passed to crypt() and the implementation in glibc (which is the libc of Debian/Ubuntu) extracts the method and number of rounds from this.

What happens inside crypt() is more complicated than just computing a single hash and the result is base64 encoded in the end. That's why the result you showed contains all kinds of characters after the last '$' and not only [0-9a-f] as in the typical hex string of a SHA-512 hash. The algorithm is described in detail in the already mentioned SHA-Crypt document.

Comments

[Ben Voigt]

C++ doesn't even have a restrict keyword, so it is following C99 rather than C++. But generally, potential aliasing kills optimization.

[Ben Voigt]

Also, if the data is updated from other threads, shouldn't you have made it volatile or include a memory barrier in some other way? That's certainly required in CPU-side C and C++.

[user0002128]

@BenVoigt: the data is modified by other threads but wont be modified at and during the time of this kernel call, I dont think I need volatile here (In my case, there is a global flag across all threads to control this, but thats CPU side stuff), as for restrict, it is a NVCC recongized keyword.

[njuffa]

const means "treat this object as read-only in this scope". Compilers can usually detect usage that conflicts with this modifier. restrict means "this pointer is the only path through which the pointed to object will be accessed in this scope". It is an assertion the programmer makes and that a compiler in general cannot verify. If the programmer breaks their promise (i.e. there is in fact aliasing) the code is bound not to work as the programmer desired. Since C++ (even in its 2011 incarnation) does not support restricted pointers, CUDA follows C99 but uses the keyword restrict.

[Admin]

mkpasswd is generating hashes in the "modular crypt format": pythonhosted.org/passlib/modular_crypt_format.html