CUDA学习（三十六）

内置的矢量类型：
char, short, int, long, longlong, float, double：
这些是从基本整数和浮点类型派生的矢量类型。 它们是结构，第一，第二，第三和第四个组件分别可以通过字段x，y，z和w访问。 它们都带有一个构造函数，形式为make_ <type name>; 例如:

int2 make_int2(int x, int y);

它创建一个类型为int2的向量，其值为（x，y）
表3列出了矢量类型的对齐要求:


dim3:
此类型是基于uint3的整数矢量类型，用于指定尺寸。 当定义类型为dim3的变量时，任何未指定的组件都被初始化为1。
内置变量:
内置变量指定网格和块的尺寸以及块和线索的索引。 它们仅在设备上执行的功能内有效;
gridDim:
此变量的类型为dim3（请参阅dim3）并包含网格(grid)的尺寸
blockIdx:
这个变量的类型是uint3（参见char，short，int，long，longlong，float，double），并在网格中包含块索引。
blockDim:
该变量的类型为dim3（请参阅dim3）并包含块的尺寸
threadIdx:
这个变量的类型为uint3（参见char，short，int，long，longlong，float，double）并包含块内的线程索引
warpSize:
此变量的类型为int，并且包含线程中的warp大小（有关warp的定义，请参见SIMT体系结构）。
记忆栅栏功能:
CUDA编程模型假设一个设备的内存模型是弱顺序的，即CUDA线程将数据写入共享内存，全局内存，页面锁定主机内存或对等设备内存的顺序不一定是 观察数据的顺序是由另一个CUDA或主机线程写入的。
例如，如果线程1执行writeXY（），并且线程2执行readXY（），如以下代码示例中所定义:

__device__ volatile int X = 1, Y = 2;
__device__ void writeXY()
{
    X = 10;
    Y = 20;
}
__device__ void readXY()
{
    int A = X;
    int B = Y;
}

有可能B最终等于20，A对于线程2等于1.在一个强排序的内存模型中，唯一的可能性是：

• A equal to 1 and B equal to 2,
• A equal to 10 and B equal to 2,
• A equal to 10 and B equal to 20

内存围栏功能可以用来执行内存访问的一些排序。 内存栅栏功能在执行顺序的范围上有所不同，但它们独立于所访问的内存空间（共享内存，全局内存，页锁内存和对等设备的内存）。
void __threadfence_block();
确保:

• 在对__threadfence_block（）的调用之前，调用线程的所有线程都会观察到所有写入由调用线程所做的所有内存的全部内存，而不是在调用__threadfence_block（）之后全部写入到调用线程所做的所有内存之前。
• 在对__threadfence_block（）的调用之前，调用线程所做的所有内存读取都是在调用__threadfence_block（）之后的调用线程所有内存读取之前进行的。
  void __threadfence();

对调用线程的块中的所有线程都起着__threadfence_block（）的作用，并且还确保在调用__threadfence（）之后调用线程所做的所有内存都不会被设备中的任何线程观察到， 在调用__threadfence（）之前调用线程所做的所有内存。 请注意，为了确保此顺序保证为真，观察线程必须真正观察内存，而不是其缓存版本; 这通过使用挥发性限定符中详细说明的volatile关键字来保证。

void __threadfence_system(); 

对调用线程的块中的所有线程都起着__threadfence_block（）的作用，并确保在调用__threadfence_system（）之前由设备中的所有线程，主线程和所有线程 在调用__threadfence_system（）之后，在对等设备中的所有线程写入由调用线程创建的所有内存之前，
__threadfence_system（）仅受计算能力2.x和更高的设备的支持
在前面的代码示例中，在X = 10之间插入栅栏函数调用; 和Y = 20; 和int A = X之间; int B = Y; 将确保对于线程2，如果B等于20，则A总是等于10.如果线程1和2属于同一个块，则使用__threadfence_block（）就足够了。 如果线程1和2不属于同一个块，则必须使用__threadfence（）（如果它们是来自同一设备的CUDA线程，并且必须使用__threadfence_system（）（如果它们是来自两个不同设备的CUDA线程）。
一个常见的用例是线程消耗其他线程产生的一些数据，如以下内核代码示例所示，该内核计算一个调用中N个数字的和。 每个块首先将数组的一个子集相加并将结果存储在全局内存中。 当所有块完成时，最后完成的块从全局内存中读取这些部分和，并将它们相加以获得最终结果。 为了确定哪个块最后完成，每个块以原子方式递增一个计数器，以通过计算并存储其部分和（参见关于原子函数的原子函数）来表明它已完成。 最后一个块是接收等于gridDim.x-1的计数器值的块。 如果在存储部分总和和递增计数器之间没有放置栅栏，计数器可能会在部分总和存储之前递增，因此可能会达到gridDim.x-1，并且让最后一个块在实际更新之前开始读取部分总和 在记忆中。
内存围栏功能只影响线程对内存操作的排序; 它们不确保这些内存操作对其他线程可见（如__syncthreads（）对块内的线程所做的操作（请参阅同步函数））。 在下面的代码示例中，通过将结果变量声明为volatile来确保内存操作对结果变量的可见性（请参阅“易失性限定符”）。

__device__ unsigned int count = 0;
__shared__ bool isLastBlockDone;
__global__ void sum(const float* array, unsigned int N,
    volatile float* result)
{
    // Each block sums a subset of the input array.
    float partialSum = calculatePartialSum(array, N);
    if (threadIdx.x == 0) {
        // Thread 0 of each block stores the partial sum
        // to global memory. The compiler will use
        // a store operation that bypasses the L1 cache
        // since the "result" variable is declared as
        // volatile. This ensures that the threads of
        // the last block will read the correct partial
        // sums computed by all other blocks.
        result[blockIdx.x] = partialSum;
        // Thread 0 makes sure that the incrementation
        // of the "count" variable is only performed after
        // the partial sum has been written to global memory.
        __threadfence();
        // Thread 0 signals that it is done.
        unsigned int value = atomicInc(&count, gridDim.x);
        // Thread 0 determines if its block is the last
        // block to be done.
        isLastBlockDone = (value == (gridDim.x - 1));
    }
    // Synchronize to make sure that each thread reads
    // the correct value of isLastBlockDone.
    __syncthreads();
    if (isLastBlockDone) {
        // The last block sums the partial sums
        // stored in result[0 .. gridDim.x-1]
        float totalSum = calculateTotalSum(result);
        if (threadIdx.x == 0) {
            // Thread 0 of last block stores the total sum
            // to global memory and resets the count
            // varialble, so that the next kernel call
            // works properly.
            result[0] = totalSum;
            count = 0;
        }
    }
}