 }}}
 
 
+

            bitwise_un_op_two_loads<uint, 4, opid><<<grid, threads>>>(rows, cols, src, dst, mask); 
            break;
        }
        if (stream == 0) cudaSafeCall(cudaThreadSynchronize());
    }



        dim3 grid(divUp(cols * elem_size, threads.x * sizeof(uint)), 
                  divUp(rows, threads.y));
        bitwise_un_op<opid><<<grid, threads>>>(rows, cols * elem_size, src, dst);
        if (stream == 0) 
            cudaSafeCall(cudaThreadSynchronize());
    }


