Skip to content

Tutorial: Poseidon2 Hash Function

This tutorial walks you through building a Poseidon2 hash application using cuPQC-Hash. You'll learn how to use GPU-accelerated Poseidon2, a zero-knowledge-friendly hash function designed for efficient use in ZK proof systems.

API reference

Configurations & constants: Supported parameter combinations (field, width, capacity) and the round constants and MDS matrix for BabyBear and KoalaBear are in the Poseidon2 operators section of the cuPQC-Hash API docs.

Step 1: Project Setup

Clone the cuPQC repository to get the example files:

git clone https://github.com/NVIDIA/cuPQC.git
cd cuPQC/examples/hash

Step 2: Include Required Headers

Start by including the necessary cuPQC SDK headers:

#include <vector>
#include <iomanip>
#include <iostream>

#include <hash.hpp>

using namespace cupqc;

The hash.hpp header provides access to Poseidon2 and other cryptographic hash primitives.

Step 3: Define Hash Function Type

Define the Poseidon2 hash function type, specifying the field and width:

using HASH = decltype(POSEIDON2_KB_8_16() + Thread());

POSEIDON2_KB_8_16() chooses Poseidon2 with KoalaBear field, capacity 8, and width 16.
Thread() specifies that the operator will execute independently for each thread.

Step 4: Create Hash Kernel

Create a CUDA kernel to compute the Poseidon2 hash:

__global__ void hash_poseidon2_kernel(
    uint32_t* digest, 
    const uint32_t* msg, 
    size_t inbuf_len, 
    size_t out_len)
{
    HASH hash {};
    hash.reset();
    hash.update(msg, inbuf_len);
    hash.finalize();
    hash.digest(digest, out_len);
}

hash.reset(): Initialize the hash state to begin a new hash computation
hash.update(msg, inbuf_len): Add input message data (field elements) to be hashed, can be called multiple times for streaming data
hash.finalize(): Complete the hash process and prepare for digest extraction
hash.digest(digest, out_len): Extract the final hash value (digest) into the output buffer, using out_len to specify the output read size
Only single thread 0 runs it

The hash computation follows a standard pattern: reset the state, update with input message data (field elements), finalize the computation, and then extract the digest. Note that Poseidon2 works with field elements (uint32_t) rather than bytes.

Step 5: Implement Host Function

Create a host function to manage GPU memory and launch the kernel. Only thread 0 performs the hash (single message per block):

void hash_poseidon2(
    std::vector<uint32_t>& digest, 
    std::vector<uint32_t>& msg, 
    size_t out_len)
{
    uint32_t* d_msg;
    uint32_t* d_digest;

    // Allocate device memory
    cudaMalloc(reinterpret_cast<void**>(&d_msg), msg.size() * sizeof(uint32_t));
    cudaMalloc(reinterpret_cast<void**>(&d_digest), digest.size() * sizeof(uint32_t));

    // Copy input message to device
    cudaMemcpy(d_msg, msg.data(), msg.size() * sizeof(uint32_t), cudaMemcpyHostToDevice);

    // Launch kernel (1 block, 1 thread)
    hash_poseidon2_kernel<<<1, 1>>>(d_digest, d_msg, msg.size(), out_len);

    // Copy result back to host
    cudaMemcpy(digest.data(), d_digest, 
               digest.size() * sizeof(uint32_t), 
               cudaMemcpyDeviceToHost);

    // Cleanup
    cudaFree(d_msg);
    cudaFree(d_digest);
}

Step 6: Create the Main Application

Implement the main function with example usage. It will take a sample message (field elements), compute the Poseidon2 hash, and display the result:

int main(int argc, char* argv[]) {
    std::cout << "================================================================\n";
    std::cout << "Poseidon2 Hash Function Example\n";
    std::cout << "================================================================\n\n";

    std::cout << "This example demonstrates Poseidon2 hashing using cuPQC SDK.\n";
    std::cout << "Poseidon2 is a zero-knowledge-friendly hash function designed for "
              << "efficient use in zero-knowledge proof systems. It uses arithmetic "
              << "operations native to ZK circuits.\n\n";
    std::cout << "Configuration: Capacity 8, Width 16, KoalaBear field\n\n";

    constexpr size_t in_len = 64;
    constexpr size_t out_len = 16;
    std::vector<uint32_t> msg(in_len, 0);

    // Generate sample input data (field elements)
    for (size_t i = 0; i < in_len; i++) {
        msg[i] = i % cupqc_common::BabyBearPrime;
    }

    std::cout << "========================================\n";
    std::cout << "Input Data\n";
    std::cout << "========================================\n";
    std::cout << "Input size: " << in_len << " field elements (" 
              << (in_len * sizeof(uint32_t)) << " bytes)\n";
    std::cout << "Input data (first 8 values):\n";
    for (size_t i = 0; i < std::min(in_len, size_t(8)); i++) {
        std::cout << "  Input[" << i << "]: 0x" << std::hex 
                  << std::setw(8) << std::setfill('0') << msg[i] 
                  << std::dec << " (" << msg[i] << ")\n";
    }
    if (in_len > 8) {
        std::cout << "  ... (showing first 8 of " << in_len << " values)\n";
    }
    std::cout << "\n";

    // Compute hash
    std::cout << "========================================\n";
    std::cout << "Computing Poseidon2 Hash\n";
    std::cout << "========================================\n";
    std::vector<uint32_t> digest(out_len, 0);
    hash_poseidon2(digest, msg, out_len);

    // Display results
    std::cout << "Computed Hash (" << out_len << " field elements):\n";
    std::cout << "  ";
    for (size_t i = 0; i < digest.size(); i++) {
        std::cout << "0x" << std::hex << std::setw(8) 
                  << std::setfill('0') << digest[i] << std::dec;
        if (i < digest.size() - 1) {
            std::cout << " ";
        }
    }
    std::cout << "\n\n";

    std::cout << "Hash values (decimal):\n";
    std::cout << "  ";
    for (size_t i = 0; i < digest.size(); i++) {
        std::cout << digest[i];
        if (i < digest.size() - 1) {
            std::cout << " ";
        }
    }
    std::cout << "\n\n";

    std::cout << "========================================\n";
    std::cout << "Hash Information\n";
    std::cout << "========================================\n";
    std::cout << "Hash size: " << out_len << " field elements (" 
              << (out_len * sizeof(uint32_t)) << " bytes)\n";
    std::cout << "Field: KoalaBear\n";
    std::cout << "Poseidon2 parameters: Capacity=8, Width=16\n\n";

    std::cout << "Example completed successfully.\n";
    return 0;
}

Step 7: Build and Run

The Makefile will build all examples in the folder. Run the Poseidon2 example:

make
./example_poseidon2

Step 8: Understanding the Output

Expected output: 

================================================================
Poseidon2 Hash Function Example
================================================================

This example demonstrates Poseidon2 hash using cuPQC SDK.
Poseidon2 is a zero-knowledge-friendly hash function designed for
efficient use in zero-knowledge proof systems. It uses arithmetic
operations native to ZK circuits.

Configuration: Capacity 8, Width 16, KoalaBear field

========================================
Input Data
========================================
Input size: 64 field elements (256 bytes)
Input data (first 8 values):
  Input[0]: 0x00000000 (0)
  Input[1]: 0x00000001 (1)
  Input[2]: 0x00000002 (2)
  Input[3]: 0x00000003 (3)
  Input[4]: 0x00000004 (4)
  Input[5]: 0x00000005 (5)
  Input[6]: 0x00000006 (6)
  Input[7]: 0x00000007 (7)
  ... (showing first 8 of 64 values)

========================================
Computing Poseidon2 Hash
========================================
Computed Hash (16 field elements):
  0x17fe7d2c 0x0e1f2486 0x26230fb8 0x40272398 0x27f4ab68 0x72080cbe 0x1bbb0dbf 0x06732192
  0x49661043 0x5f2bfe0e 0x3dd644f7 0x5b5fb4fa 0x49266af5 0x670d04aa 0x2a6edc71 0x47526886

Hash values (decimal):
  402554156 236921990 639831992 1076306840 670346088 1913130174 465243583 108208530
  1231425603 1596718606 1037452535 1532998906 1227254517 1728906410 711908465 1196583046

========================================
Hash Information
========================================
Hash size: 16 field elements (64 bytes)
Field: KoalaBear
Poseidon2 parameters: Capacity=8, Width=16

Example completed successfully.

Customization Tips

Use Different Fields and Parameters

Poseidon2 supports different prime fields (BabyBear, KoalaBear), widths (16, 24), and capacities (8). You can change these parameters in the HASH descriptor:

// BabyBear field, Capacity 8, Width 24
using HASH_BB_24 = decltype(POSEIDON2_BB_8_24() + Thread());

// KoalaBear field, Capacity 8, Width 16 (as used in this tutorial)
using HASH_KB_16 = decltype(POSEIDON2_KB_8_16() + Thread());
A complete list of supported Poseidon2 configurations is available in the Hash Operators Documentation.

Learn More