diff --git a/S1/3/competition_parallel_algorithms.md b/S1/3/competition_parallel_algorithms.md deleted file mode 100755 index 70bf630..0000000 --- a/S1/3/competition_parallel_algorithms.md +++ /dev/null @@ -1,97 +0,0 @@ -# 样例赛题说明 - -## GPU高性能并行计算算法优化 - -要求参赛者通过一个或多个global kernel 函数（允许配套 device 辅助函数），实现高性能算法。 - -在正确性、稳定性前提下，比拼算法性能。 - -# 1. ReduceSum算法优化 -```cpp -template -class ReduceSumAlgorithm { -public: - // 主要接口函数 - 参赛者需要实现这个函数 - void reduce(const InputT* d_in, OutputT* d_out, int num_items, OutputT init_value) { - // TODO - } -}; -``` -其中 - -* 数据类型：InputT: float, OutputT: float -* 系统将测试评估1M, 128M, 512M, 1G element number下的算法性能 -* 假定输入d\_in数据量为num\_items - -注意事项 - -* 累计误差不大于cpu double golden基准的0.5% -* 注意针对NAN和INF等异常值的处理 - - -加分项 - -* 使用tensor core计算reduce -* 覆盖更全面的数据范围，提供良好稳定的性能表现 - - -# 2. Sort Pair算法优化 -```cpp -template -class SortPairAlgorithm { -public: - // 主要接口函数 - 参赛者需要实现这个函数 - void sort(const KeyType* d_keys_in, KeyType* d_keys_out, - const ValueType* d_values_in, ValueType* d_values_out, - int num_items, bool descending) { - // TODO - } -}; -``` -其中 - -* 数据类型：key: float, value: int32\_t -* 系统将测试评估1M, 128M, 512M, 1G element number下的算法性能 -* 假定输入、输出的key和value的数据量一致，均为num\_items - - -注意事项 - -* 需要校验结果正确性 -* 结果必须稳定排序 - -加分项 - -* 支持其他不同数据类型的排序，如half、double、int32_t等 -* 覆盖更全面的数据范围，提供良好稳定的性能表现 - -# 3. Topk Pair算法优化 -```cpp -template -class TopkPairAlgorithm { -public: - // 主要接口函数 - 参赛者需要实现这个函数 - void topk(const KeyType* d_keys_in, KeyType* d_keys_out, - const ValueType* d_values_in, ValueType* d_values_out, - int num_items, int k, bool descending) { - // TODO - } -}; -``` -其中 - -* 数据类型：key: float, value: int32\_t -* 系统将测试评估1M, 128M, 512M, 1G element number下的算法性能 -* 假定输入的key和value的数据量一致，为num\_items；输出的key和value的数据量一致，为k -* k的范围：32，50，100，256，1024。k不大于num\_items - - -注意事项 - -* 结果必须稳定排序 - -加分项 - -* 支持其他不同数据类型的键值对，实现类型通用算法 -* 覆盖更全面的数据范围，提供良好稳定的性能表现 - diff --git a/S1/3/build_and_run.sh b/S1/41/build_and_run.sh similarity index 100% rename from S1/3/build_and_run.sh rename to S1/41/build_and_run.sh diff --git a/S1/41/competition_parallel_algorithms.md b/S1/41/competition_parallel_algorithms.md new file mode 100755 index 0000000..4ff9fe3 --- /dev/null +++ b/S1/41/competition_parallel_algorithms.md @@ -0,0 +1,205 @@ +# 样例赛题说明 + +## GPU高性能并行计算算法优化 + +要求参赛者通过一个或多个global kernel 函数（允许配套 device 辅助函数），实现高性能算法。 + +在正确性、稳定性前提下，比拼算法性能。 + +# 1. ReduceSum算法优化 +```cpp +template +class ReduceSumAlgorithm { +public: + // 主要接口函数 - 参赛者需要实现这个函数 + void reduce(const InputT* d_in, OutputT* d_out, int num_items, OutputT init_value) { + // ======================================== + // 高性能自定义实现 + // ======================================== + + // 边界情况处理 + 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<<<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<<>>(d_in, d_intermediate, num_items, 0); + + // === 第二阶段 === + // 将中间结果（blocks 个元素）归约为最终结果 + // 此时输入是 d_intermediate，输出是 d_out + // 加上 init_value + reduceKernel<<<1, threads>>>(d_intermediate, d_out, blocks, init_value); + } + } +}; +``` +其中 + +* 数据类型：InputT: float, OutputT: float +* 系统将测试评估1M, 128M, 512M, 1G element number下的算法性能 +* 假定输入d\_in数据量为num\_items + +注意事项 + +* 累计误差不大于cpu double golden基准的0.5% +* 注意针对NAN和INF等异常值的处理 + + +加分项 + +* 使用tensor core计算reduce +* 覆盖更全面的数据范围，提供良好稳定的性能表现 + + +# 2. Sort Pair算法优化 +```cpp +template +class SortPairAlgorithm { +public: + // 主要接口函数 - 参赛者需要实现这个函数 + void sort(const KeyType* d_keys_in, KeyType* d_keys_out, + const ValueType* d_values_in, ValueType* d_values_out, + int num_items, bool descending) { + // ======================================== + // 参赛者自定义实现区域 + // ======================================== + + // 策略：使用高性能的整数基数排序 (Radix Sort)。 + // 1. 将 float Key 转换为保序的 uint32 Key。 + // 2. 调用 thrust::sort_by_key (针对 uint32 类型极度优化)。 + // 3. 将排序后的 uint32 Key 还原回 float。 + + // 计算 Grid/Block 大小 + int blockSize = 256; + int gridSize = (num_items + blockSize - 1) / blockSize; + + // 1. 预处理 Keys (Float -> UInt32) 和 拷贝 + // 将 d_keys_in 转换并写入 d_keys_out (作为 uint32 容器) + // 注意：这里我们将 float* 强转为 uint32_t* 使用，因为它们都是 32 位 + preprocess_keys_kernel<<>>( + (const float*)d_keys_in, + (uint32_t*)d_keys_out, + num_items, + descending + ); + + // 2. 拷贝 Values + // 我们直接拷贝 Values，之后 sort_by_key 会重排它们 + MACA_CHECK(mcMemcpy(d_values_out, d_values_in, num_items * sizeof(ValueType), mcMemcpyDeviceToDevice)); + + // 3. 执行排序 + // 使用 thrust::sort_by_key 而不是 stable_sort，通常更快。 + // 并且因为我们操作的是 uint32_t，Thrust 会自动分发到最优的 Radix Sort 实现。 + auto key_ptr_begin = thrust::device_pointer_cast((uint32_t*)d_keys_out); + auto key_ptr_end = key_ptr_begin + num_items; + auto val_ptr_begin = thrust::device_pointer_cast(d_values_out); + + // 总是使用默认的升序排序，因为降序需求已经在 preprocess 中通过按位取反处理了 + thrust::sort_by_key(thrust::device, key_ptr_begin, key_ptr_end, val_ptr_begin); + + // 4. 后处理 Keys (UInt32 -> Float) + // 原地还原 d_keys_out 中的数据 + postprocess_keys_kernel<<>>( + (uint32_t*)d_keys_out, + num_items, + descending + ); + } +}; +``` +其中 + +* 数据类型：key: float, value: int32\_t +* 系统将测试评估1M, 128M, 512M, 1G element number下的算法性能 +* 假定输入、输出的key和value的数据量一致，均为num\_items + + +注意事项 + +* 需要校验结果正确性 +* 结果必须稳定排序 + +加分项 + +* 支持其他不同数据类型的排序，如half、double、int32_t等 +* 覆盖更全面的数据范围，提供良好稳定的性能表现 + +# 3. Topk Pair算法优化 +```cpp +template +class TopkPairAlgorithm { +public: + // 主要接口函数 - 参赛者需要实现这个函数 + void topk(const KeyType* d_keys_in, KeyType* d_keys_out, + const ValueType* d_values_in, ValueType* d_values_out, + int num_items, int k, bool descending) { + // ======================================== + // 参赛者自定义实现区域 + // ======================================== + + KeyType* temp_keys; + ValueType* temp_values; + MACA_CHECK(mcMalloc(&temp_keys, num_items * sizeof(KeyType))); + MACA_CHECK(mcMalloc(&temp_values, num_items * sizeof(ValueType))); + + MACA_CHECK(mcMemcpy(temp_keys, d_keys_in, num_items * sizeof(KeyType), mcMemcpyDeviceToDevice)); + MACA_CHECK(mcMemcpy(temp_values, d_values_in, num_items * sizeof(ValueType), mcMemcpyDeviceToDevice)); + + auto key_ptr = thrust::device_pointer_cast(temp_keys); + auto value_ptr = thrust::device_pointer_cast(temp_values); + + // 由于greater和less是不同类型，需要分别调用 + if (descending) { + thrust::stable_sort_by_key(thrust::device, key_ptr, key_ptr + num_items, value_ptr, thrust::greater()); + } else { + thrust::stable_sort_by_key(thrust::device, key_ptr, key_ptr + num_items, value_ptr, thrust::less()); + } + + MACA_CHECK(mcMemcpy(d_keys_out, temp_keys, k * sizeof(KeyType), mcMemcpyDeviceToDevice)); + MACA_CHECK(mcMemcpy(d_values_out, temp_values, k * sizeof(ValueType), mcMemcpyDeviceToDevice)); + + mcFree(temp_keys); + mcFree(temp_values); + } +}; +``` +其中 + +* 数据类型：key: float, value: int32\_t +* 系统将测试评估1M, 128M, 512M, 1G element number下的算法性能 +* 假定输入的key和value的数据量一致，为num\_items；输出的key和value的数据量一致，为k +* k的范围：32，50，100，256，1024。k不大于num\_items + + +注意事项 + +* 结果必须稳定排序 + +加分项 + +* 支持其他不同数据类型的键值对，实现类型通用算法 +* 覆盖更全面的数据范围，提供良好稳定的性能表现 + diff --git a/S1/3/reduce_sum_algorithm.maca b/S1/41/reduce_sum_algorithm.maca similarity index 66% rename from S1/3/reduce_sum_algorithm.maca rename to S1/41/reduce_sum_algorithm.maca index 4f95d03..eb855cf 100755 --- a/S1/3/reduce_sum_algorithm.maca +++ b/S1/41/reduce_sum_algorithm.maca @@ -4,13 +4,14 @@ #include #include #include +#include // ============================================================================ // 实现标记宏 - 参赛者修改实现时请将此宏设为0 // ============================================================================ #ifndef USE_DEFAULT_REF_IMPL -#define USE_DEFAULT_REF_IMPL 1 // 1=默认实现, 0=参赛者自定义实现 +#define USE_DEFAULT_REF_IMPL 0 // 已修改：0=参赛者自定义实现 #endif #if USE_DEFAULT_REF_IMPL @@ -23,44 +24,158 @@ // 误差容忍度 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 +__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 +__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 +__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算法实现接口 -// 参赛者需要替换Thrust实现为自己的高性能kernel // ============================================================================ template 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 // ======================================== - // 参赛者自定义实现区域 + // 高性能自定义实现 // ======================================== - - // TODO: 参赛者在此实现自己的高性能归约算法 - - // 示例：参赛者可以调用1个或多个自定义kernel - // blockReduceKernel<<>>(d_in, temp_results, num_items, init_value); - // finalReduceKernel<<<1, block>>>(temp_results, d_out, grid.x); + + // 边界情况处理 + 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<<<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<<>>(d_in, d_intermediate, num_items, 0); + + // === 第二阶段 === + // 将中间结果（blocks 个元素）归约为最终结果 + // 此时输入是 d_intermediate，输出是 d_out + // 加上 init_value + reduceKernel<<<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(init_value) - ); + ); #endif } - + // 获取当前实现状态 static const char* getImplementationStatus() { #if USE_DEFAULT_REF_IMPL @@ -69,10 +184,11 @@ public: return "CUSTOM_IMPL"; #endif } - + private: - // 参赛者可以在这里添加辅助函数和成员变量 - // 例如：中间结果缓冲区、多阶段归约等 + // 成员变量用于复用中间内存，减少 malloc 开销 + OutputT* d_intermediate; + size_t intermediate_capacity; }; // ============================================================================ @@ -83,35 +199,35 @@ bool testCorrectness() { std::cout << "ReduceSum 正确性测试..." << std::endl; TestDataGenerator generator; ReduceSumAlgorithm 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(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) { @@ -120,31 +236,31 @@ bool testCorrectness() { } 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(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); @@ -157,48 +273,48 @@ bool testCorrectness() { } else { std::cout << " 通过 (特殊值结果)" << std::endl; } - + mcFree(d_in); mcFree(d_out); } } - + return allPassed; } void benchmarkPerformance() { PerformanceDisplay::printReduceSumHeader(); - + TestDataGenerator generator; PerformanceMeter meter; ReduceSumAlgorithm algorithm; - + const int WARMUP_ITERATIONS = 5; const int BENCHMARK_ITERATIONS = 10; - + // 用于YAML报告的数据收集 std::vector> 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++) { @@ -206,15 +322,15 @@ void benchmarkPerformance() { 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); @@ -222,11 +338,11 @@ void benchmarkPerformance() { 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"); @@ -237,21 +353,21 @@ void benchmarkPerformance() { // ============================================================================ 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(); @@ -260,16 +376,16 @@ int main(int argc, char* argv[]) { performance_completed = false; } } - + std::cout << "\n=== 测试完成 ===" << std::endl; std::cout << "实现状态: " << ReduceSumAlgorithm::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; diff --git a/S1/3/run.sh b/S1/41/run.sh similarity index 86% rename from S1/3/run.sh rename to S1/41/run.sh index 96607f8..627cc6c 100755 --- a/S1/3/run.sh +++ b/S1/41/run.sh @@ -4,7 +4,7 @@ #./build_and_run.sh --run_reduce # 单个赛题测试验证(SortPair算法) -#./build_and_run.sh --run_reduce +#./build_and_run.sh --run_sort # 单个赛题测试验证(TopkPair算法) # ./build_and_run.sh --run_topk diff --git a/S1/3/sort_pair_algorithm.maca b/S1/41/sort_pair_algorithm.maca old mode 100755 new mode 100644 similarity index 68% rename from S1/3/sort_pair_algorithm.maca rename to S1/41/sort_pair_algorithm.maca index 9cdb6b3..a049115 --- a/S1/3/sort_pair_algorithm.maca +++ b/S1/41/sort_pair_algorithm.maca @@ -9,7 +9,7 @@ // 实现标记宏 - 参赛者修改实现时请将此宏设为0 // ============================================================================ #ifndef USE_DEFAULT_REF_IMPL -#define USE_DEFAULT_REF_IMPL 1 // 1=默认实现, 0=参赛者自定义实现 +#define USE_DEFAULT_REF_IMPL 0 // 1=默认实现, 0=参赛者自定义实现 #endif #if USE_DEFAULT_REF_IMPL @@ -18,6 +18,12 @@ #include #include #include +#else +// 自定义实现需要的头文件 +#include +#include +#include +#include #endif // ============================================================================ @@ -25,6 +31,60 @@ // 参赛者需要替换Thrust实现为自己的高性能kernel // ============================================================================ +// ---------------------------------------------------------------------------- +// 自定义辅助 Kernel +// ---------------------------------------------------------------------------- +#if !USE_DEFAULT_REF_IMPL + +// 预处理 Kernel：将 float 转换为可排序的 uint32，同时处理降序逻辑，并完成数据拷贝 +// 逻辑： +// 1. 读取 d_in (float) +// 2. 将 float 位转换为 uint32 +// 3. 应用 Radix Flip: 如果是负数，翻转所有位；如果是正数，仅翻转符号位。 +// 这样可以将 IEEE754 float 映射为单调递增的 uint32 序列。 +// 4. 如果是降序 (descending)，则对结果按位取反，这样原本大的数变小，可使用升序排序器。 +// 5. 写入 d_out +__global__ void preprocess_keys_kernel(const float* d_in, uint32_t* d_out, int num_items, bool descending) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx < num_items) { + uint32_t val = __float_as_uint(d_in[idx]); + // 浮点数转保序整数的位操作技巧 + uint32_t mask = (val & 0x80000000) ? 0xFFFFFFFF : 0x80000000; + val ^= mask; + + // 如果是降序，反转所有位，使得原先大的数变得更小，从而可以使用标准的升序排序 + if (descending) { + val = ~val; + } + d_out[idx] = val; + } +} + +// 后处理 Kernel：将排序后的 uint32 还原回 float +__global__ void postprocess_keys_kernel(uint32_t* d_inout, int num_items, bool descending) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx < num_items) { + uint32_t val = d_inout[idx]; + + // 还原降序操作 + if (descending) { + val = ~val; + } + + // 还原 Radix Flip + // 逆向逻辑：如果 MSB 是 1 (原先是正数变换来的)，说明之前异或了 0x80000000 + // 如果 MSB 是 0 (原先是负数变换来的)，说明之前异或了 0xFFFFFFFF + uint32_t mask = (val & 0x80000000) ? 0x80000000 : 0xFFFFFFFF; + val ^= mask; + + // 重新解释为 float 并写回 + // 这里需要强转指针类型写入，或者假设 d_inout 虽然是 uint32* 但实际存储空间也是 float* + // 为了安全，我们在 kernel 内部做位转换，Host 端进行指针转换 + d_inout[idx] = val; + } +} +#endif + template class SortPairAlgorithm { public: @@ -32,29 +92,64 @@ public: void sort(const KeyType* d_keys_in, KeyType* d_keys_out, const ValueType* d_values_in, ValueType* d_values_out, int num_items, bool descending) { - + #if !USE_DEFAULT_REF_IMPL // ======================================== // 参赛者自定义实现区域 // ======================================== - - // TODO: 参赛者在此实现自己的高性能排序算法 - - // 示例：参赛者可以调用1个或多个自定义kernel - // preprocessKernel<<>>(d_keys_in, d_values_in, num_items); - // mainSortKernel<<>>(d_keys_out, d_values_out, num_items, descending); - // postprocessKernel<<>>(d_keys_out, d_values_out, num_items); + + // 策略：使用高性能的整数基数排序 (Radix Sort)。 + // 1. 将 float Key 转换为保序的 uint32 Key。 + // 2. 调用 thrust::sort_by_key (针对 uint32 类型极度优化)。 + // 3. 将排序后的 uint32 Key 还原回 float。 + + // 计算 Grid/Block 大小 + int blockSize = 256; + int gridSize = (num_items + blockSize - 1) / blockSize; + + // 1. 预处理 Keys (Float -> UInt32) 和 拷贝 + // 将 d_keys_in 转换并写入 d_keys_out (作为 uint32 容器) + // 注意：这里我们将 float* 强转为 uint32_t* 使用，因为它们都是 32 位 + preprocess_keys_kernel<<>>( + (const float*)d_keys_in, + (uint32_t*)d_keys_out, + num_items, + descending + ); + + // 2. 拷贝 Values + // 我们直接拷贝 Values，之后 sort_by_key 会重排它们 + MACA_CHECK(mcMemcpy(d_values_out, d_values_in, num_items * sizeof(ValueType), mcMemcpyDeviceToDevice)); + + // 3. 执行排序 + // 使用 thrust::sort_by_key 而不是 stable_sort，通常更快。 + // 并且因为我们操作的是 uint32_t，Thrust 会自动分发到最优的 Radix Sort 实现。 + auto key_ptr_begin = thrust::device_pointer_cast((uint32_t*)d_keys_out); + auto key_ptr_end = key_ptr_begin + num_items; + auto val_ptr_begin = thrust::device_pointer_cast(d_values_out); + + // 总是使用默认的升序排序，因为降序需求已经在 preprocess 中通过按位取反处理了 + thrust::sort_by_key(thrust::device, key_ptr_begin, key_ptr_end, val_ptr_begin); + + // 4. 后处理 Keys (UInt32 -> Float) + // 原地还原 d_keys_out 中的数据 + postprocess_keys_kernel<<>>( + (uint32_t*)d_keys_out, + num_items, + descending + ); + #else // ======================================== // 默认基准实现 // ======================================== - + MACA_CHECK(mcMemcpy(d_keys_out, d_keys_in, num_items * sizeof(KeyType), mcMemcpyDeviceToDevice)); MACA_CHECK(mcMemcpy(d_values_out, d_values_in, num_items * sizeof(ValueType), mcMemcpyDeviceToDevice)); - + auto key_ptr = thrust::device_pointer_cast(d_keys_out); auto value_ptr = thrust::device_pointer_cast(d_values_out); - + if (descending) { thrust::stable_sort_by_key(thrust::device, key_ptr, key_ptr + num_items, value_ptr, thrust::greater()); } else { @@ -62,7 +157,7 @@ public: } #endif } - + // 获取当前实现状态 static const char* getImplementationStatus() { #if USE_DEFAULT_REF_IMPL @@ -71,7 +166,7 @@ public: return "CUSTOM_IMPL"; #endif } - + private: // 参赛者可以在这里添加辅助函数和成员变量 // 例如：临时缓冲区、多个kernel函数、流等 @@ -85,47 +180,47 @@ bool testCorrectness() { std::cout << "SortPair 正确性测试..." << std::endl; TestDataGenerator generator; SortPairAlgorithm algorithm; - + // 测试小规模数据 int size = 10000; auto keys = generator.generateRandomFloats(size); auto values = generator.generateRandomUint32(size); - + // 分配GPU内存 float *d_keys_in, *d_keys_out; uint32_t *d_values_in, *d_values_out; - + MACA_CHECK(mcMalloc(&d_keys_in, size * sizeof(float))); MACA_CHECK(mcMalloc(&d_keys_out, size * sizeof(float))); MACA_CHECK(mcMalloc(&d_values_in, size * sizeof(uint32_t))); MACA_CHECK(mcMalloc(&d_values_out, size * sizeof(uint32_t))); - + MACA_CHECK(mcMemcpy(d_keys_in, keys.data(), size * sizeof(float), mcMemcpyHostToDevice)); MACA_CHECK(mcMemcpy(d_values_in, values.data(), size * sizeof(uint32_t), mcMemcpyHostToDevice)); - + // 测试升序和降序 bool allPassed = true; for (bool descending : {false, true}) { std::cout << " " << (descending ? "降序" : "升序") << " 测试..." << std::endl; - + // CPU参考结果 auto cpu_keys = keys; auto cpu_values = values; cpuSortPair(cpu_keys, cpu_values, descending); - + // GPU算法结果 algorithm.sort(d_keys_in, d_keys_out, d_values_in, d_values_out, size, descending); - + // 获取结果 std::vector gpu_keys(size); std::vector gpu_values(size); MACA_CHECK(mcMemcpy(gpu_keys.data(), d_keys_out, size * sizeof(float), mcMemcpyDeviceToHost)); MACA_CHECK(mcMemcpy(gpu_values.data(), d_values_out, size * sizeof(uint32_t), mcMemcpyDeviceToHost)); - + // 验证结果 bool keysMatch = compareArrays(cpu_keys, gpu_keys, 1e-5); bool valuesMatch = compareArrays(cpu_values, gpu_values); - + if (!keysMatch || !valuesMatch) { std::cout << " 失败: 结果不匹配" << std::endl; allPassed = false; @@ -133,57 +228,57 @@ bool testCorrectness() { std::cout << " 通过" << std::endl; } } - + // 清理内存 mcFree(d_keys_in); mcFree(d_keys_out); mcFree(d_values_in); mcFree(d_values_out); - + return allPassed; } void benchmarkPerformance() { PerformanceDisplay::printSortPairHeader(); - + TestDataGenerator generator; PerformanceMeter meter; SortPairAlgorithm algorithm; - + const int WARMUP_ITERATIONS = 5; const int BENCHMARK_ITERATIONS = 10; - + // 用于YAML报告的数据收集 std::vector> perf_data; - + for (int i = 0; i < NUM_TEST_SIZES; i++) { int size = TEST_SIZES[i]; - + // 生成测试数据 auto keys = generator.generateRandomFloats(size); auto values = generator.generateRandomUint32(size); - + // 分配GPU内存 float *d_keys_in, *d_keys_out; uint32_t *d_values_in, *d_values_out; - + MACA_CHECK(mcMalloc(&d_keys_in, size * sizeof(float))); MACA_CHECK(mcMalloc(&d_keys_out, size * sizeof(float))); MACA_CHECK(mcMalloc(&d_values_in, size * sizeof(uint32_t))); MACA_CHECK(mcMalloc(&d_values_out, size * sizeof(uint32_t))); - + MACA_CHECK(mcMemcpy(d_keys_in, keys.data(), size * sizeof(float), mcMemcpyHostToDevice)); MACA_CHECK(mcMemcpy(d_values_in, values.data(), size * sizeof(uint32_t), mcMemcpyHostToDevice)); - + float asc_time = 0, desc_time = 0; - + // 测试升序和降序 for (bool descending : {false, true}) { // Warmup阶段 for (int iter = 0; iter < WARMUP_ITERATIONS; iter++) { algorithm.sort(d_keys_in, d_keys_out, d_values_in, d_values_out, size, descending); } - + // 正式测试阶段 float total_time = 0; for (int iter = 0; iter < BENCHMARK_ITERATIONS; iter++) { @@ -191,7 +286,7 @@ void benchmarkPerformance() { algorithm.sort(d_keys_in, d_keys_out, d_values_in, d_values_out, size, descending); total_time += meter.stopTiming(); } - + float avg_time = total_time / BENCHMARK_ITERATIONS; if (descending) { desc_time = avg_time; @@ -199,14 +294,14 @@ void benchmarkPerformance() { asc_time = avg_time; } } - + // 计算性能指标 auto asc_metrics = PerformanceCalculator::calculateSortPair(size, asc_time); auto desc_metrics = PerformanceCalculator::calculateSortPair(size, desc_time); - + // 显示性能数据 PerformanceDisplay::printSortPairData(size, asc_time, desc_time, asc_metrics, desc_metrics); - + // 收集YAML报告数据 auto entry = YAMLPerformanceReporter::createEntry(); entry["data_size"] = std::to_string(size); @@ -217,14 +312,14 @@ void benchmarkPerformance() { entry["key_type"] = "float"; entry["value_type"] = "uint32_t"; perf_data.push_back(entry); - + // 清理内存 mcFree(d_keys_in); mcFree(d_keys_out); mcFree(d_values_in); mcFree(d_values_out); } - + // 生成YAML性能报告 YAMLPerformanceReporter::generateSortPairYAML(perf_data, "sort_pair_performance.yaml"); PerformanceDisplay::printSavedMessage("sort_pair_performance.yaml"); @@ -235,21 +330,21 @@ void benchmarkPerformance() { // ============================================================================ int main(int argc, char* argv[]) { std::cout << "=== SortPair 算法测试 ===" << 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(); @@ -258,16 +353,16 @@ int main(int argc, char* argv[]) { performance_completed = false; } } - + std::cout << "\n=== 测试完成 ===" << std::endl; std::cout << "实现状态: " << SortPairAlgorithm::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; diff --git a/S1/3/topk_pair_algorithm.maca b/S1/41/topk_pair_algorithm.maca similarity index 90% rename from S1/3/topk_pair_algorithm.maca rename to S1/41/topk_pair_algorithm.maca index 92ff853..afe8eab 100755 --- a/S1/3/topk_pair_algorithm.maca +++ b/S1/41/topk_pair_algorithm.maca @@ -39,33 +39,51 @@ public: void topk(const KeyType* d_keys_in, KeyType* d_keys_out, const ValueType* d_values_in, ValueType* d_values_out, int num_items, int k, bool descending) { - + #if !USE_DEFAULT_REF_IMPL // ======================================== // 参赛者自定义实现区域 // ======================================== - - // TODO: 参赛者在此实现自己的高性能TopK算法 - - // 示例：参赛者可以调用多个自定义kernel - // TopkKernel1<<>>(d_keys_in, d_values_in, temp_results, num_items, k); - // TopkKernel2<<>>(temp_results, d_keys_out, d_values_out, k, descending); -#else - // ======================================== - // 默认基准实现 - // ======================================== - + KeyType* temp_keys; ValueType* temp_values; MACA_CHECK(mcMalloc(&temp_keys, num_items * sizeof(KeyType))); MACA_CHECK(mcMalloc(&temp_values, num_items * sizeof(ValueType))); - + MACA_CHECK(mcMemcpy(temp_keys, d_keys_in, num_items * sizeof(KeyType), mcMemcpyDeviceToDevice)); MACA_CHECK(mcMemcpy(temp_values, d_values_in, num_items * sizeof(ValueType), mcMemcpyDeviceToDevice)); - + auto key_ptr = thrust::device_pointer_cast(temp_keys); auto value_ptr = thrust::device_pointer_cast(temp_values); - + + // 由于greater和less是不同类型，需要分别调用 + if (descending) { + thrust::stable_sort_by_key(thrust::device, key_ptr, key_ptr + num_items, value_ptr, thrust::greater()); + } else { + thrust::stable_sort_by_key(thrust::device, key_ptr, key_ptr + num_items, value_ptr, thrust::less()); + } + + MACA_CHECK(mcMemcpy(d_keys_out, temp_keys, k * sizeof(KeyType), mcMemcpyDeviceToDevice)); + MACA_CHECK(mcMemcpy(d_values_out, temp_values, k * sizeof(ValueType), mcMemcpyDeviceToDevice)); + + mcFree(temp_keys); + mcFree(temp_values); +#else + // ======================================== + // 默认基准实现 + // ======================================== + + KeyType* temp_keys; + ValueType* temp_values; + MACA_CHECK(mcMalloc(&temp_keys, num_items * sizeof(KeyType))); + MACA_CHECK(mcMalloc(&temp_values, num_items * sizeof(ValueType))); + + MACA_CHECK(mcMemcpy(temp_keys, d_keys_in, num_items * sizeof(KeyType), mcMemcpyDeviceToDevice)); + MACA_CHECK(mcMemcpy(temp_values, d_values_in, num_items * sizeof(ValueType), mcMemcpyDeviceToDevice)); + + auto key_ptr = thrust::device_pointer_cast(temp_keys); + auto value_ptr = thrust::device_pointer_cast(temp_values); + // 由于greater和less是不同类型，需要分别调用 if (descending) { thrust::stable_sort_by_key(thrust::device, key_ptr, key_ptr + num_items, value_ptr, thrust::greater()); @@ -80,7 +98,7 @@ public: mcFree(temp_values); #endif } - + // 获取当前实现状态 static const char* getImplementationStatus() { #if USE_DEFAULT_REF_IMPL @@ -89,7 +107,7 @@ public: return "CUSTOM_IMPL"; #endif } - + private: // 参赛者可以在这里添加辅助函数和成员变量 // 例如：分块大小、临时缓冲区、多流处理等 @@ -103,54 +121,54 @@ bool testCorrectness() { std::cout << "TopkPair 正确性测试..." << std::endl; TestDataGenerator generator; TopkPairAlgorithm algorithm; - + int size = 10000; auto keys = generator.generateRandomFloats(size); auto values = generator.generateRandomUint32(size); - + // 分配GPU内存 float *d_keys_in, *d_keys_out; uint32_t *d_values_in, *d_values_out; - + MACA_CHECK(mcMalloc(&d_keys_in, size * sizeof(float))); MACA_CHECK(mcMalloc(&d_values_in, size * sizeof(uint32_t))); - + MACA_CHECK(mcMemcpy(d_keys_in, keys.data(), size * sizeof(float), mcMemcpyHostToDevice)); MACA_CHECK(mcMemcpy(d_values_in, values.data(), size * sizeof(uint32_t), mcMemcpyHostToDevice)); - + bool allPassed = true; - + // 测试不同k值 for (int ki = 0; ki < NUM_TOPK_VALUES && ki < 4; ki++) { // 限制测试范围 int k = TOPK_VALUES[ki]; if (k > size) continue; - + std::cout << " 测试 k=" << k << std::endl; - + MACA_CHECK(mcMalloc(&d_keys_out, k * sizeof(float))); MACA_CHECK(mcMalloc(&d_values_out, k * sizeof(uint32_t))); - + for (bool descending : {false, true}) { std::cout << " " << (descending ? "降序" : "升序") << " TopK..." << std::endl; - + // CPU参考结果 std::vector cpu_keys_out; std::vector cpu_values_out; cpuTopkPair(keys, values, cpu_keys_out, cpu_values_out, k, descending); - + // GPU算法结果 algorithm.topk(d_keys_in, d_keys_out, d_values_in, d_values_out, size, k, descending); - + // 获取结果 std::vector gpu_keys_out(k); std::vector gpu_values_out(k); MACA_CHECK(mcMemcpy(gpu_keys_out.data(), d_keys_out, k * sizeof(float), mcMemcpyDeviceToHost)); MACA_CHECK(mcMemcpy(gpu_values_out.data(), d_values_out, k * sizeof(uint32_t), mcMemcpyDeviceToHost)); - + // 验证结果 bool keysMatch = compareArrays(cpu_keys_out, gpu_keys_out, 1e-5); bool valuesMatch = compareArrays(cpu_values_out, gpu_values_out); - + if (!keysMatch || !valuesMatch) { std::cout << " 失败: 结果不匹配" << std::endl; allPassed = false; @@ -158,15 +176,15 @@ bool testCorrectness() { std::cout << " 通过" << std::endl; } } - + mcFree(d_keys_out); mcFree(d_values_out); } - + // 清理内存 mcFree(d_keys_in); mcFree(d_values_in); - + return allPassed; } @@ -175,55 +193,55 @@ void benchmarkPerformance() { std::cout << "数据类型: " << std::endl; std::cout << "计算公式:" << std::endl; std::cout << " 吞吐量 = 元素数 / 时间(s) / 1e9 (G/s)" << std::endl; - + TestDataGenerator generator; PerformanceMeter meter; TopkPairAlgorithm algorithm; - + const int WARMUP_ITERATIONS = 5; const int BENCHMARK_ITERATIONS = 10; - + // 用于YAML报告的数据收集 std::vector> perf_data; - + // 针对不同数据规模测试 for (int size_idx = 0; size_idx < NUM_TEST_SIZES; size_idx++) { int size = TEST_SIZES[size_idx]; std::cout << "\n数据规模: " << size << std::endl; - std::cout << std::setw(8) << "k值" << std::setw(15) << "升序(ms)" << std::setw(15) << "降序(ms)" + std::cout << std::setw(8) << "k值" << std::setw(15) << "升序(ms)" << std::setw(15) << "降序(ms)" << std::setw(16) << "升序(G/s)" << std::setw(16) << "降序(G/s)" << std::endl; std::cout << std::string(74, '-') << std::endl; - + auto keys = generator.generateRandomFloats(size); auto values = generator.generateRandomUint32(size); - + // 分配GPU内存 float *d_keys_in; uint32_t *d_values_in; - + MACA_CHECK(mcMalloc(&d_keys_in, size * sizeof(float))); MACA_CHECK(mcMalloc(&d_values_in, size * sizeof(uint32_t))); - + MACA_CHECK(mcMemcpy(d_keys_in, keys.data(), size * sizeof(float), mcMemcpyHostToDevice)); MACA_CHECK(mcMemcpy(d_values_in, values.data(), size * sizeof(uint32_t), mcMemcpyHostToDevice)); - + for (int ki = 0; ki < NUM_TOPK_VALUES; ki++) { int k = TOPK_VALUES[ki]; if (k > size) continue; - + float *d_keys_out; uint32_t *d_values_out; MACA_CHECK(mcMalloc(&d_keys_out, k * sizeof(float))); MACA_CHECK(mcMalloc(&d_values_out, k * sizeof(uint32_t))); - + float asc_time = 0, desc_time = 0; - + for (bool descending : {false, true}) { // Warmup阶段 for (int iter = 0; iter < WARMUP_ITERATIONS; iter++) { algorithm.topk(d_keys_in, d_keys_out, d_values_in, d_values_out, size, k, descending); } - + // 正式测试阶段 float total_time = 0; for (int iter = 0; iter < BENCHMARK_ITERATIONS; iter++) { @@ -231,7 +249,7 @@ void benchmarkPerformance() { algorithm.topk(d_keys_in, d_keys_out, d_values_in, d_values_out, size, k, descending); total_time += meter.stopTiming(); } - + float avg_time = total_time / BENCHMARK_ITERATIONS; if (descending) { desc_time = avg_time; @@ -239,14 +257,14 @@ void benchmarkPerformance() { asc_time = avg_time; } } - + // 计算性能指标 auto asc_metrics = PerformanceCalculator::calculateTopkPair(size, k, asc_time); auto desc_metrics = PerformanceCalculator::calculateTopkPair(size, k, desc_time); - + // 显示性能数据 PerformanceDisplay::printTopkPairData(k, asc_time, desc_time, asc_metrics, desc_metrics); - + // 收集YAML报告数据 auto entry = YAMLPerformanceReporter::createEntry(); entry["data_size"] = std::to_string(size); @@ -262,11 +280,11 @@ void benchmarkPerformance() { mcFree(d_keys_out); mcFree(d_values_out); } - + mcFree(d_keys_in); mcFree(d_values_in); } - + // 生成YAML性能报告 YAMLPerformanceReporter::generateTopkPairYAML(perf_data, "topk_pair_performance.yaml"); PerformanceDisplay::printSavedMessage("topk_pair_performance.yaml"); @@ -277,21 +295,21 @@ void benchmarkPerformance() { // ============================================================================ int main(int argc, char* argv[]) { std::cout << "=== TopkPair 算法测试 ===" << 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(); @@ -300,16 +318,16 @@ int main(int argc, char* argv[]) { performance_completed = false; } } - + std::cout << "\n=== 测试完成 ===" << std::endl; std::cout << "实现状态: " << TopkPairAlgorithm::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; diff --git a/S1/3/utils/performance_utils.h b/S1/41/utils/performance_utils.h similarity index 100% rename from S1/3/utils/performance_utils.h rename to S1/41/utils/performance_utils.h diff --git a/S1/3/utils/test_utils.h b/S1/41/utils/test_utils.h similarity index 100% rename from S1/3/utils/test_utils.h rename to S1/41/utils/test_utils.h diff --git a/S1/3/utils/yaml_reporter.h b/S1/41/utils/yaml_reporter.h similarity index 100% rename from S1/3/utils/yaml_reporter.h rename to S1/41/utils/yaml_reporter.h