Blog - 0x00

A CUDA SHAllenge

2025/01/01

Tabla de contenidos

Hace unos meses, el usuario “quirino” publicó en HackerNews (HN) un interesante shallenge. Encuentra el valor de hash SHA256 más bajo para una cadena con un formato determinado:

Envía una cadena con el formato “{username}/{nonce}”, donde:

• username: de 1 a 32 caracteres de a-zA-Z0-9_-
• nonce: de 1 a 64 caracteres en Base64 (a-zA-Z0-9+/)

Cuando se publicó esto, quise participar, competir con otros usuarios, hacer que mi código funcionara lo más rápido posible y ver si tenía suerte encontrando un valor de hash bajo.

En esta entrada del blog describiré los diferentes intentos que hice para realizar el hash más rápido, cómo pasé de 1,5 megahashes/s a 18,8 gigahashes/s y compartiré mi solución final.

CPU #

Mi primer intento fue usando mi CPU. Para ser sincero, ni siquiera pensé en usar una tarjeta gráfica (GPU), ya que nunca lo había hecho antes y, bueno, se llama tarjeta gráfica, no tarjeta de hash.

Para dar algo de contexto, la CPU que estoy usando en esta parte es una Intel i5-8250U (4c, 8t), una CPU móvil.

Un solo hilo #

Comencé escribiendo una primera implementación en Go. Es muy simple: genera una cadena aleatoria, calcula su valor de hash usando el paquete crypto y comprueba si es menor que el valor de hash almacenado.

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33

package main

import (
	"crypto/sha256"
	"encoding/hex"
	"fmt"
	"math/rand"
)

const charset = "ABCEDFGHIJKLMNOPQRSTUVWXYZabcdefghijklmnopqrstuvwxyz0123456789+/"

func generate_random_string() string {
	b := make([]byte, 18)
	for i := range b {
		b[i] = charset[rand.Int63()%int64(64)]
	}
	return string(b)
}

func main() {
	lowest_hash := "00000000ffffffffffffffffffffffffffffffffffffffffffffffffffffffff"
	username := "0x00cl/Hello+HN/"
	for {
		random_string := generate_random_string()
		string_to_hash := username + random_string
		hash := sha256.Sum256([]byte(string_to_hash))
		hash_string := hex.EncodeToString(hash[:])
		if hash_string < lowest_hash {
			lowest_hash = hash_string
			fmt.Println(lowest_hash)
		}
	}
}

Ejecutar esto me dio unos 1,5 megahashes/s (MHs). Cuanto más larga era la cadena aleatoria, más bajo era el rendimiento de hash. En ese momento, pensé que no había mucho que mejorar en el código, ya que era bastante directo.

Volví a la publicación de HN para revisar los comentarios y ver si alguien había publicado su solución y cuán cerca estaba la mía de las suyas, pero muy pocos usuarios habían compartido la suya; la mayoría afirmaba tener tasas de hash mucho más altas que las que yo podía alcanzar. Al mismo tiempo, era difícil comparar diferentes CPUs, especialmente porque yo estaba usando una CPU móvil más antigua y los usuarios de HN usaban CPUs de escritorio más recientes.

Al leer los comentarios, se me ocurrió la idea de crear una solución multihilo, lo que debería mejorar la tasa de hash. Aunque nunca había intentado crear un programa multihilo y no quería molestarme en modificar el código, pensé en un truco mucho más simple: lanzar múltiples instancias del programa.

$ for i in {1..4}; do ./shallenge >> hash.log &; done
[2] 12275
[3] 12276
[4] 12277
[5] 12282

en teoría, la cantidad de hashes ahora sería x4 (6 MHs). Probé la tasa de hash para ver si había mejorado 4 veces como pensaba, pero solo mejoró muy poco: estaba obteniendo 2,1 MHs en total. Aunque parecía haber un cuello de botella en algún lugar, no estoy seguro de si tenía que ver con la función generadora de números aleatorios o con algo más del sistema.

Esto me pareció un truco y no un código realmente escrito para multihilo, así que decidí intentar escribir una versión multihilo para ver si mejoraba la tasa de hash.

Multihilo #

