stl源码剖析之priority_queue

发表于 2014-12-29 分类于编程语言， C++ ， STL 阅读次数： Valine：
本文字数： 2.2k 阅读时长 ≈ 2 分钟

暑假时候好好阅读过，stl源码剖析这本书，之后又看了c++标准程序库，effective c++，more
effecttive c++。虽然每次都如醍醐灌顶般，明白了很多知识概念，有种拨开迷雾的感觉，但是终究抵不过
遗忘，所以了还是随便写写吧。
首先从实现代码差不多最短的优先队列开始，虽然简单的代码后面依赖很庞大的东西，比如泛型
迭代器，还有泛型二叉堆算法。我是这样理解优先队列的。首先，数据是存在一个泛型序列容器(线性容器)，
这个可以指定，默认是vector，当然指定list和deque也是可以的。另外还需要给数据类型提供个小于比较仿
函数，有些类型有默认的less泛型仿函数。所以了，对模版比较熟悉的同学，很自然就猜到了，priority_queue
的外观->一个模版（类型，线性容器，less仿函数）。
优先队列支持哪些操作了？什么empty，size，top，直接转交给底层容器就行了，所谓的适配器模式。
只需要思考，如何实现push和pop操作。对算法熟悉的同学，都知道底层数据的组织是用二叉堆的，而且是最大
堆。问题是，我们还需不需要实现个堆操作了。不需要了，stl做了所有的事情。push_heap就是将输入范围的最
后一个元素作为在原堆后面加入的新元素，再重构堆。pop_heap则是将堆首元素交换到堆末，再重构缩小之后的堆。
至于二叉堆的原理，就不仔细介绍了。其实也比较简单，如果不去思考一些复杂度的证明。
到最后，priority_queue的代码可以很清爽的出来了。尤其注意下面的那些typedef，可以说是模版
trick的常见手法，习惯了就好。这种写法还有个用法是为了使内嵌模版声明可见，记得effect c++提到过。以下
源码基本来自stl源码剖析随书奉送代码，做了些修改。

#include <algorithm>

template <typename T, typename Sequence = vector<T>,
    typename Compare = less<typename Sequence::value_type >>
class  priority_queue
{
public:
    typedef typename Sequence::value_type value_type;
    typedef typename Sequence::size_type size_type;
    typedef typename Sequence::reference reference;
    typedef typename Sequence::const_reference const_reference; //为了使内嵌模版声明可见

protected:
    Sequence c; // 底层容器
    Compare comp;   // 元素大小比较仿函数

    priority_queue() : c() {}
    explicit priority_queue(const Compare x) : c(), comp(x) {}

    template <class InputIterator>
    priority_queue(InputIterator first, InputIterator last, const Compare x)
        : c(first, last), comp(x) {
        make_heap(c.begin(), c.end(), comp);
    }

    template <class InputIterator>
    priority_queue(InputIterator first, InputIterator last)
        : c(first, last) {
        make_heap(c.begin(), c.end(), comp);
    }

    bool empty() const { return c.empty(); }
    size_type size() const { return c.size(); }
    const_reference top() const { return c.front(); }

    void push(const value_type x)
    {
        try
        {
            c.push_back(x);
            push_heap(c.begin(), c.end(), comp);
        }
        catch (...)
        {
            c.clear();
            throw;
        }
    }

    void pop()
    {
        try
        {
            pop_heap(c.begin(), c.end(), comp);
            c.pop_back();
        }
        catch (...)
        {
            c.clear();
            throw;
        }
    }
};

再啰嗦几句，刚看源码剖析的时候，对某些模版写法很头疼，但是看完书之后，另外看了几本其它的书，
就对这些写法习以为常了，觉得不这么写才怪了。有点艺术的感觉。

任意点到二维直线或者三维平面上的投影点

发表于 2014-12-09 分类于图形学阅读次数： Valine：
本文字数： 454 阅读时长 ≈ 1 分钟

