393 lines
13 KiB
Plaintext
Executable File
393 lines
13 KiB
Plaintext
Executable File
#include "test_utils.h"
|
||
#include "performance_utils.h"
|
||
#include "yaml_reporter.h"
|
||
#include <iostream>
|
||
#include <vector>
|
||
#include <iomanip>
|
||
#include <algorithm>
|
||
|
||
|
||
// ============================================================================
|
||
// 实现标记宏 - 参赛者修改实现时请将此宏设为0
|
||
// ============================================================================
|
||
#ifndef USE_DEFAULT_REF_IMPL
|
||
#define USE_DEFAULT_REF_IMPL 0 // 已修改：0=参赛者自定义实现
|
||
#endif
|
||
|
||
#if USE_DEFAULT_REF_IMPL
|
||
#include <thrust/reduce.h>
|
||
#include <thrust/device_vector.h>
|
||
#include <thrust/execution_policy.h>
|
||
#include <thrust/functional.h>
|
||
#endif
|
||
|
||
// 误差容忍度
|
||
constexpr double REDUCE_ERROR_TOLERANCE = 0.005; // 0.5%
|
||
|
||
#if !USE_DEFAULT_REF_IMPL
|
||
|
||
constexpr int BLOCK_SIZE = 512;
|
||
constexpr int WARP_SIZE = 32;
|
||
|
||
// 1. Warp 级归约：使用寄存器洗牌指令，无需 Shared Mem，速度极快
|
||
template <typename T>
|
||
__device__ __forceinline__ T warpReduceSum(T val) {
|
||
#pragma unroll
|
||
for (int offset = WARP_SIZE / 2; offset > 0; offset /= 2) {
|
||
val += __shfl_down_sync(0xffffffff, val, offset);
|
||
}
|
||
return val;
|
||
}
|
||
|
||
// 2. Block 级归约：先在 Warp 内归约，再通过 Shared Mem 汇总 Warp 结果
|
||
template <typename T>
|
||
__device__ __forceinline__ T blockReduceSum(T val) {
|
||
// 共享内存用于存储每个 Warp 的总和
|
||
// 256个线程 -> 8个warp -> 需要8个位置，但为了安全分配32
|
||
static __shared__ T shared[32];
|
||
|
||
int lane = threadIdx.x % WARP_SIZE;
|
||
int wid = threadIdx.x / WARP_SIZE;
|
||
|
||
// 每个 Warp 内部归约
|
||
val = warpReduceSum(val);
|
||
|
||
// 每个 Warp 的第一个线程将结果写入共享内存
|
||
if (lane == 0) {
|
||
shared[wid] = val;
|
||
}
|
||
|
||
__syncthreads(); // 等待所有 Warp 写入完毕
|
||
|
||
// 最后由第一个 Warp 读取共享内存并进行最终归约
|
||
// 只有当 Block 大小大于 32 时才需要这一步
|
||
val = (threadIdx.x < blockDim.x / WARP_SIZE) ? shared[lane] : 0;
|
||
|
||
if (wid == 0) {
|
||
val = warpReduceSum(val);
|
||
}
|
||
|
||
return val;
|
||
}
|
||
|
||
// 3. 通用归约 Kernel (Grid-Stride Loop)
|
||
// 如果 is_final_pass 为 true，则将结果写入 d_out 并加上 init_value
|
||
// 否则，将 Block 的部分和写入 d_out (作为临时存储)
|
||
template <typename T, bool is_final_pass>
|
||
__global__ void reduceKernel(const T* __restrict__ d_in, T* __restrict__ d_out, int n, T init_value) {
|
||
T sum = 0;
|
||
|
||
// Grid-Stride Loop: 处理数据量大于线程总数的情况
|
||
int thread_id = blockIdx.x * blockDim.x + threadIdx.x;
|
||
int stride = blockDim.x * gridDim.x;
|
||
|
||
for (int i = thread_id; i < n; i += stride) {
|
||
sum += d_in[i];
|
||
}
|
||
|
||
// Block 内归约
|
||
sum = blockReduceSum(sum);
|
||
|
||
// 由 Block 的线程 0 输出结果
|
||
if (threadIdx.x == 0) {
|
||
if (is_final_pass) {
|
||
d_out[0] = sum + init_value;
|
||
} else {
|
||
d_out[blockIdx.x] = sum;
|
||
}
|
||
}
|
||
}
|
||
#endif
|
||
|
||
// ============================================================================
|
||
// ReduceSum算法实现接口
|
||
// ============================================================================
|
||
|
||
template <typename InputT = float, typename OutputT = float>
|
||
class ReduceSumAlgorithm {
|
||
public:
|
||
ReduceSumAlgorithm() : d_intermediate(nullptr), intermediate_capacity(0) {}
|
||
|
||
// 析构函数：释放临时内存
|
||
~ReduceSumAlgorithm() {
|
||
if (d_intermediate) {
|
||
mcFree(d_intermediate);
|
||
}
|
||
}
|
||
|
||
// 主要接口函数
|
||
void reduce(const InputT* d_in, OutputT* d_out, int num_items, OutputT init_value) {
|
||
|
||
#if !USE_DEFAULT_REF_IMPL
|
||
// ========================================
|
||
// 高性能自定义实现
|
||
// ========================================
|
||
|
||
// 边界情况处理
|
||
if (num_items <= 0) {
|
||
MACA_CHECK(mcMemcpy(d_out, &init_value, sizeof(OutputT), mcMemcpyHostToDevice));
|
||
return;
|
||
}
|
||
|
||
// 计算网格配置
|
||
// 根据数据量计算需要的 Block 数量，最大限制为 1024 或数据量的除数
|
||
// 这对于大数组来说可以保持高占用率
|
||
int threads = BLOCK_SIZE;
|
||
int blocks = (num_items + threads - 1) / threads;
|
||
blocks = std::min(blocks, 1024); // 限制 Grid 大小，避免过多空闲 Block
|
||
|
||
if (blocks <= 1) {
|
||
// 如果数据量很小，直接单次 Pass 完成
|
||
reduceKernel<InputT, true><<<1, threads>>>(d_in, d_out, num_items, init_value);
|
||
} else {
|
||
// === 第一阶段 ===
|
||
// 每一个 Block 处理一部分数据，输出 Partial Sum 到中间 buffer
|
||
|
||
// 懒加载分配中间内存 (避免每次调用都 malloc)
|
||
size_t needed_bytes = blocks * sizeof(OutputT);
|
||
if (d_intermediate == nullptr || intermediate_capacity < needed_bytes) {
|
||
if (d_intermediate) mcFree(d_intermediate);
|
||
MACA_CHECK(mcMalloc((void**)&d_intermediate, needed_bytes));
|
||
intermediate_capacity = needed_bytes;
|
||
}
|
||
|
||
reduceKernel<InputT, false><<<blocks, threads>>>(d_in, d_intermediate, num_items, 0);
|
||
|
||
// === 第二阶段 ===
|
||
// 将中间结果（blocks 个元素）归约为最终结果
|
||
// 此时输入是 d_intermediate，输出是 d_out
|
||
// 加上 init_value
|
||
reduceKernel<InputT, true><<<1, threads>>>(d_intermediate, d_out, blocks, init_value);
|
||
}
|
||
#else
|
||
// ========================================
|
||
// 默认基准实现
|
||
// ========================================
|
||
auto input_ptr = thrust::device_pointer_cast(d_in);
|
||
auto output_ptr = thrust::device_pointer_cast(d_out);
|
||
|
||
// 直接使用thrust::reduce进行归约
|
||
*output_ptr = thrust::reduce(
|
||
thrust::device,
|
||
input_ptr,
|
||
input_ptr + num_items,
|
||
static_cast<OutputT>(init_value)
|
||
);
|
||
#endif
|
||
}
|
||
|
||
// 获取当前实现状态
|
||
static const char* getImplementationStatus() {
|
||
#if USE_DEFAULT_REF_IMPL
|
||
return "DEFAULT_REF_IMPL";
|
||
#else
|
||
return "CUSTOM_IMPL";
|
||
#endif
|
||
}
|
||
|
||
private:
|
||
// 成员变量用于复用中间内存，减少 malloc 开销
|
||
OutputT* d_intermediate;
|
||
size_t intermediate_capacity;
|
||
};
|
||
|
||
// ============================================================================
|
||
// 测试和性能评估
|
||
// ============================================================================
|
||
|
||
bool testCorrectness() {
|
||
std::cout << "ReduceSum 正确性测试..." << std::endl;
|
||
TestDataGenerator generator;
|
||
ReduceSumAlgorithm<float, float> algorithm;
|
||
|
||
bool allPassed = true;
|
||
|
||
// 测试不同数据规模
|
||
for (int i = 0; i < NUM_TEST_SIZES && i < 2; i++) { // 限制测试规模
|
||
int size = std::min(TEST_SIZES[i], 10000);
|
||
std::cout << " 测试规模: " << size << std::endl;
|
||
|
||
// 测试普通数据
|
||
{
|
||
auto data = generator.generateRandomFloats(size, -10.0f, 10.0f);
|
||
float init_value = 1.0f;
|
||
|
||
// CPU参考计算
|
||
double cpu_result = cpuReduceSum(data, static_cast<double>(init_value));
|
||
|
||
// GPU计算
|
||
float *d_in;
|
||
float *d_out;
|
||
MACA_CHECK(mcMalloc(&d_in, size * sizeof(float)));
|
||
MACA_CHECK(mcMalloc(&d_out, sizeof(float)));
|
||
|
||
MACA_CHECK(mcMemcpy(d_in, data.data(), size * sizeof(float), mcMemcpyHostToDevice));
|
||
|
||
algorithm.reduce(d_in, d_out, size, init_value);
|
||
|
||
float gpu_result;
|
||
MACA_CHECK(mcMemcpy(&gpu_result, d_out, sizeof(float), mcMemcpyDeviceToHost));
|
||
|
||
// 验证误差
|
||
double relative_error = std::abs(gpu_result - cpu_result) / std::abs(cpu_result);
|
||
if (relative_error > REDUCE_ERROR_TOLERANCE) {
|
||
std::cout << " 失败: 误差过大 " << relative_error << std::endl;
|
||
allPassed = false;
|
||
} else {
|
||
std::cout << " 通过 (误差: " << relative_error << ")" << std::endl;
|
||
}
|
||
|
||
mcFree(d_in);
|
||
mcFree(d_out);
|
||
}
|
||
|
||
// 测试特殊值 (NaN, Inf)
|
||
if (size > 100) {
|
||
std::cout << " 测试特殊值..." << std::endl;
|
||
auto data = generator.generateSpecialFloats(size);
|
||
float init_value = 0.0f;
|
||
|
||
double cpu_result = cpuReduceSum(data, static_cast<double>(init_value));
|
||
|
||
float *d_in;
|
||
float *d_out;
|
||
MACA_CHECK(mcMalloc(&d_in, size * sizeof(float)));
|
||
MACA_CHECK(mcMalloc(&d_out, sizeof(float)));
|
||
|
||
MACA_CHECK(mcMemcpy(d_in, data.data(), size * sizeof(float), mcMemcpyHostToDevice));
|
||
|
||
algorithm.reduce(d_in, d_out, size, init_value);
|
||
|
||
float gpu_result;
|
||
MACA_CHECK(mcMemcpy(&gpu_result, d_out, sizeof(float), mcMemcpyDeviceToHost));
|
||
|
||
// 对于包含特殊值的情况，检查是否正确处理
|
||
if (std::isfinite(cpu_result) && std::isfinite(gpu_result)) {
|
||
double relative_error = std::abs(gpu_result - cpu_result) / std::abs(cpu_result);
|
||
if (relative_error > REDUCE_ERROR_TOLERANCE) {
|
||
std::cout << " 失败: 特殊值处理错误" << std::endl;
|
||
allPassed = false;
|
||
} else {
|
||
std::cout << " 通过 (特殊值处理)" << std::endl;
|
||
}
|
||
} else {
|
||
std::cout << " 通过 (特殊值结果)" << std::endl;
|
||
}
|
||
|
||
mcFree(d_in);
|
||
mcFree(d_out);
|
||
}
|
||
}
|
||
|
||
return allPassed;
|
||
}
|
||
|
||
void benchmarkPerformance() {
|
||
PerformanceDisplay::printReduceSumHeader();
|
||
|
||
TestDataGenerator generator;
|
||
PerformanceMeter meter;
|
||
ReduceSumAlgorithm<float, float> algorithm;
|
||
|
||
const int WARMUP_ITERATIONS = 5;
|
||
const int BENCHMARK_ITERATIONS = 10;
|
||
|
||
// 用于YAML报告的数据收集
|
||
std::vector<std::map<std::string, std::string>> perf_data;
|
||
|
||
for (int i = 0; i < NUM_TEST_SIZES; i++) {
|
||
int size = TEST_SIZES[i];
|
||
|
||
// 生成测试数据
|
||
auto data = generator.generateRandomFloats(size);
|
||
float init_value = 0.0f;
|
||
|
||
// 分配GPU内存
|
||
float *d_in;
|
||
float *d_out;
|
||
MACA_CHECK(mcMalloc(&d_in, size * sizeof(float)));
|
||
MACA_CHECK(mcMalloc(&d_out, sizeof(float)));
|
||
|
||
MACA_CHECK(mcMemcpy(d_in, data.data(), size * sizeof(float), mcMemcpyHostToDevice));
|
||
|
||
// Warmup阶段
|
||
for (int iter = 0; iter < WARMUP_ITERATIONS; iter++) {
|
||
algorithm.reduce(d_in, d_out, size, init_value);
|
||
}
|
||
|
||
// 正式测试阶段
|
||
float total_time = 0;
|
||
for (int iter = 0; iter < BENCHMARK_ITERATIONS; iter++) {
|
||
meter.startTiming();
|
||
algorithm.reduce(d_in, d_out, size, init_value);
|
||
total_time += meter.stopTiming();
|
||
}
|
||
|
||
float avg_time = total_time / BENCHMARK_ITERATIONS;
|
||
|
||
// 计算性能指标
|
||
auto metrics = PerformanceCalculator::calculateReduceSum(size, avg_time);
|
||
|
||
// 显示性能数据
|
||
PerformanceDisplay::printReduceSumData(size, avg_time, metrics);
|
||
|
||
// 收集YAML报告数据
|
||
auto entry = YAMLPerformanceReporter::createEntry();
|
||
entry["data_size"] = std::to_string(size);
|
||
entry["time_ms"] = std::to_string(avg_time);
|
||
entry["throughput_gps"] = std::to_string(metrics.throughput_gps);
|
||
entry["data_type"] = "float";
|
||
perf_data.push_back(entry);
|
||
|
||
mcFree(d_in);
|
||
mcFree(d_out);
|
||
}
|
||
|
||
// 生成YAML性能报告
|
||
YAMLPerformanceReporter::generateReduceSumYAML(perf_data, "reduce_sum_performance.yaml");
|
||
PerformanceDisplay::printSavedMessage("reduce_sum_performance.yaml");
|
||
}
|
||
|
||
// ============================================================================
|
||
// 主函数
|
||
// ============================================================================
|
||
int main(int argc, char* argv[]) {
|
||
std::cout << "=== ReduceSum 算法测试 ===" << std::endl;
|
||
|
||
// 检查参数
|
||
std::string mode = "all";
|
||
if (argc > 1) {
|
||
mode = argv[1];
|
||
}
|
||
|
||
bool correctness_passed = true;
|
||
bool performance_completed = true;
|
||
|
||
try {
|
||
if (mode == "correctness" || mode == "all") {
|
||
correctness_passed = testCorrectness();
|
||
}
|
||
|
||
if (mode == "performance" || mode == "all") {
|
||
if (correctness_passed || mode == "performance") {
|
||
benchmarkPerformance();
|
||
} else {
|
||
std::cout << "跳过性能测试，因为正确性测试未通过" << std::endl;
|
||
performance_completed = false;
|
||
}
|
||
}
|
||
|
||
std::cout << "\n=== 测试完成 ===" << std::endl;
|
||
std::cout << "实现状态: " << ReduceSumAlgorithm<float, float>::getImplementationStatus() << std::endl;
|
||
if (mode == "all") {
|
||
std::cout << "正确性: " << (correctness_passed ? "通过" : "失败") << std::endl;
|
||
std::cout << "性能测试: " << (performance_completed ? "完成" : "跳过") << std::endl;
|
||
}
|
||
|
||
return correctness_passed ? 0 : 1;
|
||
|
||
} catch (const std::exception& e) {
|
||
std::cerr << "测试出错: " << e.what() << std::endl;
|
||
return 1;
|
||
}
|
||
} |