我将该仓库中的check1api.cu及error.cuh代码直接拷贝在ubuntu 22.04 LTS的VS code中并保存在同一文件夹下，后用命令nvcc check1api.cu，编译结果是no error。GPU是GeForce RTX 4090。不知道具体原因为何，希望查看解答一下。

Why Chinese people don't write English. It becomes totally hard to translate each sentence.

行12处应该为 N/TILE_DIM，因为74页描述就是，N/tile_dim=128/32=4

为什么我跑出来的结果是0 1 2 3或者0 0 0 0
这是为什么？

你好，非常nice的一本书，但是平时习惯了看电子书，且电子书做笔记方便，所以想问一下，有没有pdf版的呢？

97页（第9章 原子函数的合理使用， 9.3.1 C++版本的开发）cutoff_square 取为 $1.9^2 A^2$。 应为 $ 1.9^2$:

如题。个人GPU为GeForce RTX 4090。
在个人所练习的CUDA 12.5版本中，cuda-memcheck已被废弃，P45【关于CUDA-MEMCHECK的更多内容，参见https://docs.nvidia.com/cuda/cuda-memcheck。】这段文字已完全失效！我看目前的勘误表里尚未订正，特地声明。

这一行有问题，__host__运行在主机端

老师您好，我非常期待您的CUDA编程书籍，迫不及待想问问您何时能出版，谢谢！

这应该是“其性能基本上与数组长度N成正比”吧.

针对最新的勘误表核对了下2022年10月第5次印刷的书籍。
大部分已经该改正。除了以下几处。

1. 书中Listing 12.4代码与现GitHub项目中的oversubscription2.cu代码差别较大，与之前CUDA-Programming-v1.0的代码也有出入。至少“size 应改为 size / sizeof(uint64_t)”改的不成功，改成了“size) / sizeof(uint64_t)”。至少是笔误了。

• 2022年10月第5次印刷的书中代码：
  

• 最新oversubscription2.cu代码：

int main(void)
{
    for (int n = 1; n <= N; ++n)
    {
        const size_t memory_size = size_t(n) * 1024 * 1024 * 1024;
        const size_t data_size = memory_size / sizeof(uint64_t);
        uint64_t *x;
        CHECK(cudaMallocManaged(&x, memory_size));
        gpu_touch<<<(data_size - 1) / 1024 + 1, 1024>>>(x, data_size);
        CHECK(cudaGetLastError());
        CHECK(cudaDeviceSynchronize());
        CHECK(cudaFree(x));
        printf("Allocated %d GB unified memory with GPU touch.\n", n);
    }
    return 0;
}

• 老版本oversubscription2.cu代码：

int main(void)
{
    for (int n = 1; n <= N; ++n)
    {
        const size_t size = size_t(n) * 1024 * 1024 * 1024;
        uint64_t *x;
        CHECK(cudaMallocManaged(&x, size));
        gpu_touch<<<size / sizeof(uint64_t) / 1024, 1024>>>(x, size);
        CHECK(cudaGetLastError());
        CHECK(cudaDeviceSynchronize());
        CHECK(cudaFree(x));
        printf("Allocated %d GB unified memory with GPU touch.\n", n);
    }
    return 0;
}

2. 核函数的笔误还没改。（虽然没啥影响）

3. 这个认知错误，2022年10月书中好像没看到说明。（或者转述了没看出来？）

老师好，
P48页，”第一次计算时，机器（无论是CPU还是GPU）都可能处于预热状态，测得的时间往往偏大”
第一次运行CUDA程序，或者运行其他深度学习算法，模型载入后，第一次计算相比之后确实会非常耗时。请问第一次运行程序时，具体都是因为哪些原因导致计算较慢呢？

我查到可能的原因包括：即时编译、将内核传输到GPU内存，缓存内容
https://www.javaroad.cn/questions/137499#

请问此外还包含哪些原因吗？

在本书的第129页，Listing 11.2 中列举了本章程序 kernel-kernel.cu 的部分内容。

另外，在书本的第130页，第二段的末尾，您也说明，“为了计时方便，核函数中故意做了10^6次加法运算。”
但是，在Github相应的代码库中(CUDA-Programming/src/11-stream/kernel-kernel.cu)，我们可以发现，其中的add函数仅仅增加了10^5次，也就是说，您展示于库中的代码，与书本展示的代码并不匹配。

void __global__ add(const real *d_x, const real *d_y, real *d_z)
{
    const int n = blockDim.x * blockIdx.x + threadIdx.x;
    if (n < N1)
    {
        for (int i = 0; i < 100000; ++i)
        {
            d_z[n] = d_x[n] + d_y[n];
        }
    }
}

Page 63关于6.2.4节最后一段第三行，“该表只包含了少数几个计算能力，更完整的列表见《CUDA C++ Programming Guide》的附录H”，这一句话是不对的。因为该在线网站已经没有附录了。其本该对应的资料已经转移到https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#compute-capabilities中，这是对应于其第16章【计算能力】的表21中的部分数据。此处需要更新改正！

