CUDA 内核模板实例化导致编译错误

Question

我正在尝试为图像上的逻辑操作定义一个模板 CUDA 内核。代码如下所示：

#define AND 1
#define OR 2
#define XOR 3
#define SHL  4
#define SHR 5 

template<typename T, int opcode> 
__device__ inline T operation_lb(T a, T b)
{
    switch(opcode)
    {
    case AND:
        return a & b;
    case OR:
        return a | b;
    case XOR:
        return a ^ b;
    case SHL:
        return a << b;
    case SHR:
        return a >> b;
    default:
        return 0;
    }
}

//Logical Operation With A Constant
template<typename T, int channels, int opcode> 
__global__ void kernel_logical_constant(T* src, const T val, T* dst, int width, int height, int pitch)
{
    const int xIndex = blockIdx.x * blockDim.x + threadIdx.x;
    const int yIndex = blockIdx.y * blockDim.y + threadIdx.y;

    if(xIndex >= width || yIndex >= height) return;

    unsigned int tid = yIndex * pitch + (channels * xIndex);

    #pragma unroll
    for(int i=0; i<channels; i++)
        dst[tid + i] = operation_lb<T,opcode>(src[tid + i],val);
}

问题是，当我实例化内核进行位移位时，出现以下编译错误

错误 1 错误：Ptx 组装因错误而中止

内核瞬间是这样的：

template __global__ void kernel_logical_constant<unsigned char,1,SHL>(unsigned char*,unsigned char,unsigned char*,int,int,int);

对于

unsigned char

、

unsigned short

、1 和 3 通道以及所有逻辑运算，还有 19 个这样的时刻。但只有位移实例，即

SHL

和

SHR

才会导致错误。当我删除这些实例时，代码可以完美编译并运行。 如果我用

operation_lb

设备函数内的任何其他操作替换位移位，该代码也可以工作。 我想知道这是否与由于内核的许多不同实例而生成的 ptx 代码量有关。

我使用的是 CUDA 5.5、Visual Studio 2010、Windows 8 x64。编译

compute_1x, sm_1x

。

如有任何帮助，我们将不胜感激。

Answer 1

最初的问题指定发帖者使用的是

compute_20, sm_20

。这样，我无法使用代码here重现错误。然而，评论中指出，实际上正在使用

sm_10

。当我切换到编译

sm_10

时，我能够重现该错误。

它似乎是编译器中的一个错误。我这么说只是因为我不相信编译器应该生成汇编器无法处理的代码。然而除此之外，我不知道根本原因。我已向 NVIDIA 提交了错误报告。

在我有限的测试中，它似乎只发生在

unsigned char

而不是

int

。

作为一种可能的解决方法，对于 cc2.0 及更新版本的设备，请在编译时指定

-arch=sm_20

。