Welcome to A GPU calculation!
The practice of calculating hashes on video cards, also known as GPU (Graphics Processing Unit) hashing or GPU mining, gained popularity with the rise of cryptocurrencies, particularly Bitcoin. Bitcoin mining involves solving complex mathematical problems to validate transactions and secure the network. Initially, Bitcoin mining was performed on central processing units (CPUs), but as the difficulty of mining increased, CPUs became inefficient and were replaced by GPUs.
People started calculating hashes on video cards around 2010 to mine cryptocurrencies like Bitcoin more efficiently. Video cards, or GPUs, are highly parallel processors that can perform many calculations simultaneously, making them ideal for the computationally intensive nature of mining. Using GPUs allowed miners to process more hash calculations per second, increasing their chances of earning mining rewards.
The utilization of GPUs extends beyond their parallel processing capabilities. These robust computing devices possess a distinctive combination of characteristics that render them highly proficient in executing specific tasks more efficiently than CPUs:
-
Parallelism: CPUs, limited in their cores, engage in multitasking. In contrast, GPUs, with their multitude of agile cores, orchestrate a grand symphony of parallel processing, enabling simultaneous execution of a wide range of tasks with finesse.
-
Specialization: GPUs are architecturally optimized for specific tasks, with some models featuring specialized components like tensor cores for machine learning or ray-tracing units for realistic rendering. These purpose-built designs offer notable performance advantages over general-purpose CPUs.
-
Load Balancing: By offloading compute-intensive tasks to GPUs, CPUs can focus on their strengths, such as managing system processes and user input. This harmonious division of labor results in a more responsive and efficient overall system.
-
Energy Efficiency: GPUs excel at tasks suited to their design, achieving a higher number of calculations per watt of power compared to CPUs. This energy efficiency is particularly valuable in environments like large-scale data centers or high-performance computing facilities where energy consumption is a prominent concern.
-
Data Processing: With enhanced memory bandwidth, GPUs adeptly handle larger datasets, delivering an advantage in tasks such as image processing, simulations, and extensive data analysis.
-
Data Locality: GPUs possess dedicated memory (VRAM), promoting improved data locality and reduced latency. This dedicated memory enhances performance in specific calculations.
-
Software Libraries: Developers can leverage optimized software libraries and frameworks such as CUDA for general-purpose GPU computing, cuDNN for deep learning, and OpenCL for heterogeneous computing, enabling seamless utilization of GPU power.
-
Heterogeneous Computing: Task offloading to GPUs facilitates a seamless integration of CPU and GPU capabilities, resulting in more efficient and high-performing systems.
-
Scalability: For tasks that derive substantial performance gains from GPUs, such as machine learning or simulations, expanding GPU capabilities can be more cost-effective and scalable than increasing the number of CPU cores.
Preparing the environment
Step 1: Install Visual Studio
-
If you haven't already, download and install Visual Studio from the official Visual Studio website (https://visualstudio.microsoft.com/).
-
Make sure to install the "Desktop development with C++" workload during the Visual Studio installation, as CUDA development requires C++ development tools.
Step 2: Install CUDA Toolkit
-
Go to the NVIDIA CUDA website (https://developer.nvidia.com/cuda-toolkit) and download the latest version of CUDA Toolkit that is compatible with your GPU and operating system.
-
Run the CUDA Toolkit installer and follow the on-screen instructions to install CUDA Toolkit on your system.
Step 3: Configure Visual Studio for CUDA
-
Open Visual Studio and go to "Extensions" > "Manage Extensions".
-
Search for "CUDA" in the Extensions and Updates dialog box, and install the "NVIDIA CUDA Toolkit" extension.
-
Restart Visual Studio after the extension is installed.
-
After restarting, go to "CUDA" > "NVIDIA Nsight" > "Options" in the Visual Studio menu to open the "NVIDIA Nsight" options page.
-
In the "CUDA" tab, specify the path to the CUDA Toolkit installation folder that you installed in Step 2.
-
Click "OK" to save the settings.
Step 4: Create a CUDA Project
-
In Visual Studio, go to "File" > "New" > "Project" to create a new project.
-
Select "CUDA" under "Installed" > "Templates" > "Visual C++" > "NVIDIA" in the New Project dialog box.
-
Choose a CUDA project template, such as "CUDA Runtime Project" or "CUDA Driver Project", and click "Next".
-
Specify the project name, location, and other settings as desired, and click "Create" to create the CUDA project.
Step 5: Write and Run CUDA Code
1. In the CUDA project, you can write CUDA code in the ".cu" source files, which can be compiled and executed on the GPU.
2. To build and run the CUDA project, select the desired configuration (e.g., "Debug" or "Release") and click the "Local Windows Debugger" button in the Visual Studio toolbar.
3. Visual Studio will build and run the CUDA project, and you can view the output and debug CUDA code using the Visual Studio debugger.
Step 6. Setup CUDA debugger
-
Open Visual Studio: Launch Visual Studio on your system.
-
Create a new CUDA project or open an existing one: Either create a new project by selecting "File" -> "New" -> "Project" -> "CUDA" in Visual Studio or open an existing CUDA project.
-
Set project properties: Right-click on your CUDA project in the Solution Explorer and select "Properties" from the context menu.
-
Select the Debug configuration: In the project properties window, navigate to the "Configuration Properties" section and select the "Debug" configuration.
-
Configure CUDA Debugger settings:
-
Ensure that the "CUDA Debugger" tab is selected in the left-hand pane.
-
In the "Debugger Type" dropdown, select "NVIDIA CUDA Debugger".
-
Verify that the "Debugger Path" points to the correct location of the CUDA debugger executable (e.g., "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\vX.Y\extras\Visual Studio Integration\cuda_debugger.exe"). Adjust the path if necessary, based on your CUDA Toolkit installation directory and version.
-
-
Set breakpoints and start debugging: Place breakpoints in your CUDA code where you want the debugger to halt. Then, press F5 or select "Debug" -> "Start Debugging" to launch the CUDA debugger and start debugging your CUDA project.
-
Debugging process:
-
The debugger will halt at the breakpoints you set in your CUDA code, allowing you to inspect variables, step through the code, and analyze the program's behavior.
-
You can use the various debugging features provided by Visual Studio, such as stepping over lines, stepping into functions, inspecting variables, and viewing call stacks.
-
Note: Ensure that you have installed the appropriate CUDA Toolkit version and that your GPU supports debugging. Also, make sure you have the necessary CUDA project settings and configurations properly configured for debugging.
Warm-up: GPU-based paralleled 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;
}
Workout: GPU-based paralleled 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;
}
Conclusion
You can find the code of the entire project on our github
We appreciate your support and look forward to your continued engagement in our community
Any questions of the authors of the article can be sent to the e-mail: articles@stofu.io
Thank you for your attention and have a nice day!
