Look at the following code:
I am not getting the expected performance boost from gpu acceleration.
# o1 mini - Segmented Sieve of Euler with GPU acceleration
import time
import pyopencl as cl
import math

def popcount32(arr, bits_to_count):
    """
    Perform a population count (number of set bits) on a NumPy array of uint32,
    considering only up to bits_to_count bits.

    Parameters:
        arr (np.ndarray): NumPy array with dtype=np.uint32.
        bits_to_count (int): Total number of bits to count.

    Returns:
        int: Total number of set bits.
    """
    # Copy the array to avoid modifying the original
    x = arr.copy()

    # Step 1: Subtract adjacent bits
    x = x - ((x >> 1) & 0x55555555)

    # Step 2: Sum pairs of bits
    x = (x & 0x33333333) + ((x >> 2) & 0x33333333)

    # Step 3: Sum nibbles (4 bits)
    x = (x + (x >> 4)) & 0x0F0F0F0F

    # Step 4: Multiply by a magic number and extract byte counts
    x = x * 0x01010101
    counts = (x >> 24) & 0xFF

    # Initial total count
    total = counts.sum()

    # Handle remaining bits if bits_to_count is not a multiple of 32
    remainder = bits_to_count % 32
    if remainder != 0:
        # Number of full 32-bit words
        full_words = bits_to_count // 32
        # Extract the last relevant word
        if full_words < len(arr):
            last_word = arr[full_words]
        else:
            last_word = 0
        # Create a mask for the relevant bits
        mask = (1 << remainder) - 1
        # Apply the mask
        last_word_masked = last_word & mask
        # Recompute popcount for the masked last word
        x_last = last_word_masked
        x_last = x_last - ((x_last >> 1) & 0x55555555)
        x_last = (x_last & 0x33333333) + ((x_last >> 2) & 0x33333333)
        x_last = (x_last + (x_last >> 4)) & 0x0F0F0F0F
        x_last = x_last * 0x01010101
        count_last = (x_last >> 24) & 0xFF

        # Adjust the total count
        if full_words > 0:
            total = counts[:full_words].sum()
        else:
            total = 0
        total += count_last

    return int(total)

def sieve_euler(limit):
    """
    Perform the Sieve of Euler (Linear Sieve) up to 'limit'.
    Returns an array of prime numbers up to 'limit'.
    Includes the prime number 2.

    Parameters:
        limit (int): The upper limit to sieve primes up to.

    Returns:
        np.ndarray: Array of prime numbers up to 'limit'.
    """
    if limit < 2:
        return np.array([], dtype=np.uint32)

    smallest_prime = np.zeros(limit + 1, dtype=np.int32)
    primes = []

    for num in range(2, limit + 1):
        if smallest_prime[num] == 0:
            smallest_prime[num] = num
            primes.append(num)
        for p in primes:
            if p > smallest_prime[num] or p * num > limit:
                break
            smallest_prime[p * num] = p

    primes_np = np.array(primes, dtype=np.uint32)
    return primes_np

