#include <immintrin.h> // AVX2 头文件

void simd_add(float* a, float* b, float* c, int n) {
    // 每次处理 8 个 float (256-bit 寄存器)
    int chunks = n / 8;

    for (int i = 0; i < chunks; i++) {
        // 加载数据到 SIMD 寄存器
        __m256 va = _mm256_loadu_ps(a + i*8);
        __m256 vb = _mm256_loadu_ps(b + i*8);

        // 执行向量加法
        __m256 vc = _mm256_add_ps(va, vb);

        // 存储结果
        _mm256_storeu_ps(c + i*8, vc);
    }

    // 处理剩余元素 (标量处理)
    for (int i = chunks*8; i < n; i++) {
        c[i] = a[i] + b[i];
    }
}

__global__ void simt_add(float* a, float* b, float* c, int n) {
    // 计算全局线程索引
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    // 检查边界
    if (i < n) {
        // 每个线程处理一个元素
        c[i] = a[i] + b[i];
    }
}

// 主机调用代码示例
void launch_kernel(float* d_a, float* d_b, float* d_c, int n) {
    // 配置线程块和网格
    int blockSize = 256;
    int gridSize = (n + blockSize - 1) / blockSize;

    // 启动核函数
    simt_add<<<gridSize, blockSize>>>(d_a, d_b, d_c, n);
}

#include <immintrin.h>

void simd_conditional(float* a, float* b, int n) {
    int chunks = n / 8;
    __m256 threshold = _mm256_set1_ps(5.0f);
    __m256 add_val = _mm256_set1_ps(10.0f);

    for (int i = 0; i < chunks; i++) {
        __m256 va = _mm256_loadu_ps(a + i*8);

        // 1. 创建掩码 (a > 5.0)
        __m256 mask = _mm256_cmp_ps(va, threshold, _CMP_GT_OQ);

        // 2. 条件为真时加10，否则保持原值
        __m256 vb = _mm256_add_ps(va, add_val);
        __m256 vc = _mm256_blendv_ps(va, vb, mask);

        _mm256_storeu_ps(b + i*8, vc);
    }

    // 处理尾部元素
    for (int i = chunks*8; i < n; i++) {
        b[i] = (a[i] > 5.0f) ? a[i] + 10.0f : a[i];
    }
}

__global__ void simt_conditional(float* a, float* b, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    if (i < n) {
        // 直接使用 if 语句
        if (a[i] > 5.0f) {
            b[i] = a[i] + 10.0f;
        } else {
            b[i] = a[i];
        }
    }
}

// 注意：实际使用时需考虑分支发散问题

#include <cuda.h>
#include <cuda_fp16.h>
#include <cuda_runtime.h>
#include <cublas_v2.h>
#include <wmma.h>

using namespace nvcuda;

__global__ void matrix_mult_tensorcore(
    half* a, half* b, float* c,
    int M, int N, int K) {

    // 声明矩阵分块
    wmma::fragment<wmma::matrix_a, 16, 16, 16, half, wmma::row_major> a_frag;
    wmma::fragment<wmma::matrix_b, 16, 16, 16, half, wmma::col_major> b_frag;
    wmma::fragment<wmma::accumulator, 16, 16, 16, float> c_frag;

    // 初始化累加器
    wmma::fill_fragment(c_frag, 0.0f);

    // 计算网格中的分块位置
    int warpM = (blockIdx.y * blockDim.y + threadIdx.y) / warpSize;
    int warpN = blockIdx.x * blockDim.x + threadIdx.x;

    // 分块矩阵乘法
    for (int i = 0; i < K; i += 16) {
        // 加载分块
        wmma::load_matrix_sync(a_frag, a + warpM * M * 16 + i, M);
        wmma::load_matrix_sync(b_frag, b + i * N + warpN * 16, N);

        // Tensor Core 矩阵乘加
        wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);
    }

    // 存储结果
    wmma::store_matrix_sync(c + warpM * M * 16 + warpN * 16, c_frag, M, wmma::mem_row_major);
}