msw.cuh File Reference

CUDA-accelerated Most Significant Word finding algorithms. More...

#include <cstddef>
#include <cstdint>

Include dependency graph for msw.cuh:

This graph shows which files directly or indirectly include this file:

Go to the source code of this file.

Functions
size_t	seq_find_mismatch (const uintmax_t *const a, const uintmax_t *const b, size_t n)
	Sequential implementation of finding the first mismatching word between two arrays.
size_t	par_find_mismatch (const uintmax_t *const a, const uintmax_t *const b, size_t n)
	GPU-accelerated implementation of finding the first mismatching word.
size_t	par_find_mismatch (const uintmax_t *const d_a, const uintmax_t *const b, uintmax_t *d_large_block, size_t n)
	GPU-accelerated mismatch finding with pre-allocated device memory.
uintmax_t *	alloc_large_block_to_device (size_t max_n)
	Allocates a large block of device memory for parallel mismatch operations.

Detailed Description

CUDA-accelerated Most Significant Word finding algorithms.

Provides functions for finding the most significant word (MSW) and the first mismatching word between two arrays of words. The implementation uses GPU-accelerated algorithms with a square root decomposition approach for efficient parallel processing.

Definition in file msw.cuh.

Function Documentation

alloc_large_block_to_device()

uintmax_t * alloc_large_block_to_device

(

size_t

n

)

Allocates a large block of device memory for parallel mismatch operations.

Allocates enough memory for the square root decomposition algorithm, which requires space for the second array, section information, and √n additional elements.

Parameters

max_n

Maximum number of elements to be processed

Returns

Pointer to allocated device memory

Calculates the required memory size based on the input size n and allocates a contiguous block of memory on the device. The size is calculated to accommodate:

• A copy of the second array (n elements)
• Storage for the result (1 element)
• The signature array for chunk identification (sqrt(n) elements)
• An extra element for safety (1 element)

Parameters

n	Number of elements to be processed

Returns

Pointer to allocated device memory of size n + sqrt(n-1) + 2

Definition at line 256 of file msw.cu.

{
    size_t large_block_size = n + std::sqrt(n - 1) + 2;
 
    return alloc_to_device<uintmax_t>(large_block_size);
}

par_find_mismatch() [1/2]

size_t par_find_mismatch	(	const uintmax_t *const	a,
		const uintmax_t *const	b,
		size_t	n
	)

GPU-accelerated implementation of finding the first mismatching word.

Allocates device memory, copies data, and uses parallel algorithms to find the first mismatch. Uses a square root decomposition approach that divides the arrays into √n chunks for efficient parallel processing.

Parameters

a	First input array (host memory)
b	Second input array (host memory)
n	Number of elements in the arrays

Returns

Index of the first mismatching word, or n if arrays are identical

GPU-accelerated implementation of finding the first mismatching word.

This function handles all the device memory management, including:

• Allocating device memory for both arrays
• Copying the first array to the device
• Calling the core implementation
• Freeing all allocated device memory

This provides a simpler interface for callers who don't need to manage device memory themselves or reuse memory across multiple calls.

Parameters

a	First input array (host memory)
b	Second input array (host memory)
n	Number of elements in the arrays

Returns

Index of the first mismatching word, or n if arrays are identical

Definition at line 280 of file msw.cu.

{
    uintmax_t *d_a = copy_to_device(a, n);
    uintmax_t *d_large_block = alloc_large_block_to_device(n);
 
    size_t result = par_find_mismatch(d_a, b, d_large_block, n);
 
    device_free(d_a);
    device_free(d_large_block);
 
    return result;
}

par_find_mismatch() [2/2]

size_t par_find_mismatch	(	const uintmax_t *const	d_a,
		const uintmax_t *const	b,
		uintmax_t *	d_large_block,
		size_t	n
	)

GPU-accelerated mismatch finding with pre-allocated device memory.

Uses a square root decomposition approach: first divides the input into √n chunks, identifies the first chunk with a mismatch in O(n) time and constant span, then finds the exact mismatch within that chunk, also in O(n) time and constant span.

Parameters

d_a	First input array (already in device memory)
b	Second input array (host memory, will be copied to device)
d_large_block	Pre-allocated device memory block for temporary storage
n	Number of elements in the arrays

Returns

Index of the first mismatching word, or n if arrays are identical

GPU-accelerated mismatch finding with pre-allocated device memory.

This function implements a two-phase square root decomposition algorithm:

1. Divide the arrays into sqrt(n) chunks and identify which chunk contains the first mismatch
2. Within the identified chunk, find the exact position of the first mismatch

The algorithm uses three CUDA kernels:

