$ cat t243.cu
#include <stdio.h>

#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);
}

template __global__ void kernel_logical_constant<unsigned char,1,SHL>(unsigned char*,unsigned char,unsigned char*,int,int,int);
template __global__ void kernel_logical_constant<unsigned char,1,SHR>(unsigned char*,unsigned char,unsigned char*,int,int,int);
template __global__ void kernel_logical_constant<unsigned char,1,AND>(unsigned char*,unsigned char,unsigned char*,int,int,int);
template __global__ void kernel_logical_constant<unsigned char,1,OR>(unsigned char*,unsigned char,unsigned char*,int,int,int);
template __global__ void kernel_logical_constant<unsigned char,1,XOR>(unsigned char*,unsigned char,unsigned char*,int,int,int);
template __global__ void kernel_logical_constant<unsigned char,3,SHL>(unsigned char*,unsigned char,unsigned char*,int,int,int);
template __global__ void kernel_logical_constant<unsigned char,3,SHR>(unsigned char*,unsigned char,unsigned char*,int,int,int);
template __global__ void kernel_logical_constant<unsigned char,3,AND>(unsigned char*,unsigned char,unsigned char*,int,int,int);
template __global__ void kernel_logical_constant<unsigned char,3,OR>(unsigned char*,unsigned char,unsigned char*,int,int,int);
template __global__ void kernel_logical_constant<unsigned char,3,XOR>(unsigned char*,unsigned char,unsigned char*,int,int,int);
template __global__ void kernel_logical_constant<unsigned short,1,SHL>(unsigned short*,unsigned short,unsigned short*,int,int,int);
template __global__ void kernel_logical_constant<unsigned short,1,SHR>(unsigned short*,unsigned short,unsigned short*,int,int,int);
template __global__ void kernel_logical_constant<unsigned short,1,AND>(unsigned short*,unsigned short,unsigned short*,int,int,int);
template __global__ void kernel_logical_constant<unsigned short,1,OR>(unsigned short*,unsigned short,unsigned short*,int,int,int);
template __global__ void kernel_logical_constant<unsigned short,1,XOR>(unsigned short*,unsigned short,unsigned short*,int,int,int);
template __global__ void kernel_logical_constant<unsigned short,3,SHL>(unsigned short*,unsigned short,unsigned short*,int,int,int);
template __global__ void kernel_logical_constant<unsigned short,3,SHR>(unsigned short*,unsigned short,unsigned short*,int,int,int);
template __global__ void kernel_logical_constant<unsigned short,3,AND>(unsigned short*,unsigned short,unsigned short*,int,int,int);
template __global__ void kernel_logical_constant<unsigned short,3,OR>(unsigned short*,unsigned short,unsigned short*,int,int,int);
template __global__ void kernel_logical_constant<unsigned short,3,XOR>(unsigned short*,unsigned short,unsigned short*,int,int,int);

int main(){

  int *int_src, int_val, *int_dst, h_width, h_height, h_pitch;
  unsigned char *uc_src, uc_val, *uc_dst;
  cudaMalloc((void **)&int_src, sizeof(int));
  cudaMalloc((void **)&int_dst, sizeof(int));
  cudaMalloc((void **)&uc_src, sizeof(unsigned char));
  cudaMalloc((void **)&uc_dst, sizeof(unsigned char));
  h_width = 0;
  h_height = 0;
  h_pitch = 0;
  int_val = 0;
  uc_val = 0;

  kernel_logical_constant<int, 3, SHL><<<1,1>>>(int_src, int_val, int_dst, h_width, h_height, h_pitch);

  kernel_logical_constant<unsigned char, 1, SHL><<<1,1>>>(uc_src, uc_val, uc_dst, h_width, h_height, h_pitch);

  return 0;
}
$ nvcc -arch=sm_20 -o t243 t243.cu
$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2013 NVIDIA Corporation
Built on Wed_Jul_17_18:36:13_PDT_2013
Cuda compilation tools, release 5.5, V5.5.0
$