reduce sum，sort_pair和topk_pair算子优化 #12
|
|
@ -1,97 +0,0 @@
|
|||
# 样例赛题说明
|
||||
|
||||
## GPU高性能并行计算算法优化
|
||||
|
||||
要求参赛者通过一个或多个global kernel 函数（允许配套 device 辅助函数），实现高性能算法。
|
||||
|
||||
在正确性、稳定性前提下，比拼算法性能。
|
||||
|
||||
# 1. ReduceSum算法优化
|
||||
```cpp
|
||||
template <typename InputT = float, typename OutputT = float>
|
||||
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 <typename KeyType, typename ValueType>
|
||||
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 <typename KeyType, typename ValueType>
|
||||
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
|
||||
|
||||
|
||||
注意事项
|
||||
|
||||
* 结果必须稳定排序
|
||||
|
||||
加分项
|
||||
|
||||
* 支持其他不同数据类型的键值对，实现类型通用算法
|
||||
* 覆盖更全面的数据范围，提供良好稳定的性能表现
|
||||
|
||||
|
|
@ -0,0 +1,205 @@
|
|||
# 样例赛题说明
|
||||
|
||||
## GPU高性能并行计算算法优化
|
||||
|
||||
要求参赛者通过一个或多个global kernel 函数（允许配套 device 辅助函数），实现高性能算法。
|
||||
|
||||
在正确性、稳定性前提下，比拼算法性能。
|
||||
|
||||
# 1. ReduceSum算法优化
|
||||
```cpp
|
||||
template <typename InputT = float, typename OutputT = float>
|
||||
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<InputT, true><<<1, threads>>>(d_in, d_out, num_items, init_value);
|
||||
} else {
|
||||
// === 第一阶段 ===
|
||||
// 每一个 Block 处理一部分数据，输出 Partial Sum 到中间 buffer
|
||||
|
||||
// 懒加载分配中间内存 (避免每次调用都 malloc)
|
||||
size_t needed_bytes = blocks * sizeof(OutputT);
|
||||
if (d_intermediate == nullptr || intermediate_capacity < needed_bytes) {
|
||||
if (d_intermediate) mcFree(d_intermediate);
|
||||
MACA_CHECK(mcMalloc((void**)&d_intermediate, needed_bytes));
|
||||
intermediate_capacity = needed_bytes;
|
||||
}
|
||||
|
||||
reduceKernel<InputT, false><<<blocks, threads>>>(d_in, d_intermediate, num_items, 0);
|
||||
|
||||
// === 第二阶段 ===
|
||||
// 将中间结果（blocks 个元素）归约为最终结果
|
||||
// 此时输入是 d_intermediate，输出是 d_out
|
||||
// 加上 init_value
|
||||
reduceKernel<InputT, true><<<1, threads>>>(d_intermediate, d_out, blocks, init_value);
|
||||
}
|
||||
}
|
||||
};
|
||||
```
|
||||
其中
|
||||
|
||||
* 数据类型：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 <typename KeyType, typename ValueType>
|
||||
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<<<gridSize, blockSize>>>(
|
||||
(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<<<gridSize, blockSize>>>(
|
||||
(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 <typename KeyType, typename ValueType>
|
||||
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<KeyType>());
|
||||
} else {
|
||||
thrust::stable_sort_by_key(thrust::device, key_ptr, key_ptr + num_items, value_ptr, thrust::less<KeyType>());
|
||||
}
|
||||
|
||||
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
|
||||
|
||||
|
||||
注意事项
|
||||
|
||||
* 结果必须稳定排序
|
||||
|
||||
加分项
|
||||
|
||||
* 支持其他不同数据类型的键值对，实现类型通用算法
|
||||
* 覆盖更全面的数据范围，提供良好稳定的性能表现
|
||||
|
||||
|
|
@ -4,13 +4,14 @@
|
|||
#include <iostream>
|
||||
#include <vector>
|
||||
#include <iomanip>
|
||||
#include <algorithm>
|
||||
|
||||
|
||||
// ============================================================================
|
||||
// 实现标记宏 - 参赛者修改实现时请将此宏设为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 <typename T>
|
||||
__device__ __forceinline__ T warpReduceSum(T val) {
|
||||
#pragma unroll
|
||||
for (int offset = WARP_SIZE / 2; offset > 0; offset /= 2) {
|
||||
val += __shfl_down_sync(0xffffffff, val, offset);
|
||||
}
|
||||
return val;
|
||||
}
|
||||
|
||||
// 2. Block 级归约：先在 Warp 内归约，再通过 Shared Mem 汇总 Warp 结果
|
||||
template <typename T>
|
||||
__device__ __forceinline__ T blockReduceSum(T val) {
|
||||
// 共享内存用于存储每个 Warp 的总和
|
||||
// 256个线程 -> 8个warp -> 需要8个位置，但为了安全分配32
|
||||
static __shared__ T shared[32];
|
||||
|
||||
int lane = threadIdx.x % WARP_SIZE;
|
||||
int wid = threadIdx.x / WARP_SIZE;
|
||||
|
||||
// 每个 Warp 内部归约
|
||||
val = warpReduceSum(val);
|
||||
|
||||
// 每个 Warp 的第一个线程将结果写入共享内存
|
||||
if (lane == 0) {
|
||||
shared[wid] = val;
|
||||
}
|
||||
|
||||
__syncthreads(); // 等待所有 Warp 写入完毕
|
||||
|
||||
// 最后由第一个 Warp 读取共享内存并进行最终归约
|
||||
// 只有当 Block 大小大于 32 时才需要这一步
|
||||
val = (threadIdx.x < blockDim.x / WARP_SIZE) ? shared[lane] : 0;
|
||||
|
||||
if (wid == 0) {
|
||||
val = warpReduceSum(val);
|
||||
}
|
||||
|
||||
return val;
|
||||
}
|
||||
|
||||
// 3. 通用归约 Kernel (Grid-Stride Loop)
|
||||
// 如果 is_final_pass 为 true，则将结果写入 d_out 并加上 init_value
|
||||
// 否则，将 Block 的部分和写入 d_out (作为临时存储)
|
||||
template <typename T, bool is_final_pass>
|
||||
__global__ void reduceKernel(const T* __restrict__ d_in, T* __restrict__ d_out, int n, T init_value) {
|
||||
T sum = 0;
|
||||
|
||||
// Grid-Stride Loop: 处理数据量大于线程总数的情况
|
||||
int thread_id = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int stride = blockDim.x * gridDim.x;
|
||||
|
||||
for (int i = thread_id; i < n; i += stride) {
|
||||
sum += d_in[i];
|
||||
}
|
||||
|
||||
// Block 内归约
|
||||
sum = blockReduceSum(sum);
|
||||
|
||||
// 由 Block 的线程 0 输出结果
|
||||
if (threadIdx.x == 0) {
|
||||
if (is_final_pass) {
|
||||
d_out[0] = sum + init_value;
|
||||
} else {
|
||||
d_out[blockIdx.x] = sum;
|
||||
}
|
||||
}
|
||||
}
|
||||
#endif
|
||||
|
||||
// ============================================================================
|
||||
// ReduceSum算法实现接口
|
||||
// 参赛者需要替换Thrust实现为自己的高性能kernel
|
||||
// ============================================================================
|
||||
|
||||
template <typename InputT = float, typename OutputT = float>
|
||||
class ReduceSumAlgorithm {
|
||||
public:
|
||||
// 主要接口函数 - 参赛者需要实现这个函数
|
||||
ReduceSumAlgorithm() : d_intermediate(nullptr), intermediate_capacity(0) {}
|
||||
|
||||
// 析构函数：释放临时内存
|
||||
~ReduceSumAlgorithm() {
|
||||
if (d_intermediate) {
|
||||
mcFree(d_intermediate);
|
||||
}
|
||||
}
|
||||
|
||||
// 主要接口函数
|
||||
void reduce(const InputT* d_in, OutputT* d_out, int num_items, OutputT init_value) {
|
||||
|
||||
|
||||
#if !USE_DEFAULT_REF_IMPL
|
||||
// ========================================
|
||||
// 参赛者自定义实现区域
|
||||
// 高性能自定义实现
|
||||
// ========================================
|
||||
|
||||
// TODO: 参赛者在此实现自己的高性能归约算法
|
||||
|
||||
// 示例：参赛者可以调用1个或多个自定义kernel
|
||||
// blockReduceKernel<<<grid, block>>>(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<InputT, true><<<1, threads>>>(d_in, d_out, num_items, init_value);
|
||||
} else {
|
||||
// === 第一阶段 ===
|
||||
// 每一个 Block 处理一部分数据，输出 Partial Sum 到中间 buffer
|
||||
|
||||
// 懒加载分配中间内存 (避免每次调用都 malloc)
|
||||
size_t needed_bytes = blocks * sizeof(OutputT);
|
||||
if (d_intermediate == nullptr || intermediate_capacity < needed_bytes) {
|
||||
if (d_intermediate) mcFree(d_intermediate);
|
||||
MACA_CHECK(mcMalloc((void**)&d_intermediate, needed_bytes));
|
||||
intermediate_capacity = needed_bytes;
|
||||
}
|
||||
|
||||
reduceKernel<InputT, false><<<blocks, threads>>>(d_in, d_intermediate, num_items, 0);
|
||||
|
||||
// === 第二阶段 ===
|
||||
// 将中间结果（blocks 个元素）归约为最终结果
|
||||
// 此时输入是 d_intermediate，输出是 d_out
|
||||
// 加上 init_value
|
||||
reduceKernel<InputT, true><<<1, threads>>>(d_intermediate, d_out, blocks, init_value);
|
||||
}
|
||||
#else
|
||||
// ========================================
|
||||
// 默认基准实现
|
||||
// ========================================
|
||||
auto input_ptr = thrust::device_pointer_cast(d_in);
|
||||
auto output_ptr = thrust::device_pointer_cast(d_out);
|
||||
|
||||
|
||||
// 直接使用thrust::reduce进行归约
|
||||
*output_ptr = thrust::reduce(
|
||||
thrust::device,
|
||||
input_ptr,
|
||||
input_ptr + num_items,
|
||||
static_cast<OutputT>(init_value)
|
||||
);
|
||||
);
|
||||
#endif
|
||||
}
|
||||
|
||||
|
||||
// 获取当前实现状态
|
||||
static const char* getImplementationStatus() {
|
||||
#if USE_DEFAULT_REF_IMPL
|
||||
|
|
@ -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<float, float> algorithm;
|
||||
|
||||
|
||||
bool allPassed = true;
|
||||
|
||||
|
||||
// 测试不同数据规模
|
||||
for (int i = 0; i < NUM_TEST_SIZES && i < 2; i++) { // 限制测试规模
|
||||
int size = std::min(TEST_SIZES[i], 10000);
|
||||
std::cout << " 测试规模: " << size << std::endl;
|
||||
|
||||
|
||||
// 测试普通数据
|
||||
{
|
||||
auto data = generator.generateRandomFloats(size, -10.0f, 10.0f);
|
||||
float init_value = 1.0f;
|
||||
|
||||
|
||||
// CPU参考计算
|
||||
double cpu_result = cpuReduceSum(data, static_cast<double>(init_value));
|
||||
|
||||
|
||||
// GPU计算
|
||||
float *d_in;
|
||||
float *d_out;
|
||||
MACA_CHECK(mcMalloc(&d_in, size * sizeof(float)));
|
||||
MACA_CHECK(mcMalloc(&d_out, sizeof(float)));
|
||||
|
||||
|
||||
MACA_CHECK(mcMemcpy(d_in, data.data(), size * sizeof(float), mcMemcpyHostToDevice));
|
||||
|
||||
|
||||
algorithm.reduce(d_in, d_out, size, init_value);
|
||||
|
||||
|
||||
float gpu_result;
|
||||
MACA_CHECK(mcMemcpy(&gpu_result, d_out, sizeof(float), mcMemcpyDeviceToHost));
|
||||
|
||||
|
||||
// 验证误差
|
||||
double relative_error = std::abs(gpu_result - cpu_result) / std::abs(cpu_result);
|
||||
if (relative_error > REDUCE_ERROR_TOLERANCE) {
|
||||
|
|
@ -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<double>(init_value));
|
||||
|
||||
|
||||
float *d_in;
|
||||
float *d_out;
|
||||
MACA_CHECK(mcMalloc(&d_in, size * sizeof(float)));
|
||||
MACA_CHECK(mcMalloc(&d_out, sizeof(float)));
|
||||
|
||||
|
||||
MACA_CHECK(mcMemcpy(d_in, data.data(), size * sizeof(float), mcMemcpyHostToDevice));
|
||||
|
||||
|
||||
algorithm.reduce(d_in, d_out, size, init_value);
|
||||
|
||||
|
||||
float gpu_result;
|
||||
MACA_CHECK(mcMemcpy(&gpu_result, d_out, sizeof(float), mcMemcpyDeviceToHost));
|
||||
|
||||
|
||||
// 对于包含特殊值的情况，检查是否正确处理
|
||||
if (std::isfinite(cpu_result) && std::isfinite(gpu_result)) {
|
||||
double relative_error = std::abs(gpu_result - cpu_result) / std::abs(cpu_result);
|
||||
|
|
@ -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<float, float> algorithm;
|
||||
|
||||
|
||||
const int WARMUP_ITERATIONS = 5;
|
||||
const int BENCHMARK_ITERATIONS = 10;
|
||||
|
||||
|
||||
// 用于YAML报告的数据收集
|
||||
std::vector<std::map<std::string, std::string>> perf_data;
|
||||
|
||||
|
||||
for (int i = 0; i < NUM_TEST_SIZES; i++) {
|
||||
int size = TEST_SIZES[i];
|
||||
|
||||
|
||||
// 生成测试数据
|
||||
auto data = generator.generateRandomFloats(size);
|
||||
float init_value = 0.0f;
|
||||
|
||||
|
||||
// 分配GPU内存
|
||||
float *d_in;
|
||||
float *d_out;
|
||||
MACA_CHECK(mcMalloc(&d_in, size * sizeof(float)));
|
||||
MACA_CHECK(mcMalloc(&d_out, sizeof(float)));
|
||||
|
||||
|
||||
MACA_CHECK(mcMemcpy(d_in, data.data(), size * sizeof(float), mcMemcpyHostToDevice));
|
||||
|
||||
|
||||
// Warmup阶段
|
||||
for (int iter = 0; iter < WARMUP_ITERATIONS; iter++) {
|
||||
algorithm.reduce(d_in, d_out, size, init_value);
|
||||
}
|
||||
|
||||
|
||||
// 正式测试阶段
|
||||
float total_time = 0;
|
||||
for (int iter = 0; iter < BENCHMARK_ITERATIONS; iter++) {
|
||||
|
|
@ -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<float, float>::getImplementationStatus() << std::endl;
|
||||
if (mode == "all") {
|
||||
std::cout << "正确性: " << (correctness_passed ? "通过" : "失败") << std::endl;
|
||||
std::cout << "性能测试: " << (performance_completed ? "完成" : "跳过") << std::endl;
|
||||
}
|
||||
|
||||
|
||||
return correctness_passed ? 0 : 1;
|
||||
|
||||
|
||||
} catch (const std::exception& e) {
|
||||
std::cerr << "测试出错: " << e.what() << std::endl;
|
||||
return 1;
|
||||
|
|
@ -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
|
||||
|
|
@ -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 <thrust/execution_policy.h>
|
||||
#include <thrust/iterator/zip_iterator.h>
|
||||
#include <thrust/tuple.h>
|
||||
#else
|
||||
// 自定义实现需要的头文件
|
||||
#include <thrust/sort.h>
|
||||
#include <thrust/device_ptr.h>
|
||||
#include <thrust/execution_policy.h>
|
||||
#include <cstdint>
|
||||
#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 <typename KeyType, typename ValueType>
|
||||
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<<<grid, block>>>(d_keys_in, d_values_in, num_items);
|
||||
// mainSortKernel<<<grid, block>>>(d_keys_out, d_values_out, num_items, descending);
|
||||
// postprocessKernel<<<grid, block>>>(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<<<gridSize, blockSize>>>(
|
||||
(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<<<gridSize, blockSize>>>(
|
||||
(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<KeyType>());
|
||||
} 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<float, uint32_t> 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<float> gpu_keys(size);
|
||||
std::vector<uint32_t> 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<float, uint32_t> algorithm;
|
||||
|
||||
|
||||
const int WARMUP_ITERATIONS = 5;
|
||||
const int BENCHMARK_ITERATIONS = 10;
|
||||
|
||||
|
||||
// 用于YAML报告的数据收集
|
||||
std::vector<std::map<std::string, std::string>> perf_data;
|
||||
|
||||
|
||||
for (int i = 0; i < NUM_TEST_SIZES; i++) {
|
||||
int size = TEST_SIZES[i];
|
||||
|
||||
|
||||
// 生成测试数据
|
||||
auto 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<float, uint32_t>::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;
|
||||
|
|
@ -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<<<grid, block>>>(d_keys_in, d_values_in, temp_results, num_items, k);
|
||||
// TopkKernel2<<<grid, block>>>(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<KeyType>());
|
||||
} else {
|
||||
thrust::stable_sort_by_key(thrust::device, key_ptr, key_ptr + num_items, value_ptr, thrust::less<KeyType>());
|
||||
}
|
||||
|
||||
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<KeyType>());
|
||||
|
|
@ -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<float, uint32_t> 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<float> cpu_keys_out;
|
||||
std::vector<uint32_t> 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<float> gpu_keys_out(k);
|
||||
std::vector<uint32_t> 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 << "数据类型: <float, uint32_t>" << std::endl;
|
||||
std::cout << "计算公式:" << std::endl;
|
||||
std::cout << " 吞吐量 = 元素数 / 时间(s) / 1e9 (G/s)" << std::endl;
|
||||
|
||||
|
||||
TestDataGenerator generator;
|
||||
PerformanceMeter meter;
|
||||
TopkPairAlgorithm<float, uint32_t> algorithm;
|
||||
|
||||
|
||||
const int WARMUP_ITERATIONS = 5;
|
||||
const int BENCHMARK_ITERATIONS = 10;
|
||||
|
||||
|
||||
// 用于YAML报告的数据收集
|
||||
std::vector<std::map<std::string, std::string>> 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<float, uint32_t>::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;
|
||||
Loading…
Reference in New Issue