def count_primes_with_gpu(n):
    """
    Count the number of primes up to 'n' using the Sieve of Euler
    with GPU acceleration and advanced optimizations.

    Parameters:
        n (int): The upper limit to count primes up to.

    Returns:
        int: Number of primes up to 'n'.
    """
    # Step 1: Find small primes up to sqrt(n) using Sieve of Euler
    sqrt_n = int(math.isqrt(n))
    small_primes = sieve_euler(sqrt_n)
    print(f"Number of small primes up to sqrt(n)={sqrt_n}: {len(small_primes)}")

    # Step 2: Initialize OpenCL
    platforms = cl.get_platforms()
    if not platforms:
        raise RuntimeError("No OpenCL platforms found.")

    # Select the first platform with a GPU device
    device = None
    for platform in platforms:
        devices = platform.get_devices(device_type=cl.device_type.GPU)
        if devices:
            device = devices[0]
            break

    if device is None:
        raise RuntimeError("No GPU device found on any platform.")

    ctx = cl.Context([device])
    # Create a command queue with profiling enabled
    queue = cl.CommandQueue(ctx, properties=cl.command_queue_properties.PROFILING_ENABLE)

    # Step 3: Define the OpenCL kernel with atomic operations
    kernel_code = """
    __kernel void sieve_of_eratosthenes(__global unsigned int *segment,
                                        __global const unsigned int *small_primes,
                                        const unsigned int num_small_primes,
                                        const unsigned int segment_start,
                                        const unsigned int segment_size,
                                        const unsigned int max_n)
    {
        int idx = get_global_id(0);
        if (idx >= num_small_primes) return;

        unsigned int p = small_primes[idx];

        // Find the first multiple of p within the segment
        unsigned int start = (p * p > segment_start) ? p * p : segment_start + ((p - (segment_start % p)) % p);
        
        // Ensure start is odd
        if (start % 2 == 0)
            start += p;

        // Iterate over multiples of p within the segment
        for(unsigned int multiple = start; multiple < segment_start + segment_size && multiple <= max_n; multiple += p * 2)
        {
            if (multiple < segment_start)
                continue;
            unsigned int index = (multiple - segment_start) / 2;
            if(index >= (segment_size + 1) / 2)
                break;
            
            // Calculate word and bit indices
            unsigned int word_idx = index / 32;
            unsigned int bit_idx = index % 32;
            unsigned int mask = ~(1u << bit_idx);
            
            // Atomic AND operation to clear the bit
            atomic_and(&segment[word_idx], mask);
        }
    }
    """

    try:
        program = cl.Program(ctx, kernel_code).build()
    except cl.RuntimeError as e:
        # Retrieve and print the build log directly from the program
        try:
            build_log = program.get_build_info(device, cl.program_build_info.LOG)
        except:
            build_log = "No build log available."
        print("Build log:")
        print(build_log)
        raise

    sieve_kernel = program.sieve_of_eratosthenes

    # Step 4: Segmented Sieve Parameters
    segment_size = 10_000_000  # 10 million numbers per segment
    remaining_numbers = n - sqrt_n
    num_segments = (remaining_numbers + segment_size -1) // segment_size
    print(f"Total segments to process: {num_segments}")

    # Step 5: Allocate Odd Small Primes Buffer (exclude 2)
    # Exclude 2 to optimize since even numbers are already excluded in sieve segments
    odd_small_primes = small_primes[small_primes != 2]
    print(f"Number of odd small primes (excluding 2): {len(odd_small_primes)}")
    small_primes_np = np.array(odd_small_primes, dtype=np.uint32)
    small_primes_buffer = cl.Buffer(ctx, cl.mem_flags.READ_ONLY | cl.mem_flags.COPY_HOST_PTR, hostbuf=small_primes_np)

    # Step 6: Initialize Total Prime Count (includes all small primes, including 2)
    total_primes = len(small_primes)  # 'small_primes' includes 2

    # Step 7: Process Each Segment
    for seg_idx in range(num_segments):
        low = sqrt_n + 1 + seg_idx * segment_size  # Start strictly above sqrt_n
        high = min(low + segment_size, n + 1)

        # Ensure low is odd and >=3
        if low < 3:
            low = 3
        if low % 2 == 0:
            low += 1
        current_segment_size = high - low
        sieve_bits = (current_segment_size + 1) // 2  # Only odd numbers
        sieve_words = (sieve_bits + 31) // 32  # Ceiling division for 32-bit words

        # Initialize the segment sieve: all bits set to 1 (each bit represents an odd number)
        sieve = np.full(sieve_words, 0xFFFFFFFF, dtype=np.uint32)

        # Create buffer for the segment sieve
        sieve_buffer = cl.Buffer(ctx, cl.mem_flags.READ_WRITE | cl.mem_flags.COPY_HOST_PTR, hostbuf=sieve)

        # Set global work size
        global_work_size = len(odd_small_primes)

        # Set kernel arguments
        sieve_kernel.set_args(
            sieve_buffer,
            small_primes_buffer,
            np.uint32(len(odd_small_primes)),  # arg#3: num_small_primes
            np.uint32(low),                    # arg#4: segment_start
            np.uint32(current_segment_size),   # arg#5: segment_size
            np.uint32(n)                       # arg#6: max_n
        )

        # Enqueue the kernel
        try:
            event = cl.enqueue_nd_range_kernel(
                queue,
                sieve_kernel,
                (global_work_size,),
                None  # Let OpenCL decide the local work size
            )
        except cl.RuntimeError as e:
            print(f"Error enqueueing kernel on segment {seg_idx}: {e}")
            build_log = program.get_build_info(device, cl.program_build_info.LOG)
            print("Build log:")
            print(build_log)
            raise

        # Read back the sieve data after kernel execution
        sieve_host = np.empty_like(sieve)
        try:
            cl.enqueue_copy(queue, sieve_host, sieve_buffer, is_blocking=True, wait_for=[event])
        except cl.RuntimeError as e:
            print(f"Error copying data from GPU on segment {seg_idx}: {e}")
            raise

        # Calculate the maximum number in this segment
        max_num = high - 1  # since range is [low, high)
        max_num = min(max_num, n)

        # Calculate the number of bits to count
        bits_to_count = (max_num - low) // 2 + 1  # Only odd numbers

        # Ensure bits_to_count does not exceed sieve_bits
        bits_to_count = min(bits_to_count, sieve_bits)

        # Count the number of set bits (primes) in the segment using popcount32
        primes_in_segment = popcount32(sieve_host, bits_to_count)

        # Add to the total prime count
        total_primes += primes_in_segment

        # Optional: Progress update
        progress = (seg_idx + 1) / num_segments * 100
        print(f"Processed segment [{low}, {high}): {primes_in_segment} primes found. Progress: {progress:.2f}%")

    return total_primes

# Example Usage
if __name__ == "__main__":
    test_values = [100, 1000, 10000, 100000, 100000000]  # Add more values as needed

    for n in test_values:
        print("\n" + "-"*50)
        print(f"Counting primes up to {n}")
        start_time = time.time()
        prime_count = count_primes_with_gpu(n)
        end_time = time.time()
        print(f"Total number of primes up to {n}: {prime_count}")
        print(f"Execution time: {end_time - start_time:.2f} seconds")