Bluegill

0%

发表于： 更新于： 分类于：

再次考虑 dot 计算

在第 5 章的 dot 计算中，我们在每个块上做完 reduction 之后就将数据拷贝回到 CPU 了，然后让 CPU 做最后的加法。

为什么在 compute capability 2.0 之前，atomicAdd 只支持整数？因为原子加法是不能指定计算的发生顺序的，因而每个计算都必须遵守结合律，也就是 $(A+B)+C$ 必须等于 $A+(B+C)$。但是浮点数因为中间结果的舍入问题，并不能保证这一点！！

本节接下来是讲解用原子操作实现一个忙等待的 mutex，用于同步多个 CUDA 线程（因为在写书的时候浮点数的原子加法还没有受到设备的广泛支持）。atomicCAS 就是 CUDA 上的比较并交换。

// mutex 的类型是 int *，而且是分配在 GPU 上的
__device__ void lock( void ) {
  while( atomicCAS( mutex, 0, 1 ) != 0 );
}

__device__ void unlock( void ) {
  atomicExch( mutex, 1 );
}

Note

这里的 unlock 方法并不是直接对 *mutex 赋值为 1，因为 CUDA 中原子操作和普通的内存访问经过的路径不同，所以应该在 unlock 的时候也统一使用原子操作。

发表于： 更新于： 分类于：

源码可以参考 https://github.com/yottaawesome/cuda-by-example/ ，官网的源码链接挂了。

书中的代码有些需要用 opengl 来跑。安装了 freeglut3-dev 和 mesa-utils。（不确定 libgl1-mesa-dev 是否是必要的。）然后 cmake 规则中要 link 对应的库：

cmake_minimum_required(VERSION 3.20.1)
project(chapter3 LANGUAGES CXX CUDA)
set(CMAKE_CUDA_STANDARD 17)
set(CMAKE_CUDA_STANDARD_REQUIRED ON)
set(CMAKE_CUDA_EXTENSIONS OFF)
add_executable(ray ray_global.cu)
target_link_libraries(ray GL glut)
#                         ^^^^^^^

第 3 章 CUDA 源文件

用 nvcc 编译时不需要为 cuda 内置函数额外包含头文件。这些头文件是在 host 端才需要的。

https://stackoverflow.com/questions/6302695/difference-between-cuda-h-cuda-runtime-h-cuda-runtime-api-h

发表于： 更新于： 分类于：

第 6 章 Constant Memory and Events

常量内存

常量内存是在全局区域声明的。如果漏掉了 __constant__ 关键字，就会将其定义在全局内存区域，尽管存储方式、分配的时机和用 cudaMalloc 申请的内存有一些差异。

__constant__ Sphere s[SPHERES];

常量内存的内存拷贝方法比较特殊：

HANDLE_ERROR( cudaMemcpyToSymbol( s, temp_s,
                                  sizeof(Sphere) * SPHERES) );

CUDA 线程对常量内存是只读的，也就是只有 host 能操作常量内存。通过将反复读取的数据移动到常量内存区域而不是全局内存，可以加速。但是要注意常量内存的大小非常有限（）。书中的例子只是对 20 个球体做光线追踪。

发表于： 更新于： 分类于：

第 9 章 原子操作

You should know that atomic operations on global memory are supported only on GPUs of compute capability 1.1 or higher. Furthermore, atomic operations on shared memory require a GPU of compute capability 1.2 or higher.

指定计算能力：

nvcc -arch=sm_11

这样就指定了计算能力是 1.1，当有些指令是只有 1.1 才能编译时加这个参数可以确保编译。同时，有了更加精确的生成目标，nvcc 可以执行一些和硬件相关的优化手段，这些优化手段在更早的架构上可能没有。

但是一个硬件上不一定支持给定的计算能力。通过 -arch-ls 可以列出设备支持的计算能力：