Bienvenue dans un calcul GPU !
La pratique du calcul de hash sur des cartes vidéo, également connue sous le nom de hachage GPU (Graphics Processing Unit) ou minage GPU, a gagné en popularité avec l'essor des cryptomonnaies, particulièrement Bitcoin. Le minage de Bitcoin consiste à résoudre des problèmes mathématiques complexes pour valider des transactions et sécuriser le réseau. Initialement, le minage de Bitcoin était effectué sur des unités de traitement central (CPU), mais à mesure que la difficulté de minage augmentait, les CPU devenaient inefficaces et étaient remplacés par des GPU.
Les gens ont commencé à calculer des hashes sur des cartes vidéo vers 2010 pour miner des cryptomonnaies telles que Bitcoin plus efficacement. Les cartes vidéo, ou GPU, sont des processeurs hautement parallèles qui peuvent effectuer de nombreuses calculs simultanément, ce qui les rend idéaux pour la nature computationnellement intensive du minage. L'utilisation de GPU a permis aux mineurs de traiter plus de calculs de hash par seconde, augmentant ainsi leurs chances de gagner des récompenses de minage.
L'utilisation des GPU s'étend au-delà de leurs capacités de traitement parallèle. Ces dispositifs de calcul robustes possèdent une combinaison distinctive de caractéristiques qui les rendent extrêmement compétents pour exécuter certaines tâches plus efficacement que les CPU :
Parallélisme : Les CPU, limités dans leurs cœurs, se livrent au multitâche. En revanche, les GPU, dotés de leur multitude de cœurs agiles, orchestrent une grande symphonie de traitement parallèle, permettant l'exécution simultanée d'une large gamme de tâches avec finesse.
Spécialisation : Les GPU sont optimisés architecturalement pour des tâches spécifiques, certains modèles comportant des composants spécialisés comme les cœurs tensor pour l'apprentissage automatique ou les unités de ray-tracing pour un rendu réaliste. Ces conceptions dédiées offrent des avantages de performance notables par rapport aux CPU à usage général
Équilibrage de charge : En déléguant les tâches intensives en calcul aux GPU, les CPU peuvent se concentrer sur leurs points forts, tels que la gestion des processus systèmes et des entrées utilisateur. Cette division du travail harmonieuse résulte en un système global plus réactif et efficace.
Efficacité énergétique : Les GPU excellent dans les tâches adaptées à leur conception, réalisant un nombre plus élevé de calculs par watt de puissance comparé aux CPU. Cette efficacité énergétique est particulièrement précieuse dans des environnements comme les centres de données à grande échelle ou les installations de calcul haute performance où la consommation d'énergie est une préoccupation majeure.
Traitement des données : Avec une bande passante mémoire améliorée, les GPU gèrent habilement de plus grands ensembles de données, offrant un avantage dans des tâches telles que le traitement d'images, les simulations et l'analyse de données extensive.
Localité des données : Les GPU possèdent une mémoire dédiée (VRAM), favorisant une meilleure localité des données et réduisant la latence. Cette mémoire dédiée améliore les performances dans des calculs spécifiques.
Bibliothèques logicielles : Les développeurs peuvent tirer parti de bibliothèques logicielles et de cadres optimisés tels que CUDA pour le calcul GPU généraliste, cuDNN pour l'apprentissage profond et OpenCL pour le calcul hétérogène, permettant une utilisation sans heurts de la puissance des GPU.
Calcul hétérogène : La délégation de tâches aux GPU facilite une intégration transparente des capacités des CPU et des GPU, résultant en des systèmes plus efficaces et performants.
Scalabilité : Pour les tâches qui bénéficient de gains de performance substantiels grâce aux GPU, comme l'apprentissage automatique ou les simulations, augmenter les capacités des GPU peut être plus rentable et plus évolutif que d'augmenter le nombre de cœurs CPU.
Préparation de l'environnement
Étape 1 : Installer Visual Studio
Si vous ne l'avez pas déjà fait, téléchargez et installez Visual Studio depuis le site officiel de Visual Studio (https://visualstudio.microsoft.com/).
Assurez-vous d'installer la charge de travail "Développement de bureau avec C++" lors de l'installation de Visual Studio, car le développement CUDA nécessite des outils de développement C++.
Étape 2 : Installer le CUDA Toolkit
-
Rendez-vous sur le site NVIDIA CUDA (https://developer.nvidia.com/cuda-toolkit) et téléchargez la dernière version du CUDA Toolkit compatible avec votre GPU et votre système d'exploitation.
-
Exécutez l'installateur du CUDA Toolkit et suivez les instructions à l'écran pour installer le CUDA Toolkit sur votre système.
Étape 3 : Configurer Visual Studio pour CUDA
Ouvrez Visual Studio et allez dans "Extensions" > "Gérer les extensions".
Recherchez "CUDA" dans la boîte de dialogue Extensions et Mises à jour, et installez l'extension "NVIDIA CUDA Toolkit".
Redémarrez Visual Studio après l'installation de l'extension
Après le redémarrage, allez dans "CUDA" > "NVIDIA Nsight" > "Options" dans le menu de Visual Studio pour ouvrir la page des options de "NVIDIA Nsight".
Dans l'onglet "CUDA", spécifiez le chemin vers le dossier d'installation de CUDA Toolkit que vous avez installé à l'étape 2.
Cliquez sur "OK" pour enregistrer les paramètres
Étape 4 : Créer un projet CUDA
Dans Visual Studio, allez dans "Fichier" > "Nouveau" > "Projet" pour créer un nouveau projet.
Sélectionnez "CUDA" sous "Installé" > "Modèles" > "Visual C++" > "NVIDIA" dans la boîte de dialogue Nouveau Projet.
Choisissez un modèle de projet CUDA, tel que "Projet d'exécution CUDA" ou "Projet de pilote CUDA", et cliquez sur "Suivant".
Spécifiez le nom du projet, l'emplacement et les autres paramètres selon vos souhaits, et cliquez sur "Créer" pour créer le projet CUDA.
Ouvrez Visual Studio : Lancez Visual Studio sur votre système.
Créez un nouveau projet CUDA ou ouvrez-en un existant : Créez un nouveau projet en sélectionnant "Fichier" -> "Nouveau" -> "Projet" -> "CUDA" dans Visual Studio, ou ouvrez un projet CUDA existant.
Définissez les propriétés du projet : Cliquez droit sur votre projet CUDA dans l'Explorateur de solutions et sélectionnez "Propriétés" dans le menu contextuel.
Sélectionnez la configuration de débogage : Dans la fenêtre des propriétés du projet, naviguez jusqu'à la section "Propriétés de configuration" et sélectionnez la configuration "Debug".
Configurez les paramètres du débogueur CUDA :
Assurez-vous que l'onglet "CUDA Debugger" est sélectionné dans le volet gauche.
Dans le menu déroulant "Type de débogueur", sélectionnez "NVIDIA CUDA Debugger".
Vérifiez que le "Chemin du débogueur" pointe vers l'emplacement correct de l'exécutable du débogueur CUDA (par exemple, "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\vX.Y\extras\Visual Studio Integration\cuda_debugger.exe"). Ajustez le chemin si nécessaire, en fonction du répertoire d'installation et de la version de votre Toolkit CUDA.
Placez des points d'arrêt et commencez le débogage : Placez des points d'arrêt dans votre code CUDA là où vous souhaitez que le débogueur s'arrête. Ensuite, appuyez sur F5 ou sélectionnez "Déboguer" -> "Démarrer le débogage" pour lancer le débogueur CUDA et commencer le débogage de votre projet CUDA.
Processus de débogage :
Le débogueur s'arrêtera aux points d'arrêt que vous avez définis dans votre code CUDA, vous permettant d'inspecter les variables, de parcourir le code et d'analyser le comportement du programme.
Vous pouvez utiliser les différentes fonctionnalités de débogage fournies par Visual Studio, comme passer au-dessus des lignes, entrer dans les fonctions, inspecter les variables, et visualiser les piles d'appels.
Étape 5 : Écrire et exécuter le code CUDA
1. Dans le projet CUDA, vous pouvez écrire du code CUDA dans les fichiers sources ".cu", qui peuvent être compilés et exécutés sur le GPU.
2. Pour construire et exécuter le projet CUDA, sélectionnez la configuration désirée (par exemple, "Debug" ou "Release") et cliquez sur le bouton "Local Windows Debugger" dans la barre d'outils de Visual Studio.
3. Visual Studio va compiler et exécuter le projet CUDA, et vous pouvez voir le résultat et déboguer le code CUDA en utilisant le débogueur de Visual Studio.
Étape 6. Configuration du débogueur CUDA
Note : Assurez-vous d'avoir installé la version appropriée du CUDA Toolkit et que votre GPU prend en charge le débogage. Assurez-vous également que vous disposez des paramètres et des configurations de projet CUDA nécessaires correctement configurés pour le débogage.
Échauffement : CRC32 parallélisé basé sur GPU
/*
* FUNCTION: __device__ __host__ uint32_t crc32
*
* ARGS:
* const uint8_t* buffer - Input buffer containing data for CRC32 calculation.
* int size - Size of the input buffer.
*
* DESCRIPTION:
* This function calculates the CRC32 checksum for a given input buffer on both CPU and GPU
* devices.
* The CRC32 calculation algorithm used is the standard CRC32 polynomial with initial value of
* 0xFFFFFFFF and final XOR of 0xFFFFFFFF.
* The function iterates through each byte in the input buffer using a for loop, performing bitwise XOR
* and shift operations to calculate the CRC32 value.
* The calculated CRC32 value is then bitwise negated (~crc) and returned as the final result.
*
* RETURN VALUE:
* uint32_t - The calculated CRC32 checksum for the input buffer.
* This function returns a 32-bit unsigned integer representing the CRC32 checksum value.
*/
__device__ __host__ uint32_t crc32(const uint8_t* buffer, int size)
{
uint32_t crc = 0xFFFFFFFF;
for (int i = 0; i < size; ++i)
{
crc ^= buffer[i];
for (int j = 0; j < 8; ++j)
crc = (crc >> 1) ^ (0xEDB88320 & (-(crc & 1)));
}
return ~crc;
}
/*
* FUNCTION: __global__ void crc32Kernel
*
* ARGS:
* In uint8_t* buffers - Input buffer containing data for CRC32 calculation.
* Out uint32_t* crcResults - Output buffer to store CRC32 results.
* int numBuffers - Number of input buffers.
* int bufferSize - Size of each input buffer.
*
* DESCRIPTION:
* This is a CUDA kernel function for calculating CRC32 checksums in parallel on a GPU device.
* Each thread in the GPU grid corresponds to a unique thread identifier (tid) calculated from blockIdx.x and blockDim.x.
* The bufferIndex is calculated based on tid and bufferSize to determine the starting index of the current buffer to be processed.
* The function performs CRC32 calculation on each buffer by iterating through each byte in the buffer using a for loop.
* The calculated CRC32 value is then saved to the crcResults array at the corresponding tid index.
*
* RETURN VALUE: void
* This function does not return a value.
*/
__global__ void crc32Kernel(_In_ const uint8_t* buffers, _Out_ uint32_t* crcResults, int numBuffers, int bufferSize)
{
/* Calculate unique thread identifier */
int tid = blockIdx.x * blockDim.x + threadIdx.x;
/* Calculate index of the current buffer */
int bufferIndex = tid * bufferSize;
/* Check if buffer index is within valid range */
if (bufferIndex < numBuffers * bufferSize)
/* Call crc32 function to calculate CRC32 for the current buffer */
crcResults[tid] = crc32(buffers + bufferIndex, bufferSize);
}
/*
* FUNCTION: std::vector< uint32_t > testCRC32CPU
*
* ARGS:
* std::vector< std::vector< uint8_t >> const& buffers - A vector of input buffers to calculate CRC32 checksums.
*
* DESCRIPTION:
* This function calculates the CRC32 checksum for buffers of random data on the GPU using CUDA.
* It dynamically аllocates memory on the device (GPU) for the buffers and CRC32 results.
* The function launches a CUDA kernel on the device to calculate the CRC32 checksum for each buffer in parallel.
* It then copies the results back from the device to the host and frees the allocated memory.
*
* RETURN VALUE: std::vector
* CRC32 checksums for each buffer in the input vector on the GPU.
*/
std::vector< uint32_t > testCRC32GPU(std::vector< std::vector< uint8_t >> const& buffers)
{
const int numBuffers = buffers.size();
const int bufferSize = buffers[0].size();
/* Dynamic memory allocation on the device */
unsigned char* d_buffers;
uint32_t* d_crcResults;
cudaMalloc(reinterpret_cast< void** >(&d_buffers), numBuffers * bufferSize * sizeof(unsigned char));
cudaMalloc(reinterpret_cast< void** >(&d_crcResults), numBuffers * sizeof(uint32_t));
/* Copy data from host to device using cudaMemcpy2D */
for (int i = 0; i < numBuffers; ++i)
cudaMemcpy(d_buffers + i * bufferSize, buffers[i].data(), bufferSize * sizeof(unsigned char), cudaMemcpyHostToDevice);
/* Calculate number of blocks and threads per block for the kernel launch */
const int blockSize = 256;
const int numBlocks = (numBuffers + blockSize - 1) / blockSize;
/* Launch the kernel on the device indicating the number of blocks(numBlocks) and block size(blockSize) that will be used for parallel execution of calculations on the GPU. */
crc32::crc32Kernel << < numBlocks, blockSize >> > (d_buffers, d_crcResults, numBuffers, bufferSize);
/* Copy results from device to host directly into a vector without intermediate buffer */
std::vector< uint32_t > checksums(numBuffers);
cudaMemcpy(checksums.data(), d_crcResults, numBuffers * sizeof(uint32_t), cudaMemcpyDeviceToHost);
/* Free device memory */
cudaFree(d_buffers);
cudaFree(d_crcResults);
/* Return the CRC32 checksums as a vector */
return checksums;
}
Entraînement : SHA512 parallélisé basé sur GPU !
__device__ static const uint64_t K[80] = {
UINT64_C(0x428a2f98d728ae22), UINT64_C(0x7137449123ef65cd),
UINT64_C(0xb5c0fbcfec4d3b2f), UINT64_C(0xe9b5dba58189dbbc),
UINT64_C(0x3956c25bf348b538), UINT64_C(0x59f111f1b605d019),
UINT64_C(0x923f82a4af194f9b), UINT64_C(0xab1c5ed5da6d8118),
UINT64_C(0xd807aa98a3030242), UINT64_C(0x12835b0145706fbe),
UINT64_C(0x243185be4ee4b28c), UINT64_C(0x550c7dc3d5ffb4e2),
UINT64_C(0x72be5d74f27b896f), UINT64_C(0x80deb1fe3b1696b1),
UINT64_C(0x9bdc06a725c71235), UINT64_C(0xc19bf174cf692694),
UINT64_C(0xe49b69c19ef14ad2), UINT64_C(0xefbe4786384f25e3),
UINT64_C(0x0fc19dc68b8cd5b5), UINT64_C(0x240ca1cc77ac9c65),
UINT64_C(0x2de92c6f592b0275), UINT64_C(0x4a7484aa6ea6e483),
UINT64_C(0x5cb0a9dcbd41fbd4), UINT64_C(0x76f988da831153b5),
UINT64_C(0x983e5152ee66dfab), UINT64_C(0xa831c66d2db43210),
UINT64_C(0xb00327c898fb213f), UINT64_C(0xbf597fc7beef0ee4),
UINT64_C(0xc6e00bf33da88fc2), UINT64_C(0xd5a79147930aa725),
UINT64_C(0x06ca6351e003826f), UINT64_C(0x142929670a0e6e70),
UINT64_C(0x27b70a8546d22ffc), UINT64_C(0x2e1b21385c26c926),
UINT64_C(0x4d2c6dfc5ac42aed), UINT64_C(0x53380d139d95b3df),
UINT64_C(0x650a73548baf63de), UINT64_C(0x766a0abb3c77b2a8),
UINT64_C(0x81c2c92e47edaee6), UINT64_C(0x92722c851482353b),
UINT64_C(0xa2bfe8a14cf10364), UINT64_C(0xa81a664bbc423001),
UINT64_C(0xc24b8b70d0f89791), UINT64_C(0xc76c51a30654be30),
UINT64_C(0xd192e819d6ef5218), UINT64_C(0xd69906245565a910),
UINT64_C(0xf40e35855771202a), UINT64_C(0x106aa07032bbd1b8),
UINT64_C(0x19a4c116b8d2d0c8), UINT64_C(0x1e376c085141ab53),
UINT64_C(0x2748774cdf8eeb99), UINT64_C(0x34b0bcb5e19b48a8),
UINT64_C(0x391c0cb3c5c95a63), UINT64_C(0x4ed8aa4ae3418acb),
UINT64_C(0x5b9cca4f7763e373), UINT64_C(0x682e6ff3d6b2b8a3),
UINT64_C(0x748f82ee5defb2fc), UINT64_C(0x78a5636f43172f60),
UINT64_C(0x84c87814a1f0ab72), UINT64_C(0x8cc702081a6439ec),
UINT64_C(0x90befffa23631e28), UINT64_C(0xa4506cebde82bde9),
UINT64_C(0xbef9a3f7b2c67915), UINT64_C(0xc67178f2e372532b),
UINT64_C(0xca273eceea26619c), UINT64_C(0xd186b8c721c0c207),
UINT64_C(0xeada7dd6cde0eb1e), UINT64_C(0xf57d4f7fee6ed178),
UINT64_C(0x06f067aa72176fba), UINT64_C(0x0a637dc5a2c898a6),
UINT64_C(0x113f9804bef90dae), UINT64_C(0x1b710b35131c471b),
UINT64_C(0x28db77f523047d84), UINT64_C(0x32caab7b40c72493),
UINT64_C(0x3c9ebe0a15c9bebc), UINT64_C(0x431d67c49c100d4c),
UINT64_C(0x4cc5d4becb3e42b6), UINT64_C(0x597f299cfc657e2a),
UINT64_C(0x5fcb6fab3ad6faec), UINT64_C(0x6c44198c4a475817)
};
/* Various logical functions for calculating sha-512 hash on GPU */
#define ROR64c(x, y) \
( ((((x)&UINT64_C(0xFFFFFFFFFFFFFFFF))>>((uint64_t)(y)&UINT64_C(63))) | \
((x)<<((uint64_t)(64-((y)&UINT64_C(63)))))) & UINT64_C(0xFFFFFFFFFFFFFFFF))
#define STORE64H(x, y) \
{ (y)[0] = (unsigned char)(((x)>>56)&255); (y)[1] = (unsigned char)(((x)>>48)&255); \
(y)[2] = (unsigned char)(((x)>>40)&255); (y)[3] = (unsigned char)(((x)>>32)&255); \
(y)[4] = (unsigned char)(((x)>>24)&255); (y)[5] = (unsigned char)(((x)>>16)&255); \
(y)[6] = (unsigned char)(((x)>>8)&255); (y)[7] = (unsigned char)((x)&255); }
#define LOAD64H(x, y) \
{ x = (((uint64_t)((y)[0] & 255))<<56)|(((uint64_t)((y)[1] & 255))<<48) | \
(((uint64_t)((y)[2] & 255))<<40)|(((uint64_t)((y)[3] & 255))<<32) | \
(((uint64_t)((y)[4] & 255))<<24)|(((uint64_t)((y)[5] & 255))<<16) | \
(((uint64_t)((y)[6] & 255))<<8)|(((uint64_t)((y)[7] & 255))); }
#define Ch(x,y,z) (z ^ (x & (y ^ z)))
#define Maj(x,y,z) (((x | y) & z) | (x & y))
#define S(x, n) ROR64c(x, n)
#define R(x, n) (((x) &UINT64_C(0xFFFFFFFFFFFFFFFF))>>((uint64_t)n))
#define Sigma0(x) (S(x, 28) ^ S(x, 34) ^ S(x, 39))
#define Sigma1(x) (S(x, 14) ^ S(x, 18) ^ S(x, 41))
#define Gamma0(x) (S(x, 1) ^ S(x, 8) ^ R(x, 7))
#define Gamma1(x) (S(x, 19) ^ S(x, 61) ^ R(x, 6))
#ifndef MIN
#define MIN(x, y) ( ((x)<(y))?(x):(y) )
#endif
/*
* FUNCTION: static int __device__ __host__ sha512_compress
*
* ARGS:
* sha512_context* md - Pointer to the SHA-512 context structure.
* unsigned char* buf - Pointer to the buffer containing the data to be compressed.
*
* DESCRIPTION:
* This function performs the compression step of the SHA-512 algorithm on a block of data.
* It performs the following steps:
* - Copies the current state values from the SHA-512 context (md) into local variables (S).
* - Copies the input data block (buf) into an array of 80 64-bit unsigned integers (W).
* - Fills the remaining elements of W[16..79] using bitwise operations and additions as per the SHA-512 algorithm.
* - Performs a series of 80 rounds of SHA-512 operations (RND macro) on the state variables (S) and elements of W.
* - Updates the state variables (md->state) by adding the values from the local variables (S).
* This function is marked as static, which means it can only be accessed within the same source file. It can be called from both device (GPU) and host (CPU) code, as denoted by the __device__ and __host__ qualifiers.
*
* RETURN VALUE: int
* Returns 0 on success, and a non-zero value if any error occurs (currently not used in the function).
*/
static int __device__ __host__ sha512_compress(sha512_context* md, unsigned char* buf)
{
uint64_t S[8], W[80], t0, t1;
int i;
/* copy state into S */
for (i = 0; i < 8; i++)
S[i] = md->state[i];
/* copy the state into 1024-bits into W[0..15] */
for (i = 0; i < 16; i++)
LOAD64H(W[i], buf + (8 * i));
/* fill W[16..79] */
for (i = 16; i < 80; i++)
W[i] = Gamma1(W[i - 2]) + W[i - 7] + Gamma0(W[i - 15]) + W[i - 16];
/* Compress */
#define RND(a,b,c,d,e,f,g,h,i) \
t0 = h + Sigma1(e) + Ch(e, f, g) + K[i] + W[i]; \
t1 = Sigma0(a) + Maj(a, b, c);\
d += t0; \
h = t0 + t1;
for (i = 0; i < 80; i += 8) {
RND(S[0], S[1], S[2], S[3], S[4], S[5], S[6], S[7], i + 0);
RND(S[7], S[0], S[1], S[2], S[3], S[4], S[5], S[6], i + 1);
RND(S[6], S[7], S[0], S[1], S[2], S[3], S[4], S[5], i + 2);
RND(S[5], S[6], S[7], S[0], S[1], S[2], S[3], S[4], i + 3);
RND(S[4], S[5], S[6], S[7], S[0], S[1], S[2], S[3], i + 4);
RND(S[3], S[4], S[5], S[6], S[7], S[0], S[1], S[2], i + 5);
RND(S[2], S[3], S[4], S[5], S[6], S[7], S[0], S[1], i + 6);
RND(S[1], S[2], S[3], S[4], S[5], S[6], S[7], S[0], i + 7);
}
#undef RND
for (i = 0; i < 8; i++)
md->state[i] = md->state[i] + S[i];
return 0;
}
/*
* FUNCTION: int __device__ __host__ sha512_init
*
* ARGS:
* sha512_context* md - Pointer to the SHA-512 context structure.
*
* DESCRIPTION:
* This function initializes the SHA-512 context by setting the initial state values for the SHA-512 hash calculation.
* It performs the following steps:
* - Checks for a NULL pointer for the input SHA-512 context, which is an error condition.
* - Sets the buffer length (curlen) and original message length (length) in the context to 0.
* - Sets the initial state values (8 64-bit unsigned integers) in the context as per the SHA-512 algorithm specifications.
* This function can be called from both device (GPU) and host (CPU) code, as denoted by the __device__ and __host__ qualifiers.
*
* RETURN VALUE: int
* Returns 0 on success, and a non-zero value if any error occurs (e.g., NULL pointer for the input context).
*/
int __device__ __host__ sha512_init(sha512_context* md)
{
if (md == NULL) return 1;
md->curlen = 0;
md->length = 0;
md->state[0] = UINT64_C(0x6a09e667f3bcc908);
md->state[1] = UINT64_C(0xbb67ae8584caa73b);
md->state[2] = UINT64_C(0x3c6ef372fe94f82b);
md->state[3] = UINT64_C(0xa54ff53a5f1d36f1);
md->state[4] = UINT64_C(0x510e527fade682d1);
md->state[5] = UINT64_C(0x9b05688c2b3e6c1f);
md->state[6] = UINT64_C(0x1f83d9abfb41bd6b);
md->state[7] = UINT64_C(0x5be0cd19137e2179);
return 0;
}
/*
* FUNCTION: int __device__ __host__ sha512_update
*
* ARGS:
* sha512_context* md - Pointer to the SHA-512 context structure.
* const uint8_t* in - Pointer to the input message buffer.
* size_t inlen - Length of the input message buffer.
*
* DESCRIPTION:
* This function updates the SHA-512 hash calculation with additional input data. It processes the input data in blocks of 128 bytes and updates the SHA-512 context accordingly.
* It performs the following steps:
* - Checks for NULL pointers for the input SHA-512 context and input message buffer.
* - Checks if the current length of the message buffer in the context is greater than the size of the buffer, which is an error condition.
* - Processes the input data in blocks of 128 bytes:
* - If the current length of the message buffer in the context is 0 and the input data length is greater than or equal to 128 bytes, it directly compresses the input data using sha512_compress() function, updates the length of the original message, and advances the input data buffer and length.
* - Otherwise, it copies the input data to the message buffer in the context until the buffer is full (128 bytes):
* - If the buffer is full, it compresses the buffer using sha512_compress() function, updates the length of the original message, and resets the buffer length.
* - Continues this process until all the input data is processed.
* This function can be called from both device (GPU) and host (CPU) code, as denoted by the __device__ and __host__ qualifiers.
*
* RETURN VALUE: int
* Returns 0 on success, and a non-zero value if any error occurs.
*/
int __device__ __host__ sha512_update(sha512_context* md, const uint8_t* in, size_t inlen)
{
size_t n;
int err;
/* Check if input parameters are valid */
if (md == NULL) return 1;
if (in == NULL) return 1;
if (md->curlen > sizeof(md->buf)) return 1;
/* Process input data in blocks of HASH_SIZE bytes */
while (inlen > 0)
{
/* If there is enough input data and buffer is empty, directly compress the input data */
if (md->curlen == 0 && inlen >= HASH_SIZE)
{
if ((err = sha512_compress(md, (unsigned char*)in)) != 0) return err;
md->length += HASH_SIZE * 8;
in += HASH_SIZE;
inlen -= HASH_SIZE;
}
else
{
/* Copy input data to buffer until it is full or input data is exhausted */
n = MIN(inlen, (HASH_SIZE - md->curlen));
for (size_t i = 0; i < n; ++i)
md->buf[i + md->curlen] = in[i];
md->curlen += n;
in += n;
inlen -= n;
/* If buffer is full, compress it */
if (md->curlen == HASH_SIZE) {
if ((err = sha512_compress(md, md->buf)) != 0) return err;
md->length += 8 * HASH_SIZE;
md->curlen = 0;
}
}
}
return 0;
}
/*
* FUNCTION: int __device__ __host__ sha512_final
*
* ARGS:
* sha512_context* md - Pointer to the SHA-512 context structure.
* uint8_t* out - Pointer to the output buffer for storing the final SHA-512 hash.
*
* DESCRIPTION:
* This function finalizes the SHA-512 hash calculation by padding the input message and storing the calculated hash in the output buffer.
* It performs the following steps:
* - Checks for NULL pointers for the input SHA-512 context and output buffer.
* - Appends the `1` bit to the message buffer.
* - If the length of the message buffer is greater than 112 bytes, it appends zeros and compresses the buffer.
* - Appends zeros to the message buffer until it reaches a length of 120 bytes.
* - Stores the length of the original message in big-endian format in the last 8 bytes of the buffer.
* - Performs the final compression using sha512_compress() function.
* - Copies the resulting hash from the SHA-512 context to the output buffer.
* This function can be called from both device (GPU) and host (CPU) code, as denoted by the __device__ and __host__ qualifiers.
*
* RETURN VALUE: int
* Returns 0 on success, and a non-zero value if any error occurs.
*/
int __device__ __host__ sha512_final(sha512_context* md, uint8_t* out)
{
/* Check if input parameters are valid */
if (md == NULL) return 1;
if (out == NULL) return 1;
if (md->curlen >= sizeof(md->buf)) return 1;
/* increase the length of the message */
md->length += md->curlen * UINT64_C(8);
/* append the '1' bit */
md->buf[md->curlen++] = (unsigned char)0x80;
/* if the length is currently above 112 bytes append zeros then compress. Then can fall back to padding zeros and length encoding like normal */
if (md->curlen > 112) {
while (md->curlen < HASH_SIZE)
md->buf[md->curlen++] = (unsigned char)0;
sha512_compress(md, md->buf);
md->curlen = 0;
}
while (md->curlen < 120)
md->buf[md->curlen++] = (unsigned char)0;
/* store length */
STORE64H(md->length, md->buf + 120);
sha512_compress(md, md->buf);
/* copy output */
for (int i = 0; i < 8; i++)
STORE64H(md->state[i], out + (8 * i));
return 0;
}
/*
* FUNCTION: int __device__ __host__ sha512
*
* ARGS:
* const uint8_t* message - Pointer to the input message whose SHA-512 hash needs to be calculated.
* size_t length - Length of the input message.
* uint8_t* out - Pointer to the output buffer for storing the calculated SHA-512 hash.
*
* DESCRIPTION:
* This function calculates the SHA-512 hash for the input message using the sha512_context structure and associated functions.
* It initializes the sha512_context using sha512_init() function, updates the context with the input message using sha512_update() function, and finalizes the context to obtain the SHA-512 hash using sha512_final() function.
* The calculated hash is stored in the output buffer pointed to by `out`.
* This function can be called from both device (GPU) and host (CPU) code, as denoted by the __device__ and __host__ qualifiers.
*
* RETURN VALUE: int
* Returns the status of the SHA-512 calculation, where 0 indicates success, and any other value indicates an error.
*/
int __device__ __host__ sha512(const uint8_t* message, size_t length, uint8_t* out)
{
sha512_context ctx;
int status;
if ((status = sha512_init(&ctx))) return status;
if ((status = sha512_update(&ctx, message, length))) return status;
if ((status = sha512_final(&ctx, out))) return status;
return status;
}
/*
* FUNCTION: std::string __host__ sha512
*
* ARGS:
* const uint8_t* message - Pointer to the input message whose SHA-512 hash needs to be calculated.
* size_t length - Length of the input data.
*
* DESCRIPTION:
* This function calculates the SHA-512 hash of the input data using a GPU-based implementation.
* It performs the following steps:
* - Initializes a SHA-512 context structure (ctx) from the sha512GPU namespace.
* - Updates the context with the input data using sha512GPU::sha512_update() function.
* - Finalizes the hash calculation and stores the resulting digest in a local buffer (digest) using sha512GPU::sha512_final() function.
* - Converts the digest from binary to hexadecimal representation and stores it in a string buffer (buf) using sprintf() function.
* - Returns the calculated SHA-512 hash as a string.
* This function is marked with __host__ qualifier, which means it can be called from host (CPU) code, but not from device (GPU) code.
*
* RETURN VALUE: std::vector< uint8_t >
* Returns the calculated std::vector< uint8_t > as a hexadecimal bytes.
*/
std::vector< uint8_t > __host__ sha512(const uint8_t* message, size_t length)
{
std::vector< uint8_t > digest(DIGEST_SIZE);
hashes::sha512_context ctx;
int status;
if ((status = hashes::sha512_init(&ctx))) return digest;
if ((status = hashes::sha512_update(&ctx, message, length))) return digest;
if ((status = hashes::sha512_final(&ctx, digest.data()))) return digest;
return digest;
}
/*
* FUNCTION: void __global__ sha512Kernel
*
* ARGS:
* char* inputs - Pointer to the input buffers in GPU memory.
* int numInputs - Number of input buffers to process.
* uint8_t* outputs - Pointer to the output buffer in GPU memory for storing the calculated SHA-512 hashes.
* size_t bufferSize - Size of each input buffer.
* int bufferLength - Length of each input buffer.
* This function is meant to be called from host code and executed on the GPU using CUDA.
*
* DESCRIPTION:
* This CUDA kernel function is launched on the GPU to calculate the SHA-512 hashes for the input buffers in parallel.
* It calculates the global thread ID using blockIdx.x and threadIdx.x, and checks if the thread ID is within bounds of the number of input buffers.
* If the thread ID is within bounds, it calls the sha512() function to calculate the SHA-512 hash for the corresponding input buffer, and stores the result in the output buffer in GPU memory.
*/
void __global__ sha512Kernel(char* inputs, int numInputs, uint8_t* outputs, int bufferLength)
{
/* Calculate global thread ID */
int index = blockIdx.x * blockDim.x + threadIdx.x;
/* Check if thread ID is within bounds and call SHA-512 function */
if (index < numInputs)
sha512((uint8_t*)(inputs + index * bufferLength), bufferLength, outputs + index * DIGEST_SIZE);
}
/*
* TEST FUNCTION: std::vector< std::vector< uint8_t >> sha512BuffersGPU
*
* ARGS:
* const std::vector< std::vector< uint8_t >>& buffers - A vector of input buffers to calculate SHA-512 hashes.
*
* DESCRIPTION:
* This function calculates the SHA-512 hash for a vector of input buffers on the GPU using CUDA parallel processing.
* It allocates GPU memory for input and output buffers, copies input buffers from host to GPU memory, and launches a CUDA kernel function to perform the hash calculation.
* The results are then copied back from GPU to host memory using CUDA streams for faster copying.
* Finally, the function converts the hash results from binary to hexadecimal string format and returns them as a vector of strings.
*
* RETURN VALUE: std::vector< std::vector< uint8_t >>
* A vector of SHA-512 hashes for the input buffers.
*/
std::vector< std::vector< uint8_t >> testSHA512GPU(const std::vector< std::vector< uint8_t >>& buffers)
{
int numInputs = buffers.size();
/* Size of each input buffer (assuming all strings have the same size) */
size_t bufferSize = buffers[0].size();
int bufferLength = static_cast< int >(bufferSize);
/* Create and copy input buffers to GPU memory */
char* d_inputs;
cudaMalloc((void**)&d_inputs, numInputs * bufferLength);
for (int i = 0; i < numInputs; ++i)
cudaMemcpy(d_inputs + i * bufferLength, buffers[i].data(), bufferLength, cudaMemcpyHostToDevice);
unsigned char* d_outputs;
/* 128 - size of SHA-512 hash in bytes */
cudaMalloc((void**)&d_outputs, numInputs * hashes::DIGEST_SIZE);
/* Calculate grid size and block size for CUDA threads */
const int blockSize = 256;
const int gridSize = (numInputs + blockSize - 1) / blockSize;
/* Call the sha512Kernel CUDA kernel function on GPU to calculate hashes for each input buffer and save results into output buffer */
hashes::sha512Kernel << < gridSize, blockSize >> > (d_inputs, numInputs, d_outputs, bufferLength);
/* Allocate memory on host for results */
std::vector< std::vector< uint8_t >> results(numInputs);
/* Allocate memory on host for output buffer */
std::vector< unsigned char > h_outputs(numInputs * hashes::DIGEST_SIZE);
/* Create CUDA stream for faster copying */
cudaStream_t stream;
cudaStreamCreate(&stream);
/* Copy results using CUDA stream */
cudaMemcpyAsync(h_outputs.data(), d_outputs, numInputs * hashes::DIGEST_SIZE, cudaMemcpyDeviceToHost, stream);
/* Synchronize CUDA stream to complete copying */
cudaStreamSynchronize(stream);
/* Copy results to vector of vectors */
for (int i = 0; i < numInputs; ++i) {
results[i].resize(hashes::DIGEST_SIZE);
memcpy(results[i].data(), h_outputs.data() + i * hashes::DIGEST_SIZE, hashes::DIGEST_SIZE);
}
/* Free GPU memory */
cudaFree(d_inputs);
cudaFree(d_outputs);
/* Destroy CUDA stream */
cudaStreamDestroy(stream);
return results;
}
Conclusion
Vous pouvez trouver le code du projet entier sur notre github
Nous apprécions votre soutien et attendons avec impatience votre engagement continu dans notre communauté
Toutes les questions adressées aux auteurs de l'article peuvent être envoyées par e-mail à : articles@stofu.io
Merci de votre attention et passez une bonne journée !