El código de la solución que usa multihilo está aquí en Codeberg (git). No quiero copiar y pegar un gran bloque de código, así que solo mencionaré una parte del código en este blog.

Como nunca había probado el multihilo antes, pensé que sería más complicado, pero afortunadamente las goroutines de Go lo hacen mucho más fácil. Al principio necesité un poco de ayuda y usé modelos de lenguaje grandes (LLM) para empezar, pero una vez que entendí cómo funcionaban, pude modificarlo por mi cuenta.

Antes de que comentes algo sobre que las goroutines no son verdaderamente paralelas sino concurrentes, soy consciente de esto sobre las goroutines.

“Una goroutine tiene un modelo simple: es una función que se ejecuta de forma concurrente con otras goroutines en el mismo espacio de direcciones… Las goroutines se multiplexan en múltiples hilos del sistema operativo, de modo que si una se bloquea, como mientras espera una E/S, las demás continúan ejecutándose. Su diseño oculta muchas de las complejidades de la creación y gestión de hilos.”

Fuente: https://go.dev/doc/effective_go#goroutines

De nuevo, no hay nada especial en el código; la única diferencia es que, en lugar de usar print cuando encuentra un valor de hash más bajo, lo envía a través de un canal para que se imprima con una marca de tiempo. Los cambios en el código se centraron principalmente en el manejo de las goroutines.

58
59
60
61
62
63
64
65
66
67
68

	for {
		random_string := generate_random_string()
		string_to_hash := username + random_string
		hash := sha256.Sum256([]byte(string_to_hash))
		hash_string := hex.EncodeToString(hash[:])
		if hash_string < *lowest_hash {
			*lowest_hash = hash_string
			t := time.Now()
			result_chan <- "[" + t.Format("2006-01-02T15:04:05 -070000") + "] " + hash_string + " (" + string_to_hash + ")"
		}
	}

Aunque solo fuera concurrente, pensé que obtendría una tasa de hash más alta y quizás algo cercano a lo que había calculado en teoría, pero me decepcioné cuando lo probé y vi que, como mucho, alcanzaba 4,1 MHs. Noté que 3 goroutines era el número con el que obtenía la tasa de hash más alta; usar un valor mayor o menor la reducía.

Una vez más, aunque lograra paralelizar realmente el proceso de hash, no obtendría una tasa mucho mejor que la de otros usuarios, principalmente debido a mi CPU, así que más o menos renuncié a optimizarlo y lo dejé ejecutándose durante unas horas para ver si podía obtener un resultado suficientemente bueno.

Logré encontrar un valor de hash que comenzaba con 9 ceros para la cadena “0x00cl/i5+8250U/Hi/HN/OinJeJj/eam3VRRI9d” (00000000 03b98cfc f91b9b3e 63735491 6d9cb960 63837375 8d810723 49bcb778).

GPU (CUDA) #

Al ver la tabla de líderes, se podía notar que la mayoría de los usuarios en los primeros puestos usaban GPUs para obtener sus hashes, y los usuarios en la publicación de HN también se dieron cuenta de esto. En ese momento, no tenía mucho interés en CUDA, especialmente porque probablemente estaba haciendo otras cosas y, aunque tengo una tarjeta gráfica con CUDA habilitado (GTX 650 Ti BOOST (768 núcleos CUDA)), ni siquiera la consideré porque era una tarjeta antigua con la que había intentado hacer cosas en el pasado, solo para descubrir que no era compatible y que necesitaba una tarjeta gráfica más moderna. Por ejemplo, el GPU Passthrough, para poder usar una distribución de Linux y jugar en una VM de Windows, no era compatible; o prácticamente cualquier software que usara CUDA requería una tarjeta gráfica más moderna para poder utilizarse.

Después de unos meses sin tocar/ejecutar el proyecto, recordé el shallenge y quise intentarlo de nuevo, ya que había bajado algunas posiciones en la tabla de líderes. Este desarrollo se hizo en Windows, usando la misma tarjeta gráfica para la pantalla.

Configuración del kit de herramientas CUDA #

