下面一段代码
cudaDeviceProp deviceProp;
cudaGetDeviceProperties(&deviceProp, 0); //assuming current device ID is 0
收集设备的性能进deviceProp
。您可以看到here,在成功拨打cudaGetDeviceProperties
后,您将可以访问具有所需设备属性的deviceProp
成员。例如,deviceProp.maxThreadsPerMultiProcessor
表示每多处理器的最大线程数,deviceProp.maxThreadsPerBlock
表示每块的最大线程数等
每块与块的整体数量的线程适当数量与您通话大多取决于您的设备属性的功能和你的程序。您调用的每个块都占用SM的一部分。多少取决于您的块请求的资源:线程,寄存器和共享内存。
考虑这个例子。假设您的设备SM最多可以有2048个线程,48 KB的共享内存和64 KB的寄存器。如果您的块需要512个线程,并且同时使用SM可用的所有共享内存和寄存器,则不可能在SM中具有另一个具有相同特征的块。因此,如果不能使用2048减512个潜在SM线程,则可以将最大占用率降低到25%。现在,如果您通过将块中的线程数增加到1024来设计块,则可以使用相同数量的寄存器和共享内存,将占用率提高了一倍,达到50%。
通常不推荐使用大量的块。 GPU将新块安排到可用的SM中。如果所有的SM都被占用,它将对该块进行排队,直到SM有足够的空闲资源用于该块。调度新块对GPU有开销(尽管很小)。在找到最佳块大小之后,通过SM计算(或剖析)块占用率,然后调用尽可能多的块,因为它占用了所有GPU SM。如果您需要更多的块,您可以重用已完成作业的块的线程。
例如转换,其中
GPU_kernel<<<1024,512>>>();
__global__ void GPU_kernel(void){
unsigned int tid = (blockIdx.x * blockDim.x) + threadIdx.x;
//rest of code
}
成其中
GPU_kernel<<<(number_of_SMs*number_of_blocks_per_SM),512>>>();
__global__ void GPU_kernel(void){
unsigned int tid = (blockIdx.x * blockDim.x) + threadIdx.x;
for (; tid < 1024*512; tid += blockIdx.x* gridDim.x) {
//rest of code
}
}
通常导致更好的perfor曼斯。
另外请注意,在上述代码片段中,我没有包含适当的CUDA错误检查。请运用您自己的方法来处理可能的错误。说明here。