• par_sig_sqrt: Marks chunks containing mismatches
• naive_leftmost_prisoner: Eliminates all but the leftmost non-zero element
• par_get_section_diff: Identifies the index of the first non-zero element

Time complexity: O(n) work with O(1) span when sufficient threads are available Space complexity: O(n + sqrt(n)) for temporary storage

Precondition

d_a is already populated on the device

d_large_block has size at least n + sqrt(n) + 2 to hold b and signature arrays

Parameters

d_a	First input array (already in device memory)
b	Second input array (host memory, will be copied to device)
d_large_block	Pre-allocated device memory block for temporary storage
n	Number of elements in the arrays

Returns

Index of the first mismatching word, or n if arrays are identical

Definition at line 178 of file msw.cu.

{
    // Copy the second array to device memory using pre-allocated buffer
    uintmax_t *d_b = copy_to_device(d_large_block, b, n);
 
    // Set up memory pointers within the large block:
    // - d_section: single value to store result
    // - d_sqrt: array of size sqrt(n) to track which chunks have mismatches
    uintmax_t *d_section = d_b + n;
    uintmax_t *d_sqrt = d_section + 1;
    size_t num_sqrt = std::sqrt(n - 1) + 1;  // Calculate number of sqrt-sized chunks
 
    // Initialize the section result to max value (no match found yet)
    memset_within_device(d_section, 1, 0xFF);
 
    // Initialize the sqrt array to zeros (no mismatches found yet)
    memset_within_device(d_sqrt, num_sqrt, 0);
 
    // PHASE 1: Find which sqrt(n)-sized chunk has the first mismatch
    // For each element, compute XOR of the two arrays and mark its chunk in d_sqrt if mismatch
    par_sig_sqrt<<<BLOCKS, THREADS>>>(d_a, d_b, d_sqrt, n);
 
    // Find the first non-zero chunk (chunk with a mismatch) in parallel
    naive_leftmost_prisoner<<<BLOCKS, THREADS>>>(d_sqrt, num_sqrt);
 
    // Extract the index of the first chunk with a mismatch
    par_get_section_diff<<<BLOCKS, THREADS>>>(d_sqrt, d_section, num_sqrt);
 
    // Copy the result back to host memory
    uintmax_t *h_section = copy_from_device(d_section, 1);
    size_t section = *h_section;
 
    // If no mismatch found, return n (indicating arrays are identical)
    if (section == std::numeric_limits<uintmax_t>::max())
    {
        delete[] h_section;
        return n;
    }
 
    // Reset the section result for the second phase
    memset_within_device(d_section, 1, 0xFF);
 
    // PHASE 2: Find the exact mismatch position within the identified chunk
    // Calculate the offset and size of the identified chunk
    size_t section_offset = section * num_sqrt;
    size_t section_end = std::min(n, section_offset + num_sqrt);
    uintmax_t *d_sig_section = d_b + section_offset;
    size_t num_words = section_end - section_offset;
 
    // Find the first non-zero word within the chunk (containing XOR results)
    naive_leftmost_prisoner<<<BLOCKS, THREADS>>>(d_sig_section, num_words);
 
    // Extract the index of the first mismatching word within the chunk
    par_get_section_diff<<<BLOCKS, THREADS>>>(d_sig_section, d_section, num_words);
 
    // Copy the result back to host memory
    h_section = copy_from_device(d_section, 1);
    size_t result = *h_section;
 
    delete[] h_section;
 
    // Return the global index of the first mismatch (chunk offset + position within chunk)
    return section_offset + result;
}

seq_find_mismatch()

size_t seq_find_mismatch	(	const uintmax_t *const	arr1,
		const uintmax_t *const	arr2,
		size_t	size
	)

Sequential implementation of finding the first mismatching word between two arrays.

Compares two arrays element by element until finding the first mismatch. This is the baseline CPU implementation.

Parameters

a	First input array (host memory)
b	Second input array (host memory)
n	Number of elements in the arrays

Returns

Index of the first mismatching word, or n if arrays are identical

This is a simple linear scan through both arrays, comparing corresponding elements until a mismatch is found. This function serves as a baseline for comparison with the parallel implementation and is used for validation.

Time complexity: O(n) in the worst case Space complexity: O(1)

Parameters

arr1	First input array
arr2	Second input array
size	Number of elements in the arrays

Returns

Index of the first mismatching word, or size if arrays are identical

Definition at line 141 of file msw.cu.

{
    for (size_t i = 0; i < size; ++i)
    {
        if (arr1[i] != arr2[i])
        {
            return i;
        }
    }
 
    return size;
}