nvcc
编译器未检测到本地可用的GPU,默认情况下始终针对SM 1.0。否则,在建立不同系统时,它可能会引入一些令人困惑的行为。
要编译可用设备,您要么需要要求用户指定SM版本,要么在构建时间内运行一些检测代码。我坚信将硬件检测代码放入运行时更容易,然后根据需要配置内核启动。
题
我有一个内核,表现出最高的性能 不同的 在Kepler和Fermi硬件上运行时的块大小。我希望在编译时检查当前的体系结构目标并定义一个 THREADS_PER_BLOCK
宏到i)启动内核; ii)确定必要的块数; iii)在内核中静态设置共享内存大小。
以下证明了我正在尝试做什么。假设我针对GK104硬件,因此使用 nvcc -arch=sm_30
. 。这仍然会导致 THREADS_PER_BLOCK = 256
自从 __CUDA_ARCH__
未针对主机代码编译定义。 (我了解,例如 这个答案, ,为什么它不能这样工作。)
#if __CUDA_ARCH__ >= 300
#define THREADS_PER_BLOCK 512
#else
#define THREADS_PER_BLOCK 256
#endif
__global__ void some_kernel(int* a, int* b) {
__shared__ sm_data[THREADS_PER_BLOCK];
// Do something.
}
int main(void) {
// Initialize data.
// Calculate blocks based on THREADS_PER_BLOCK, problem size and some max.
some_kernel<<blocks, THREADS_PER_BLOCK>>>(d_a, d_b)
return 0;
}
我可以在运行时检查设备属性并使用动态共享内存,但是想知道是否可以在编译时间进行硬编码,而无需手动添加-dfermi或-dkepler和设置 THREADS_PER_BLOCK
基于此。 NB:此代码的任何用户都将自己编译, 几乎 当然,对于一个架构,这不是一个不合理的选择。鉴于通过 -arch=
旗帜。
解决方案
nvcc
编译器未检测到本地可用的GPU,默认情况下始终针对SM 1.0。否则,在建立不同系统时,它可能会引入一些令人困惑的行为。
要编译可用设备,您要么需要要求用户指定SM版本,要么在构建时间内运行一些检测代码。我坚信将硬件检测代码放入运行时更容易,然后根据需要配置内核启动。