性能优化技术

// ❌ 非合并访问 - 步长为32
__global__ void strided_access(float *data, float *out) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    out[tid] = data[tid * 32];  // 每个线程访问不同cacheline
}

// ✅ 合并访问 - 连续访问
__global__ void coalesced_access(float *data, float *out) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    out[tid] = data[tid];  // 连续访问，1次事务
}

// ✅ 向量化加载 - float4一次加载128位
__global__ void vectorized_load(float *data, float *out) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    float4 temp = reinterpret_cast<float4*>(data)[tid];
    temp.x *= 2.0f;
    temp.y *= 2.0f;
    temp.z *= 2.0f;
    temp.w *= 2.0f;
    reinterpret_cast<float4*>(out)[tid] = temp;
}

// ❌ Bank Conflict示例
__shared__ float shared[32][33];  // 已经padding了
float val = shared[threadIdx.x][0];  // 每行32字节=8个float，跨越多个Bank

// ❌ 转置导致冲突
__shared__ float tile[32][32];
float val = tile[0][threadIdx.x];  // 每个线程访问同一行不同列
// 如果列索引%32相同，则冲突

// ✅ Padding消除冲突
__shared__ float tile[32][33];  // 多一列，错开Bank
float val = tile[0][threadIdx.x];  // 现在无冲突

// Stream并发示例 - 计算传输重叠
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

// Stream1: 数据传输 + 计算
cudaMemcpyAsync(d_A, h_A, size, cudaMemcpyHostToDevice, stream1);
matmul_kernel<<<grid, block, 0, stream1>>>(d_A, d_B, d_C1);

// Stream2: 数据传输 + 计算（与Stream1重叠）
cudaMemcpyAsync(d_B2, h_B2, size, cudaMemcpyHostToDevice, stream2);
matmul_kernel<<<grid, block, 0, stream2>>>(d_A, d_B2, d_C2);

// 等待所有Stream完成
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);

// ❌ 两个Kernel，多次Global Memory访问
relu_kernel<<<...>>>(input, temp);   // 读input，写temp
softmax_kernel<<<...>>>(temp, output); // 读temp，写output
// 总共: 2次Global Read + 2次Global Write

// ✅ 融合为一个Kernel
__global__ void relu_softmax_kernel(float *input, float *output) {
    __shared__ float smem[256];
    
    // 加载到Shared Memory
    smem[threadIdx.x] = input[threadIdx.x + blockIdx.x * blockDim.x];
    __syncthreads();
    
    // ReLU + Softmax在Shared Memory中完成
    float val = fmaxf(smem[threadIdx.x], 0.0f);  // ReLU
    // ... Softmax计算 ...
    
    output[threadIdx.x + blockIdx.x * blockDim.x] = result;
}
// 总共: 1次Global Read + 1次Global Write

// 循环展开示例
__global__ void unrolled_kernel(float *data, int n) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = blockDim.x * gridDim.x;
    
    // 4路循环展开
    for (int i = tid; i < n; i += stride * 4) {
        if (i < n)         data[i] *= 2;
        if (i + stride < n)   data[i + stride] *= 2;
        if (i + stride*2 < n) data[i + stride*2] *= 2;
        if (i + stride*3 < n) data[i + stride*3] *= 2;
    }
}