【发布时间】：2013-09-25 18:42:50
【问题描述】：
我正在尝试为图像上的逻辑操作定义模板 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编译。
任何帮助将不胜感激。
【问题讨论】：
-
a << int(b)之类的东西有帮助吗？ -
@Eric... 不。仍然没有编译。
-
非模板版本的 SHR 可以工作吗？或 SHR 的模板特化而不是
switch。 -
能否请您提供编译问题的完整重现者？当我尝试创建自己的时，如 here 所示，它编译正确。是的，我意识到我使用的是 linux 而不是 windows。由于问题似乎出在设备编译器方面，我不确定 linux 与 windows 是否会有所不同。如果我发现你所说的复制器在 linux 下编译干净，我会切换到 windows。
-
@sgar91：所有支持 CUDA 的设备都支持位移，并且有硬件指令可以做到这一点。您可能遇到了编译器错误。 sm_1x 平台使用 Open64 作为编译器前端，而后期平台的编译器前端是 NVVM（基于 LLVM），因此观察到不同的行为。我建议提交一份错误报告，并附上一个独立的复制案例。错误报告表单链接自注册的开发者网站。