老师好，
P59页“全局内存的生命周期不是由核函数决定的，而是由主机端决定的”，第10章代码，reduce2static.cu中，声明静态全局变量如下，
device real static_y[GRID_SIZE];
但该数组并未释放。请问静态全局内存变量的生命周期？采用cudaFree(static_y)释放时报错，请问该如何释放该变量呢？

这里的cudaGetErrorString后是否应该添加反斜杠 \ （github中的代码并没有分两行写）
我在vscode里键入同样的代码如果不加反斜杠会报错？

P65 6.3.2节第一段第三行，“获得100%的占有率并不是获得高性能的必要或充分条件”，这句话不就是指100%占有率和获得高性能毫无关系么？感觉和后面矛盾啊。最多也是就说，“获得100%的占有率并不是获得高性能的充分条件”，这样比较合适？还是说有什么更深层次的理解？这里说明力度也不够啊？我觉得读起来很困惑。

你好，有一个问题需要请教
在8.3中，通过__shared__ real S[TILE_DIM][TILE_DIM + 1]来解决bank冲突的问题。原先S[TILE_DIM][TILE_DIM]的数据刚好可以存进32个bank（具有32层）。现在数组变大了，那多出来的那部分数据被存到哪里？
有一种解释是，每个bank的层数并不是固定的32，是硬件决定的。所以多出来那部分数据在第33层？

如果这个说法正确的话，那么书上可能要明确写出bank层数不是32，及其层数计算方法。就我自己而言，看到图8.1的时候，就认为bank只会是32层。

希望采纳～

书中34页第四段，“此时，在主机中调用该核函时所用的grid_size为”
这里是不是应该是核函数？

08-shared-memory章节中的reduce2gpu.cu中的 50、75、100行中的__syncthreads();可以放在循环外，可以减少运行时间。

您好，最近看了您的作品收获很多，并使用了您书中提到的一些技巧。在使用过程中我遇到了如下问题。
我在PyTorch框架中编写了一些自定义CUDA算子，其中使用了error.cuh中的CHECK宏进行返回值检测，有如下几个现象。

  int *pts_mask = NULL;
  cudaMalloc(&pts_mask, boxes_num * pts_num * sizeof(int));  // (N, M)
  CHECK(cudaMemset(pts_mask, -1, boxes_num * pts_num * sizeof(int)));

在程序运行过程中（已经稳定运行了很久，cudaMemset函数已经被执行了很多次），被CHECK的这行代码报如下错误：

CUDA Error:
    File:       /torch_custom_ext/torchex/src/sparse_roi_voxelization/sparse_roiaware_pool3d_kernel.cu
    Line:       264
    Error code: 1
    Error text: invalid argument

在某些时候（随机发生），错误会发生在其他自定义算子中，并且提示out of memory错误。

如果将代码改为

  int *pts_mask = NULL;
  CHECK(cudaMalloc(&pts_mask, boxes_num * pts_num * sizeof(int)));  // (N, M)
  cudaMemset(pts_mask, -1, boxes_num * pts_num * sizeof(int));

错误码会变成: out of memory

根据以上现象我初步判断真正的错误原因就是显存不足（我的运行模型确实也比较大），我的问题是为什么在某些情况下会报出上面提到的invalid argument的错误。是否是因为分配内存失败过后将NULL指针传给了cudaMenset?

樊老师，您好！我已经阅读完了您的书籍，现在在使用CUDA开发一种高性能的算法，能比NV的某官方库更快。不知道您是否有类似于读者群之类的HPC交流群？我是某985大学的硕士，应该也可以为社群做出贡献。多谢！

樊老师好，读完您的著作，对cuda编程有了很多新的认识。读书过程中发现了一处问题，P108提到”如果想要在循环内去掉对线程号的约束，又要避免出现读-写竞争，可以将相关代码改写如下：

real v = 0；
for (int offset = 16; offset > 0; offset >>= 1)
{
      v += s_y[tid + offset];
      __syncwarp();
      s_y[tid] = v;
      __syncwarp();
}

”
问题在于v的初值赋为0是有问题的，比如tid=0的线程，第一次迭代后其值变为第s_y[16]，而不是s_y[0]+s_y[16]。
根据我的理解以及与同事的讨论，该处代码改成如下：

real v = s_y[tid]；
for (int offset = 16; offset > 0; offset >>= 1)
{
      v += s_y[tid + offset];
      __syncwarp();
      s_y[tid] = v;
      __syncwarp();
}

才能得到正确结果。
为了便于理解，改成如下更妥：

for (int offset = 16; offset > 0; offset >>= 1)
{
      v = s_y[tid + offset];
      __syncwarp();
      s_y[tid] += v;
      __syncwarp();
}

不知我理解是否有问题，请不吝赐教。

此处的FMA只提到了 d=a*b+c，具体是怎么算的？因为这种公式可能可以引用到论文中，或许下一版可以增加一点介绍~比如专门增加一章讲解，CUDA程序性能衡量指标，以及对应的公式，工具。
书写的很好，感谢作者！