这里统一考虑二维直线和三维平面的情况。
假设法线（二维直线和三维平面的法线）为n，其上的任意一点为p0。那么，
可以得到方程为：(p - p0) * n = 0。展开得到p * n = p0 * n。令p0 * n = d，由此可以得到
方程可以简单的表示为p * n = d，即(x, y, z) * (nx ,ny, nz) = d，二维情况直线的为(x, y) * (nx, ny) = d。
至此，我们可以方便的用(n,d)代表任意的二维直线或者三维平面了。
下面，开始应用这个表示。
比如，求点到直线或者平面的距离，然后求投影点，再求对称点。
最关键的是如何方便求出距离。假设我们要求q到直线的距离。我们现在的直线方程为
p * n = d。那么经过q与当前直线的平行直线的法线也为n。由于q在经过q的平行直线上面，有q * n = d’。
有符号距离为 d - d’ = d - q * n。通过这个式子，也可以知道d其实就是原点到直线或者平面的有符号距离。
那么投影点q’ = q + (d - d’) = q + (d - q * n)。
对称点则为 q’’ = q’ + (d - d’) = q + 2 * (d - q * n)。
利用这种表示方法，基本上可以直接给出答案。

delete NULL是安全的

发表于 2014-11-27 分类于编程语言， C++ 阅读次数： Valine：
本文字数： 286 阅读时长 ≈ 1 分钟

一直以来都默认删除指针前，都得判断下是否是空指针。直到最近在看More Effective C++
的时候，作者明确指出了C++在语言层次上保证删除空指针是安全的，我才意识到这个问题。
我用vs2013实践了下，发现没有运行错误。所以了，这大概是从学习C语言起遗留下来的思维
定势吧。以前用C调用malloc申请内存的时候，对应的是free释放，free肯定是不能释放NULL指针的，所以
自然而然就觉得delete也不能处理NULL指针。
虽然说，这件事情无伤大雅，但是确实没必要在delete之前判断下是否为NULL了，毕竟自己
判断和编译器判断都是一样的结果，也不存在什么效率问题。

vs2013下配置cuda6.5和opencv2.4.10的gpu版本

发表于 2014-11-13 分类于图像处理， CUDA 阅读次数： Valine：
本文字数： 978 阅读时长 ≈ 1 分钟

首先，当然是安装好vs2013。第二步是配置最新的cuda6.5版本，这个有相关的教程。
这里要讲的是如何生成相应的opencv的gpu版本。
首先，当然是去opencv官方网站下载opencv，然后解压到指定的目录，最好不要出现
中文路径，因为后面需要用到cmake。2.4.10的下载地址，已经不好找了。最新的是3.0beta版本，
由于模块分类完全变了，所以还是别去吃螃蟹了。
下载好解压之后，会出现两个子目录，build和source。如果，你只使用cpu版本，那么build
里面的内容就满足你的要求了。否则，打开source目录，你会发现有一个CMakeLists.txt文件。
那么，肯定是需要使用cmake生成相应的工程了。
第二步，下载cmake，安装好cmake。
第三步，用cmake生成工程。操作方法：假如opencv在f盘根目录下面，那么设置好cmake的
source目录和build目录。然后点击configure按钮，选择对应的平台如visual studio 12。注意，现在
记得选择WITH_CUDA，WITH_CUFFT，WITH_CUBLAS等相关的选项。
然后，点击generate按钮就能生成对应的工程。如果，出现配置错误，可以根据cmake的
提示寻找原因。基本上是些环境变量的设置问题，或者是平台不匹配等。
如图所示：

注意build目录不用和source目录设置为同一个，否则可能引起问题，cmake也会出警告。
第四步，就是打开工程生成文件了。用vs2013或者你的使用版本，打开build目录下的OpenCV.sln
文件。最简单的办法，就是批生成->全选->重新生成。
关键的问题出现，如果你只是这样做，可能有几个dll无法编译出来，比如说opencv_gpu2410.dll。
问题在哪里了。我试验了好几次，每次几个小时，都是这样的结果。顿时觉得很郁闷，
只能去分析vs的输出内容了。突然发现有个错误提示：大致内容是NCV.cu文件中的std::max没有
定义。这个cu文件在..\opencv\sources\modules\gpu\src\nvidia\core下面。
你最后需要做的就是打开这个文件，然后包含algorithm头文件就行了，然后再重新编译吧。
PS：我只生成过32平台下的gpu版本了，64位的就没有去尝试了。

opencv gpu或者cuda加速第一次调用

发表于 2014-11-11 分类于图像处理， CUDA 阅读次数： Valine：
本文字数： 450 阅读时长 ≈ 1 分钟

