Please fix the kernel code because we get an error around the atomic function: import time
import numpy as np
import argparse
import psutil
import csv
import os
import pyopencl as cl
import sys

def get_gpu_info():
    """
    Retrieve GPU information using PyOpenCL.
    """
    devices_info = []
    platforms = cl.get_platforms()
    for platform in platforms:
        for device in platform.get_devices(device_type=cl.device_type.GPU):
            devices_info.append({
                'platform': platform,
                'device': device,
                'name': device.name,
                'memory': device.global_mem_size / 1e9  # Convert to GB
            })
    return devices_info

def create_kernels(context, size):
    """
    Create OpenCL kernels for floating-point and integer operations.
    Removes atomic_add from the floating-point kernel due to unsupported data types.
    """
    # Floating-point operations kernel
    fp_operation_kernel = """
    #pragma OPENCL EXTENSION cl_khr_fp64 : enable
    __kernel void benchmark_fp_operation(
        __global const double* x_fp,
        __global const double* y_fp,
        const unsigned int num_iterations,
        const int operation,
        __global double* total) // 'total' will not be used for FP operations
    {{
        int gid = get_global_id(0);
        int stride = get_global_size(0);
        double partial_total = 0.0;

        for (int i = gid; i < num_iterations; i += stride)
        {{
            int index = i % {size};
            if (operation == 0) {{ // add
                partial_total += x_fp[index] + y_fp[index];
            }}
            else if (operation == 1) {{ // sub
                partial_total += x_fp[index] - y_fp[index];
            }}
            else if (operation == 2) {{ // mul
                partial_total += x_fp[index] * y_fp[index];
            }}
            else if (operation == 3) {{ // div
                partial_total += x_fp[index] / y_fp[index];
            }}
        }}

        // Dummy operation to prevent optimization
        total[0] += partial_total * 0.0;
    }}
    """.format(size=size)

    # Integer operations kernel remains unchanged
    int_operation_kernel = """
    __kernel void benchmark_int_operation(
        __global const long* x_int,
        __global const long* y_int,
        const unsigned int num_iterations,
        const int operation,
        __global long* total)
    {{
        int gid = get_global_id(0);
        int stride = get_global_size(0);
        long partial_total = 0;

        for (int i = gid; i < num_iterations; i += stride)
        {{
            int index = i % {size};
            if (operation == 0) {{ // add
                partial_total += x_int[index] + y_int[index];
            }}
            else if (operation == 1) {{ // sub
                partial_total += x_int[index] - y_int[index];
            }}
            else if (operation == 2) {{ // mul
                partial_total += x_int[index] * y_int[index];
            }}
            else if (operation == 3) {{ // div
                partial_total += x_int[index] / y_int[index];
            }}
        }}

        // Atomic add to prevent race conditions
        atomic_add(&total[0], partial_total);
    }}
    """.format(size=size)

    try:
        # Build programs
        prg_fp = cl.Program(context, fp_operation_kernel).build()
        prg_int = cl.Program(context, int_operation_kernel).build()
    except Exception as e:
        # If there's a build error, print the build log
        print("Error building OpenCL programs:")
        print(e)
        for program in [fp_operation_kernel, int_operation_kernel]:
            print(program)  # Optionally print the kernel code for debugging
        sys.exit(1)

    return prg_fp, prg_int

