c++ - NVIDIA NVCC 在使用模板特征类型时更改编译时间常量

Question

在使用 C++ 模板时，我发现 NVIDIA NVCC(经过 CUDA 4.0 和 4.1 测试)有一个奇怪的行为。我将其简化为一个演示行为的简单示例。

这已经处于错误报告状态。然而，我把它卡在这里，因为这个网站是一个越来越可靠的错误和修复来源。因此，我会不断更新此页面。

代码:

#include"stdio.h"

#define PETE_DEVICE __device__

template<class T, int N>  class ILattice;
template<class T>         class IScalar;
template<class T, int IL> struct AddILattice {};

template<class T>
PETE_DEVICE
void printType() {
  printf("%s\n",__PRETTY_FUNCTION__);
}

template<class T> class IScalar {
  T F;
};

template<class T, int N> class ILattice {
  T F[N];
};

template<class T, int N>
struct AddILattice<IScalar<T> , N> {
  typedef ILattice< T , N > Type_t;
};

#define IL 16

__global__ void kernel()
{
  printf("IL=%d\n",IL);  // Here IL==16

  typedef typename AddILattice<IScalar<float> ,IL>::Type_t Tnew;

  // This still works fine. Output:
  // void printType() [with T = ILattice<float, 16>]
  //
  printType<Tnew>();

  // Now problems begin: Output:
  // T=4 Tnew=0 IL=64
  // Here IL should still be 16
  // sizeof(Tnew) should be 16*sizeof(float)
  //
  printf("T=%d Tnew=%d IL=%d\n",sizeof(IScalar<float> ),sizeof(Tnew),IL);   
}   

int main()
{
    dim3  blocksPerGrid( 1 , 1 , 1 );
    dim3  threadsPerBlock( 1 , 1, 1);
    kernel<<< blocksPerGrid , threadsPerBlock , 48*1024 >>>( );

    cudaDeviceSynchronize();
    cudaError_t kernel_call = cudaGetLastError();
    printf("call: %s\n",cudaGetErrorString(kernel_call));

}

编译器将 IL 从 16 更改为 64 的任何想法？？

Answer 1

可能是因为你使用了错误的 printf 转换。 %d表示输出一个int，但是sizeof返回的不是int，而是一个size_t。另外使用 size_t 长度修饰符(并使其无符号)，即将 %d 替换为 %zu。

printf 无法知道(由于 var-args 列表)真正传递的类型是什么，因此不会发生类型转换，它只能知道格式字符串的类型。所以你必须在那里传递正确的参数。当您在 size_t 与 int 大小相同的系统上时，您的代码可以工作(例如，许多 32 位系统)。但您不能依赖这个事实，使用正确的转换将帮助您实现这一目标。

(所以不是编译器改变了你的常量，而是你输出错了)