找回密码
 立即注册
xiao88yan 该用户已被删除
发表于 2013-7-25 10:46:01
10972
本帖最后由 xiao88yan 于 2013-7-25 11:39 编辑

我在做关于体绘制方面的内容
我的代码是从CUDA示例程序中的volume render中加以改动而来,但是发现改动以后,kernel函数的运行时间增加到很多,从3ms增加到了186ms,严重影响了程序的运行流畅度。

比较之后发现,我的代码有几处嫌疑最大,比如多了两个纹理读取语句:
float4 startPosition = tex2D(cudafronttex,u,v);
float4 endPostion = tex2D(cudabacktex,u,v);

还多一句求根运算:
depth = sqrt((endPostion.x-pos.x)*(endPostion.x-pos.x) + (endPostion.y-pos.y)*(endPostion.y-pos.y) + (endPostion.z-pos.z)*(endPostion.z-pos.z));
以上是我认为的最有可能的原因。




下面贴出两个代码,并给出我的初步分析,请大家帮我进一步分析一下,非常感谢:


CUDA示例代码,profiler测试运行时间为3ms:
__global__ void
d_render(uint *d_output, uint imageW, uint imageH,
         float density, float brightness,
         float transferOffset, float transferScale)
{
    const int maxSteps = 500;
    const float tstep = 0.01f;
    const float opacityThreshold = 0.95f;

    const float3 boxMin = make_float3(-1.0f, -1.0f, -1.0f);
    const float3 boxMax = make_float3(1.0f, 1.0f, 1.0f);

    uint x = blockIdx.x*blockDim.x + threadIdx.x;
    uint y = blockIdx.y*blockDim.y + threadIdx.y;


    if ((x >= imageW) || (y >= imageH)) return;

    float u = (x / (float) imageW)*2.0f-1.0f;
    float v = (y / (float) imageH)*2.0f-1.0f;

    // calculate eye ray in world space
    Ray eyeRay;
    eyeRay.o = make_float3(mul(c_invViewMatrix, make_float4(0.0f, 0.0f, 0.0f, 1.0f)));
    eyeRay.d = normalize(make_float3(u, v, -2.0f));
    eyeRay.d = mul(c_invViewMatrix, eyeRay.d);

    // find intersection with box
    float tnear, tfar;
    int hit = intersectBox(eyeRay, boxMin, boxMax, &tnear, &tfar);
    if (!hit) return;

    if (tnear < 0.0f) tnear = 0.0f;     // clamp to near plane

    // march along ray from front to back, accumulating color
    float4 sum = make_float4(0.0f);
    float t = tnear;
    float3 pos = eyeRay.o + eyeRay.d*tnear;
    float3 step = eyeRay.d*tstep;

    for (int i=0; i<maxSteps; i++)
    {
        float sample = tex3D(tex, pos.x*0.5f+0.5f, pos.y*0.5f+0.5f, pos.z*0.5f+0.5f);

        // lookup in transfer function texture
        float4 col = tex1D(transferTex, (sample-transferOffset)*transferScale);
        col.w *= density;

        // "under" operator for back-to-front blending
        //sum = lerp(sum, col, col.w);

        // pre-multiply alpha
        col.x *= col.w;
        col.y *= col.w;
        col.z *= col.w;
        // "over" operator for front-to-back blending
        sum = sum + col*(1.0f - sum.w);

        // exit early if opaque
        if (sum.w > opacityThreshold)
            break;

        t += tstep;
        if (t > tfar) break;

        pos += step;
    }

    sum *= brightness;

    // write output color
    d_output[y*imageW + x] = rgbaFloatToInt(sum);

}

我的代码,profiler测试运行时间186ms:
__global__ void
d_render(uint *d_output, uint imageW, uint imageH,
         float density, float brightness,
         float transferOffset, float transferScale)
{
    const int maxSteps = 500;
    const float tstep = 0.01f;
    const float opacityThreshold = 0.95f;

        uint x = blockIdx.x*blockDim.x + threadIdx.x;
    uint y = blockIdx.y*blockDim.y + threadIdx.y;
    if ((x >= imageW) || (y >= imageH)) return;


        float u = (x / (float) imageW);
        float v = (y / (float) imageH);

    // calculate eye ray in world space
    Ray eyeRay;
    eyeRay.o = make_float3(mul(c_invViewMatrix, make_float4(0.0f, 0.0f, 0.0f, 1.0f)));
        eyeRay.d = normalize(make_float3(u-0.5f, v-0.5f, -1.0f));
    eyeRay.d = mul(c_invViewMatrix, eyeRay.d);

    float4 sum = make_float4(0.0f);

        float3 step = eyeRay.d*tstep;

        float4 startPosition = tex2D(cudafronttex,u,v);
        float4 endPostion = tex2D(cudabacktex,u,v);
        float3 pos;
        pos.x = startPosition.x;
        pos.y = startPosition.y;
        pos.z = startPosition.z;

        depth = sqrtf(pow((endPostion.x-startPosition.x),2) + pow((endPostion.y-startPosition.y),2) + pow((endPostion.z-startPosition.z),2));
        if (depth<1e-7) return;

        float t = 0;

    for(int i=0; i<maxSteps; i++)
    {
                float sample = tex3D(tex, pos.x, pos.y, pos.z);

        // lookup in transfer function texture
        float4 col = tex1D(transferTex, (sample-transferOffset)*transferScale);
        col.w *= density;

        // "under" operator for back-to-front blending

        // pre-multiply alpha
        col.x *= col.w;
        col.y *= col.w;
        col.z *= col.w;
        // "over" operator for front-to-back blending
        sum = sum + col*(1.0f - sum.w);

        // exit early if opaque
        if (sum.w > opacityThreshold)  break;


        t += tstep;
                if (t > depth) break;

        pos += step;

    }
    sum *= brightness;

    // write output color
    d_output[y*imageW + x] = rgbaFloatToInt(sum);

}

1、为方便查看,我把以上基本相同的代码标记为蓝色,红色为差异代码。从两个代码上看,在循环处理上基本没差别,在循环之外,我的代码多了两个读取纹理的操作,读取之后有一个求根操作,而CUDA示例代码多了一个 intersectBox函数,但从效果看,intersectBox耗时可以忽略不计。


2、我很纳闷的是,即使是读取纹理和求根操作比较耗时,但那也是在循环之外,每个线程只读取一次,也不至于时间消耗增加那么多,从3ms增加到180多ms


3、那么与我的纹理的动态变化有无关系?我的两个纹理都是512*512,而且是实时动态变化的,而且kernel 函数是循环执行的,kernel 函数在两次执行过程 中,同样的纹理坐标,取到的值是不一样的。我看了资料,CUDA纹理是有缓存的,这种不断动态变化的纹理内容会不会对缓存造成影响,导致缓存摇摆,影响速度?

4、再说明一下,我的两个纹理是从opengl中实时生成,并映射到CUDA中进行使用的,应该也是在设备的global memory中吧,有没有影响?




使用道具 举报 回复
发表于 2013-7-25 15:42:08
版主们来帮帮忙呀
使用道具 举报 回复 支持 反对
发表于 2013-7-25 16:19:52
各位朋友,版主好,我好像找到原因了,是从这个贴子里找到的启发http://cudazone.nvidia.cn/forum/ ... &extra=page%3D1

原因是将CU文件的属性Generate GPU Debug Informantion设置为是了

不过具体原因,还请各位高手及版主稍稍解释下,谢谢
使用道具 举报 回复 支持 反对
发新帖
您需要登录后才可以回帖 登录 | 立即注册