def single_operation_benchmark_opencl(op_type, operation, num_iterations, queue, prg_fp, prg_int, 
                                      x_fp_buf, y_fp_buf, x_int_buf, y_int_buf, 
                                      total_fp_buf, total_int_buf):
    """
    Perform a benchmark for a specific operation type on the GPU using OpenCL.
    
    Parameters:
    - op_type: 'FP' for floating-point or 'INT' for integer operations.
    - operation: 'add', 'sub', 'mul', 'div'.
    - num_iterations: Number of iterations to perform.
    - queue: OpenCL command queue.
    - prg_fp: Compiled OpenCL program for floating-point operations.
    - prg_int: Compiled OpenCL program for integer operations.
    - x_fp_buf, y_fp_buf, x_int_buf, y_int_buf: OpenCL buffers.
    - total_fp_buf, total_int_buf: OpenCL buffers for results.
    
    Returns:
    - gops: Giga Operations Per Second.
    - elapsed_time_ms: Elapsed time in milliseconds.
    - total: Accumulated result to prevent loop optimization or 'N/A' for FP operations.
    """
    # Mapping operation string to integer code
    operation_map = {'add': 0, 'sub': 1, 'mul': 2, 'div': 3}
    op_code = operation_map.get(operation)
    if op_code is None:
        raise ValueError("Invalid operation")
    
    # Select the appropriate program and buffer
    if op_type == 'FP':
        kernel = prg_fp.benchmark_fp_operation
        total_buf = total_fp_buf
        dtype = np.float64
        x_buf = x_fp_buf
        y_buf = y_fp_buf
    elif op_type == 'INT':
        kernel = prg_int.benchmark_int_operation
        total_buf = total_int_buf
        dtype = np.int64
        x_buf = x_int_buf
        y_buf = y_int_buf
    else:
        raise ValueError("Invalid operation type")
    
    if op_type == 'FP':
        # For FP operations, initialize 'total' buffer (dummy operation)
        zero = np.array([0.0], dtype=dtype) 
        cl.enqueue_copy(queue, total_buf, zero)
    else:
        # For INT operations, initialize 'total' buffer to zero
        zero = np.array([0], dtype=dtype)
        cl.enqueue_copy(queue, total_buf, zero)
    
    # Define global work size
    global_work_size = (256,)
    
    # Warm-up run
    kernel(queue, global_work_size, None,
           x_buf, y_buf,
           np.uint32(1000),
           np.int32(op_code),
           total_buf).wait()
    
    # Ensure all commands are completed before timing
    queue.finish()
    
    # Start timing
    start_time = time.perf_counter()
    
    # Launch the kernel
    kernel(queue, global_work_size, None,
           x_buf, y_buf,
           np.uint32(num_iterations),
           np.int32(op_code),
           total_buf).wait()
    
    # End timing
    end_time = time.perf_counter()
    
    if op_type == 'FP':
        # For FP operations, 'total' is not meaningful
        total = 'N/A'
    else:
        # Retrieve the total for INT operations
        total = np.empty(1, dtype=dtype)
        cl.enqueue_copy(queue, total, total_buf)
        total = total[0]
    
    # Calculate elapsed time in milliseconds
    elapsed_time_ms = (end_time - start_time) * 1e3
    
    # Calculate Operations Per Second
    ops = num_iterations / (elapsed_time_ms / 1e3)  # Operations per second
    gops = ops / 1e9  # Giga Operations Per Second
    
    return gops, elapsed_time_ms, total

def initialize_csv(file_name):
    """
    Initialize the CSV file with headers if it doesn't exist.
    """
    if not os.path.exists(file_name):
        with open(file_name, mode='w', newline='') as csv_file:
            fieldnames = ['Run', 'Configuration', 'Operation_Type', 'Operation', 'GOPS', 'Elapsed_Time_ms', 'Total', 'Device_Name', 'Device_Memory_GB']
            writer = csv.DictWriter(csv_file, fieldnames=fieldnames)
            writer.writeheader()

def append_to_csv(file_name, data_row):
    """
    Append a single row of data to the CSV file.
    """
    with open(file_name, mode='a', newline='') as csv_file:
        fieldnames = ['Run', 'Configuration', 'Operation_Type', 'Operation', 'GOPS', 'Elapsed_Time_ms', 'Total', 'Device_Name', 'Device_Memory_GB']
        writer = csv.DictWriter(csv_file, fieldnames=fieldnames)
        writer.writerow(data_row)

