A*寻路算法的核心是通过开放列表维护待探索节点，不断选取代价最小的节点扩展，直到找到目标节点。传统C++串行实现中，开放列表的节点选取、扩展操作都是顺序执行，当地图规模增大、节点数量增多时，性能会明显下降。CUDA作为NVIDIA推出的并行计算平台，能够同时调度成千上万个线程并行处理任务，非常适合优化这类存在大量可并行计算逻辑的算法。
串行A*算法的核心逻辑
在移植之前，先回顾下标准C++串行A*算法的核心实现，方便后续对比并行适配的改动点。下面的代码是简化版的串行A*实现，仅保留核心逻辑：
#include <vector>
#include <queue>
#include <unordered_map>
#include <cmath>
// 节点结构体，记录坐标、代价、父节点信息
struct Node {
int x;
int y;
float g; // 起点到当前节点的实际代价
float h; // 当前节点到目标节点的预估代价
float f; // 总代价 g + h
Node* parent;
Node(int _x, int _y) : x(_x), y(_y), g(0), h(0), f(0), parent(nullptr) {}
// 重载小于运算符，用于优先队列排序
bool operator<(const Node& other) const {
return f > other.f; // 小顶堆，f值小的优先
}
};
// 计算曼哈顿距离作为启发函数
float heuristic(Node* a, Node* b) {
return std::abs(a->x - b->x) + std::abs(a->y - b->y);
}
// 串行A*寻路实现
std::vector<Node*> astar_serial(Node* start, Node* end, const std::vector<std::vector<int>>& map) {
std::priority_queue<Node> open_list;
std::unordered_map<int, Node*> open_map; // 快速查找开放列表中的节点
std::unordered_map<int, Node*> closed_map; // 已探索节点
start->h = heuristic(start, end);
start->f = start->g + start->h;
open_list.push(*start);
open_map[start->x * map[0].size() + start->y] = start;
// 四个方向的移动偏移
int dx[4] = {1, -1, 0, 0};
int dy[4] = {0, 0, 1, -1};
while (!open_list.empty()) {
Node current = open_list.top();
open_list.pop();
Node* current_ptr = open_map[current.x * map[0].size() + current.y];
// 到达目标节点，回溯路径
if (current.x == end->x && current.y == end->y) {
std::vector<Node*> path;
Node* node = current_ptr;
while (node != nullptr) {
path.push_back(node);
node = node->parent;
}
std::reverse(path.begin(), path.end());
return path;
}
// 将当前节点移入关闭列表
closed_map[current.x * map[0].size() + current.y] = current_ptr;
open_map.erase(current.x * map[0].size() + current.y);
// 扩展相邻节点
for (int i = 0; i < 4; i++) {
int nx = current.x + dx[i];
int ny = current.y + dy[i];
// 检查边界和障碍物
if (nx < 0 || nx >= map.size() || ny < 0 || ny >= map[0].size() || map[nx][ny] == 1) {
continue;
}
int key = nx * map[0].size() + ny;
// 如果节点已在关闭列表，跳过
if (closed_map.find(key) != closed_map.end()) {
continue;
}
float new_g = current.g + 1.0f;
// 如果节点不在开放列表，或者新的g值更小，更新节点
if (open_map.find(key) == open_map.end() || new_g < open_map[key]->g) {
Node* neighbor = open_map.find(key) == open_map.end() ? new Node(nx, ny) : open_map[key];
neighbor->g = new_g;
neighbor->h = heuristic(neighbor, end);
neighbor->f = neighbor->g + neighbor->h;
neighbor->parent = current_ptr;
if (open_map.find(key) == open_map.end()) {
open_list.push(*neighbor);
open_map[key] = neighbor;
}
}
}
}
return {}; // 未找到路径
}
A*算法的并行性分析
要将A*移植到CUDA，首先需要分析算法中哪些部分可以并行执行。串行A*的性能瓶颈主要在两个地方：一是每次从开放列表选取最小f值节点的操作，二是扩展当前节点时对每个相邻节点的检查逻辑。其中，相邻节点的检查是完全独立的，没有数据依赖，非常适合用GPU多线程并行处理。而开放列表的维护由于需要全局排序和去重，并行实现复杂度较高，我们可以先采用折中方案：CPU维护开放列表，GPU并行处理节点扩展的计算任务，这也是很多轻量级移植方案的常用思路。
CUDA移植的核心步骤
1. 数据结构的适配
CUDA中无法直接高效使用C++的STL容器，因此需要将节点数据、地图数据转换为C风格数组或者CUDA支持的结构。首先定义设备端可用的节点结构体，和主机端结构体保持一致，去掉不必要的运算符重载：
// 设备端节点结构体
struct DeviceNode {
int x;
int y;
float g;
float h;
float f;
int parent_idx; // 用索引代替指针，避免设备端指针问题
bool valid; // 标记节点是否有效
};
// 设备端地图和辅助数据
struct DeviceMapData {
int* map; // 地图数据，0可走1障碍
int width;
int height;
int start_x;
int start_y;
int end_x;
int end_y;
};
2. 节点扩展的CUDA内核函数
核心的并行逻辑是节点扩展：每个线程处理一个相邻方向的检查，计算新节点的代价，判断是否符合加入开放列表的条件。下面的内核函数实现并行扩展单个节点的四个相邻方向：
__global__ void expand_node_kernel(DeviceNode* current_node, DeviceMapData map_data, DeviceNode* neighbors, float* min_f, int* min_idx) {
int tid = threadIdx.x;
if (tid >= 4) return;
// 四个方向的偏移
int dx[4] = {1, -1, 0, 0};
int dy[4] = {0, 0, 1, -1};
int nx = current_node->x + dx[tid];
int ny = current_node->y + dy[tid];
neighbors[tid].valid = false;
// 边界检查
if (nx < 0 || nx >= map_data.height || ny < 0 || ny >= map_data.width) {
return;
}
// 障碍物检查
if (map_data.map[nx * map_data.width + ny] == 1) {
return;
}
// 计算新节点的代价
neighbors[tid].x = nx;
neighbors[tid].y = ny;
neighbors[tid].g = current_node->g + 1.0f;
// 曼哈顿距离启发函数
neighbors[tid].h = fabsf(nx - map_data.end_x) + fabsf(ny - map_data.end_y);
neighbors[tid].f = neighbors[tid].g + neighbors[tid].h;
neighbors[tid].parent_idx = current_node->x * map_data.width + current_node->y;
neighbors[tid].valid = true;
// 原子操作更新最小f值（简化版，实际场景需要更复杂的同步）
float old_min = atomicMinf(min_f, neighbors[tid].f);
if (neighbors[tid].f < old_min) {
atomicExch(min_idx, tid);
}
}
3. 主机端调度逻辑
主机端负责维护开放列表、调用CUDA内核、处理数据在主机和设备之间的拷贝。核心流程如下：
#include <cuda_runtime.h>
#include <vector>
#include <queue>
#include <unordered_map>
// 原子操作辅助函数，用于浮点数最小值更新
__device__ float atomicMinf(float* addr, float value) {
float old = *addr;
while (value < old) {
old = atomicCAS((unsigned int*)addr, __float_as_uint(old), __float_as_uint(value));
}
return old;
}
// CUDA版本A*寻路主函数
std::vector<Node*> astar_cuda(Node* start, Node* end, const std::vector<std::vector<int>>& map) {
int width = map[0].size();
int height = map.size();
int node_num = width * height;
// 1. 准备主机端数据
std::vector<int> host_map(node_num);
for (int i = 0; i < height; i++) {
for (int j = 0; j < width; j++) {
host_map[i * width + j] = map[i][j];
}
}
// 2. 分配设备内存
int* dev_map;
cudaMalloc(&dev_map, node_num * sizeof(int));
cudaMemcpy(dev_map, host_map.data(), node_num * sizeof(int), cudaMemcpyHostToDevice);
DeviceMapData host_map_data;
host_map_data.map = dev_map;
host_map_data.width = width;
host_map_data.height = height;
host_map_data.start_x = start->x;
host_map_data.start_y = start->y;
host_map_data.end_x = end->x;
host_map_data.end_y = end->y;
DeviceMapData* dev_map_data;
cudaMalloc(&dev_map_data, sizeof(DeviceMapData));
cudaMemcpy(dev_map_data, &host_map_data, sizeof(DeviceMapData), cudaMemcpyHostToDevice);
// 设备端节点数组，存储所有生成的节点
DeviceNode* dev_nodes;
cudaMalloc(&dev_nodes, node_num * sizeof(DeviceNode));
cudaMemset(dev_nodes, 0, node_num * sizeof(DeviceNode));
// 初始化起点
DeviceNode start_node;
start_node.x = start->x;
start_node.y = start->y;
start_node.g = 0;
start_node.h = fabsf(start->x - end->x) + fabsf(start->y - end->y);
start_node.f = start_node.g + start_node.h;
start_node.parent_idx = -1;
start_node.valid = true;
cudaMemcpy(dev_nodes + (start->x * width + start->y), &start_node, sizeof(DeviceNode), cudaMemcpyHostToDevice);
// 主机端开放列表和关闭列表
std::priority_queue<Node> open_list;
std::unordered_map<int, Node*> open_map;
std::unordered_map<int, Node*> closed_map;
start->h = start_node.h;
start->f = start_node.f;
open_list.push(*start);
open_map[start->x * width + start->y] = start;
// 临时存储扩展结果的设备内存
DeviceNode* dev_neighbors;
cudaMalloc(&dev_neighbors, 4 * sizeof(DeviceNode));
float* dev_min_f;
cudaMalloc(&dev_min_f, sizeof(float));
int* dev_min_idx;
cudaMalloc(&dev_min_idx, sizeof(int));
// 寻路主循环
while (!open_list.empty()) {
Node current = open_list.top();
open_list.pop();
Node* current_ptr = open_map[current.x * width + current.y];
int current_key = current.x * width + current.y;
if (current.x == end->x && current.y == end->y) {
// 找到目标，回溯路径
std::vector<Node*> path;
Node* node = current_ptr;
while (node != nullptr) {
path.push_back(node);
node = node->parent;
}
std::reverse(path.begin(), path.end());
// 释放设备内存
cudaFree(dev_map);
cudaFree(dev_map_data);
cudaFree(dev_nodes);
cudaFree(dev_neighbors);
cudaFree(dev_min_f);
cudaFree(dev_min_idx);
return path;
}
closed_map[current_key] = current_ptr;
open_map.erase(current_key);
// 准备当前节点的设备端数据
DeviceNode dev_current;
dev_current.x = current.x;
dev_current.y = current.y;
dev_current.g = current.g;
dev_current.h = current.h;
dev_current.f = current.f;
dev_current.parent_idx = current.parent ? current.parent->x * width + current.parent->y : -1;
dev_current.valid = true;
// 初始化最小f值为最大值
float init_min_f = INFINITY;
int init_min_idx = -1;
cudaMemcpy(dev_min_f, &init_min_f, sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(dev_min_idx, &init_min_idx, sizeof(int), cudaMemcpyHostToDevice);
// 调用CUDA内核并行扩展节点
expand_node_kernel<<<1, 4>>>(dev_current, *dev_map_data, dev_neighbors, dev_min_f, dev_min_idx);
cudaDeviceSynchronize();
// 将扩展结果拷贝回主机
DeviceNode host_neighbors[4];
cudaMemcpy(host_neighbors, dev_neighbors, 4 * sizeof(DeviceNode), cudaMemcpyDeviceToHost);
// 处理扩展得到的相邻节点，更新开放列表
for (int i = 0; i < 4; i++) {
if (!host_neighbors[i].valid) continue;
int key = host_neighbors[i].x * width + host_neighbors[i].y;
if (closed_map.find(key) != closed_map.end()) continue;
float new_g = host_neighbors[i].g;
if (open_map.find(key) == open_map.end() || new_g < open_map[key]->g) {
Node* neighbor = open_map.find(key) == open_map.end() ? new Node(host_neighbors[i].x, host_neighbors[i].y) : open_map[key];
neighbor->g = new_g;
neighbor->h = host_neighbors[i].h;
neighbor->f = host_neighbors[i].f;
neighbor->parent = current_ptr;
if (open_map.find(key) == open_map.end()) {
open_list.push(*neighbor);
open_map[key] = neighbor;
}
// 更新设备端节点数组
DeviceNode dev_neighbor;
dev_neighbor.x = neighbor->x;
dev_neighbor.y = neighbor->y;
dev_neighbor.g = neighbor->g;
dev_neighbor.h = neighbor->h;