我们有一个长度为100,000的整型数组，需要用10个GPU线程并行地对这个数组进行自增计算，每个数组元素需要自增100次，可知每个线程负责10,000个数组元素。下面有两种方式。

#define ARRAY_SIZE 100000
#define THREAD_SIZE 10

方式1：Kernel函数内循环

每个GPU线程计算出数组元素的起始地址和尾地址，依次对每个元素进行100次自增操作。

Kernel函数：

__global__ void kernel_1(int* array, int len, int threadCount)
{
    int tid = threadIdx.x;
    int start = tid * (len + threadCount-1) / threadCount;
    int end = (tid+1) * (len + threadCount-1) / threadCount < len ? (tid+1) * (len + threadCount-1) / threadCount: len;
    for(int i = start; i<end; i++)
    {
        for(int j=0; j<100; j++)
        {
            array[i]++;
        }
    }
}

外部调用：

kernel_1<<< 1, THREAD_SIZE >>> (d_array, ARRAY_SIZE, THREAD_SIZE);

方式2：外部循环

GPU线程每次只对一个数组元素进行100次自增操作，数组元素地址由外部给定。外部使用循环体多次调用Kernel函数，完成所有数组元素的自增任务。

Kernel函数：

__global__ void kernel_2(int* array, int firstPos)
{
    int tid = threadIdx.x;
    int pos = firstPos + tid;
    for(int j=0; j<100; j++)
    {
        array[pos]++;
    }
}

外部调用：

for(int firstPos=0; firstPos<ARRAY_SIZE; firstPos+=THREAD_SIZE)
    {
        kernel_2<<< 1, THREAD_SIZE >>> (d_array, firstPos);
    }

我们统计了两种方式的执行时间，实验结果如下：

方式1：

方式2：

可见两种方式的执行时间接近，方式1略优于方式2。这是因为方式2中GPU需要多次地与CPU进行中断交互。事实证明for循环可以正常地应用于GPU的Kernel函数中。