cuda学习2-block与thread数量的选取

　　由上一节可知，在main函数中，cuda程序的并行能力是在add<<<N,1>>>( dev_a, dev_b, dev_c )函数中体现的，这里面设置的是由N个block的构成的计算网络即grid，每一个block里面有1个thread存在。那么这种选取有什么用意呢，如何针对自己的计算问题设置计算网络呢？

　　首先要说明这两个数的选取没有固定的方法，完全是根据自身需求。其实它的完整形式是Kernel<<<Dg,Db, Ns, S>>>(param list);<<<>>>运算符内是核函数的执行参数，告诉编译器运行时如何启动核函数，用于说明内核函数中的线程数量，以及线程是如何组织的。

　　参数Dg用于定义整个grid的维度和尺寸，即一个grid有多少个block。为dim3类型。Dim3 Dg(Dg.x, Dg.y, 1)表示grid中每行有Dg.x个block，每列有Dg.y个block，第三维恒为1。整个grid中共有Dg.x*Dg.y个block，其中Dg.x和Dg.y最大值为65535。

　　参数Db用于定义一个block的维度和尺寸，即一个block有多少个thread。为dim3类型。Dim3 Db(Db.x, Db.y, Db.z)表示整个block中每行有Db.x个thread，每列有Db.y个thread，高度为Db.z。Db.x和Db.y最大值为512，Db.z最大值为62。 一个block中共有Db.x*Db.y*Db.z个thread。计算能力为1.0,1.1的硬件该乘积的最大值为768，计算能力为1.2,1.3的硬件支持的最大值为1024。

　　参数Ns是一个可选参数，用于设置每个block除了静态分配的shared Memory（以后会学习到）以外，最多能动态分配的shared memory大小，单位为byte。不需要动态分配时该值为0或省略不写。

　　参数S是一个cudaStream_t类型的可选参数，初始值为零，表示该核函数处在哪个流（以后会学习到）之中。

　　在这个例子中，由于计算很简单，就选了一个<<<N,1>>>这种搭配。现在我们看一个复杂一点的例子。

　　这个例子是说要计算两个任意长的向量的加法，可能会比比65535长，超过了block数的最大范围，甚至于比65535×512（thread上限）还长，应该怎么办呢？下面就用

<<<128,128>>>的计算网络来搞定。

　　核函数改为如下：

__global__ void add( int *a, int *b, int *c ) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    while (tid < N) {
        c[tid] = a[tid] + b[tid];
        tid += blockDim.x * gridDim.x;
    }
}

　　这段代码的精髓就在于它是一个循环，当编号为tid = threadIdx.x + blockIdx.x * blockDim.x的线程进行加法运算之后，tid += blockDim.x * gridDim.x;如果tid<N，则这个线程再做一次加法，依次循环下去。因为计算网络只有blockDim.x * gridDim.x这么大（次例为128×128），那么那些大于blockDim.x * gridDim.x并且小于N的数组分量的相加任务就需要继续分配给各个线程，如上就是用循环来分配的。

　　任意长度向量相加完整代码：

/*
 * Copyright 1993-2010 NVIDIA Corporation.  All rights reserved.
 *
 * NVIDIA Corporation and its licensors retain all intellectual property and 
 * proprietary rights in and to this software and related documentation. 
 * Any use, reproduction, disclosure, or distribution of this software 
 * and related documentation without an express license agreement from
 * NVIDIA Corporation is strictly prohibited.
 *
 * Please refer to the applicable NVIDIA end user license agreement (EULA) 
 * associated with this source code for terms and conditions that govern 
 * your use of this NVIDIA software.
 * 
 */


#include "../common/book.h"

#define N   (33 * 1024)

__global__ void add( int *a, int *b, int *c ) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    while (tid < N) {
        c[tid] = a[tid] + b[tid];
        tid += blockDim.x * gridDim.x;
    }
}

int main( void ) {
    int *a, *b, *c;
    int *dev_a, *dev_b, *dev_c;

    // allocate the memory on the CPU
    a = (int*)malloc( N * sizeof(int) );
    b = (int*)malloc( N * sizeof(int) );
    c = (int*)malloc( N * sizeof(int) );

    // allocate the memory on the GPU
    HANDLE_ERROR( cudaMalloc( (void**)&dev_a, N * sizeof(int) ) );
    HANDLE_ERROR( cudaMalloc( (void**)&dev_b, N * sizeof(int) ) );
    HANDLE_ERROR( cudaMalloc( (void**)&dev_c, N * sizeof(int) ) );

    // fill the arrays 'a' and 'b' on the CPU
    for (int i=0; i<N; i++) {
        a[i] = i;
        b[i] = 2 * i;
    }

    // copy the arrays 'a' and 'b' to the GPU
    HANDLE_ERROR( cudaMemcpy( dev_a, a, N * sizeof(int),
                              cudaMemcpyHostToDevice ) );
    HANDLE_ERROR( cudaMemcpy( dev_b, b, N * sizeof(int),
                              cudaMemcpyHostToDevice ) );

    add<<<128,128>>>( dev_a, dev_b, dev_c );

    // copy the array 'c' back from the GPU to the CPU
    HANDLE_ERROR( cudaMemcpy( c, dev_c, N * sizeof(int),
                              cudaMemcpyDeviceToHost ) );

    // verify that the GPU did the work we requested
    bool success = true;
    for (int i=0; i<N; i++) {
        if ((a[i] + b[i]) != c[i]) {
            printf( "Error:  %d + %d != %d\n", a[i], b[i], c[i] );
            success = false;
        }
    }
    if (success)    printf( "We did it!\n" );

    // free the memory we allocated on the GPU
    HANDLE_ERROR( cudaFree( dev_a ) );
    HANDLE_ERROR( cudaFree( dev_b ) );
    HANDLE_ERROR( cudaFree( dev_c ) );

    // free the memory we allocated on the CPU
    free( a );
    free( b );
    free( c );

    return 0;
}

总结：我们通常选取一定数量的线程来解决问题，通常都选2的倍数。是由grid，block，thread，这种三级结构实现的。一般的程序的计算量都会超过线程数量，因此要合理的把计算量尽量平均分配给各个线程来计算。感觉上来说，编写核函数的精髓就是如何利用线程的序号（索引值）来分配计算任务。