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) and widths (16, 24). 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