def main():
    parser = argparse.ArgumentParser(description="GPU Benchmarking Tool with PyOpenCL and CSV Output")
    parser.add_argument("--iterations", type=int, default=100_000_000, help="Number of iterations per benchmark")
    parser.add_argument("--runs", type=int, default=3, help="Number of times to run each benchmark configuration")
    parser.add_argument("--output", type=str, default="gpu_benchmark_results.csv", help="CSV file to store benchmark results")
    args = parser.parse_args()

    num_iterations = args.iterations
    num_runs = args.runs
    output_file = args.output

    # Initialize CSV file
    initialize_csv(output_file)

    # Retrieve GPU devices
    devices_info = get_gpu_info()
    if not devices_info:
        print("No OpenCL-capable GPU found.")
        return

    print("Detected OpenCL-capable GPUs:")
    for gpu in devices_info:
        print(f"  Device: {gpu['name']} with {gpu['memory']:.2f} GB memory")

    # Pre-generating random data for benchmarking
    print("\nPre-generating random data for benchmarks...")
    np.random.seed(42)  # For reproducibility
    size = 1_000_000  # Size of arrays (adjusted for better cache performance)

    # Floating-point data
    x_fp = np.random.uniform(low=1.0, high=100.0, size=size).astype(np.float64)
    y_fp = np.random.uniform(low=1.0, high=100.0, size=size).astype(np.float64)

    # Integer data
    x_int = np.random.randint(low=1, high=1_000_000, size=size).astype(np.int64)
    y_int = np.random.randint(low=1, high=1_000_000, size=size).astype(np.int64)

    # Define operations to benchmark
    operations = ['add', 'sub', 'mul', 'div']
    op_types = ['FP', 'INT']

    for run in range(1, num_runs + 1):
        print(f"\n===== Run {run} =====\n")
        for gpu in devices_info:
            platform = gpu['platform']
            device = gpu['device']
            device_name = gpu['name']
            device_memory_gb = gpu['memory']

            print(f"--- Benchmarking on Device: {device_name} ---")

            # Create context and queue for the device
            context = cl.Context([device])
            queue = cl.CommandQueue(context)

            # Create OpenCL programs
            prg_fp, prg_int = create_kernels(context, size)

            # Transfer data to GPU memory
            mf = cl.mem_flags
            x_fp_buf = cl.Buffer(context, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=x_fp)
            y_fp_buf = cl.Buffer(context, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=y_fp)

            x_int_buf = cl.Buffer(context, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=x_int)
            y_int_buf = cl.Buffer(context, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=y_int)

            # Allocate buffer for results
            total_fp_buf = cl.Buffer(context, mf.WRITE_ONLY, np.float64().nbytes)
            total_int_buf = cl.Buffer(context, mf.WRITE_ONLY, np.int64().nbytes)

            # Perform benchmarks
            results = {op_type: {} for op_type in op_types}

            for op_type in op_types:
                for operation in operations:
                    print(f"Running {op_type} - {operation}...")
                    if op_type == "FP":
                        gops, elapsed_ms, total = single_operation_benchmark_opencl(
                            op_type, operation, num_iterations, queue, prg_fp, prg_int, 
                            x_fp_buf, y_fp_buf, x_int_buf, y_int_buf, 
                            total_fp_buf, total_int_buf
                        )
                    else:
                        gops, elapsed_ms, total = single_operation_benchmark_opencl(
                            op_type, operation, num_iterations, queue, prg_fp, prg_int, 
                            x_fp_buf, y_fp_buf, x_int_buf, y_int_buf, 
                            total_fp_buf, total_int_buf
                        )
                    results[op_type][operation] = {
                        'GOPS': gops,
                        'Elapsed_Time_ms': elapsed_ms,
                        'Total': total
                    }
                    print(f"  {operation.capitalize()} {op_type} GOPS: {gops:.3f} GOPS in {elapsed_ms:.2f} ms\n")

            # Store results
            for op_type in op_types:
                for operation in operations:
                    data_row = {
                        'Run': run,
                        'Configuration': f"{device_name}",
                        'Operation_Type': op_type,
                        'Operation': operation,
                        'GOPS': f"{results[op_type][operation]['GOPS']:.3f}",
                        'Elapsed_Time_ms': f"{results[op_type][operation]['Elapsed_Time_ms']:.2f}",
                        'Total': f"{results[op_type][operation]['Total']}" if op_type == 'INT' else "N/A",
                        'Device_Name': device_name,
                        'Device_Memory_GB': f"{device_memory_gb:.2f}"
                    }
                    append_to_csv(output_file, data_row)

            print(f"Completed benchmarking on Device: {device_name}.\n")
            print("="*50 + "\n")

    print(f"Benchmarking complete. Results saved to {output_file}")

if __name__ == "__main__":
    main()