I try to explain that understand
-
Declare the variable as volatile and initialize
-
Copy the content of the variable from host to device with cudaMemcpyToSymbolAsync()
But, how use the variable to stop all threads?
This is my Kernel code:
__global__ void brute_force(unsigned char *texto_dev, int lt, unsigned char *clave_dev, unsigned char *descifrado_dev, unsigned char *tid_dev) {
unsigned char alfabeto[56] = {'a','b','c','d','e','f','g','h','i','j','k','l','m','n','o','p','q','r','s','t','u','v','w','x','y','z',
'A','B','C','D','E','F','G','H','I','J','K','L','M','N','O','P','Q','R','S','T','U','V','W','X','Y','Z',' ',',','.','
global void brute_force(unsigned char *texto_dev, int lt, unsigned char *clave_dev, unsigned char *descifrado_dev, unsigned char *tid_dev) {
unsigned char alfabeto[56] = {'a','b','c','d','e','f','g','h','i','j','k','l','m','n','o','p','q','r','s','t','u','v','w','x','y','z',
‘A’,‘B’,‘C’,‘D’,‘E’,‘F’,‘G’,‘H’,‘I’,‘J’,‘K’,‘L’,‘M’,‘N’,‘O’,‘P’,‘Q’,‘R’,‘S’,‘T’,‘U’,‘V’,‘W’,‘X’,‘Y’,‘Z’,’ ‘,’,‘,’.‘,’\0’};
int y, c1, c2, lc = 5, ii, gen_key = 0, valor;
int id = threadIdx.x;
int N_combinaciones = 10000;
int key_inicial = id * N_combinaciones;
int key_final = key_inicial + N_combinaciones;
/* KEY >> */
for (ii = 0; ii < 5; ii++) {
clave_dev[257 * id + ii] = 'X';
}
clave_dev[257 * id + 5] = ‘\0’;
/* KEY << */
for (gen_key = key_inicial; gen_key < key_final && !encontrado; gen_key++) {
valor = gen_key;
for (ii = 0; ii < 5; ii++) {
clave_dev[257 * id + 4 - ii] = '0' + (valor % 10);
valor = valor / 10;
}
clave_dev[257 * id + lc] = '\0';
rc4_init(id, &clave_dev[257 * id], lc);
for (y = 0; y < lt; y++) {
descifrado_dev[1025 * id + y] = texto_dev[y] ^ rc4_output(id);
}
for (c1 = 0; c1 < lt; c1++) {
for (c2 = 0; c2 < 56; c2++) {
if (descifrado_dev[1025 * id + c1] == alfabeto[c2]) {
break;
}
}
if (c2 == 56) {
break;
}
}
if (c1 == (lt-1)) {
tid_dev[0] = (unsigned char) ('0' + id);
encontrado = 1;
break;
//gen_key = 100000;
}
}
// tid_dev[0] = (unsigned char) (‘0’ + id);
}
'};
int y, c1, c2, lc = 5, ii, gen_key = 0, valor;
int id = threadIdx.x;
int N_combinaciones = 10000;
int key_inicial = id * N_combinaciones;
int key_final = key_inicial + N_combinaciones;
/* KEY >> */
for (ii = 0; ii < 5; ii++) {
clave_dev[257 * id + ii] = 'X';
}
clave_dev[257 * id + 5] = '
global void brute_force(unsigned char *texto_dev, int lt, unsigned char *clave_dev, unsigned char *descifrado_dev, unsigned char *tid_dev) {
unsigned char alfabeto[56] = {'a','b','c','d','e','f','g','h','i','j','k','l','m','n','o','p','q','r','s','t','u','v','w','x','y','z',
‘A’,‘B’,‘C’,‘D’,‘E’,‘F’,‘G’,‘H’,‘I’,‘J’,‘K’,‘L’,‘M’,‘N’,‘O’,‘P’,‘Q’,‘R’,‘S’,‘T’,‘U’,‘V’,‘W’,‘X’,‘Y’,‘Z’,’ ‘,’,‘,’.‘,’\0’};
int y, c1, c2, lc = 5, ii, gen_key = 0, valor;
int id = threadIdx.x;
int N_combinaciones = 10000;
int key_inicial = id * N_combinaciones;
int key_final = key_inicial + N_combinaciones;
/* KEY >> */
for (ii = 0; ii < 5; ii++) {
clave_dev[257 * id + ii] = 'X';
}
clave_dev[257 * id + 5] = ‘\0’;
/* KEY << */
for (gen_key = key_inicial; gen_key < key_final && !encontrado; gen_key++) {
valor = gen_key;
for (ii = 0; ii < 5; ii++) {
clave_dev[257 * id + 4 - ii] = '0' + (valor % 10);
valor = valor / 10;
}
clave_dev[257 * id + lc] = '\0';
rc4_init(id, &clave_dev[257 * id], lc);
for (y = 0; y < lt; y++) {
descifrado_dev[1025 * id + y] = texto_dev[y] ^ rc4_output(id);
}
for (c1 = 0; c1 < lt; c1++) {
for (c2 = 0; c2 < 56; c2++) {
if (descifrado_dev[1025 * id + c1] == alfabeto[c2]) {
break;
}
}
if (c2 == 56) {
break;
}
}
if (c1 == (lt-1)) {
tid_dev[0] = (unsigned char) ('0' + id);
encontrado = 1;
break;
//gen_key = 100000;
}
}
// tid_dev[0] = (unsigned char) (‘0’ + id);
}
';
/* KEY << */
for (gen_key = key_inicial; gen_key < key_final && !encontrado; gen_key++) {
valor = gen_key;
for (ii = 0; ii < 5; ii++) {
clave_dev[257 * id + 4 - ii] = '0' + (valor % 10);
valor = valor / 10;
}
clave_dev[257 * id + lc] = '
global void brute_force(unsigned char *texto_dev, int lt, unsigned char *clave_dev, unsigned char *descifrado_dev, unsigned char *tid_dev) {
unsigned char alfabeto[56] = {'a','b','c','d','e','f','g','h','i','j','k','l','m','n','o','p','q','r','s','t','u','v','w','x','y','z',
‘A’,‘B’,‘C’,‘D’,‘E’,‘F’,‘G’,‘H’,‘I’,‘J’,‘K’,‘L’,‘M’,‘N’,‘O’,‘P’,‘Q’,‘R’,‘S’,‘T’,‘U’,‘V’,‘W’,‘X’,‘Y’,‘Z’,’ ‘,’,‘,’.‘,’\0’};
int y, c1, c2, lc = 5, ii, gen_key = 0, valor;
int id = threadIdx.x;
int N_combinaciones = 10000;
int key_inicial = id * N_combinaciones;
int key_final = key_inicial + N_combinaciones;
/* KEY >> */
for (ii = 0; ii < 5; ii++) {
clave_dev[257 * id + ii] = 'X';
}
clave_dev[257 * id + 5] = ‘\0’;
/* KEY << */
for (gen_key = key_inicial; gen_key < key_final && !encontrado; gen_key++) {
valor = gen_key;
for (ii = 0; ii < 5; ii++) {
clave_dev[257 * id + 4 - ii] = '0' + (valor % 10);
valor = valor / 10;
}
clave_dev[257 * id + lc] = '\0';
rc4_init(id, &clave_dev[257 * id], lc);
for (y = 0; y < lt; y++) {
descifrado_dev[1025 * id + y] = texto_dev[y] ^ rc4_output(id);
}
for (c1 = 0; c1 < lt; c1++) {
for (c2 = 0; c2 < 56; c2++) {
if (descifrado_dev[1025 * id + c1] == alfabeto[c2]) {
break;
}
}
if (c2 == 56) {
break;
}
}
if (c1 == (lt-1)) {
tid_dev[0] = (unsigned char) ('0' + id);
encontrado = 1;
break;
//gen_key = 100000;
}
}
// tid_dev[0] = (unsigned char) (‘0’ + id);
}
';
rc4_init(id, &clave_dev[257 * id], lc);
for (y = 0; y < lt; y++) {
descifrado_dev[1025 * id + y] = texto_dev[y] ^ rc4_output(id);
}
for (c1 = 0; c1 < lt; c1++) {
for (c2 = 0; c2 < 56; c2++) {
if (descifrado_dev[1025 * id + c1] == alfabeto[c2]) {
break;
}
}
if (c2 == 56) {
break;
}
}
if (c1 == (lt-1)) {
tid_dev[0] = (unsigned char) ('0' + id);
encontrado = 1;
break;
//gen_key = 100000;
}
}
// tid_dev[0] = (unsigned char) ('0' + id);
}
My variable encontrado (means “found” in english) is that determine if one thread found the key or not, and is my flag to stop the all others threads.