La documentación es mala. Era la primera vez que usaba CUDA y, por supuesto, descargué la última versión de CUDA desde su sitio web (12.6). La instalé e intenté compilar usando nvcc, pero obtuve un error que decía que cl.exe (MSVC) faltaba en el PATH, lo cual fue culpa mía, ya que la documentación sí mencionaba que MSVC era un requisito. Así que instalé Visual Studio 2022 Community Edition, descargué las herramientas de C++ y compilé este programa simple:

#include <stdio.h>

__global__ void hello_kernel() {
	printf("Hello, GPU!");
}

int main() {
	hello_kernel<<<1,1>>>();
	return 0;
}

Se compiló sin errores, generó un ejecutable y lo ejecuté, pero por alguna razón, cuando corrí el programa, no imprimía nada. Probé con otro código simple, pero a veces lanzaba un SIGSEGV (error de segmentación).

Intenté varias cosas, pero ninguna funcionó. Por un momento pensé que mi GPU simplemente tenía algún tipo de error que impedía la ejecución de programas CUDA. Seguí intentando averiguar cuál podría ser el problema. Revisé las opciones de nvcc para la depuración y una de ellas llamó mi atención: -arch. Tenía una opción “nativa” que detectaba tu GPU y generaba código para ella, así que intenté compilar con -arch=native, pero me dio un error indicando que la capacidad de cómputo (CC) de mi GPU no era compatible.

Ese era el problema. Por defecto, compilaba para tarjetas gráficas con CC 5.2 y, al ejecutarlo, no mostraba ningún error. Supuse que el CUDA Toolkit sería compatible con cualquier tarjeta gráfica habilitada para CUDA, pero las diferentes versiones del CUDA Toolkit solo admiten un conjunto específico de CC. En el caso de la versión 12.6, es compatible con CC de 5.0 a 9.0, algo que no se menciona explícitamente en ningún lugar de la documentación ni como un requisito. La única forma de saber si es compatible con tu CC es buscando en la documentación y revisando las tablas donde se mencionan múltiples CC, como en esta sección de instrucciones aritméticas.

Como la CC de mi tarjeta es 3.0, no es compatible con CUDA Toolkit 12.6, por lo que tuve que consultar la documentación de versiones anteriores, donde sí se mencionaba la compatibilidad con CC 3.0. La versión más reciente de CUDA Toolkit que admite la CC de mi tarjeta es la 10.2. Así que descargué e instalé esa versión, compilé el programa y me dio un error: la versión de Visual Studio no era compatible (la 2022 era demasiado nueva); solo se admitían versiones entre 2015 y 2019 (suspiro). Entonces fui a descargar VS 2019, pero Microsoft lo había puesto detrás de una suscripción a MSDN para poder descargarlo. Afortunadamente, encontré esta respuesta en StackOverflow donde Microsoft todavía mantenía un enlace no listado para descargar VS 2019 desde sus servidores. Lo instalé, compilé y ejecuté el programa CUDA, y finalmente logré que imprimiera algo.

Primera solución #

En mi primer intento, realmente no estaba seguro de lo que estaba haciendo. Simplemente veía la GPU como una CPU: en lugar de tener 4 núcleos, ahora tenía 768.

250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293

// Used this implementation of SHA256
// https://github.com/mochimodev/cuda-hashing-algos/blob/master/sha256.cu

// For RNG I used this as __device__ code
// https://www.pcg-random.org/

