如何避免CUDA中隐式使用本地内存？性能瓶颈排查求助

我正在开发一款CUDA路径追踪渲染器，目前遇到L1TEX本地加载/存储访问模式优化不佳的问题。通过NVIDIA Nsight Compute（NCU）分析，性能瓶颈集中在无栈线性BVH遍历的核心代码中，相关指标为L2 Theoretical Sectors Local。以下是瓶颈代码（已标记本地内存访问位置）：

__device__ float ray_intersect_bvh(
    const Ray& ray,
    const cudaTextureObject_t bvh_leaves,
    const cudaTextureObject_t node_fronts,
    const cudaTextureObject_t node_backs,
    ConstF4Ptr cached_nodes,
    const PrecomputedArray& verts,
    int& min_index,
    int& min_obj_idx,
    float& prim_u,
    float& prim_v,
    const int node_num,
    const int cache_num,
    float min_dist
) {
    bool valid_cache = false;         // whether we find valid node-intersection in cached nodes
    int node_idx     = 0;             // BVH tree node index
    float aabb_tmin  = 0;             // minimum intersection time of AABB (positive)
    // The following while lobe checks cached (in shared memory) node intersection
    // near root layers of the BVH tree are cached for faster access
    while (node_idx < cache_num && !valid_cache) {
        const LinearNode node(
            cached_nodes[node_idx],
            cached_nodes[node_idx + cache_num]
        );
        bool intersect_node = node.aabb.intersect(ray, aabb_tmin) && aabb_tmin < min_dist;
        int all_offset = node.aabb.base(), gmem_index = node.aabb.prim_cnt();
        int increment = (!intersect_node) * all_offset + (intersect_node && all_offset != 1) * 1;

        node_idx += increment;

        if (intersect_node && all_offset == 1) {
            valid_cache = true;
            node_idx = gmem_index;
        }
    }
    // if we find a valid intersection in cached nodes, we continue traversal in texture memory
    if (valid_cache) {
        while (node_idx < node_num) {
            const LinearNode node(tex1Dfetch<float4>(node_fronts, node_idx),
                            tex1Dfetch<float4>(node_backs, node_idx));
            bool intersect_node = node.aabb.intersect(ray, aabb_tmin) && aabb_tmin < min_dist;
            int beg_idx = 0, end_idx = 0;
            node.get_range(beg_idx, end_idx);
            
            int increment = (!intersect_node) * (end_idx < 0 ? -end_idx : 1) + int(intersect_node);
            if (intersect_node && end_idx > 0) {
                end_idx += beg_idx;
                // For BVH leaf node: traverse all the triangles within the range
                for (int idx = beg_idx; idx < end_idx; idx ++) {   /* Up to 44%.12 L2 theoretical sectors are accessed for this line */
                    /* SASS for the above line: LDL.LU, load 32 bit. */
                    bool valid  = false;
                    int2 obj_prim_idx = tex1Dfetch<int2>(bvh_leaves, idx);
                    float it_u = 0, it_v = 0, dist = Primitive::intersect(ray, verts, obj_prim_idx.y, it_u, it_v, obj_prim_idx.x >= 0);
                    valid = dist > EPSILON && dist < min_dist;
                    min_dist = valid ? dist : min_dist;
                    prim_u   = valid ? it_u : prim_u;
                    prim_v   = valid ? it_v : prim_v;
                    min_index = valid ? obj_prim_idx.y : min_index;
                    min_obj_idx = valid ? obj_prim_idx.x : min_obj_idx;
                }
            }
            node_idx += increment;
        }
    }
    return min_dist;
}

此外，AABB相交测试函数中存在更严重的本地内存存储模式问题，NCU显示该函数的本地内存访问效率极低：

__device__ bool intersect(const Ray& ray, float& t_near) const {
    // Vec3 is a simple encapsulation of float3
    Vec3 invDir = ray.d.rcp();
    Vec3 t1s = (mini - ray.o) * invDir;             // long scoreboard
    Vec3 t2s = (maxi - ray.o) * invDir;

    float tmin = t1s.minimize(t2s).max_elem();
    float tmax = t1s.maximize(t2s).min_elem();
    t_near = tmin;
    return (tmax > tmin) && (tmax > 0);             /* SASS: STL, store 32 bit. local memory access problem */
}

上下文信息 #

• 编译参数：使用-maxrregcount=56，设备为RTX3060 Laptop；NCU显示上述两个函数的最大活跃寄存器数为52，寄存器资源充足，理论上不应出现意外寄存器溢出。
• 无动态本地数组索引，仅使用共享内存、全局内存和纹理内存访问。
• NCU数据详情：平均每个线程仅利用了每个扇区32字节传输数据中的1.1字节（存储）|15.8字节（加载），利用率极低。

疑问与需求 #

1. 为何在寄存器充足的情况下，编译器仍会生成隐式本地内存访问？
2. 有哪些有效的编译器提示或代码优化方法可以避免这种情况？（已知C++的register关键字基本无效）

内容的提问来源于stack exchange，提问作者Enigmatisms