diff --git a/icicle/src/hash/blake2s/Makefile b/icicle/src/hash/blake2s/Makefile
index b979f3cd..9f58e80e 100644
--- a/icicle/src/hash/blake2s/Makefile
+++ b/icicle/src/hash/blake2s/Makefile
@@ -3,6 +3,7 @@ test_blake2s: test_blake2s.cu blake2s.cu
 
 test_blake2s_hasher: test_blake2s_hasher.cu blake2s.cu 
 	nvcc -o test_blake2s_hasher -I. -I../../../include test_blake2s_hasher.cu -g
+	./test_blake2s_hasher
 
 test_blake2s_tree: test_tree.cu blake2s.cu ../../merkle-tree/merkle.cu
 	nvcc -DMERKLE_DEBUG -o test_blake2s_tree -I../../../include test_tree.cu
diff --git a/icicle/src/hash/blake2s/blake2s.cu b/icicle/src/hash/blake2s/blake2s.cu
index ccf9b7b6..bf75f21a 100644
--- a/icicle/src/hash/blake2s/blake2s.cu
+++ b/icicle/src/hash/blake2s/blake2s.cu
@@ -225,11 +225,15 @@ __global__ void kernel_blake2s_hash(const BYTE *indata, WORD inlen, BYTE *outdat
     if (thread >= n_batch) {
         return;
     }
+    BYTE key[32] = "";  // Null key
+    WORD keylen = 0;
+    CUDA_BLAKE2S_CTX blake_ctx;
     const BYTE *in = indata + thread * inlen;
     BYTE *out = outdata + thread * BLAKE2S_BLOCK_SIZE;
-    CUDA_BLAKE2S_CTX ctx = c_CTX;
-    cuda_blake2s_update(&ctx, in, inlen);
-    cuda_blake2s_final(&ctx, out);
+    
+    cuda_blake2s_init(&blake_ctx, key, keylen, (BLAKE2S_BLOCK_SIZE << 3));
+    cuda_blake2s_update(&blake_ctx, in, inlen);
+    cuda_blake2s_final(&blake_ctx, out);
 }
 
 extern "C" {
@@ -269,22 +273,11 @@ cudaError_t Blake2s::run_hash_many_kernel(
     WORD        output_len,
     const device_context::DeviceContext& ctx) const
     {
-        BYTE key[32] = "";  // Null key
-        WORD keylen = strlen((char *)key);
         const WORD BLAKE2S_BLOCK_SIZE = (output_len >> 3);
-        CUDA_BLAKE2S_CTX blake_ctx;
-        cpu_blake2s_init(&blake_ctx, key, keylen, output_len);
-        cudaMemcpyToSymbol(c_CTX, &blake_ctx, sizeof(CUDA_BLAKE2S_CTX), 0, cudaMemcpyHostToDevice);
-    
         WORD thread = 256;
         WORD block = (number_of_states + thread - 1) / thread;
 
         kernel_blake2s_hash<<<block, thread, 0, ctx.stream>>>(input, input_len, output, number_of_states, BLAKE2S_BLOCK_SIZE);
-        // cudaDeviceSynchronize();
-        // cudaError_t error = cudaGetLastError();
-        // if (error != cudaSuccess) {
-        //     printf("Error cuda blake2s hash: %s \n", cudaGetErrorString(error));
-        // }
 
         CHK_IF_RETURN(cudaPeekAtLastError());
         return CHK_LAST();