经测试，第一次gpu调用，无论是用opencv的gpu模块或者cuda都会比较耗时，可能将近1s。
额，这也不是我一个人碰到的问题，确实有这回事。stackoverflow上面有个帖子就是关于这个的。
帖子地址：关于速度慢，关于解决办法。
如果，看了上面的帖子话，这篇文章也不用看下去了。因为我要讲的就是这件事情。
我的观点也是，在第一次调用之前先建立cuda context。为什么需要做这样的事情了？额，假如，
你只调用一次gpu处理，总不能太慢了吧。那样就看不到速度了，所以先调用次垃圾操作初始化cuda环境。
方法是调用cudaFree(0)，在这之前最好调用先调用cudaSetDevice(0)。记住包含cuda的头文件，
如果只有opencv的头文件，这2个函数是找不到的。还有，建立的工程是cuda runtime模版。
在我的L0Smooth代码里，这样的处理之后，初次L0Smooth调用能减少1s左右，从3s变成了2s。
其余的耗时，基本都在gpu版本的dft和idft，还得继续寻找加快速度的办法。

opencv的gpu部分实现dft

发表于 2014-11-10 分类于图像处理， OpenCV 阅读次数： Valine：
本文字数： 1.6k 阅读时长 ≈ 1 分钟

最近打算使用gpu优化L0Smooth的代码，但是不熟悉opencv的gpu部分的使用，
直接用cuda觉得太麻烦，还是继续试试opencv吧。
首先，从网站上下载下来的默认版本是不支持gpu的，必须下载代码下来，用cmake生成工程，注意
选择支持cuda等选项，具体参考教程。这样生成的版本才能支持gpu运算。
那么如何测试自己的opencv是否支持gpu计算了，以及自己的硬件是否符合要求？
使用这句代码打印cuda设备个数：
printf(“Device Num:%d\n”, cv::gpu::getCudaEnabledDeviceCount());
如果，个数大于0，那么说明你的显卡是支持cuda的，并且你的opencv版本支持gpu运算了。
下面是我测试成功的gpumat实现的dft和idft函数，输入和输出的都是cpu上的mat。

void Dft(cv::Mat in, cv::Mat out)
{
    cv::gpu::GpuMat gpuIn0(in.size(), CV_32FC1);
    gpuIn0.upload(in);

    std::vector<cv::gpu::GpuMat> planes;

    planes.push_back(gpuIn0);

    cv::Mat zero = cv::Mat::zeros(in.size(), CV_32FC1);
    cv::gpu::GpuMat gpuZero(in.size(), CV_32FC1);
    gpuZero.upload(zero);

    planes.push_back(gpuZero);

    cv::gpu::GpuMat gpuIn(in.size(), CV_32FC2);
    cv::gpu::merge(planes, gpuIn);

    cv::gpu::GpuMat gpuOut(gpuIn.size(), CV_32FC2);

    cv::gpu::dft(gpuIn, gpuOut, gpuIn.size(), 0);

    out.create(in.size(), CV_32FC2);
    gpuOut.download(out);
}

void IDft(cv::Mat in, cv::Mat out)
{
    cv::gpu::GpuMat gpuIn(in.size(), CV_32FC2);
    cv::gpu::GpuMat gpuOut(in.size(), CV_32FC2);

    gpuIn.upload(in);
    cv::gpu::dft(gpuIn, gpuOut, gpuIn.size(), cv::DFT_INVERSE);

    cv::gpu::GpuMat splitter[2];
    cv::gpu::split(gpuOut, splitter);

    out.create(in.size(), CV_32FC1);

    double minV, maxV;
    cv::gpu::minMax(splitter[0], minV, maxV);
    splitter[0].convertTo(splitter[0], CV_32F, 255.0 / maxV);
    splitter[0].download(out);
}

从代码上，可以看出，要尽可能把运算放到gpu上去，包括很多辅助运算，比如说merge,convert等可以加快速度。
这个版本的代码，dft的输入和输出都是双通道的，也就是一个通道实数，一个是复数的矩阵。通过merge生成和split分离。
这两个函数在分离通道单独处理时候非常有用，比如说可以分离彩色图像的三个通道，单独进行dft处理。最后再将结果合并。

opencv实现快速傅立叶变换和逆变换

发表于 2014-11-07 分类于图像处理， OpenCV 阅读次数： Valine：
本文字数： 3.1k 阅读时长 ≈ 3 分钟