为何对于去掉if的文件进行编译后，用cuda-memcheck工具检测错误会死循环，一直无输出

Hi bruce,
很期待看到您的新书,不知有没有前几章的电子预览版可以以供尝鲜呢?
谢谢!

您好，请教几个问题。
我的显卡是2080Ti

当我读到，109页10.3 更多线程束内的基本函数中的Listing 10.2（如下）

#include "error.cuh"
#include <stdio.h>

const unsigned WIDTH = 8;
const unsigned BLOCK_SIZE = 16;
const unsigned FULL_MASK = 0xffffffff;

void __global__ test_warp_primitives(void);

int main(int argc, char **argv)
{
    test_warp_primitives<<<1, BLOCK_SIZE>>>();
    CHECK(cudaDeviceSynchronize());
    return 0;
}

void __global__ test_warp_primitives(void)
{
    int tid = threadIdx.x;
    int lane_id = tid % WIDTH;

    if (tid == 0) printf("threadIdx.x: ");
    printf("%2d ", tid);
    if (tid == 0) printf("\n");

    if (tid == 0) printf("lane_id:     ");
    printf("%2d ", lane_id);
    if (tid == 0) printf("\n");

    unsigned mask1 = __ballot_sync(FULL_MASK, tid > 0);
    unsigned mask2 = __ballot_sync(FULL_MASK, tid == 0);
    if (tid == 0) printf("FULL_MASK = %x\n", FULL_MASK);
    if (tid == 1) printf("mask1     = %x\n", mask1);
    if (tid == 0) printf("mask2     = %x\n", mask2);

    int result = __all_sync(FULL_MASK, tid);
    if (tid == 0) printf("all_sync (FULL_MASK): %d\n", result);

    result = __all_sync(mask1, tid);
    if (tid == 1) printf("all_sync     (mask1): %d\n", result);

    result = __any_sync(FULL_MASK, tid);
    if (tid == 0) printf("any_sync (FULL_MASK): %d\n", result);

    result = __any_sync(mask2, tid);
    if (tid == 0) printf("any_sync     (mask2): %d\n", result);

    int value = __shfl_sync(FULL_MASK, tid, 2, WIDTH);
    if (tid == 0) printf("shfl:      ");
    printf("%2d ", value);
    if (tid == 0) printf("\n");

    value = __shfl_up_sync(FULL_MASK, tid, 1, WIDTH);
    if (tid == 0) printf("shfl_up:   ");
    printf("%2d ", value);
    if (tid == 0) printf("\n");

    value = __shfl_down_sync(FULL_MASK, tid, 1, WIDTH);
    if (tid == 0) printf("shfl_down: ");
    printf("%2d ", value);
    if (tid == 0) printf("\n");

    value = __shfl_xor_sync(FULL_MASK, tid, 1, WIDTH);
    if (tid == 0) printf("shfl_xor:  ");
    printf("%2d ", value);
    if (tid == 0) printf("\n");
}

输出是

threadIdx.x:  0  1  2  3  4  5  6  7  8  9 10 11 12 13 14 15 
lane_id:      0  1  2  3  4  5  6  7  0  1  2  3  4  5  6  7 
FULL_MASK = ffffffff
mask1     = fffe
mask2     = 1
all_sync (FULL_MASK): 0
all_sync     (mask1): 1
any_sync (FULL_MASK): 1
any_sync     (mask2): 0
shfl:       2  2  2  2  2  2  2  2 10 10 10 10 10 10 10 10 
shfl_up:    0  0  1  2  3  4  5  6  8  8  9 10 11 12 13 14 
shfl_down:  1  2  3  4  5  6  7  7  9 10 11 12 13 14 15 15 
shfl_xor:   1  0  3  2  5  4  7  6  9  8 11 10 13 12 15 14 

其中的

if (tid == 0) printf("threadIdx.x: ");
printf("%2d ", tid);

想请问：

1. 对于线程束中的不同线程，为什么第二行一定在第一行之后执行？
2. 对于线程束中的不同线程，为什么会顺序执行第二行？
   我尝试加入如下代码

    int k = 0;
    if (tid == 0) {
        for ( int i = 0; i < 100000000; i++ ) {
            k = k + 1;
        }
    }
    if (tid == 0) printf("threadIdx.x: ");
    printf("%2d ", tid);
    if (tid == 0) printf("\n");

   但是仍然如下输出，我期待线程0应该是最晚执行完成的。

    threadIdx.x:  0  1  2  3  4  5  6  7  8  9 10 11 12 13 14 15 
   

书中106页讲从伏特架构开始，引入了独立线程调度（independent thread scheduling）机制。 每个线程有自己的程序计数器。我的理解是线程之间的执行顺序应该是随机的。这也正是线程束内同步函数 ＿syncwarp()存在的意义。

在 08-shared-memory/reduce2gpu.cu文件中有

extern __shared__ real s_y[];

请问这里为什么可以使用extern呢？该数组是在哪里定义的呢？