CUDA学习笔记（4）光线追踪（常量内存）

最新推荐文章于 2025-07-13 12:26:11 发布

原创

最新推荐文章于 2025-07-13 12:26:11 发布 · 1.3k 阅读

3 ·

CC 4.0 BY-SA版权

文章标签：

#cuda #线程 #nvidia #硬件 #gpu

根据书上内容，本节通过一个实验“光线追踪”来了解和测试常量内存带来的效果提升。

常量内存是NVIDIA提供的一个64KB大小的内存空间，它的处理方式和普通的全局内存和共享内存都不一样，是有cuda专门提供的。

线程束的概念：线程束是指一个包含32个线程的集合，在程序中的每一行，线程束中的每个线程都将在不同的数据上执行相同的指令。

因此，常量内存的作用是，能够将单次内存的读取操作广播到每个半线程束（即16个线程），所以如果在半线程束中的每个线程都从常量内存的相同地址上读取数据，那么GPU只会产生一次读取请求，并将其广播，显而易见，这种方式的内存流量只是使用全局内存流量的1/16。这是常量内存的第一个好处，第二个好处则是由于这块内存的内容是不会发生变化的，因此硬件将主动把这个常量内存数据缓存到GPU上，这样第一次从敞亮内存的某个地址上读取后，其他半线程束请求同一个地址时，将直接在GPU上命中缓存，因此也减少了额外的内存流量。

使用常量内存只需加上:__constant__修饰符，当从主机内存复制内存到GPU上常量内存时，不用cudaMemcpy()而用cudaMemcpyToSymbol()，这样就复制到常量内存里了。

在测试程序中涉及到了cuda事件，我们通过这个API去记录程序运行的时间，以此来将是否用到常量内存这两个程序进行比较。这部分代码大致为：

cudaEvent_t start,stop;

cudaEventCreat(&start);

cudaEventCreat(&stop);

cudaEventRecord(start,0);

//执行工作

cudaEventRecord(stop,0);

cudaEventSynchronize(stop);

大致含义就是建立两个事件，并且记录第一个start时间，此时开始记录时间，工作执行完之后，再记录结束时间。然而我们为了让GPU完全执行完语句，再记录Stop里的准确时间，我们不得不加入cudaEventSynchronize()这个函数，否则将得到不可靠的时间结果。

以下代码为没有使用常量内存的光线追踪代码，测试球面为30个

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "H:\cuda_by_example\common\book.h"
#include "H:\cuda_by_example\common\cpu_bitmap.h"
#include "device_functions.h"
#include <stdio.h>
  
#define DIM 1024  
  
#define rnd( x ) (x * rand() / RAND_MAX)  
#define INF     2e10f  
  
//数据结构对球面建模
struct Sphere {  
    float   r,b,g;  
    float   radius;  
    float   x,y,z;  
	//hit方法，计算光线是否与球面相交，若相交则返回光线到命中球面处的距离
    __device__ float hit( float ox, float oy, float *n ) {  
        float dx = ox - x;  
        float dy = oy - y;  
        if (dx*dx + dy*dy < radius*radius) {  
            float dz = sqrtf( radius*radius - dx*dx - dy*dy );  
            *n = dz / sqrtf( radius * radius );  
            return dz + z;