说实话觉得网上很多人转载的文章的挺坑的，全部是opencv文档程序的翻译，看来看去都是那一
篇，真的没啥意思。文档的地址。
本来opencv实现dft就是一个函数的事情，但是很少有关于逆变换使用的资料。我这几天在翻译
matlab版本的L0Smooth到opencv上面，就碰到这样一件很坑爹的事情。
首先，很少有人说清楚这个函数的使用方法。还有，根据教程，dft之前最好扩充原矩阵到合适的尺
寸(2,3,5的倍数)，再调用dft会加快速度。那么，idft的时候了？如何恢复原有的尺寸？
在我的L0Smooth代码里，就碰到这样的事情了。如果，图片尺寸是2，3，5的倍数，那么能够得到
正确结果。否则得到是全黑的图片。如果，我不扩张矩阵，那么就能正确处理。
所以，到这里，我不推荐调用dft之前先扩充矩阵了。因为，我找了很久也没找到解决办法。
我数学水平有限，也分析不出原因，也没有时间去系统的学习这些了。
这里提供两个例子，说明dft和idft的使用。
例子一：类似于opencv官方文档的例子

#include "opencv2/core/core.hpp"
#include "opencv2/imgproc/imgproc.hpp"
#include "opencv2/highgui/highgui.hpp"
#include <iostream>

#ifdef _DEBUG
#pragma comment(lib, "opencv_core247d.lib")
#pragma comment(lib, "opencv_imgproc247d.lib")
#pragma comment(lib, "opencv_highgui247d.lib")
#else
#pragma comment(lib, "opencv_core247.lib")
#pragma comment(lib, "opencv_imgproc247.lib")
#pragma comment(lib, "opencv_highgui247.lib")
#endif // DEBUG

int main()
{
    // Read image from file
    // Make sure that the image is in grayscale
    cv::Mat img = cv::imread("lena.JPG",0);

    cv::Mat planes[] = {cv::Mat_<float>(img), cv::Mat::zeros(img.size(), CV_32F)};
    cv::Mat complexI;    //Complex plane to contain the DFT coefficients {[0]-Real,[1]-Img}
    cv::merge(planes, 2, complexI);
    cv::dft(complexI, complexI);  // Applying DFT

    //这里可以对复数矩阵comlexI进行处理

    // Reconstructing original imae from the DFT coefficients
    cv::Mat invDFT, invDFTcvt;
    cv::idft(complexI, invDFT, cv::DFT_SCALE | cv::DFT_REAL_OUTPUT ); // Applying IDFT
    cv::invDFT.convertTo(invDFTcvt, CV_8U); 
    cv::imshow("Output", invDFTcvt);

    //show the image
    cv::imshow("Original Image", img);

    // Wait until user press some key
    cv::waitKey(0);

    return 0;
}

代码意思很简单，dft之后再idft，注意参数额，必须有DFT_SCALE。代码中，先merge了个
复数矩阵，在例子2中可以看到，其实这一步可以去掉。
例子2：

#include "opencv2/core/core.hpp"
#include "opencv2/imgproc/imgproc.hpp"
#include "opencv2/highgui/highgui.hpp"
#include <iostream>

#ifdef _DEBUG
#pragma comment(lib, "opencv_core247d.lib")
#pragma comment(lib, "opencv_imgproc247d.lib")
#pragma comment(lib, "opencv_highgui247d.lib")
#else
#pragma comment(lib, "opencv_core247.lib")
#pragma comment(lib, "opencv_imgproc247.lib")
#pragma comment(lib, "opencv_highgui247.lib")
#endif // DEBUG

int main()
{
    // Read image from file
    // Make sure that the image is in grayscale
    cv:;Mat img = cv::imread("lena.JPG",0);

    cv::Mat dftInput1, dftImage1, inverseDFT, inverseDFTconverted;
    cv::img.convertTo(dftInput1, CV_32F);
    cv::dft(dftInput1, dftImage1, cv::DFT_COMPLEX_OUTPUT);    // Applying DFT

    // Reconstructing original imae from the DFT coefficients
    cv::idft(dftImage1, inverseDFT, cv::DFT_SCALE | cv::DFT_REAL_OUTPUT ); // Applying IDFT
    cv::inverseDFT.convertTo(inverseDFTconverted, CV_8U);
    cv::imshow("Output", inverseDFTconverted);

    //show the image
    cv::imshow("Original Image", img);

    // Wait until user press some key
    waitKey(0);
    return 0;
}

