CUDA Samples 之 lineOfSight：基于 Thrust 的视线可见性分析示例深入解析

CUDA Samples 之 lineOfSight基于 Thrust 的视线可见性分析示例深入解析【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples本篇技术指南聚焦 CUDA Samples 仓库中的lineOfSight示例位于 cpp/4_CUDA_Libraries/lineOfSight它实现了一个经典的数据并行算法——视线分析Line of Sight给定一张高度图height map和一条从观测点出发的射线计算射线上所有点中哪些点对观测点可见。该示例以开源 CUDA Thrust 库为核心实现并配套纹理对象采样、前缀扫描max-scan与黄金参考解验证。读完本文你将掌握该算法的数学原理、两个 CUDA 核函数与 Thrust 库的组合方式以及如何在当前仓库中构建、运行并验证该示例。一、示例概览与算法背景根据 README.mdlineOfSight的核心任务是给定一张高度图与一条从某观测点出发的射线计算射线上所有从观测点可见的点。实现基于Thrust Library而算法思想来源于 Guy E. Blelloch.Vector models for>// Boolean typedef unsigned char Bool; enum { False 0, True 1 }; // 2D height field二维高度场 struct HeightField { int width; float *height; }; // Ray射线 struct Ray { float3 origin; // 观测点3D 坐标含高度 z float2 dir; // 水平方向上的前进方向 int length; // 射线上的采样点数量 float oneOverLength; // 1 / length用于插值定位 };Bool用单字节unsigned char表示布尔值便于在设备端紧凑存储可见性结果HeightField仅记录宽度与数据指针采用行主序row-major存储height width * y x即二维坐标(x, y)处的高度Ray中origin.z表示观测点高度dir是水平面上的 2D 方向向量oneOverLength预计算用于把整数步长 i 归一化到[0, 1]区间做线性插值。三、核心算法两个 CUDA 核函数 Thrust 扫描算法在主机端runTest中被组织为三阶段流水线lineOfSight.cufor (uint i 0; i numIterations; i) { // 1) 计算射线上每个点的仰角 computeAngles_kernelgrid, block(ray, thrust::raw_pointer_cast(d_angles[0]), heightFieldTex); // 2) 对仰角数组执行 inclusive max-scan thrust::inclusive_scan(d_angles.begin(), d_angles.end(), d_scannedAngles.begin(), thrust::maximumfloat()); // 3) 依据仰角与其扫描结果判定可见性 computeVisibilities_kernelgrid, block(thrust::raw_pointer_cast(d_angles[0]), thrust::raw_pointer_cast(d_scannedAngles[0]), ray.length, thrust::raw_pointer_cast(d_visibilities[0])); }3.1 computeAngles_kernel并行计算仰角__global__ void computeAngles_kernel(const Ray ray, float *angles, cudaTextureObject_t HeightFieldTex) { uint i blockDim.x * blockIdx.x threadIdx.x; if (i ray.length) { float2 location getLocation(ray, i 1); // 射线上的第 i1 个采样点坐标 float height tex2Dfloat(HeightFieldTex, location.x, location.y); // 纹理采样取高度 float angle getAngle(ray, location, height); // 计算仰角 angles[i] angle; } }关键点每个线程对应射线上的一个采样点通过经典的blockDim.x * blockIdx.x threadIdx.x扁平化索引做数据并行高度图通过 CUDA 纹理对象texture object采样而不是直接全局内存读取。这一设计让 GPU 纹理硬件完成二维插值与边界处理getLocation与getAngle均为__device__ __host__双修饰函数可在主机与设备两端复用lineOfSight.cu__device__ __host__ float2 getLocation(const Ray ray, int i) { float step i * ray.oneOverLength; return make_float2(ray.origin.x, ray.origin.y) ray.dir * step; } __device__ __host__ float getAngle(const Ray ray, float2 location, float height) { float2 dir location - make_float2(ray.origin.x, ray.origin.y); return atanf((height - ray.origin.z) / length(dir)); }其中仰角定义为atanf((height - origin.z) / horizontal_distance)即观测点与该地形点连线相对水平面的夹角length(dir)即水平距离。3.2 Thrust inclusive_scan前缀最大值扫描仰角数组计算完毕后调用一次 Thrust 的包含式扫描thrust::inclusive_scan(d_angles.begin(), d_angles.end(), d_scannedAngles.begin(), thrust::maximumfloat());二元操作符为thrust::maximumfloat()因此这是一个max-scan前缀最大值结果数组d_scannedAngles[i]表示d_angles[0..i]中的最大值即到当前位置为止视线需要跨越的最大仰角数组由thrust::device_vectorfloat管理lineOfSight.cu在两次核函数之间无需显式分配或释放显存这也是 Thrust 容器在本示例中的核心价值。3.3 computeVisibilities_kernel可见性判定__global__ void computeVisibilities_kernel(const float *angles, const float *scannedAngles, int numAngles, Bool *visibilities) { uint i blockDim.x * blockIdx.x threadIdx.x; if (i numAngles) { visibilities[i] scannedAngles[i] angles[i]; } }判定规则极其简洁当该点自身仰角不小于前方最大仰角时可见。由于scannedAngles[i]包含了angles[i]自身保证了与观测点连线平齐不被前方任何点高出的点同样可见。最终可见性结果存于thrust::device_vectorBool d_visibilities再通过thrust::copy拷回主机。3.4 黄金参考解gold reference为保证结果正确示例在主机端实现了一个串行参考版本lineOfSight_goldlineOfSight.cu沿射线顺序遍历用angleMax记录已见最大仰角仅当当前点仰角超过angleMax时才标记可见。这与设备端max-scan 逐点比较的并行写法在数学上完全等价构成可交叉验证的对照实现。四、纹理对象配置从高度图到可采样纹理高度图数据在主机端用thrust::host_vectorfloat生成一个任意构造的正弦曲面见 lineOfSight.cu随后经历CUDA 数组 → 纹理对象的标准流程lineOfSight.cu// 1) 创建 32 位单通道 float 通道描述符 cudaChannelFormatDesc channelDesc cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat); cudaArray *heightFieldArray; checkCudaErrors(cudaMallocArray(heightFieldArray, channelDesc, dim.x, dim.y)); // 2) 2D 拷贝主机高度数据 - CUDA 数组 checkCudaErrors(cudaMemcpy2DToArray(heightFieldArray, 0, 0, heightField.height, dim.x * sizeof(float), dim.x * sizeof(float), dim.y, cudaMemcpyHostToDevice)); // 3) 配置资源描述与纹理描述 cudaTextureObject_t heightFieldTex; cudaResourceDesc texRes; memset(texRes, 0, sizeof(cudaResourceDesc)); texRes.resType cudaResourceTypeArray; texRes.res.array.array heightFieldArray; cudaTextureDesc texDescr; memset(texDescr, 0, sizeof(cudaTextureDesc)); texDescr.normalizedCoords false; // 使用像素坐标非归一化坐标 texDescr.filterMode cudaFilterModePoint; // 点采样不做双线性插值 texDescr.addressMode[0] cudaAddressModeClamp; // x 方向越界钳制 texDescr.addressMode[1] cudaAddressModeClamp; // y 方向越界钳制 texDescr.readMode cudaReadModeElementType; // 按元素类型原样读取 // 4) 创建纹理对象 checkCudaErrors(cudaCreateTextureObject(heightFieldTex, texRes, texDescr, NULL));配置要点解读cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat)声明单通道 32 位浮点通道对应 README 中列出的 Runtime APIcudaCreateChannelDesc、cudaMallocArrayfilterMode cudaFilterModePoint表示采样时不做滤波直接取最近像素的高度值与黄金参考解中floorf取整读数的行为保持一致lineOfSight.cuaddressMode cudaAddressModeClamp在射线端点处出现越界采样时把坐标钳制到纹理边界保证核函数不会访问越界数据采样点在核函数内通过tex2Dfloat(HeightFieldTex, location.x, location.y)完成lineOfSight.cu。该示例也同时覆盖了 README 中列出的另外三个 Runtime APIcudaFreeArray释放高度图数组lineOfSight.cu、cudaDeviceSynchronize等待全部核函数执行完毕lineOfSight.cu与cudaCreateTextureObject。五、执行配置、计时与结果验证5.1 执行配置核函数的网格配置由射线长度决定lineOfSight.cudim3 block(256); dim3 grid((uint)ceil(ray.length / (double)block.x));示例构造的测试场景为一个10000 x 100的高度场dim make_int2(10000, 100)射线从(0, 0, 2.0f)出发、沿高度场对角线方向前进dir (9999, 99)因此ray.length max(9999, 99) 9999个采样点lineOfSight.cu向上取整得到 40 个线程块。5.2 计时主机端通过sdkCreateTimer/sdkStartTimer/sdkStopTimer/sdkGetTimerValue来自 Common/helper_functions.h对 100 次迭代numIterations 100计时最后输出单次平均耗时printf(Average time: %f ms\n\n, sdkGetTimerValue(timer) / numIterations);5.3 与参考解比对设备端结果经thrust::copy拷回主机后调用compareData定义于 Common/helper_image.h与黄金参考解逐元素比对bool res compareData(thrust::raw_pointer_cast(h_visibilitiesRef[0]), thrust::raw_pointer_cast(h_visibilities[0]), ray.length, 0.0f, 0.0f);compareData的语义是当threshold 0.0f时要求两数组逐元素完全一致允许的浮点差为 epsilon此处传 0可见性结果为离散布尔值因此可以实现严格零误差比对。main函数依据runTest的返回值打印Test passed或Test failedlineOfSight.cu。5.4 设备选择runTest首先调用findCudaDevice(argc, (const char **)argv)lineOfSight.cu该工具函数位于 Common/helper_cuda.h支持通过命令行参数-deviceN指定目标 GPU否则自动选择计算能力最高Gflops/s 最大的设备。六、构建与运行该示例隶属于4_CUDA_Libraries目录已通过 cpp/4_CUDA_Libraries/CMakeLists.txt 中的add_subdirectory(lineOfSight)接入整体 CMake 构建。其独立的 CMakeLists.txt 要点如下project(lineOfSight LANGUAGES CUDA)find_package(CUDAToolkit REQUIRED)声明 CUDA 语言支持并定位 Toolkitinclude_directories(../../../Common)引入仓库公共头文件目录用于解析helper_cuda.h、helper_functions.h、helper_math.hset(CMAKE_CUDA_ARCHITECTURES 75 80 86 87 89 90 100 110 120)默认生成多个现代架构对应 README 中列出的 SM 7.5 至 SM 9.0 支持范围README 同时声明支持 SM 5.0/5.2/5.3/6.0/6.1/7.0/7.2 等更早架构可按需覆盖该变量开启CUDA_SEPARABLE_COMPILATION并启用--extended-lambda、C17 特性通过cmake/InstallSamples.cmake中的setup_samples_install()支持安装。构建方式与仓库中其他示例一致以仓库根目录为基准的典型流程mkdir build cd build cmake ../cpp/4_CUDA_Libraries/lineOfSight make -j$(nproc) ./lineOfSight运行后程序会输出Line of sight、Average time: xxx ms以及最终的Test passed验证结果。也可传入-deviceN指定 GPU 设备。七、支持环境与依赖依据 README.md 与构建文件支持的 SM 架构SM 5.0、5.2、5.3、6.0、6.1、7.0、7.2、7.5、8.0、8.6、8.7、8.9、9.0覆盖 Maxwell 至 Hopper/Blackwell 世代支持的操作系统Linux、Windows支持的 CPU 架构x86_64、armv7l前置条件安装对应平台的 CUDA Toolkit 公共头文件。八、小结数据并行视线分析的实现范式lineOfSight是一个体量精简但结构完整的教科书式示例它演示了三条可复用的 CUDA 工程范式用一次 max-scan 把顺序依赖变为数据并行视线可见性天然具有前方最大值的前缀依赖Thrust 的inclusive_scan将其转化为两阶段并行计算并行求仰角 → 并行扫描 → 并行比较纹理对象管理二维场数据cudaMallocArraycudaCreateTextureObjecttex2D的完整链路配合cudaFilterModePoint与cudaAddressModeClamp实现无越界、无插值的精确采样黄金参考解驱动的自验证串行参考实现与设备端并行实现交叉比对compareData零容差配合计时器输出平均耗时兼顾正确性与性能可观测性。该示例非常适合作为学习 CUDA 运行时 API、纹理对象与 Thrust 算法库三者协同工作的入门到进阶素材读者可直接以 lineOfSight.cu 为模板将其推广到地形可视域分析、雷达/通信链路遮挡判定等实际场景。【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考