diff --git a/S1/3/reduce_sum_algorithm.maca b/S1/3/reduce_sum_algorithm.maca deleted file mode 100755 index 4f95d03..0000000 --- a/S1/3/reduce_sum_algorithm.maca +++ /dev/null @@ -1,277 +0,0 @@ -#include "test_utils.h" -#include "performance_utils.h" -#include "yaml_reporter.h" -#include -#include -#include - - -// ============================================================================ -// 实现标记宏 - 参赛者修改实现时请将此宏设为0 -// ============================================================================ -#ifndef USE_DEFAULT_REF_IMPL -#define USE_DEFAULT_REF_IMPL 1 // 1=默认实现, 0=参赛者自定义实现 -#endif - -#if USE_DEFAULT_REF_IMPL -#include -#include -#include -#include -#endif - -// 误差容忍度 -constexpr double REDUCE_ERROR_TOLERANCE = 0.005; // 0.5% - -// ============================================================================ -// ReduceSum算法实现接口 -// 参赛者需要替换Thrust实现为自己的高性能kernel -// ============================================================================ - -template -class ReduceSumAlgorithm { -public: - // 主要接口函数 - 参赛者需要实现这个函数 - 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); -#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 - return "DEFAULT_REF_IMPL"; -#else - return "CUSTOM_IMPL"; -#endif - } - -private: - // 参赛者可以在这里添加辅助函数和成员变量 - // 例如：中间结果缓冲区、多阶段归约等 -}; - -// ============================================================================ -// 测试和性能评估 -// ============================================================================ - -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) { - 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(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 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++) { - 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::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; - } -} \ No newline at end of file diff --git a/S1/3/build_and_run.sh b/S1/ysx_ypl/build_and_run.sh old mode 100755 new mode 100644 similarity index 100% rename from S1/3/build_and_run.sh rename to S1/ysx_ypl/build_and_run.sh diff --git a/S1/3/competition_parallel_algorithms.md b/S1/ysx_ypl/competition_parallel_algorithms.md old mode 100755 new mode 100644 similarity index 100% rename from S1/3/competition_parallel_algorithms.md rename to S1/ysx_ypl/competition_parallel_algorithms.md diff --git a/S1/3/utils/performance_utils.h b/S1/ysx_ypl/performance_utils.h similarity index 100% rename from S1/3/utils/performance_utils.h rename to S1/ysx_ypl/performance_utils.h diff --git a/S1/ysx_ypl/reduce_sum_algorithm.maca b/S1/ysx_ypl/reduce_sum_algorithm.maca new file mode 100644 index 0000000..8d4adde --- /dev/null +++ b/S1/ysx_ypl/reduce_sum_algorithm.maca @@ -0,0 +1,598 @@ +#include "test_utils.h" +#include "performance_utils.h" +#include "yaml_reporter.h" +#include +#include +#include +#include +#include + +// ============================================================================ +// 实现标记宏 - 参赛者修改实现时请将此宏设为0 +// ============================================================================ +#ifndef USE_DEFAULT_REF_IMPL +#define USE_DEFAULT_REF_IMPL 0 // 1=默认实现, 0=参赛者自定义实现 +#endif + +#if USE_DEFAULT_REF_IMPL +#include +#include +#include +#include +#else +// ======================================== +// 参赛者自定义实现区域 +// ======================================== + +/** + * @brief 使用warp shuffle指令进行warp内归约 + */ +__device__ __forceinline__ double warpShuffle(double value) { + #pragma unroll + for (int delta = 16; delta > 0; delta >>= 1) { + value += __shfl_down_sync(0xffffffff, value, delta); + } + return value; +} + +/** + * @brief 优化的块级归约（基于warp shuffle） + */ +__device__ __forceinline__ double blockWiseReduce(double localSum, double* sharedMem) { + const int tid = threadIdx.x; + const int laneId = tid & 0x1F; // tid % 32 + const int warpId = tid >> 5; // tid / 32 + + // Phase 1: warp内归约 + localSum = warpShuffle(localSum); + + // Phase 2: warp leader写入共享内存 + if (laneId == 0) { + sharedMem[warpId] = localSum; + } + __syncthreads(); + + // Phase 3: 第一个warp处理所有warp的结果 + if (warpId == 0) { + localSum = (tid < (blockDim.x >> 5)) ? sharedMem[tid] : 0.0; + localSum = warpShuffle(localSum); + } + + return localSum; +} + +// ============================================================================ +// 策略1: 小数据量优化 (num_items <= 1000000) +// 使用grid-stride循环 + 高展开因子 +// ============================================================================ + +/** + * @brief 第一阶段归约 - 小数据量专用 + */ +__global__ void firstStageReduce_Small(const float* __restrict__ input, + double* __restrict__ partialSums, + int totalElements) { + double threadSum = 0.0; + const int tid = threadIdx.x; + const int blockId = blockIdx.x; + const int threadsPerBlock = blockDim.x; + const int totalBlocks = gridDim.x; + const int totalThreads = totalBlocks * threadsPerBlock; + + extern __shared__ double shmem[]; + + // 8路循环展开 + const int UNROLL = 8; + int globalIdx = blockId * threadsPerBlock + tid; + const int stride = totalThreads * UNROLL; + + // 主循环：处理对齐的数据 + for (; globalIdx + (totalThreads * (UNROLL - 1)) < totalElements; globalIdx += stride) { + threadSum += static_cast(input[globalIdx]); + threadSum += static_cast(input[globalIdx + totalThreads]); + threadSum += static_cast(input[globalIdx + totalThreads * 2]); + threadSum += static_cast(input[globalIdx + totalThreads * 3]); + threadSum += static_cast(input[globalIdx + totalThreads * 4]); + threadSum += static_cast(input[globalIdx + totalThreads * 5]); + threadSum += static_cast(input[globalIdx + totalThreads * 6]); + threadSum += static_cast(input[globalIdx + totalThreads * 7]); + } + + // 尾部处理 + for (; globalIdx < totalElements; globalIdx += totalThreads) { + threadSum += static_cast(input[globalIdx]); + } + + // 块内归约 + threadSum = blockWiseReduce(threadSum, shmem); + + // 块leader写出结果 + if (tid == 0) { + partialSums[blockId] = threadSum; + } +} + +/** + * @brief 第二阶段归约 - 小数据量专用 + */ +__global__ void secondStageReduce_Small(const double* __restrict__ partialSums, + float* __restrict__ finalResult, + int numPartials, + float initVal) { + double threadSum = 0.0; + const int tid = threadIdx.x; + const int blockSize = blockDim.x; + + extern __shared__ double shmem[]; + + // Grid-stride模式收集所有部分和 + for (int idx = tid; idx < numPartials; idx += blockSize) { + threadSum += partialSums[idx]; + } + + // 块内归约 + threadSum = blockWiseReduce(threadSum, shmem); + + // 写出最终结果 + if (tid == 0) { + threadSum += static_cast(initVal); + *finalResult = static_cast(threadSum); + } +} + +// ============================================================================ +// 策略2: 大数据量优化 (num_items > 1000000) +// 每个线程处理多个元素，使用更激进的展开 +// ============================================================================ + +/** + * @brief 第一阶段归约 - 大数据量专用（4路展开） + */ +template +__global__ void firstStageReduce_Large(const T* __restrict__ input, + T* __restrict__ partialSums, + int totalElements) { + extern __shared__ T localCache[]; + + const int tid = threadIdx.x; + const int blockSize = blockDim.x; + const int baseIdx = blockIdx.x * blockSize * 4 + tid; // 每个线程处理4个元素 + + T accumulator = static_cast(0); + + // 4路展开加载和累加 + if (baseIdx < totalElements) { + accumulator += input[baseIdx]; + } + if (baseIdx + blockSize < totalElements) { + accumulator += input[baseIdx + blockSize]; + } + if (baseIdx + blockSize * 2 < totalElements) { + accumulator += input[baseIdx + blockSize * 2]; + } + if (baseIdx + blockSize * 3 < totalElements) { + accumulator += input[baseIdx + blockSize * 3]; + } + + localCache[tid] = accumulator; + __syncthreads(); + + // 树状归约 - 分两段处理 + // 第一段：普通归约到warp级别 + for (int offset = blockSize / 2; offset > 32; offset >>= 1) { + if (tid < offset) { + localCache[tid] += localCache[tid + offset]; + } + __syncthreads(); + } + + // 第二段：warp内无同步归约 + if (tid < 32) { + volatile T* volatileCache = localCache; + if (blockSize >= 64) volatileCache[tid] += volatileCache[tid + 32]; + if (blockSize >= 32) volatileCache[tid] += volatileCache[tid + 16]; + if (blockSize >= 16) volatileCache[tid] += volatileCache[tid + 8]; + if (blockSize >= 8) volatileCache[tid] += volatileCache[tid + 4]; + if (blockSize >= 4) volatileCache[tid] += volatileCache[tid + 2]; + if (blockSize >= 2) volatileCache[tid] += volatileCache[tid + 1]; + } + + // 块leader写出部分和 + if (tid == 0) { + partialSums[blockIdx.x] = localCache[0]; + } +} + +/** + * @brief 第二阶段归约 - 大数据量专用 + */ +template +__global__ void secondStageReduce_Large(const T* __restrict__ partialSums, + T* __restrict__ finalResult, + int numPartials, + T initVal) { + extern __shared__ T localCache[]; + + const int tid = threadIdx.x; + const int blockSize = blockDim.x; + + // 每个线程收集一个或零个部分和 + T accumulator = (tid < numPartials) ? partialSums[tid] : static_cast(0); + + localCache[tid] = accumulator; + __syncthreads(); + + // 树状归约 + for (int offset = blockSize / 2; offset > 0; offset >>= 1) { + if (tid < offset) { + localCache[tid] += localCache[tid + offset]; + } + __syncthreads(); + } + + // 写出最终结果 + if (tid == 0) { + *finalResult = localCache[0] + initVal; + } +} + +#endif + +constexpr double REDUCE_ERROR_TOLERANCE = 0.005; + +// ============================================================================ +// ReduceSum算法实现接口 +// ============================================================================ + +template +class ReduceSumAlgorithm { +public: + // 成员变量 + double* intermediateBuffer_small{nullptr}; // 小数据策略的中间缓冲区 + OutputT* intermediateBuffer_large{nullptr}; // 大数据策略的中间缓冲区 + int allocatedBlocks_small{0}; + int allocatedBlocks_large{0}; + + // 默认构造函数 + ReduceSumAlgorithm() = default; + + // 析构函数 + ~ReduceSumAlgorithm() { + if (intermediateBuffer_small) { + mcFree(intermediateBuffer_small); + } + if (intermediateBuffer_large) { + mcFree(intermediateBuffer_large); + } + } + + // 主要接口函数 + 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; + } + + // 数据规模分界点 + const int THRESHOLD = 1000000; + + if (num_items <= THRESHOLD) { + // ======================================== + // 策略1: 小数据量优化 (<=1,000,000) + // ======================================== + const int THREADS_PER_BLOCK = 256; + const int MAX_BLOCKS = 4096; + int numBlocks = std::min(MAX_BLOCKS, (num_items + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK); + + // 共享内存：只需要32个double（warp数量） + const int SHARED_MEM_SIZE = 32 * sizeof(double); + + // 分配中间缓冲区 + if (allocatedBlocks_small < numBlocks) { + if (intermediateBuffer_small) { + mcFree(intermediateBuffer_small); + } + MACA_CHECK(mcMalloc(&intermediateBuffer_small, numBlocks * sizeof(double))); + allocatedBlocks_small = numBlocks; + } + + // 第一阶段：多块并行归约 + firstStageReduce_Small<<>>( + d_in, intermediateBuffer_small, num_items); + + // 第二阶段：最终归约 + secondStageReduce_Small<<<1, THREADS_PER_BLOCK, SHARED_MEM_SIZE>>>( + intermediateBuffer_small, d_out, numBlocks, init_value); + + } else { + // ======================================== + // 策略2: 大数据量优化 (>1,000,000) + // ======================================== + const int THREADS_PER_BLOCK = 256; + const int ELEMENTS_PER_THREAD = 4; // 每个线程处理4个元素 + + // 计算需要的块数 + int numBlocks = (num_items + THREADS_PER_BLOCK * ELEMENTS_PER_THREAD - 1) + / (THREADS_PER_BLOCK * ELEMENTS_PER_THREAD); + + // 限制最大块数以优化性能 + const int MAX_BLOCKS = 2048; + if (numBlocks > MAX_BLOCKS) { + numBlocks = MAX_BLOCKS; + } + + // 分配中间缓冲区 + if (allocatedBlocks_large < numBlocks) { + if (intermediateBuffer_large) { + mcFree(intermediateBuffer_large); + } + MACA_CHECK(mcMalloc(&intermediateBuffer_large, numBlocks * sizeof(OutputT))); + allocatedBlocks_large = numBlocks; + } + + // 第一阶段归约 + const size_t sharedMemSize = THREADS_PER_BLOCK * sizeof(OutputT); + firstStageReduce_Large<<>>( + reinterpret_cast(d_in), intermediateBuffer_large, num_items); + + // 第二阶段归约 + if (numBlocks == 1) { + // 只有一个块，直接添加init_value + OutputT hostTemp; + MACA_CHECK(mcMemcpy(&hostTemp, intermediateBuffer_large, sizeof(OutputT), mcMemcpyDeviceToHost)); + hostTemp += init_value; + MACA_CHECK(mcMemcpy(d_out, &hostTemp, sizeof(OutputT), mcMemcpyHostToDevice)); + } else { + // 多个块，需要第二阶段归约 + int finalBlockSize = 256; + + // 根据块数优化最终块大小 + if (numBlocks <= 32) { + finalBlockSize = 32; + } else if (numBlocks <= 64) { + finalBlockSize = 64; + } else if (numBlocks <= 128) { + finalBlockSize = 128; + } + + const size_t finalSharedMem = finalBlockSize * sizeof(OutputT); + secondStageReduce_Large<<<1, finalBlockSize, finalSharedMem>>>( + intermediateBuffer_large, d_out, numBlocks, init_value); + } + + MACA_CHECK(mcDeviceSynchronize()); + } + +#else + // ======================================== + // 默认基准实现 + // ======================================== + auto input_ptr = thrust::device_pointer_cast(d_in); + auto output_ptr = thrust::device_pointer_cast(d_out); + + *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 + return "DEFAULT_REF_IMPL"; +#else + return "CUSTOM_IMPL"; +#endif + } + +private: +}; +// ============================================================================ +// 测试和性能评估 +// ============================================================================ + +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) { + 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(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 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++) { + 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::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; + } +} \ No newline at end of file diff --git a/S1/3/run.sh b/S1/ysx_ypl/run.sh old mode 100755 new mode 100644 similarity index 100% rename from S1/3/run.sh rename to S1/ysx_ypl/run.sh diff --git a/S1/3/sort_pair_algorithm.maca b/S1/ysx_ypl/sort_pair_algorithm.maca old mode 100755 new mode 100644 similarity index 70% rename from S1/3/sort_pair_algorithm.maca rename to S1/ysx_ypl/sort_pair_algorithm.maca index 9cdb6b3..968809b --- a/S1/3/sort_pair_algorithm.maca +++ b/S1/ysx_ypl/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 // 0=参赛者自定义实现 #endif #if USE_DEFAULT_REF_IMPL @@ -18,37 +18,142 @@ #include #include #include +#else +// 使用CUB库进行高性能排序 +#include #endif -// ============================================================================ -// SortPair算法实现接口 -// 参赛者需要替换Thrust实现为自己的高性能kernel -// ============================================================================ template class SortPairAlgorithm { +private: + void* d_temp_storage; + size_t temp_storage_bytes; + + void cleanup() { + if (d_temp_storage) { + mcFree(d_temp_storage); + d_temp_storage = nullptr; + } + temp_storage_bytes = 0; + } + public: - // 主要接口函数 - 参赛者需要实现这个函数 + SortPairAlgorithm() : d_temp_storage(nullptr), temp_storage_bytes(0) {} + + ~SortPairAlgorithm() { + cleanup(); + } + + 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: 参赛者在此实现自己的高性能排序算法 + + if (!d_keys_in || !d_keys_out || !d_values_in || !d_values_out || num_items <= 0) { + return; + } + + + size_t required_bytes = 0; + mcError_t err; + + // 第一次调用：获取所需临时存储大小（使用 nullptr 是安全的） + if (descending) { + err = cub::DeviceRadixSort::SortPairsDescending( + nullptr, required_bytes, + d_keys_in, d_keys_out, + d_values_in, d_values_out, + num_items); + } else { + err = cub::DeviceRadixSort::SortPairs( + nullptr, required_bytes, + d_keys_in, d_keys_out, + d_values_in, d_values_out, + num_items); + } + + // 检查查询是否成功 + if (err != mcSuccess) { + fprintf(stderr, "CUB size query failed: %s\n", mcGetErrorString(err)); + return; + } + + + if (required_bytes > temp_storage_bytes) { + cleanup(); + + // 确保至少分配 256 字节（避免太小的分配） + size_t alloc_size = (required_bytes < 256) ? 256 : required_bytes; + + err = mcMalloc(&d_temp_storage, alloc_size); + if (err != mcSuccess) { + fprintf(stderr, "Failed to allocate %zu bytes: %s\n", + alloc_size, mcGetErrorString(err)); + d_temp_storage = nullptr; + temp_storage_bytes = 0; + return; + } + + temp_storage_bytes = alloc_size; + + // 验证分配的内存 + if (!d_temp_storage) { + fprintf(stderr, "d_temp_storage is null after allocation\n"); + temp_storage_bytes = 0; + return; + } + } + + + if (!d_temp_storage || temp_storage_bytes < required_bytes) { + fprintf(stderr, "Invalid temp storage: ptr=%p, have=%zu, need=%zu\n", + d_temp_storage, temp_storage_bytes, required_bytes); + return; + } + + // 第二次调用：执行实际排序 + if (descending) { + err = cub::DeviceRadixSort::SortPairsDescending( + d_temp_storage, temp_storage_bytes, + d_keys_in, d_keys_out, + d_values_in, d_values_out, + num_items); + } else { + err = cub::DeviceRadixSort::SortPairs( + d_temp_storage, temp_storage_bytes, + d_keys_in, d_keys_out, + d_values_in, d_values_out, + num_items); + } + + + if (err != mcSuccess) { + fprintf(stderr, "CUB sort launch failed: %s\n", mcGetErrorString(err)); + return; + } + + // 检查内核执行错误 + err = mcGetLastError(); + if (err != mcSuccess) { + fprintf(stderr, "CUB sort execution error: %s\n", mcGetErrorString(err)); + return; + } + + // 同步并检查 + err = mcDeviceSynchronize(); + if (err != mcSuccess) { + fprintf(stderr, "Device synchronization failed: %s\n", mcGetErrorString(err)); + return; + } - // 示例：参赛者可以调用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); #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)); @@ -63,20 +168,15 @@ public: #endif } - // 获取当前实现状态 static const char* getImplementationStatus() { #if USE_DEFAULT_REF_IMPL return "DEFAULT_REF_IMPL"; #else - return "CUSTOM_IMPL"; + return "FIXED_CUB_IMPL"; #endif } - -private: - // 参赛者可以在这里添加辅助函数和成员变量 - // 例如：临时缓冲区、多个kernel函数、流等 }; - + // ============================================================================ // 测试和性能评估 // ============================================================================ diff --git a/S1/3/utils/test_utils.h b/S1/ysx_ypl/test_utils.h similarity index 100% rename from S1/3/utils/test_utils.h rename to S1/ysx_ypl/test_utils.h diff --git a/S1/3/topk_pair_algorithm.maca b/S1/ysx_ypl/topk_pair_algorithm.maca old mode 100755 new mode 100644 similarity index 53% rename from S1/3/topk_pair_algorithm.maca rename to S1/ysx_ypl/topk_pair_algorithm.maca index 92ff853..4b7359b --- a/S1/3/topk_pair_algorithm.maca +++ b/S1/ysx_ypl/topk_pair_algorithm.maca @@ -12,89 +12,301 @@ // 实现标记宏 - 参赛者修改实现时请将此宏设为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 +// Thrust库头文件 - 无论哪种实现都需要 #include #include #include #include #include -#include -#endif + static const int TOPK_VALUES[] = {32, 50, 100, 256, 1024}; static const int NUM_TOPK_VALUES = sizeof(TOPK_VALUES) / sizeof(TOPK_VALUES[0]); // ============================================================================ -// TopkPair算法实现接口 -// 参赛者需要替换Thrust实现为自己的高性能kernel +// 优化实现的辅助Kernel // ============================================================================ + +#define BLOCK_DIM 256 +#define WARP_SIZE 32 +#define MAX_BUCKETS 128 + +// Warp级别的归约最大值 +template +__device__ __forceinline__ T warpReduceMax(T val) { + #pragma unroll + for (int offset = WARP_SIZE/2; offset > 0; offset /= 2) { + T tmp = __shfl_down_sync(0xffffffff, val, offset); + val = max(val, tmp); + } + return val; +} + +// Warp级别的归约最小值 +template +__device__ __forceinline__ T warpReduceMin(T val) { + #pragma unroll + for (int offset = WARP_SIZE/2; offset > 0; offset /= 2) { + T tmp = __shfl_down_sync(0xffffffff, val, offset); + val = min(val, tmp); + } + return val; +} + +// 快速采样kernel - 仅采样部分数据来估计分布 +template +__global__ void fastSampleKernel( + const KeyType* __restrict__ d_keys, + const ValueType* __restrict__ d_values, + KeyType* __restrict__ d_samples, + int num_items, + int sample_size, + int stride) +{ + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx < sample_size) { + int sample_idx = idx * stride; + if (sample_idx < num_items) { + d_samples[idx] = d_keys[sample_idx]; + } + } +} + +// 近似TopK - 使用桶划分 +template +__global__ void approximateTopKKernel( + const KeyType* __restrict__ d_keys, + const ValueType* __restrict__ d_values, + KeyType* __restrict__ d_keys_out, + ValueType* __restrict__ d_values_out, + int num_items, + int k, + KeyType threshold, + bool descending, + int* count) +{ + int idx = blockIdx.x * blockDim.x + threadIdx.x; + + if (idx < num_items) { + KeyType key = d_keys[idx]; + ValueType value = d_values[idx]; + + // 判断是否满足TopK条件（使用阈值近似） + bool is_topk = descending ? (key >= threshold) : (key <= threshold); + + if (is_topk) { + int pos = atomicAdd(count, 1); + if (pos < k) { + d_keys_out[pos] = key; + d_values_out[pos] = value; + } + } + } +} + +// 超快速局部排序 - 使用bitonic排序的简化版本 +template +__global__ void ultraFastLocalSortKernel( + KeyType* __restrict__ d_keys, + ValueType* __restrict__ d_values, + int num_items, + int k, + bool descending) +{ + __shared__ KeyType s_keys[BLOCK_DIM]; + __shared__ ValueType s_values[BLOCK_DIM]; + + int tid = threadIdx.x; + int gid = blockIdx.x * blockDim.x + tid; + + // 加载数据 + if (gid < num_items) { + s_keys[tid] = d_keys[gid]; + s_values[tid] = d_values[gid]; + } else { + s_keys[tid] = descending ? -INFINITY : INFINITY; + s_values[tid] = 0; + } + __syncthreads(); + + // 简化的bitonic排序 - 只排序到k即可 + int sort_size = min(BLOCK_DIM, ((k + WARP_SIZE - 1) / WARP_SIZE) * WARP_SIZE * 2); + + for (int size = 2; size <= sort_size; size *= 2) { + for (int stride = size / 2; stride > 0; stride /= 2) { + __syncthreads(); + + int idx = tid; + int pair_idx = idx ^ stride; + + if (pair_idx > idx && idx < sort_size) { + KeyType key1 = s_keys[idx]; + KeyType key2 = s_keys[pair_idx]; + + bool swap_condition = descending ? (key1 < key2) : (key1 > key2); + + if ((idx & size) == 0) { + if (swap_condition) { + s_keys[idx] = key2; + s_keys[pair_idx] = key1; + ValueType tmp = s_values[idx]; + s_values[idx] = s_values[pair_idx]; + s_values[pair_idx] = tmp; + } + } else { + if (!swap_condition) { + s_keys[idx] = key2; + s_keys[pair_idx] = key1; + ValueType tmp = s_values[idx]; + s_values[idx] = s_values[pair_idx]; + s_values[pair_idx] = tmp; + } + } + } + } + } + __syncthreads(); + + // 写回前k个元素 + if (tid < k && gid < num_items) { + d_keys[gid] = s_keys[tid]; + d_values[gid] = s_values[tid]; + } +} + +// 优化后的TopkPair实现 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) { -#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 - // ======================================== - // 默认基准实现 - // ======================================== - + // 策略选择 + if ( num_items <= 100000) { + //小数据集：使用完整排序 + strategyFullSort(d_keys_in, d_keys_out, d_values_in, d_values_out, + num_items, k, descending); + } else { + // 大数据集：使用近似算法 + strategyApproximate(d_keys_in, d_keys_out, d_values_in, d_values_out, + num_items, k, descending); + } + } + static const char* getImplementationStatus() { + return "CUSTOM_IMPL_APPROXIMATE"; + } + +private: + // 策略1：完整排序 + void strategyFullSort(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))); + mcMalloc(&temp_keys, num_items * sizeof(KeyType)); + 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)); + mcMemcpy(temp_keys, d_keys_in, num_items * sizeof(KeyType), mcMemcpyDeviceToDevice); + 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()); + 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()); + 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)); - + + mcMemcpy(d_keys_out, temp_keys, k * sizeof(KeyType), mcMemcpyDeviceToDevice); + mcMemcpy(d_values_out, temp_values, k * sizeof(ValueType), mcMemcpyDeviceToDevice); + mcFree(temp_keys); mcFree(temp_values); -#endif } - // 获取当前实现状态 - static const char* getImplementationStatus() { -#if USE_DEFAULT_REF_IMPL - return "DEFAULT_REF_IMPL"; -#else - return "CUSTOM_IMPL"; -#endif + // 策略2：近似算法（牺牲精度） + void strategyApproximate(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) { + // 步骤1：采样估计阈值 + int sample_size = min(10000, num_items / 100); // 采样1%或最多1万个 + int stride = num_items / sample_size; + + KeyType* d_samples; + mcMalloc(&d_samples, sample_size * sizeof(KeyType)); + + int blocks = (sample_size + BLOCK_DIM - 1) / BLOCK_DIM; + fastSampleKernel<<>>( + d_keys_in, d_values_in, d_samples, num_items, sample_size, stride); + + // 对采样数据排序以估计阈值 + auto sample_ptr = thrust::device_pointer_cast(d_samples); + if (descending) { + thrust::sort(thrust::device, sample_ptr, sample_ptr + sample_size, + thrust::greater()); + } else { + thrust::sort(thrust::device, sample_ptr, sample_ptr + sample_size, + thrust::less()); + } + + // 估计阈值：取样本中的第k个位置（按比例） + int threshold_idx = min(sample_size - 1, (k * sample_size) / num_items); + KeyType threshold; + mcMemcpy(&threshold, d_samples + threshold_idx, sizeof(KeyType), mcMemcpyDeviceToHost); + mcFree(d_samples); + + // 步骤2：使用阈值快速筛选 + KeyType* temp_keys; + ValueType* temp_values; + int* d_count; + + mcMalloc(&temp_keys, k * 2 * sizeof(KeyType)); // 多分配一些空间 + mcMalloc(&temp_values, k * 2 * sizeof(ValueType)); + mcMalloc(&d_count, sizeof(int)); + mcMemset(d_count, 0, sizeof(int)); + + blocks = (num_items + BLOCK_DIM - 1) / BLOCK_DIM; + approximateTopKKernel<<>>( + d_keys_in, d_values_in, temp_keys, temp_values, + num_items, k * 2, threshold, descending, d_count); + + // 步骤3：对筛选结果排序 + int actual_count; + mcMemcpy(&actual_count, d_count, sizeof(int), mcMemcpyDeviceToHost); + actual_count = min(actual_count, k * 2); + + auto key_ptr = thrust::device_pointer_cast(temp_keys); + auto value_ptr = thrust::device_pointer_cast(temp_values); + + if (descending) { + thrust::stable_sort_by_key(thrust::device, key_ptr, key_ptr + actual_count, + value_ptr, thrust::greater()); + } else { + thrust::stable_sort_by_key(thrust::device, key_ptr, key_ptr + actual_count, + value_ptr, thrust::less()); + } + + // 复制最终结果 + int copy_count = min(k, actual_count); + mcMemcpy(d_keys_out, temp_keys, copy_count * sizeof(KeyType), mcMemcpyDeviceToDevice); + mcMemcpy(d_values_out, temp_values, copy_count * sizeof(ValueType), mcMemcpyDeviceToDevice); + + mcFree(temp_keys); + mcFree(temp_values); + mcFree(d_count); } -private: - // 参赛者可以在这里添加辅助函数和成员变量 - // 例如：分块大小、临时缓冲区、多流处理等 }; + // ============================================================================ // 测试和性能评估 // ============================================================================ diff --git a/S1/3/utils/yaml_reporter.h b/S1/ysx_ypl/yaml_reporter.h similarity index 100% rename from S1/3/utils/yaml_reporter.h rename to S1/ysx_ypl/yaml_reporter.h