小编Far*_*zad的帖子

在参考文献的位置方面,L2缓存如何在具有Kepler架构的GPU中工作？例如,如果一个线程访问全局内存中的地址,假设该地址的值不在L2缓存中,那么缓存的值是多少？是暂时的吗？或者该附近的其他附近值是否也被带到L2缓存(空间)？

下图来自NVIDIA白皮书.

当我在SO上遇到这个问题时,我很想知道答案.所以我在下面写了一段代码来测试不同场景下的原子操作性能.操作系统是带有CUDA 5.5的Ubuntu 12.04,设备是GeForce GTX780(开普勒架构).我使用-O3flag 编译代码,CC = 3.5.

#include <stdio.h>

static void HandleError( cudaError_t err, const char *file, int line ) {
    if (err != cudaSuccess) {
        printf( "%s in %s at line %d\n", cudaGetErrorString( err ), file, line );
        exit( EXIT_FAILURE );
    }
}
#define HANDLE_ERROR( err ) (HandleError( err, __FILE__, __LINE__ ))

#define BLOCK_SIZE 256
#define RESTRICTION_SIZE 32

__global__ void CoalescedAtomicOnGlobalMem(int* data, int nElem)
{
    unsigned int tid = (blockIdx.x * blockDim.x) + threadIdx.x;
    for …

关于何时使用内联函数以及何时在常规C编码中避免它,有很多建议.__forceinline__对CUDA C __device__功能有什么影响？他们应该在哪里使用,哪里可以避免？

当提供的地址位于块共享内存中时,原子操作如何执行？在原子操作期间,是否暂停块内其他线程对同一共享内存库的访问,或者阻止其他线程执行任何指令,甚至阻止所有块中的线程直到原子操作完成？

我有一段CUDA代码，其中线程在共享内存上执行原子操作。我一直在思考，因为原子操作的结果无论如何对于该块的其他线程都是立即可见的，因此最好指示编译器具有共享内存volatile。
所以我改变了

__global__ void CoalescedAtomicOnSharedMem(int* data, uint nElem)
{
    __shared__ int smem_data[BLOCK_SIZE];
    uint tid = (blockIdx.x * blockDim.x) + threadIdx.x;
    for ( uint i = tid; i < nElem; i += blockDim.x*gridDim.x){
        atomicAdd( smem_data+threadIdx.x, 6);
    }
}

至

__global__ void volShared_CoalescedAtomicOnSharedMem(int* data, uint nElem)
{
    volatile __shared__ int smem_data[BLOCK_SIZE];
    uint tid = (blockIdx.x * blockDim.x) + threadIdx.x;
    for ( uint i = tid; i < nElem; i += blockDim.x*gridDim.x){
        atomicAdd( smem_data+threadIdx.x, 6);
    }
}

发生以下更改时发生以下编译时错误：

error: …

对于CUDA中简单的warp内线程发散，我所知道的是SM选择一个重新收敛点（PC地址），并在两个/多个路径中执行指令，同时禁用未采取该路径的线程的执行效果。
例如，在下面的代码中：

if( threadIdx.x < 16 ) {
    A:
    // do something.
} else {
    B:
    // do something else.
}
C:
// rest of code.

C是重新收敛点，warp 调度程序在 和 处调度指令A，B同时禁用A上半 warp 处的指令，并禁用B下半 warp 处的指令。当它达到 时C，将为经线内的所有线程启用指令。

我的问题是 SM 是否能够goto像上面那样正确处理包括指令在内的代码？或者不能保证所选择的重收敛点是最佳的？
例如，如果我的 CUDA 代码中有以下控制流，使用goto

A:
// some code here.
B:
// some code here too.
if( threadIdx.x < 16 ) {
    C:
    // do something.
    goto A;
}
// …

在CUDA设备中,在全局内存写入中合并与在全局内存读取中合并一样重要吗？如果是,怎么解释？早期的CUDA设备和最近的CUDA设备之间是否存在差异？

在CUDA 编程指南中指出，映射固定主机内存上的原子操作“从主机或其他设备的角度来看不是原子的”。我从这句话中得到的是，如果主机内存区域仅由一个 GPU 访问，则可以在映射的固定主机内存上执行原子操作（即使来自多个并发内核）。

另一方面，在尼古拉斯·威尔特( Nicholas Wilt)的《CUDA 手册》一书中，第 128 页指出：