Willkommen bei einer GPU-Berechnung!
Die Praxis des Berechnens von Hashes auf Grafikkarten, auch bekannt als GPU (Graphics Processing Unit) Hashing oder GPU-Mining, gewann mit dem Aufstieg von Kryptowährungen, insbesondere Bitcoin, an Popularität. Beim Bitcoin-Mining geht es darum, komplexe mathematische Probleme zu lösen, um Transaktionen zu validieren und das Netzwerk zu sichern. Anfänglich wurde das Bitcoin-Mining auf zentralen Verarbeitungseinheiten (CPUs) durchgeführt, aber als die Schwierigkeit des Minings zunahm, wurden CPUs ineffizient und durch GPUs ersetzt.
Menschen begannen um das Jahr 2010 damit, auf Grafikkarten Hashes zu berechnen, um Kryptowährungen wie Bitcoin effizienter zu minen. Grafikkarten, oder GPUs, sind hochgradig parallele Prozessoren, die gleichzeitig viele Berechnungen durchführen können, was sie ideal für die rechenintensive Natur des Minings macht. Durch den Einsatz von GPUs konnten Miner mehr Hash-Berechnungen pro Sekunde verarbeiten, was ihre Chancen erhöhte, Mining-Belohnungen zu verdienen.
Die Nutzung von GPUs geht über ihre parallelen Verarbeitungsfähigkeiten hinaus. Diese robusten Rechengeräte besitzen eine einzigartige Kombination von Eigenschaften, die sie besonders effizient in der Ausführung bestimmter Aufgaben im Vergleich zu CPUs macht:
-
Parallelismus: CPUs, begrenzt in ihren Kernen, betreiben Multitasking. Im Gegensatz dazu orchestrieren GPUs mit ihrer Vielzahl agiler Kerne eine große Symphonie der parallelen Verarbeitung und ermöglichen die gleichzeitige Ausführung einer breiten Palette von Aufgaben mit Finesse.
-
Spezialisierung: GPUs sind architektonisch für spezifische Aufgaben optimiert, wobei einige Modelle spezialisierte Komponenten wie Tensor-Kerne für maschinelles Lernen oder Ray-Tracing-Einheiten für realistisches Rendering enthalten. Diese zweckgebundenen Designs bieten deutliche Leistungsvorteile gegenüber allgemeinen CPUs.
-
Lastverteilung: Indem rechenintensive Aufgaben auf GPUs ausgelagert werden, können sich CPUs auf ihre Stärken konzentrieren, wie z.B. die Verwaltung von Systemprozessen und Benutzereingaben. Diese harmonische Arbeitsteilung führt zu einem reaktionsschnelleren und effizienteren Gesamtsystem.
-
Energieeffizienz: GPUs sind bei Aufgaben, die ihrem Design entsprechen, überlegen und erreichen eine höhere Anzahl von Berechnungen pro Watt im Vergleich zu CPUs. Diese Energieeffizienz ist besonders wertvoll in Umgebungen wie groß angelegten Rechenzentren oder Einrichtungen für Hochleistungsrechnen, wo Energieverbrauch ein wichtiges Anliegen ist.
-
Datenverarbeitung: Mit verbesserter Speicherbandbreite können GPUs größere Datensätze geschickt verarbeiten und bieten Vorteile bei Aufgaben wie Bildverarbeitung, Simulationen und umfassender Datenanalyse.
-
Datenlokalität: GPUs verfügen über dedizierten Speicher (VRAM), was eine verbesserte Datenlokalität und reduzierte Latenz fördert. Dieser dedizierte Speicher verbessert die Leistung bei spezifischen Berechnungen.
-
Softwarebibliotheken: Entwickler können optimierte Softwarebibliotheken und Frameworks wie CUDA für allgemeines GPU-Computing, cuDNN für Deep Learning und OpenCL für heterogenes Computing nutzen, um die GPU-Leistung nahtlos zu nutzen.
-
Heterogenes Computing: Die Auslagerung von Aufgaben an GPUs erleichtert eine nahtlose Integration der Fähigkeiten von CPU und GPU, was zu effizienteren und leistungsfähigeren Systemen führt.
-
Skalierbarkeit: Für Aufgaben, die erhebliche Leistungssteigerungen durch GPUs erfahren, wie maschinelles Lernen oder Simulationen, kann die Erweiterung der GPU-Fähigkeiten kosteneffektiver und skalierbarer sein als die Erhöhung der Anzahl der CPU-Kerne.
Vorbereitung der Umgebung
Schritt 1: Visual Studio installieren
-
Wenn Sie es noch nicht getan haben, laden Sie Visual Studio von der offiziellen Visual Studio-Website herunter und installieren Sie es (https://visualstudio.microsoft.com/).
-
Stellen Sie sicher, dass Sie die Arbeitslast "Desktop-Entwicklung mit C++" während der Installation von Visual Studio installieren, da die CUDA-Entwicklung C++ Entwicklungswerkzeuge erfordert.
Schritt 2: CUDA Toolkit installieren
-
Gehen Sie auf die NVIDIA CUDA-Webseite (https://developer.nvidia.com/cuda-toolkit) und laden Sie die neueste Version des CUDA Toolkits herunter, die mit Ihrer GPU und Ihrem Betriebssystem kompatibel ist.
-
Führen Sie den Installer des CUDA Toolkits aus und folgen Sie den Anweisungen auf dem Bildschirm, um das CUDA Toolkit auf Ihrem System zu installieren.
Schritt 3: Konfigurieren Sie Visual Studio für CUDA
-
Öffne Visual Studio und gehe zu „Erweiterungen“ > „Erweiterungen verwalten“.
-
Suche im Dialogfeld „Erweiterungen und Updates“ nach „CUDA“ und installiere die Erweiterung „NVIDIA CUDA Toolkit“.
-
Starte Visual Studio neu, nachdem die Erweiterung installiert wurde.
-
Nach dem Neustart gehe zu „CUDA“ > „NVIDIA Nsight“ > „Optionen“ im Visual Studio-Menü, um die Optionsseite von „NVIDIA Nsight“ zu öffnen.
-
Im „CUDA“-Reiter gib den Pfad zum Installationsordner des CUDA Toolkits an, den du in Schritt 2 installiert hast.
-
Klicke auf „OK“, um die Einstellungen zu speichern.
Schritt 4: Ein CUDA-Projekt erstellen
-
In Visual Studio gehe zu "Datei" > "Neu" > "Projekt", um ein neues Projekt zu erstellen.
-
Wähle "CUDA" unter "Installiert" > "Vorlagen" > "Visual C++" > "NVIDIA" im Dialogfeld "Neues Projekt" aus.
-
Wähle eine CUDA-Projektvorlage aus, wie zum Beispiel "CUDA Laufzeitprojekt" oder "CUDA Treiberprojekt", und klicke auf "Weiter".
-
Gib den Projektnamen, den Speicherort und andere Einstellungen nach Wunsch an und klicke auf "Erstellen", um das CUDA-Projekt zu erstellen.
Schritt 5: CUDA-Code schreiben und ausführen
1. Im CUDA-Projekt können Sie CUDA-Code in den ".cu" Quelldateien schreiben, die kompiliert und auf der GPU ausgeführt werden können.
2. Um das CUDA-Projekt zu erstellen und auszuführen, wählen Sie die gewünschte Konfiguration (z.B. "Debug" oder "Release") und klicken Sie auf den Button "Lokaler Windows-Debugger" in der Symbolleiste von Visual Studio.
3. Visual Studio wird das CUDA-Projekt erstellen und ausführen, und Sie können die Ausgabe anzeigen und den CUDA-Code mit dem Visual Studio-Debugger debuggen.
Schritt 6. CUDA-Debugger einrichten
-
Visual Studio öffnen: Starten Sie Visual Studio auf Ihrem System.
-
Erstellen Sie ein neues CUDA-Projekt oder öffnen Sie ein bestehendes: Erstellen Sie ein neues Projekt, indem Sie "Datei" -> "Neu" -> "Projekt" -> "CUDA" in Visual Studio auswählen oder ein vorhandenes CUDA-Projekt öffnen.
-
Projekteigenschaften festlegen: Klicken Sie mit der rechten Maustaste auf Ihr CUDA-Projekt im Lösungsexplorer und wählen Sie "Eigenschaften" aus dem Kontextmenü aus.
-
Wählen Sie die Debug-Konfiguration: Navigieren Sie im Fenster der Projekteigenschaften zum Abschnitt "Konfigurationseigenschaften" und wählen Sie die "Debug"-Konfiguration aus.
-
Einstellungen des CUDA-Debuggers konfigurieren:
-
Stellen Sie sicher, dass die Registerkarte "CUDA Debugger" im linken Bereich ausgewählt ist.
-
Wählen Sie im Dropdown-Menü "Debugger-Typ" den "NVIDIA CUDA Debugger" aus.
-
Überprüfen Sie, ob der "Debugger-Pfad" auf den korrekten Ort der ausführbaren CUDA-Debugger-Datei zeigt (z. B. "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\vX.Y\extras\Visual Studio Integration\cuda_debugger.exe"). Passen Sie den Pfad bei Bedarf an, basierend auf Ihrem Installationsverzeichnis und der Version des CUDA Toolkits.
-
-
Setzen Sie Haltepunkte und starten Sie das Debugging: Setzen Sie Haltepunkte in Ihrem CUDA-Code, an denen der Debugger anhalten soll. Drücken Sie dann F5 oder wählen Sie "Debuggen" -> "Debugging starten", um den CUDA-Debugger zu starten und mit dem Debugging Ihres CUDA-Projekts zu beginnen.
-
Debugging-Prozess:
-
Der Debugger wird an den von Ihnen gesetzten Haltepunkten in Ihrem CUDA-Code anhalten, sodass Sie Variablen inspizieren, den Code schrittweise durchgehen und das Verhalten des Programms analysieren können.
-
Sie können die verschiedenen Debugging-Funktionen von Visual Studio nutzen, wie z. B. Über Zeilen schrittweise weitergehen, in Funktionen eintreten, Variablen inspizieren und Aufrufstapel anzeigen.
-
Hinweis: Stellen Sie sicher, dass Sie die passende Version des CUDA Toolkits installiert haben und dass Ihre GPU das Debugging unterstützt. Außerdem vergewissern Sie sich, dass Sie die notwendigen CUDA-Projekteinstellungen und Konfigurationen korrekt für das Debugging eingerichtet haben.
Aufwärmen: GPU-basierte parallele CRC32
/*
* 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;
}
Training: GPU-basiertes paralleles SHA512!
__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;
}
Fazit
Sie können den Code des gesamten Projekts auf unserem github finden
Wir schätzen Ihre Unterstützung und freuen uns auf Ihre weiterhin aktive Beteiligung in unserer Gemeinschaft
Alle Fragen an die Autoren des Artikels können per E-Mail gesendet werden an: articles@stofu.io
Danke für Ihre Aufmerksamkeit und haben Sie einen schönen Tag!
