1.关于volatile：

对于文章中这个函数，

__global__ void reduceUnrollWarps8 (int *g_idata, int *g_odata, unsigned int n)
{// set thread IDunsigned int tid = threadIdx.x;unsigned int idx = blockIdx.x * blockDim.x * 8 + threadIdx.x;// convert global data pointer to the local pointer of this blockint *idata = g_idata + blockIdx.x * blockDim.x * 8;// unrolling 8if (idx + 7 * blockDim.x < n){int a1 = g_idata[idx];int a2 = g_idata[idx + blockDim.x];int a3 = g_idata[idx + 2 * blockDim.x];int a4 = g_idata[idx + 3 * blockDim.x];int b1 = g_idata[idx + 4 * blockDim.x];int b2 = g_idata[idx + 5 * blockDim.x];int b3 = g_idata[idx + 6 * blockDim.x];int b4 = g_idata[idx + 7 * blockDim.x];g_idata[idx] = a1 + a2 + a3 + a4 + b1 + b2 + b3 + b4;}__syncthreads();// in-place reduction in global memoryfor (int stride = blockDim.x / 2; stride > 32; stride >>= 1){if (tid < stride){idata[tid] += idata[tid + stride];}// synchronize within threadblock__syncthreads();}// unrolling warpif (tid < 32){volatile int *vmem = idata;vmem[tid] += vmem[tid + 32];vmem[tid] += vmem[tid + 16];vmem[tid] += vmem[tid +  8];vmem[tid] += vmem[tid +  4];vmem[tid] += vmem[tid +  2];vmem[tid] += vmem[tid +  1];}// write result for this block to global memif (tid == 0) g_odata[blockIdx.x] = idata[0];
}

在  //unrolling wrap 之前的for循环完成后，只剩下idata[0] 到 idata[63]（共64个元素）的数据还没有进行加和。

之前的疑问是 最后的if语句块内为什么会完成这个操作而不出问题：

因为一个线程束内的线程会进行同步，所以if里的语句是，所有线程执行第一条语句，都执行完了，再执行第二条语句。

然后关于volitale ，大家都说是禁止编译器优化，我的理解是：

如果没有volatile，各个线程写寄存器的时候是有同步的，从寄存器写到共享内存是没有同步的。(不过在我这里，去掉volatile结果也是对的。。。)

2.关于延迟隐藏

目前个人理解就是让一个线程尽可能多做事。

3.关于如何获得最大线程束 

参考：【精选】CUDA编程：笔记2_线程束_longlongqin的博客-CSDN博客

little法则给出了下面的计算公式"所需线程束 = 延迟 × 吞吐量 

带宽：一般指的是理论峰值，最大每个时钟周期能执行多少个指令；

吞吐量：是指实际操作过程中每分钟处理多少个指令。

每次想到还是看书吧。 

一个WRAP访存周期是T，刚好是Wrap处理指令时间周期t的K倍，只需要K个wrap就能隐藏访存带来的延迟。

得到最大线程束数量还要减少分支

4.选择合适的网格和块：

有个工具包（同学说不能使了，但是应该也提供了别的方法）。

4.什么是事务内存

参考：初识事务内存（Transactional Memory） - 知乎

个人理解：比u锁更好的具有原子性的互斥手段。 

5.在GPU上只有内存加载操作可以被缓存，内存存储操作不能被缓存。

内存加载操作：从全局内存或其他内存读数据加载到寄存器或缓存中。

内存存储操作：将数据从寄存器或缓存写入到内存

线程加载数据时，如果缓存中有数据，可以直接从缓存中加载。而存储操作通常不能被缓存，因为要保证数据的一致性，必须及时更新内存中的数据，而不能依赖于缓存。

6.对于固定内存和零拷贝内存

当传输大数据时，使用固定内存：cudaError_t cudaMallocHost(void **devPtr, size_t count);

cudaError_t cudaFreeHost(void *ptr);

如果想要共享主机和设备端的少量数据，使用零拷贝内存：

cudaErroer_t cudaHostAlloc(void **pHost, size_t count, unsigned int flags);

cudaFreeHost()

使用下列函数获取映射到固定内存的设备指针：

cudaError_t cudaHostGetDevicePointer(void **pDevice, void *pHost, unsigned int flags);

    // part 2: using zerocopy memory for array A and B// allocate zerocpy memoryCHECK(cudaHostAlloc((void **)&h_A, nBytes, cudaHostAllocMapped));CHECK(cudaHostAlloc((void **)&h_B, nBytes, cudaHostAllocMapped));// initialize data at host sideinitialData(h_A, nElem);initialData(h_B, nElem);memset(hostRef, 0, nBytes);memset(gpuRef,  0, nBytes);// pass the pointer to deviceCHECK(cudaHostGetDevicePointer((void **)&d_A, (void *)h_A, 0));CHECK(cudaHostGetDevicePointer((void **)&d_B, (void *)h_B, 0));// add at host side for result checkssumArraysOnHost(h_A, h_B, hostRef, nElem);// execute kernel with zero copy memorysumArraysZeroCopy<<<grid, block>>>(d_A, d_B, d_C, nElem);// copy kernel result back to host sideCHECK(cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost));// check device resultscheckResult(hostRef, gpuRef, nElem);// free  memoryCHECK(cudaFree(d_C));CHECK(cudaFreeHost(h_A));CHECK(cudaFreeHost(h_B));

有了虚拟内存寻址（UVA）后，由cudaHostAlloc()分配的固定主机内存具有相同的主机和设备指针，可以将返回的指针直接传递给核函数：

cudaHostAlloc((void **)&h_A, nBytes, cudaHostAllocMapped);
cudaHostAlloc((void **)&h_B, nBytes, cudaHostAllocMapped);initialData(h_A, nElem);
initialData(h_B, nElem);sumArrayZeroCopy<<<grid, block>>>(h_A, h_B, d_C, nElem);