从代码中可以看到，dft时候添加参数DFT_COMPLEX_OUTPUT，就可以自动得到复数矩阵了，代码更加简洁。
注意，必须先将图片对应的uchar矩阵转换为float矩阵，再进行dft，idft，最后再转换回来。

cuda结合opencv实现简单的平滑滤波

发表于 2014-11-05 分类于图像处理， CUDA 阅读次数： Valine：
本文字数： 2.4k 阅读时长 ≈ 2 分钟

这次也是使用opencv的mat加载处理图像。唯一与上次有区别的是核函数的编写。
根据cuda的线程分配模型，每一个像素是分配单独的线程处理的。那么有这样的一个疑问？
像平滑滤波这些应用，如何在每一个线程中获取周围的像素了？
其实，这个问题很好解决。因为，在核函数中，我们能够根据线程id，块id，块尺寸等计算
出当前像素的位置。那么，自然能够得到其邻域的位置。从而实现了平滑滤波。
代码如下：

#include <stdlib.h>
#include <stdio.h>
#include <opencv/cv.h>
#include <opencv/highgui.h>
#include <opencv2/opencv.hpp>

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#ifdef _DEBUG
#pragma comment(lib, "opencv_core247d.lib")
#pragma comment(lib, "opencv_imgproc247d.lib")
#pragma comment(lib, "opencv_highgui247d.lib")
#else
#pragma comment(lib, "opencv_core247.lib")
#pragma comment(lib, "opencv_imgproc247.lib")
#pragma comment(lib, "opencv_highgui247.lib")
#endif // DEBUG

__global__ void smooth_kernel(const uchar3* src, uchar3* dst, int width, int height)
{
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;

    if(x < width  y < height)
    {
        int offset = x + y * width;
        int left = offset - 1;
        if (x - 1 < 0)
        {
            left += 1;
        }
        int right = offset + 1;
        if (x + 1 >= width)
        {
            right -= 1;
        }
        int top = offset - width;
        if (y - 1 < 0)
        {
            top += width;
        }
        int bottom = offset + width;
        if (y + 1 >= height)
        {
            bottom -= width;
        }

        dst[offset].x = 0.125 * (4 * src[offset].x + src[left].x + src[right].x + src[top].x + src[bottom].x);
        dst[offset].y = 0.125 * (4 * src[offset].y + src[left].y + src[right].y + src[top].y + src[bottom].y);
        dst[offset].z = 0.125 * (4 * src[offset].z + src[left].z + src[right].z + src[top].z + src[bottom].z);
    }
}

void smooth_caller(const uchar3* src, uchar3* dst, int width, int height)
{
    dim3 threads(16, 16);
    dim3 grids((width + threads.x - 1) / threads.x, (height + threads.y - 1) / threads.y);

    smooth_kernel<< <grids, threads >> >(src, dst, width, height);
    cudaThreadSynchronize();
}

int main()
{
    cv::Mat image = cv::imread("lena.png");
    cv::imshow("src", image);

    size_t memSize = image.step * image.rows;
    uchar3* d_src = NULL;
    uchar3* d_dst = NULL;
    cudaMalloc((void**)d_src, memSize);
    cudaMalloc((void**)d_dst, memSize);
    cudaMemcpy(d_src, image.data, memSize, cudaMemcpyHostToDevice);

    smooth_caller(d_src, d_dst, image.cols, image.rows);

    cudaMemcpy(image.data, d_dst, memSize, cudaMemcpyDeviceToHost);
    cv::imshow("gpu", image);
    cv::waitKey(0);

    cudaFree(d_src);
    cudaFree(d_dst);

    return 0;
}

效果如图：

cuda和opencv混合使用

发表于 2014-11-03 分类于图像处理， CUDA 阅读次数： Valine：
本文字数： 2.4k 阅读时长 ≈ 2 分钟

这里我既不介绍opencv的基本使用，也更加不会介绍cuda的使用。推荐下cuda的一本书：GPU高性能编程CUDA实战。
opencv这么强大的工具不用肯定是浪费了，opencv也有gpu的部分，据说也是用cuda实现的，但是灵活性肯定不如直接用cuda吧。
所以，我觉得只需要使用opencv负责cpu的部分，比如加载图片，gui之类的，而cuda负责并行的处理。还有，本着方便的原则，
opencv使用cpp的版本，不想再去管内存分配释放了。虽然，Mat相对来说更难使用。
下面是一个简短的交换rb通道的cuda和opencv混合的程序。

