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;