__global__ kernel_sha256_hash()
{
	BYTE in[65];
	BYTE out[32];
	BYTE charset[65] = "ABCDEFGHIJKLMNOPQRSTUVWXYZabcdefghijklmnopqrstuvwxyz0123456789+/";
	memcpy(in, "0x00cl/", 7);

	// Setup RNG
	int rounds = 5;
	pcg32_random_t rng;
	pcg32_srandom_r(&rng, 29102, (intptr_t)&rounds);
	// Run indefinitely
	for (;;) {
		// Generate a random string and calculate hash
		for (int i = 7; i < 64; i++ {
			in[i] = charset[pcg32_boundedrand_r(&rng, 64)]
		}
		CUDA_SHA256_CTX ctx;
		cuda_sha256_init(&ctx);
		cuda_sha256_update(&ctx, in, 64);
		cuda_sha256_final(&ctx, out);

		// Check if hash starts with at least 8 zeros.
		if (out[0] == 0 && out[1] == 0 && out[2] == 0 && out[3] == 0){
			for(int j = 0; i< 32; j++) {
				printf("%02x", out[j]);
			}
			printf("\n%s\n", in);
		}
	}
}

int main()
{
	// GPU has 768 CUDA cores, lets use them all, right?
	kernel_sha256_hash<<<1, 768>>>();
	return 0;
}

Esta solución tiene varios problemas:

• Lanzar el kernel con 1 bloque y 768 hilos (<<<1, 768>>>) no es la forma correcta de hacerlo. Por lo que entiendo (a menos que el planificador de la GPU haga alguna magia detrás de escena), cada bloque se asigna a un único Multiprocesador de Transmisión (SM). Mi GPU tiene 4 SM con 192 núcleos cada uno, para un total de 768 núcleos CUDA, lo que significaba que solo se le asignaba este bloque de 768 hilos a un único SM (192 núcleos), sin aprovechar todos los núcleos, ya que había asumido que se trataba de un único bloque grande de núcleos.
• Generar una cadena aleatoria puede ser computacionalmente costoso.
• Por lo que pude encontrar, el modelo de controlador de pantalla de Windows (WDDM) incluye un temporizador de vigilancia que detecta si la pantalla se ha congelado durante más de 2 o 5 segundos; si es así, reinicia el controlador de pantalla, lo que provoca el fallo del programa CUDA. Como creé un bucle infinito dentro del kernel, la pantalla se congelaba durante unos segundos y el controlador de pantalla se reiniciaba.

El problema de la pantalla congelada y el reinicio del controlador de pantalla terminé “resolviéndolo” reemplazando el bucle infinito dentro de la función del kernel por un bucle que realizara 1000 iteraciones, lo cual fue suficiente para que el temporizador de vigilancia no intentara reiniciar el controlador de pantalla.

Aunque no conservo el código exacto, sí volví a recrear lo que hice y realicé pruebas de rendimiento, obteniendo unos 7.2 MHs con 1 bloque y 768 hilos.

En ese momento no sabía cómo funcionaban los bloques y los hilos, pero como tuve que volver a mi primera solución para poder escribir sobre ella, quise probar si usar más bloques mejoraría el rendimiento. Intenté usar 4 bloques y 768 hilos, ya que mi GPU tiene 4 SM, por lo que cada bloque debería asignarse a un SM, y definitivamente noté una mejora en la tasa de hash, alcanzando unos 27 MHs.

Solución con cudaMemcpy #

Después de leer y comprender un poco mejor la arquitectura de las GPU, y en particular las tarjetas NVIDIA, se me ocurrió copiar los datos hacia y desde el dispositivo en lugar de generar cadenas aleatorias en la GPU y verificar su valor de hash. Esto se hacía con frecuencia en los ejemplos de aplicaciones CUDA que encontré.

Mi idea era generar un array de caracteres aleatorios en mi ordenador (host) con una longitud igual al total de hilos más 64 (aunque solo usé 57 caracteres aleatorios, ¡ups!), de modo que cada hilo tuviera su propio idx y que este idx fuera único para cada hilo, variando de 0 al número total de hilos. Así, usaría el idx como un desplazamiento dentro del array de caracteres aleatorios para generar la cadena.

Por ejemplo, si uso cadenas de 5 caracteres para la parte aleatoria y mi array de caracteres aleatorios es este: ['4', 'F', 'J', 'J', '9', 'W', '7', 'd', 'S', 'i', '+', 'V', 'x', 'A'], la cadena aleatoria resultante para cada hilo sería:

• idx = 0: 4FJJ9
• idx = 1: FJJ9W
• idx = 2: JJ9W7
• idx = 3: J9W7d
• …
• idx = 9: i+VxA

250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314

// Used this implementation of SHA256
// https://github.com/mochimodev/cuda-hashing-algos/blob/master/sha256.cu

// For RNG I used this
// https://www.pcg-random.org/

__global__ kernel_sha256_hash(BYTE *rnd_chars, BYTE *rnd_hash_result)
{
	int idx = blockIdx.x * blockDim + threadIdx.x;
	BYTE text[64];
	// Creating the full string with random characters using idx as offset of the array
	memcpy(text, "0x00cl/", 7);
	memcpy(text+7, rnd_chars+idx, 57);
	CUDA_SHA256_CTX ctx;
	cuda_sha256_init(&ctx);
	cuda_sha256_update(&ctx, text, 64);
	cuda_sha256_final(&ctx, rnd_hash_result+idx*32);
}

int main()
{
    BYTE charset[65] = "ABCDEFGHIJKLMNOPQRSTUVWXYZabcdefghijklmnopqrstuvwxyz0123456789+/";
	// Setup RNG
	int rounds = 5;
	pcg32_random_t rng;
	pcg32_srandom_r(&rng, time(NULL) ^ (intptr_t)&printf, (intptr_t)&rounds);
	
	BYTE rnd_chars_h[65+BLOCKS*THREADS_PER_BLOCK];
	BYTE *rnd_chars_d;
	BYTE rnd_hash_result_h[BLOCKS*THREADS_PER_BLOCK*32];
	BYTE *rnd_hash_result_d;

	// Reserve memory and map address between host and device
	cudaMalloc((void **) &rnd_chars_d, (65+BLOCKS*THREADS_PER_BLOCK)*sizeof(BYTE));
	cudaMalloc((void **) &rnd_hash_result_d, BLOCKS*THREADS_PER_BLOCK*32*sizeof(BYTE));

	for(int i = 0; i < 100000000; i++) {
        // Generate an array of random characters as long as total threads plus 64
        // (the amount of random chars in the string)
		for (int j = 0; j < BLOCKS*THREADS_PER_BLOCK+64; i++) {
			rnd_chars_h[j] = charset[pcg32_boundedrand_r(&rng, 64)];
		}
		// Copy to device the array of random characters
		cudaMemcpy(rnd_chars_d, rnd_chars_h, (65+BLOCKS*THREADS_PER_BLOCK)*sizeof(BYTE), cudaMemcpyHostToDevice);

		kernel_sha256_hash<<<16, 384>>>(rnd_chars_h, rnd_hash_result_h);
		// Copy the results and check the number of zeros
		cudaMemcpy(rnd_hash_result_d, rnd_hash_result_h, BLOCKS*THREADS_PER_BLOCK*32*sizeof(BYTE), cudaMemcpyDeviceToHost);

		for(int k = 0; k < BLOCKS*THREADS_PER_BLOCK*32; k = k + 32){
			// Check if hash starts with at least 8 zeros.
			if (out[k] == 0 && out[k+1] == 0 && out[k+2] == 0 && out[k+3] == 0){
				for(int l = k; l < 32; l++) {
					printf("%02x", out[l]);
				}
				printf("\n");
				for(int j = k/32; j < 64+k/32; j++) {
					printf("%c", rnd_chars_d[j]);
				}
				printf("\n");
			}
		}
	}
	return 0;
}

Utilicé nvprof para perfilar y analizar el programa, y me di cuenta de que pasaba gran parte de su tiempo copiando datos desde el host al dispositivo, lo cual no es ideal.

También realicé una prueba de rendimiento de esta solución y obtuve 27,7 MHs, un resultado mejor que mi primer intento, con la diferencia de que la pantalla no se congelaba cada pocos segundos. Quizás podría haber mejorado la tasa de hash si también hubiera intentado usar un bucle dentro de la función del kernel, pero como estaba copiando datos hacia adentro y hacia afuera, habría requerido más memoria asignada; sin embargo, en ese momento no lo intenté.

Evitar cudaMemcpy #

Aun así, no estaba satisfecho con los resultados, aunque no tenía idea de si 27,7 MHs era un buen resultado o no, ya que no tenía con qué compararlo. Además, después de perfilarlo, seguía sintiendo que no era correcto pasar “mucho” tiempo copiando datos hacia o desde el dispositivo.

Así que empecé a pensar en formas de evitar cudaMemcpy(), y fue entonces cuando me di cuenta de que la cadena no necesitaba ser aleatoria. Podía optar por un enfoque secuencial y probar todas las combinaciones posibles, desde “AAAAAA”, “AAAAAB”, “AAAAAC”, …, “aaaaaA”, “aaaaaB”, …, “//////”, y todo esto podría hacerse dentro de la GPU sin tener que copiar datos externos hacia ella.

Esto resolvió en parte el problema de evitar la generación de caracteres aleatorios y la necesidad de copiarlos al dispositivo, pero ahora tenía un nuevo problema: ¿cómo hacerlo de forma secuencial si cada hilo debe obtener una secuencia de cadena única? Una variable que los hilos tienen y que es única para cada uno es idx = blockIdx.x * blockDim + threadIdx.x, y podría usar ese valor para seleccionar un carácter del conjunto de caracteres. Así, con un poco de matemática, podemos obtener una secuencia única para cada hilo.

charset[65] = "ABCDEFGHIJKLMNOPQRSTUVWXYZabcdefghijklmnopqrstuvwxyz0123456789+/";

string[3] = charset[ idx % 64];
string[2] = charset[(idx / 64) % 64];
string[1] = charset[(idx /(64*64)) % 64];
string[0] = charset[(idx /(64*64*64)) % 64];

idx 0 = "AAAA"
    string[3] = charset[ 0 % 64] -> charset[0] = "A";
	string[2] = charset[ 0 % 64] -> charset[0] = "A";
	string[1] = charset[ 0 % 64] -> charset[0] = "A";
	string[0] = charset[ 0 % 64] -> charset[0] = "A";
idx 1 = "AAAB"
	string[3] = charset[ 1 % 64] -> charset[1] = "B";
	string[2] = charset[ 0 % 64] -> charset[0] = "A";
	string[1] = charset[ 0 % 64] -> charset[0] = "A";
	string[0] = charset[ 0 % 64] -> charset[0] = "A";
---
idx 64 = "AABA"
	string[3] = charset[64 % 64] -> charset[0] = "A";
	string[2] = charset[ 1 % 64] -> charset[1] = "B";
	string[1] = charset[ 0 % 64] -> charset[0] = "A";
	string[0] = charset[ 0 % 64] -> charset[0] = "A";
idx 65 = "AABB"
	string[3] = charset[65 % 64] -> charset[1] = "B";
	string[2] = charset[ 1 % 64] -> charset[1] = "B";
	string[1] = charset[ 0 % 64] -> charset[0] = "A";
	string[0] = charset[ 0 % 64] -> charset[0] = "A";
---
idx 16777215 (64^4 - 1) = "////"
	string[3] = charset[16777215 % 64] -> charset[63] = "/";
	string[2] = charset[  262143 % 64] -> charset[63] = "/";
	string[1] = charset[    4095 % 64] -> charset[63] = "/";
	string[0] = charset[      63 % 64] -> charset[63] = "/";

ThoSin embargo, esto planteó otro problema: la cadena que uso como nonce tiene 12 caracteres, lo que significa que el número total de combinaciones posibles es 64^12 (4.722.366.482.869.645.213.696), lo que equivale a 2^72. Esto causaría un desbordamiento incluso en un entero sin signo de 64 bits (2^64) y, aunque NVIDIA añadió soporte para int128 en CUDA 11.6, mi tarjeta gráfica no lo soportaba y, como se emula mediante software, supuse que no sería la mejor opción en términos de rendimiento.

Por lo tanto, lo que hice fue dividir la secuencia de la cadena en dos partes, de modo que cada parte tuviera 64^6 (68.719.476.736) combinaciones, es decir, 2^36, lo cual cabría en un entero sin signo de 64 bits (2^64). Como dividí la secuencia de la cadena, decidí usar hilos y bloques bidimensionales al lanzar el kernel. Esto me permitiría rastrear los valores de los primeros 6 caracteres de la secuencia y los últimos 6 caracteres en dos valores distintos y con una combinación única para cada hilo: idx e idy.

135
136
137
138
139
140
141
142
143

        // Build the sequence for the string to hash.
        data[48] = charset[idx_offset & 63];
        data[54] = charset[idy_offset & 63];
        for (int j = 0; j < 5; j++) {
                idx_offset >>= 6;
                idy_offset >>= 6;
                data[47-j] = charset[idx_offset & 63];
                data[53-j] = charset[idy_offset & 63];
        }

En esta parte del código, estoy construyendo la secuencia de caracteres para el nonce. Es exactamente lo mismo que expliqué en el bloque de código anterior, con la diferencia de que ahora hay dos variables: en este caso, idx_offset e idy_offset, que son los identificadores de cada hilo con un desplazamiento. En lugar de usar operadores aritméticos como la división (/) y el módulo (%), utilizo operadores bit a bit porque, según la documentación, tienen mayor rendimiento que algunas operaciones aritméticas. Además, en su documentación de “Mejores prácticas” se recomienda usar operaciones bit a bit. Afortunadamente, como el conjunto de caracteres tiene una longitud de 64, que es una potencia de 2 (2^6), podemos usar fácilmente operaciones bit a bit, como idx_offset >>= 6, que desplaza los bits hacia la derecha 6 veces, lo que equivale a dividir entre 64. En lugar del módulo (%), se puede usar el “AND bit a bit”, que es equivalente a idx_offset % 64 = idx_offset & 63; ten en cuenta que esto solo funciona con potencias de 2.

Una vez que resolví eso, quise codificar algunos valores de forma fija, especialmente en el cálculo del hash. Sin embargo, es difícil aprovechar al máximo la codificación fija de valores cuando no se entiende cómo funciona SHA256. Pero existe un sitio web que explica y muestra cómo cambian los valores paso a paso en el algoritmo SHA256. Este sitio web definitivamente me ayudó a comprender mejor cómo se realiza.

Después de entender mejor cómo funcionaba SHA256, comencé a codificar valores de forma fija, precalculando los primeros 10 valores de m. Estos valores dependen de tu cadena, pero como solo estoy cambiando los últimos 12 caracteres de mi cadena, que tiene una longitud total de 55 caracteres, estos números permanecen iguales.

54
55
56
57
58
59
60
61

        const uint32_t m_precalculated[] = {813183024,1668034375,1415064374,892349268,1764442703,1330861103,1752462448,1932210296,808463203,1815032368};
        memcpy(m, m_precalculated, sizeof(m_precalculated));

#pragma unroll 4
        for (i = 10, j = 40; i < 14; ++i, j += 4)
                m[i] = (data[j] << 24) | (data[j + 1] << 16) | (data[j + 2] << 8) | (data[j + 3]);
        m[14] = 0;
        m[15] = 440;

Lo mismo ocurre con los valores del estado inicial (a, b, c, d, e, f, g y h) de las diferentes variables utilizadas en las 64 rondas.

Otro cambio que realicé fue que, en lugar de copiar los resultados de vuelta al host, simplemente verifiqué e imprimí los resultados en el dispositivo. Al final, agregué un bucle dentro del kernel que evita parte de la sobrecarga de sincronizar hilos, lanzar el kernel y copiar la cadena completa a hash en una variable.

Cómo funciona #

Este es un resumen de cómo funciona mi solución.

1. Host: Lanza un kernel con hilos y bloques bidimensionales (16 hilos por bloque, 32 bloques), pasando dos desplazamientos: s1, que establece los primeros 6 caracteres del nonce, y s2, que establece los últimos 6 caracteres del nonce.
2. Dispositivo: Calcula el ID de cada hilo, que es una combinación única de idx e idy más el desplazamiento de s1 y s2. En la primera iteración, estos valores van de 0 a 511.
3. Dispositivo: Copia la cadena a hash en una variable data, que ya incluye el relleno necesario: "0x00cl/GTX+650+Ti+BOOST/https+0x00+cl/2025/AAAAAAAAAAAA\x80\x00\x00\x00\x00\x00\x00\x01\xb8".
4. Dispositivo: Modifica data para incluir la secuencia de 12 caracteres usando idx_offset e idy_offset, dividiendo entre 64 para cada posición y eligiendo el carácter del array charset.
5. Dispositivo: Calcula el hash de la cadena (cuda_sha256_transform(data)). Al final, cada valor de los estados (a, b, c, d, e, f, g y h) debe sumarse a un valor de estado inicial para obtener el hash completo de la cadena. Pero como solo buscamos los primeros 11 ceros (el usuario actual en lo más alto de la tabla de clasificación encontró un hash con 15 ceros), y cada estado representa 8 dígitos hexadecimales, solo necesitamos a y b para encontrar los primeros 16 dígitos hexadecimales y verificar si comienzan con ceros. Si es así, imprimirá la cadena y los primeros 16 dígitos hexadecimales del hash.
6. Dispositivo: Repite los pasos 4 y 5, en este caso, 16*32 (512) veces, pero generando una secuencia diferente para cada hilo usando idx_offset e idy_offset en cada bucle.
7. Host: Repite desde el paso 1, lanzando nuevamente el kernel pero con un desplazamiento para s2 de s2 = s2 + 16*32*16*32.

Resultados #

Una vez hecho todo eso, comparé el rendimiento de mi solución con mi GTX 650 Ti BOOST y logré unos 211 MH, lo que supone una gran mejora en comparación con mis intentos anteriores (27,7 MH). En ese punto, parecía que ya no había más margen de optimización para el código actual, o al menos no se me ocurría nada más.

También decidí probarlo en una RTX 4090 para ver qué tal funcionaba y comprobar si podía alcanzar los 21 GH que afirmaba haber logrado el usuario del primer puesto en la clasificación. Mi solución alcanzó los 18,8 GH, lo cual sigue siendo un resultado muy bueno y cercano a los 21 GH, pero aún no llegaba del todo, y era evidente que había margen de mejora.

seletskiy, el usuario que ocupaba el primer puesto con 21 GH, comentó su solución y también la compartió. Sin embargo, por lo que pude entender de su código, estaba haciendo algo similar, aunque él codificaba los caracteres del nonce en base64, mientras que yo utilizo una variable charset que contiene todos los caracteres. Pero como no hay muchas explicaciones sobre el código, me resulta difícil saber exactamente qué atajos tomó para alcanzar los 21 GH.

Reflexiones finales #

Este ha sido, sin duda, un reto divertido en el que he disfrutado trabajando. Dediqué bastante tiempo a entender CUDA y las GPU de Nvidia solo para poder escribir un programa eficiente y aprovechar al máximo una GPU, analizando y pensando en qué enfoque adoptar para el código, realizando pequeños cambios y comprobando si mejoraban el rendimiento. Para mí, lo importante fue el proceso en sí, más que encontrar un valor de hash bajo. Al fin y al cabo, este reto es una combinación de conocimientos, suerte y disponer del dinero o el hardware necesario para realizar los cálculos. Incluso si tu programa no está completamente optimizado para el rendimiento, con una RTX 4090 probablemente funcionará más rápido que en algunas GPU más antiguas, lo que aumenta tus posibilidades de encontrar un valor de hash más bajo.

En cualquier caso, después de todo esto, veo claramente que el rendimiento de la solución con CPU podría mejorarse utilizando caracteres secuenciales en lugar de generarlos aleatoriamente, escribiendo también mi propia implementación de SHA256 con algunos valores codificados directamente y, en lugar de usar goroutines, empleando algo como C para paralelizar realmente el trabajo y utilizar todos los núcleos e hilos de la CPU.

Por otro lado, mi solución con CUDA podría mejorarse si ajustara más la cantidad de bloques en los que se ejecuta para tarjetas gráficas más potentes, aumentara el tiempo que se pasa dentro del kernel realizando cálculos (con bucles más grandes) y, finalmente, optimizara los cálculos para generar una secuencia de cadenas, donde divido enteros, lo cual es un poco más lento que la aritmética de punto flotante en una GPU, y además me falta algo relacionado con la creación de secuencias de cadenas diferentes.

En el repositorio de Git he incluido un archivo llamado shallengeBase.cu, cuyo código es el mismo que el de mi solución, excepto que algunos valores no están codificados directamente ni precalculados. Simplemente puedes modificar la cadena que aparece en la parte superior del archivo (USERNAME_NONCE), compilarlo y ejecutarlo. Este archivo está pensado para servir como punto de partida si quieres empezar a codificar algunos valores directamente para que funcione más rápido y así también puedas encontrar un valor de hash bajo.

Código fuente de mi solución con CUDA: https://codeberg.org/0x00cl/cuda-shallenge