Skip to content
This repository was archived by the owner on Jan 13, 2025. It is now read-only.

Commit 92cfbf7

Browse files
sakridgesakridge
committed
Remove chacha ctr mode
1 parent 416e642 commit 92cfbf7

File tree

6 files changed

+0
-467
lines changed

6 files changed

+0
-467
lines changed

‎src/crypt-if/chacha.h

Lines changed: 0 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -13,21 +13,9 @@ extern "C" {
1313
#define CHACHA_ROUNDS 500
1414
#define SAMPLE_SIZE 32
1515

16-
void chacha20_ctr_encrypt(const uint8_t *in, uint8_t *out, size_t in_len,
17-
const uint8_t key[CHACHA_KEY_SIZE],
18-
const uint8_t nonce[CHACHA_NONCE_SIZE],
19-
uint32_t counter);
20-
2116
void cuda_chacha20_cbc_encrypt(const uint8_t *in, uint8_t *out, size_t in_len,
2217
const uint8_t key[CHACHA_KEY_SIZE], uint8_t* ivec);
2318

24-
void chacha_ctr_encrypt_many(const unsigned char* in, unsigned char* out,
25-
size_t length,
26-
const uint8_t* keys,
27-
const uint8_t* nonces,
28-
uint32_t num_keys,
29-
float* time_us);
30-
3119
void chacha_cbc_encrypt_many(const uint8_t* in, uint8_t* out,
3220
size_t length, const uint8_t *keys,
3321
uint8_t* ivec,

‎src/cuda-crypt/chacha20_core.cu

Lines changed: 0 additions & 57 deletions
Original file line numberDiff line numberDiff line change
@@ -56,60 +56,3 @@ static void __host__ __device__ chacha20_encrypt(const u32 input[16],
5656
U32TO8_LITTLE(output + 4 * i, x[i]);
5757
}
5858
}
59-
60-
void __host__ __device__
61-
chacha20_ctr_encrypt_device(const uint8_t *in, uint8_t *out, size_t in_len,
62-
const uint8_t key[CHACHA_KEY_SIZE],
63-
const uint8_t nonce[CHACHA_NONCE_SIZE],
64-
uint32_t counter)
65-
{
66-
uint32_t input[16];
67-
uint8_t buf[64];
68-
size_t todo, i;
69-
70-
input[0] = U8TO32_LITTLE(sigma + 0);
71-
input[1] = U8TO32_LITTLE(sigma + 4);
72-
input[2] = U8TO32_LITTLE(sigma + 8);
73-
input[3] = U8TO32_LITTLE(sigma + 12);
74-
75-
input[4] = U8TO32_LITTLE(key + 0);
76-
input[5] = U8TO32_LITTLE(key + 4);
77-
input[6] = U8TO32_LITTLE(key + 8);
78-
input[7] = U8TO32_LITTLE(key + 12);
79-
80-
input[8] = U8TO32_LITTLE(key + 16);
81-
input[9] = U8TO32_LITTLE(key + 20);
82-
input[10] = U8TO32_LITTLE(key + 24);
83-
input[11] = U8TO32_LITTLE(key + 28);
84-
85-
input[12] = counter;
86-
input[13] = U8TO32_LITTLE(nonce + 0);
87-
input[14] = U8TO32_LITTLE(nonce + 4);
88-
input[15] = U8TO32_LITTLE(nonce + 8);
89-
90-
while (in_len > 0) {
91-
todo = sizeof(buf);
92-
if (in_len < todo) {
93-
todo = in_len;
94-
}
95-
96-
chacha20_encrypt(input, buf, CHACHA_ROUNDS);
97-
for (i = 0; i < todo; i++) {
98-
out[i] = in[i] ^ buf[i];
99-
}
100-
101-
out += todo;
102-
in += todo;
103-
in_len -= todo;
104-
105-
input[12]++;
106-
}
107-
}
108-
109-
void chacha20_ctr_encrypt(const uint8_t *in, uint8_t *out, size_t in_len,
110-
const uint8_t key[CHACHA_KEY_SIZE],
111-
const uint8_t nonce[CHACHA_NONCE_SIZE],
112-
uint32_t counter)
113-
{
114-
chacha20_ctr_encrypt_device(in, out, in_len, key, nonce, counter);
115-
}

‎src/cuda-crypt/chacha_cbc.cu

Lines changed: 0 additions & 123 deletions
Original file line numberDiff line numberDiff line change
@@ -131,22 +131,6 @@ __global__ void chacha20_cbc128_encrypt_sample_kernel(const uint8_t* input,
131131
}
132132

133133

134-
__global__ void chacha_ctr_encrypt_kernel(const unsigned char* input, unsigned char* output,
135-
size_t length, const uint8_t* keys,
136-
unsigned char* nonces, uint32_t num_keys,
137-
unsigned char* sha_state,
138-
uint32_t* sample_idx,
139-
uint32_t sample_len,
140-
uint32_t block_offset)
141-
{
142-
size_t i = (size_t)(blockIdx.x * blockDim.x + threadIdx.x);
143-
144-
if (i < num_keys) {
145-
chacha20_ctr_encrypt_device(input, &output[i * length], length, &keys[i * CHACHA_KEY_SIZE], &nonces[i * CHACHA_NONCE_SIZE], 0);
146-
}
147-
}
148-
149-
150134
void chacha_cbc_encrypt_many(const unsigned char *in, unsigned char *out,
151135
size_t length, const uint8_t *keys,
152136
uint8_t* ivec,
@@ -427,110 +411,3 @@ void chacha_cbc_encrypt_many_sample(const uint8_t* in,
427411
}
428412

429413

430-
431-
void chacha_ctr_encrypt_many(const unsigned char *in, unsigned char *out,
432-
size_t length,
433-
const uint8_t *keys,
434-
const uint8_t* nonces,
435-
uint32_t num_keys,
436-
float* time_us)
437-
{
438-
if (length < BLOCK_SIZE) {
439-
printf("ERROR! block size(%d) > length(%zu)\n", BLOCK_SIZE, length);
440-
return;
441-
}
442-
uint8_t* in_device = NULL;
443-
uint8_t* in_device0 = NULL;
444-
uint8_t* in_device1 = NULL;
445-
uint8_t* keys_device = NULL;
446-
uint8_t* output_device = NULL;
447-
uint8_t* output_device0 = NULL;
448-
uint8_t* output_device1 = NULL;
449-
uint8_t* nonces_device = NULL;
450-
451-
uint8_t* sha_state_device = NULL;
452-
453-
uint32_t sample_len = 0;
454-
uint32_t* samples_device = NULL;
455-
456-
CUDA_CHK(cudaMalloc(&in_device0, BLOCK_SIZE));
457-
CUDA_CHK(cudaMalloc(&in_device1, BLOCK_SIZE));
458-
459-
size_t keys_size = CHACHA_KEY_SIZE * num_keys;
460-
CUDA_CHK(cudaMalloc(&keys_device, keys_size));
461-
CUDA_CHK(cudaMemcpy(keys_device, keys, keys_size, cudaMemcpyHostToDevice));
462-
463-
size_t nonces_size = CHACHA_NONCE_SIZE * num_keys;
464-
CUDA_CHK(cudaMalloc(&nonces_device, nonces_size));
465-
CUDA_CHK(cudaMemcpy(nonces_device, nonces, nonces_size, cudaMemcpyHostToDevice));
466-
467-
size_t output_size = (size_t)num_keys * (size_t)BLOCK_SIZE;
468-
CUDA_CHK(cudaMalloc(&output_device0, output_size));
469-
CUDA_CHK(cudaMalloc(&output_device1, output_size));
470-
471-
int num_threads_per_block = 64;
472-
int num_blocks = ROUND_UP_DIV(num_keys, num_threads_per_block);
473-
474-
perftime_t start, end;
475-
476-
get_time(&start);
477-
478-
cudaStream_t stream, stream0, stream1;
479-
cudaStreamCreate(&stream0);
480-
cudaStreamCreate(&stream1);
481-
482-
ssize_t slength = length;
483-
size_t num_data_blocks = (length + BLOCK_SIZE - 1) / (BLOCK_SIZE);
484-
485-
LOG("num_blocks: %d threads_per_block: %d keys size: %zu in: %p ind0: %p ind1: %p output_size: %zu num_data_blocks: %zu\n",
486-
num_blocks, num_threads_per_block, keys_size, in, in_device0, in_device1, output_size, num_data_blocks);
487-
488-
for (uint32_t i = 0;; i++) {
489-
//if (i & 0x1) {
490-
if (0) {
491-
in_device = in_device1;
492-
output_device = output_device1;
493-
stream = stream1;
494-
} else {
495-
in_device = in_device0;
496-
output_device = output_device0;
497-
stream = stream0;
498-
}
499-
size_t size = std::min(slength, (ssize_t)BLOCK_SIZE);
500-
//printf("copying to in_device: %p in: %p size: %zu num_data_blocks: %zu\n", in_device, in, size, num_data_blocks);
501-
CUDA_CHK(cudaMemcpyAsync(in_device, in, size, cudaMemcpyHostToDevice, stream));
502-
503-
chacha_ctr_encrypt_kernel<<<num_blocks, num_threads_per_block, 0, stream>>>(
504-
in_device, output_device, size,
505-
keys_device, nonces_device, num_keys,
506-
sha_state_device,
507-
samples_device,
508-
sample_len,
509-
i * BLOCK_SIZE);
510-
#ifdef DO_COPY
511-
for (uint32_t j = 0; j < num_keys; j++) {
512-
size_t block_offset = j * length + i * BLOCK_SIZE;
513-
size_t out_offset = j * size;
514-
//printf("i: %d j: %d copy %zi b block offset: %zu output offset: %zu num_data_blocks: %zu\n",
515-
// i, j, size, block_offset, out_offset, num_data_blocks);
516-
CUDA_CHK(cudaMemcpy(&out[block_offset], &output_device[out_offset], size, cudaMemcpyDeviceToHost));
517-
}
518-
#endif
519-
520-
slength -= BLOCK_SIZE;
521-
in += BLOCK_SIZE;
522-
if (slength <= 0) {
523-
break;
524-
}
525-
}
526-
527-
#ifndef DO_COPY
528-
CUDA_CHK(cudaDeviceSynchronize());
529-
#endif
530-
531-
get_time(&end);
532-
*time_us = get_diff(&start, &end);
533-
534-
//printf("gpu time: %f us\n", get_diff(&start, &end));
535-
}
536-

‎src/cuda-crypt/modes.h

Lines changed: 0 additions & 167 deletions
Original file line numberDiff line numberDiff line change
@@ -22,10 +22,6 @@ typedef void (*cbc128_f) (const unsigned char *in, unsigned char *out,
2222
size_t len, const void *key,
2323
unsigned char ivec[16], int enc);
2424

25-
typedef void (*ctr128_f) (const unsigned char *in, unsigned char *out,
26-
size_t blocks, const void *key,
27-
const unsigned char ivec[16]);
28-
2925
typedef void (*ccm128_f) (const unsigned char *in, unsigned char *out,
3026
size_t blocks, const void *key,
3127
const unsigned char ivec[16],
@@ -39,169 +35,6 @@ void CRYPTO_cbc128_decrypt(const unsigned char *in, unsigned char *out,
3935
size_t len, const void *key,
4036
unsigned char ivec[16], block128_f block);
4137

42-
void CRYPTO_ctr128_encrypt(const unsigned char *in, unsigned char *out,
43-
size_t len, const void *key,
44-
unsigned char ivec[16],
45-
unsigned char ecount_buf[16], unsigned int *num,
46-
block128_f block);
47-
48-
void CRYPTO_ctr128_encrypt_ctr32(const unsigned char *in, unsigned char *out,
49-
size_t len, const void *key,
50-
unsigned char ivec[16],
51-
unsigned char ecount_buf[16],
52-
unsigned int *num, ctr128_f ctr);
53-
54-
void CRYPTO_ofb128_encrypt(const unsigned char *in, unsigned char *out,
55-
size_t len, const void *key,
56-
unsigned char ivec[16], int *num,
57-
block128_f block);
58-
59-
void CRYPTO_cfb128_encrypt(const unsigned char *in, unsigned char *out,
60-
size_t len, const void *key,
61-
unsigned char ivec[16], int *num,
62-
int enc, block128_f block);
63-
void CRYPTO_cfb128_8_encrypt(const unsigned char *in, unsigned char *out,
64-
size_t length, const void *key,
65-
unsigned char ivec[16], int *num,
66-
int enc, block128_f block);
67-
void CRYPTO_cfb128_1_encrypt(const unsigned char *in, unsigned char *out,
68-
size_t bits, const void *key,
69-
unsigned char ivec[16], int *num,
70-
int enc, block128_f block);
71-
72-
size_t CRYPTO_cts128_encrypt_block(const unsigned char *in,
73-
unsigned char *out, size_t len,
74-
const void *key, unsigned char ivec[16],
75-
block128_f block);
76-
size_t CRYPTO_cts128_encrypt(const unsigned char *in, unsigned char *out,
77-
size_t len, const void *key,
78-
unsigned char ivec[16], cbc128_f cbc);
79-
size_t CRYPTO_cts128_decrypt_block(const unsigned char *in,
80-
unsigned char *out, size_t len,
81-
const void *key, unsigned char ivec[16],
82-
block128_f block);
83-
size_t CRYPTO_cts128_decrypt(const unsigned char *in, unsigned char *out,
84-
size_t len, const void *key,
85-
unsigned char ivec[16], cbc128_f cbc);
86-
87-
size_t CRYPTO_nistcts128_encrypt_block(const unsigned char *in,
88-
unsigned char *out, size_t len,
89-
const void *key,
90-
unsigned char ivec[16],
91-
block128_f block);
92-
size_t CRYPTO_nistcts128_encrypt(const unsigned char *in, unsigned char *out,
93-
size_t len, const void *key,
94-
unsigned char ivec[16], cbc128_f cbc);
95-
size_t CRYPTO_nistcts128_decrypt_block(const unsigned char *in,
96-
unsigned char *out, size_t len,
97-
const void *key,
98-
unsigned char ivec[16],
99-
block128_f block);
100-
size_t CRYPTO_nistcts128_decrypt(const unsigned char *in, unsigned char *out,
101-
size_t len, const void *key,
102-
unsigned char ivec[16], cbc128_f cbc);
103-
104-
typedef struct gcm128_context GCM128_CONTEXT;
105-
106-
GCM128_CONTEXT *CRYPTO_gcm128_new(void *key, block128_f block);
107-
void CRYPTO_gcm128_init(GCM128_CONTEXT *ctx, void *key, block128_f block);
108-
void CRYPTO_gcm128_setiv(GCM128_CONTEXT *ctx, const unsigned char *iv,
109-
size_t len);
110-
int CRYPTO_gcm128_aad(GCM128_CONTEXT *ctx, const unsigned char *aad,
111-
size_t len);
112-
int CRYPTO_gcm128_encrypt(GCM128_CONTEXT *ctx,
113-
const unsigned char *in, unsigned char *out,
114-
size_t len);
115-
int CRYPTO_gcm128_decrypt(GCM128_CONTEXT *ctx,
116-
const unsigned char *in, unsigned char *out,
117-
size_t len);
118-
int CRYPTO_gcm128_encrypt_ctr32(GCM128_CONTEXT *ctx,
119-
const unsigned char *in, unsigned char *out,
120-
size_t len, ctr128_f stream);
121-
int CRYPTO_gcm128_decrypt_ctr32(GCM128_CONTEXT *ctx,
122-
const unsigned char *in, unsigned char *out,
123-
size_t len, ctr128_f stream);
124-
int CRYPTO_gcm128_finish(GCM128_CONTEXT *ctx, const unsigned char *tag,
125-
size_t len);
126-
void CRYPTO_gcm128_tag(GCM128_CONTEXT *ctx, unsigned char *tag, size_t len);
127-
void CRYPTO_gcm128_release(GCM128_CONTEXT *ctx);
128-
129-
typedef struct ccm128_context CCM128_CONTEXT;
130-
131-
void CRYPTO_ccm128_init(CCM128_CONTEXT *ctx,
132-
unsigned int M, unsigned int L, void *key,
133-
block128_f block);
134-
int CRYPTO_ccm128_setiv(CCM128_CONTEXT *ctx, const unsigned char *nonce,
135-
size_t nlen, size_t mlen);
136-
void CRYPTO_ccm128_aad(CCM128_CONTEXT *ctx, const unsigned char *aad,
137-
size_t alen);
138-
int CRYPTO_ccm128_encrypt(CCM128_CONTEXT *ctx, const unsigned char *inp,
139-
unsigned char *out, size_t len);
140-
int CRYPTO_ccm128_decrypt(CCM128_CONTEXT *ctx, const unsigned char *inp,
141-
unsigned char *out, size_t len);
142-
int CRYPTO_ccm128_encrypt_ccm64(CCM128_CONTEXT *ctx, const unsigned char *inp,
143-
unsigned char *out, size_t len,
144-
ccm128_f stream);
145-
int CRYPTO_ccm128_decrypt_ccm64(CCM128_CONTEXT *ctx, const unsigned char *inp,
146-
unsigned char *out, size_t len,
147-
ccm128_f stream);
148-
size_t CRYPTO_ccm128_tag(CCM128_CONTEXT *ctx, unsigned char *tag, size_t len);
149-
150-
typedef struct xts128_context XTS128_CONTEXT;
151-
152-
int CRYPTO_xts128_encrypt(const XTS128_CONTEXT *ctx,
153-
const unsigned char iv[16],
154-
const unsigned char *inp, unsigned char *out,
155-
size_t len, int enc);
156-
157-
size_t CRYPTO_128_wrap(void *key, const unsigned char *iv,
158-
unsigned char *out,
159-
const unsigned char *in, size_t inlen,
160-
block128_f block);
161-
162-
size_t CRYPTO_128_unwrap(void *key, const unsigned char *iv,
163-
unsigned char *out,
164-
const unsigned char *in, size_t inlen,
165-
block128_f block);
166-
size_t CRYPTO_128_wrap_pad(void *key, const unsigned char *icv,
167-
unsigned char *out, const unsigned char *in,
168-
size_t inlen, block128_f block);
169-
size_t CRYPTO_128_unwrap_pad(void *key, const unsigned char *icv,
170-
unsigned char *out, const unsigned char *in,
171-
size_t inlen, block128_f block);
172-
173-
# ifndef OPENSSL_NO_OCB
174-
typedef struct ocb128_context OCB128_CONTEXT;
175-
176-
typedef void (*ocb128_f) (const unsigned char *in, unsigned char *out,
177-
size_t blocks, const void *key,
178-
size_t start_block_num,
179-
unsigned char offset_i[16],
180-
const unsigned char L_[][16],
181-
unsigned char checksum[16]);
182-
183-
OCB128_CONTEXT *CRYPTO_ocb128_new(void *keyenc, void *keydec,
184-
block128_f encrypt, block128_f decrypt,
185-
ocb128_f stream);
186-
int CRYPTO_ocb128_init(OCB128_CONTEXT *ctx, void *keyenc, void *keydec,
187-
block128_f encrypt, block128_f decrypt,
188-
ocb128_f stream);
189-
int CRYPTO_ocb128_copy_ctx(OCB128_CONTEXT *dest, OCB128_CONTEXT *src,
190-
void *keyenc, void *keydec);
191-
int CRYPTO_ocb128_setiv(OCB128_CONTEXT *ctx, const unsigned char *iv,
192-
size_t len, size_t taglen);
193-
int CRYPTO_ocb128_aad(OCB128_CONTEXT *ctx, const unsigned char *aad,
194-
size_t len);
195-
int CRYPTO_ocb128_encrypt(OCB128_CONTEXT *ctx, const unsigned char *in,
196-
unsigned char *out, size_t len);
197-
int CRYPTO_ocb128_decrypt(OCB128_CONTEXT *ctx, const unsigned char *in,
198-
unsigned char *out, size_t len);
199-
int CRYPTO_ocb128_finish(OCB128_CONTEXT *ctx, const unsigned char *tag,
200-
size_t len);
201-
int CRYPTO_ocb128_tag(OCB128_CONTEXT *ctx, unsigned char *tag, size_t len);
202-
void CRYPTO_ocb128_cleanup(OCB128_CONTEXT *ctx);
203-
# endif /* OPENSSL_NO_OCB */
204-
20538
# ifdef __cplusplus
20639
}
20740
# endif

0 commit comments

Comments
 (0)