关于理解CUDA朴素前缀和（Naive Prefix Sum）代码的技术问询

我最近在研究NVIDIA《GPU Gems 3》第39章的示例代码，其中一段实现朴素前缀和（Naive Prefix Sum）的CUDA Kernel我完全摸不着头脑，原代码片段如下：

__global__ void scan(float *g_odata, float *g_idata, int n) {
    extern __shared__ float temp[]; // allocated on invocation
    int thid = threadIdx.x;
    int pout = 0, pin = 1;
    // Load input into shared memory.
    // This is exclusive scan, so shift right by one
    // and set first element to 0
    temp[pout...

原代码没贴完整，但仅前面这部分就有好几个地方让我困惑：

• 为什么要用extern __shared__ float temp[]这种方式声明共享内存？这种共享内存在Kernel调用时是怎么确定大小的？
• pout和pin这两个变量的作用是什么？为什么初始化要设成0和1？
• 注释里提到这是exclusive scan（排他前缀和），为什么要把输入右移一位，还把第一个元素设为0？这和普通的前缀和计算逻辑有什么区别？

有没有大佬能帮我拆解这段代码的核心逻辑，尤其是共享内存的使用方式和排他前缀和的实现思路？

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