#include <cuda_runtime.h>
#include <cuda_fp16.h>

__global__ void fused_attention_kernel(
    const half* __restrict__ q,
    const half* __restrict__ k,
    const half* __restrict__ v,
    half* __restrict__ output,
    int batch_size, int seq_len, int num_heads, int head_dim
) {
    extern __shared__ float s_data[];
    
    int tidx = threadIdx.x;
    int bidx = blockIdx.x;
    
    // Compute QK^T
    for (int i = tidx; i < seq_len; i += blockDim.x) {
        float qki = 0.0f;
        for (int j = 0; j < head_dim; j++) {
            int idx = bidx * seq_len * head_dim + i * head_dim + j;
            qki += __half2float(q[idx]) * __half2float(k[idx]);
        }
        qki /= sqrtf(head_dim);
        s_data[i] = qki;
    }
    __syncthreads();
    
    // Softmax
    float max_val = -INFINITY;
    for (int i = tidx; i < seq_len; i += blockDim.x) {
        max_val = max(max_val, s_data[i]);
    }
    __syncthreads();
    
    // Reduce max_val
    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
        if (tidx < stride) {
            s_data[tidx] = max(s_data[tidx], s_data[tidx + stride]);
        }
        __syncthreads();
    }
    max_val = s_data[0];
    
    float sum = 0.0f;
    for (int i = tidx; i < seq_len; i += blockDim.x) {
        float exp_val = expf(s_data[i] - max_val);
        s_data[i] = exp_val;
        sum += exp_val;
    }
    __syncthreads();
    
    // Reduce sum
    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
        if (tidx < stride) {
            s_data[tidx] += s_data[tidx + stride];
        }
        __syncthreads();
    }
    sum = s_data[0];
    
    // Normalize
    for (int i = tidx; i < seq_len; i += blockDim.x) {
        s_data[i] /= sum;
    }
    __syncthreads();
    
    // Compute output
    for (int i = tidx; i < seq_len * head_dim; i += blockDim.x) {
        int seq_idx = i / head_dim;
        int dim_idx = i % head_dim;
        float out_val = 0.0f;
        for (int j = 0; j < seq_len; j++) {
            out_val += s_data[j] * __half2float(v[bidx * seq_len * head_dim + j * head_dim + dim_idx]);
        }
        output[bidx * seq_len * head_dim + i] = __float2half(out_val);
    }
}

__global__ void fused_layernorm_kernel(
    const half* __restrict__ input,
    half* __restrict__ output,
    const float* __restrict__ gamma,
    const float* __restrict__ beta,
    int batch_size, int seq_len, int hidden_size
) {
    extern __shared__ float s_data[];
    
    int tidx = threadIdx.x;
    int bidx = blockIdx.x;
    
    float sum = 0.0f;
    float sq_sum = 0.0f;
    
    for (int i = tidx; i < hidden_size; i += blockDim.x) {
        int idx = bidx * hidden_size + i;
        float val = __half2float(input[idx]);
        sum += val;
        sq_sum += val * val;
    }
    
    // Reduce sum and sq_sum
    s_data[tidx] = sum;
    s_data[tidx + blockDim.x] = sq_sum;
    __syncthreads();
    
    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
        if (tidx < stride) {
            s_data[tidx] += s_data[tidx + stride];
            s_data[tidx + blockDim.x] += s_data[tidx + blockDim.x + stride];
        }
        __syncthreads();
    }
    
    if (tidx == 0) {
        float mean = s_data[0] / hidden_size;
        float var = (s_data[blockDim.x] / hidden_size) - (mean * mean);
        s_data[0] = mean;
        s_data[1] = rsqrtf(var + 1e-5f);
    }
    __syncthreads();
    
    float mean = s_data[0];
    float inv_std = s_data[1];
    
    for (int i = tidx; i < hidden_size; i += blockDim.x) {
        int idx = bidx * hidden_size + i;
        float normalized = (__half2float(input[idx]) - mean) * inv_std;
        output[idx] = __float2half(gamma[i] * normalized + beta[i]);
    }
}

extern "C" {

void launch_fused_attention(const half* q, const half* k, const half* v, half* output,
                            int batch_size, int seq_len, int num_heads, int head_dim) {
    dim3 grid(batch_size * num_heads);
    dim3 block(256);
    size_t shared_mem_size = seq_len * sizeof(float);
    fused_attention_kernel<<<grid, block, shared_mem_size>>>(q, k, v, output, batch_size, seq_len, num_heads, head_dim);
}

void launch_fused_layernorm(const half* input, half* output, const float* gamma, const float* beta,
                            int batch_size, int seq_len, int hidden_size) {
    dim3 grid(batch_size * seq_len);
    dim3 block(256);
    size_t shared_mem_size = 2 * block.x * sizeof(float);
    fused_layernorm_kernel<<<grid, block, shared_mem_size>>>(input, output, gamma, beta, batch_size, seq_len, hidden_size);
}

}  improve this code performacne signficanlty