NVCC-不同的块大小，具体取决于编译时间的拱形

Question

我有一个内核，表现出最高的性能 不同的 在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= 旗帜。

Answer 1

nvcc 编译器未检测到本地可用的GPU，默认情况下始终针对SM 1.0。否则，在建立不同系统时，它可能会引入一些令人困惑的行为。

要编译可用设备，您要么需要要求用户指定SM版本，要么在构建时间内运行一些检测代码。我坚信将硬件检测代码放入运行时更容易，然后根据需要配置内核启动。