#include <stdlib.h>
#include <stdio.h>
#include <opencv/cv.h>
#include <opencv/highgui.h>
#include <opencv2/opencv.hpp>

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#ifdef _DEBUG
#pragma comment(lib, "opencv_core247d.lib")
#pragma comment(lib, "opencv_imgproc247d.lib")
#pragma comment(lib, "opencv_highgui247d.lib")
#else
#pragma comment(lib, "opencv_core247.lib")
#pragma comment(lib, "opencv_imgproc247.lib")
#pragma comment(lib, "opencv_highgui247.lib")
#endif // DEBUG

__global__ void swap_rb_kernel(const uchar3* src, uchar3* dst, int width, int height)
{
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;

    if(x < width  y < height)
    {
        int offset = x + y * width;
        uchar3 v = src[offset];
        dst[offset].x = v.z;
        dst[offset].y = v.y;
        dst[offset].z = v.x;
    }
}

void swap_rb_caller(const uchar3* src, uchar3* dst, int width, int height)
{
    dim3 threads(16, 16);
    dim3 grids((width + threads.x - 1) / threads.x, (height + threads.y - 1) / threads.y);

    swap_rb_kernel<<<grids, threads>>>(src, dst, width, height);
    cudaThreadSynchronize();
}

int main()
{
    cv::Mat image = cv::imread("lena_1.jpg");
    cv::imshow("src", image);

    size_t memSize = image.step * image.rows;
    uchar3* d_src = NULL;
    uchar3* d_dst = NULL;
    cudaMalloc((void**)d_src, memSize);
    cudaMalloc((void**)d_dst, memSize);
    cudaMemcpy(d_src, image.data, memSize, cudaMemcpyHostToDevice);

    swap_rb_caller(d_src, d_dst, image.cols, image.rows);

    cudaMemcpy(image.data, d_dst, memSize, cudaMemcpyDeviceToHost);
    cv::imshow("gpu", image);
    cv::waitKey(0);

    cudaFree(d_src);
    cudaFree(d_dst);

    return 0;
}

运行效果：

opencv部分不用做过多解释了，cuda的那些内存操作函数也不用解释。唯一需要解释的是核函数里面的这两句：int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;千万别搞错x和y，否则效果就完全不对了。根据cuda的模型，线程是一块一块的。一个线程块里面有很多线程，那么如何索引这些线程了？
把线程块看作是是三维的（一般用到一维或者二维）数组，然后根据数组索引得到具体线程。至于blockDim指的是一共有多少线程块了，这个也是三维的，意思一个gpu格子里面，
会出现三维的线程块组合。gpu格子的大小，在cuda里面用gridDim指代。所以，从上到下就是三层模型吧，三维的grid->三维的block->三维的thread。具体的理解，参照书籍或者教程吧。

vs2013下设置cuda高亮

发表于 2014-11-03 分类于图像处理， CUDA 阅读次数： Valine：
本文字数： 572 阅读时长 ≈ 1 分钟

因为是最近才接触cuda，安装的是6.5版本，所以网上的教程都不是完全适用了。
总之，设置高亮大致分为四步，不限于vs2013，其他平台下也类似。
第一步，是在vs2013里面设置vc++文件支持.cu;cuh;文件。方法：工具->选项->文本编辑器->文件扩展名。
得到如图所示的界面：注意，在右侧可以添加vc++类型的文件扩展名，这是我的设置效果，操作就不用细说了。

第二步，是设置visual assist的目录。在va的C++directory里面，选择custom选项，然后包含你的cuda的sdk目录，效果如图：

第三步，是设置va的支持文件类型，类似于第一步。但是，这次是修改注册表的值。注册表目录：
HKEY_CURRENT_USER/Software/Whole Tomato/Visual Assist X/VANet12，修改属性ExtSource的值为：.c;.cpp;.cc;.cxx;.tli;.cu;.cuh;
意思就是添加上cuda的头文件和源文件类型，vs2010的改法类似。
第四步，完成以上步骤之后，还可能会发现一些内置变量下面是有波浪线的。怎么办了？
加上这句：#include “device_launch_parameters.h”，就行了。cuda 6.5估计把内置变量的声明放在该头文件下面了吧。
最终的效果：