我们有一个长度为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:
Starting program: /home/leee/test/threadTest
[Thread debugging using libthread_db enabled]
Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1".
[New Thread 0x7fffef7e9700 (LWP 25370)]
[New Thread 0x7fffeefe8700 (LWP 25371)]
Time:369674
[Thread 0x7fffeefe8700 (LWP 25371) exited]
[Thread 0x7fffef7e9700 (LWP 25370) exited]
[Inferior 1 (process 25338) exited with code 01]
方式2:
Starting program: /home/leee/test/threadTest
[Thread debugging using libthread_db enabled]
Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1".
[New Thread 0x7fffef7e9700 (LWP 25195)]
[New Thread 0x7fffeefe8700 (LWP 25196)]
Time:395915
[Thread 0x7fffeefe8700 (LWP 25196) exited]
[Thread 0x7fffef7e9700 (LWP 25195) exited]
[Inferior 1 (process 25164) exited with code 01]
可见两种方式的执行时间接近,方式1略优于方式2。这是因为方式2中GPU需要多次地与CPU进行中断交互。事实证明for循环可以正常地应用于GPU的Kernel函数中。
此外,如果以同步地方法执行方式2,执行时间会更长。
外部调用:
for(int firstPos=0; firstPos<ARRAY_SIZE; firstPos+=THREAD_SIZE)
{
kernel_2<<< 1, THREAD_SIZE >>> (d_array, firstPos);
err = cudaDeviceSynchronize();
}
结果:
Starting program: /home/leee/test/threadTest
[Thread debugging using libthread_db enabled]
Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1".
[New Thread 0x7fffef7e9700 (LWP 27645)]
[New Thread 0x7fffeefe8700 (LWP 27646)]
Time:597882
[Thread 0x7fffeefe8700 (LWP 27646) exited]
[Thread 0x7fffef7e9700 (LWP 27645) exited]
[Inferior 1 (process 27612) exited with code 01]