OpenGL Shader vs CUDA

问题描述:

我在使用0123AAFXAA Shader在我的OpenGL程序中进行反锯齿。现在我在CUDA中重新实现了这个代码并对其进行了测试。由此产生的图像是相同的,但CUDA版本要慢得多。 (着色器运行在60 FPS与垂直同步,而CUDA下降到〜40 FPS)OpenGL Shader vs CUDA

这里是CUDA代码:

__device__ uchar4 readChar(int x, int y){ 
    return surf2Dread<uchar4>(surfaceRead, (x)*sizeof(uchar4), (y),cudaBoundaryModeClamp); 
} 

__device__ uchar4 readFloatBilin2(float x, float y){ 
    int x1 = floor(x); 
    int y1 = floor(y); 

    uchar4 z11 = readChar(x1,y1); 
    uchar4 z12 = readChar(x1,y1+1); 
    uchar4 z21 = readChar(x1+1,y1); 
    uchar4 z22 = readChar(x1+1,y1+1); 

    float u_ratio = x - x1; 
    float v_ratio = y - y1; 
    float u_opposite = 1 - u_ratio; 
    float v_opposite = 1 - v_ratio; 
    uchar4 result = (z11 * u_opposite + z21 * u_ratio) * v_opposite + 
        (z12 * u_opposite + z22 * u_ratio) * v_ratio; 


    return result; 
} 
__device__ float fluma(const uchar4 &c){ 
    return c.x*0.299 * (1.0/255) + c.y *0.587 * (1.0/255) + c.z*0.114 * (1.0/255); 
} 
__global__ void filter_fxaa_opt(TextureData data) 
{ 

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

    if(x >= data.w || y >= data.h) 
    { 
     return; 
    } 

    uchar4 out_color; 

    const float FXAA_SPAN_MAX = 8.0; 
    const float FXAA_REDUCE_MUL = 1.0/8.0; 
    const float FXAA_REDUCE_MIN = (1.0/128.0); 



    float lumaNW = fluma(readChar(x-1,y-1)); 

    float lumaNE = fluma(readChar(x+1,y-1)); 

    float lumaSW = fluma(readChar(x-1,y+1)); 

    float lumaSE = fluma(readChar(x+1,y+1)); 

    float lumaM = fluma(readChar(x,y)); 

    float lumaMin = min(lumaM, min(min(lumaNW, lumaNE), min(lumaSW, lumaSE))); 
    float lumaMax = max(lumaM, max(max(lumaNW, lumaNE), max(lumaSW, lumaSE))); 

    float2 dir; 
    dir.x = -((lumaNW + lumaNE) - (lumaSW + lumaSE)); 
    dir.y = ((lumaNW + lumaSW) - (lumaNE + lumaSE)); 

    float dirReduce = max((lumaNW + lumaNE + lumaSW + lumaSE) * (0.25 * FXAA_REDUCE_MUL), FXAA_REDUCE_MIN); 

    float rcpDirMin = 1.0/(min(abs(dir.x), abs(dir.y)) + dirReduce); 


// float2 test = dir * rcpDirMin; 
    dir = clamp(dir * rcpDirMin,-FXAA_SPAN_MAX,FXAA_SPAN_MAX); 


    uchar4 rgbA = (
       readFloatBilin2(x+ dir.x * (1.0/3.0 - 0.5),y+ dir.y * (1.0/3.0 - 0.5))*0.5f+ 
       readFloatBilin2(x+ dir.x * (2.0/3.0 - 0.5),y+ dir.y * (2.0/3.0 - 0.5))*0.5f); 
    uchar4 rgbB = rgbA * (1.0/2.0) + (
       readFloatBilin2(x+ dir.x * (0.0/3.0 - 0.5),y+ dir.y * (0.0/3.0 - 0.5))*0.25f+ 
       readFloatBilin2(x+ dir.x * (3.0/3.0 - 0.5),y+ dir.y * (3.0/3.0 - 0.5))*0.25f); 
    float lumaB = fluma(rgbB); 


    if((lumaB < lumaMin) || (lumaB > lumaMax)){ 
     out_color=rgbA; 
    } else { 
     out_color=rgbB; 
    } 

    surf2Dwrite<uchar4>(out_color, surfaceWrite, x*sizeof(uchar4), y); 
} 

设置:

//called for the 'src' and 'dst' texture once at the beginning 
checked_cuda(cudaGraphicsGLRegisterImage(&res, gl_buffer,gl_target, cudaGraphicsRegisterFlagsSurfaceLoadStore)); 

//called for the 'src' and 'dst' texture every frame 
checked_cuda(cudaGraphicsMapResources(1, &res, 0)); 
checked_cuda(cudaGraphicsSubResourceGetMappedArray(&array, res, 0,0)); 

//kernel call every frame 
dim3 block_size(8, 8); 
dim3 grid_size; 
grid_size.x = (src->w)/(block_size.x) ; 
grid_size.y = (src->h)/(block_size.y) ; 
checked_cuda(cudaBindSurfaceToArray(surfaceRead, (cudaArray *)src->d_data)); 
checked_cuda(cudaBindSurfaceToArray(surfaceWrite, (cudaArray *)dst->d_data)); 
filter_fxaa_opt<<<grid_size, block_size>>>(*src); 

系统:

Ubuntu 14.04 
Opengl version: 4.4.0 NVIDIA 331.113 
Renderer version: GeForce GTX 760M/PCIe/SSE2 
CUDA 5.5 

问题: 什么是OpenGL Shader做得更好,为什么它速度更快?

+2

一些裸设备代码如何,没有关于如何调用的信息,OpenGL互操作如何设置以及基准测试如何让我们能够将帧速率差异解释为完全未知的着色器? – talonmies 2015-01-04 15:43:34

+0

增加了互操作代码。使用的着色器是第一句中链接的1到1个副本。 – dari 2015-01-04 16:08:19

+0

您是否在该代码中暗示您正在映射OpenGL资源并在每个框架上绑定输入和输出表面?如果是这样,为什么?这些调用不是免费的... – talonmies 2015-01-04 16:19:36

njuffa指出,主要问题是手动插值和归一化。在使用CUDA texture而不是CUDA surface之后,可以通过调用tex2D(..)而不是surf2Dread(...)来使用内插内置。

修改后的CUDA代码现在几乎与OpenGL着色器保持一致,确实性能良好。

__global__ void filter_fxaa2(TextureData data) 
{ 

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

    if(x >= data.w || y >= data.h) 
    { 
     return; 
    } 

    uchar4 out_color; 

    const float FXAA_SPAN_MAX = 8.0f; 
    const float FXAA_REDUCE_MUL = 1.0f/8.0f; 
    const float FXAA_REDUCE_MIN = (1.0f/128.0f); 

    float u = x + 0.5f; 
    float v = y + 0.5f; 

    float4 rgbNW = tex2D(texRef, u-1.0f,v-1.0f); 
    float4 rgbNE = tex2D(texRef, u+1.0f,v-1.0f); 
    float4 rgbSW = tex2D(texRef, u-1.0f,v+1.0f); 
    float4 rgbSE = tex2D(texRef, u+1.0f,v+1.0f); 
    float4 rgbM = tex2D(texRef, u,v); 

    const float4 luma = make_float4(0.299f, 0.587f, 0.114f,0.0f); 
    float lumaNW = dot(rgbNW, luma); 
    float lumaNE = dot(rgbNE, luma); 
    float lumaSW = dot(rgbSW, luma); 
    float lumaSE = dot(rgbSE, luma); 
    float lumaM = dot(rgbM, luma); 

    float lumaMin = min(lumaM, min(min(lumaNW, lumaNE), min(lumaSW, lumaSE))); 
    float lumaMax = max(lumaM, max(max(lumaNW, lumaNE), max(lumaSW, lumaSE))); 

    float2 dir; 
    dir.x = -((lumaNW + lumaNE) - (lumaSW + lumaSE)); 
    dir.y = ((lumaNW + lumaSW) - (lumaNE + lumaSE)); 

    float dirReduce = max((lumaNW + lumaNE + lumaSW + lumaSE) * (0.25f * FXAA_REDUCE_MUL), FXAA_REDUCE_MIN); 

    float rcpDirMin = 1.0f/(min(abs(dir.x), abs(dir.y)) + dirReduce); 


    float2 test = dir * rcpDirMin; 
    dir = clamp(test,-FXAA_SPAN_MAX,FXAA_SPAN_MAX); 

    float4 rgbA = (1.0f/2.0f) * (
       tex2D(texRef,u+ dir.x * (1.0f/3.0f - 0.5f),v+ dir.y * (1.0f/3.0f - 0.5f))+ 
       tex2D(texRef,u+ dir.x * (2.0f/3.0f - 0.5f),v+ dir.y * (2.0f/3.0f - 0.5f))); 
    float4 rgbB = rgbA * (1.0f/2.0f) + (1.0f/4.0f) * (
       tex2D(texRef,u+ dir.x * (0.0f/3.0f - 0.5f),v+ dir.y * (0.0f/3.0f - 0.5f))+ 
       tex2D(texRef,u+ dir.x * (3.0f/3.0f - 0.5f),v+ dir.y * (3.0f/3.0f - 0.5f))); 
    float lumaB = dot(rgbB, luma); 


    if((lumaB < lumaMin) || (lumaB > lumaMax)){ 
     out_color=toChar(rgbA); 
    } else { 
     out_color=toChar(rgbB); 
    } 


    surf2Dwrite<uchar4>(out_color, surfaceWrite, x*sizeof(uchar4), y); 
} 

更新:

性能meassured与cudaEvents

  • 旧版本:〜12.8ms
  • 新版本:〜1.2ms的

结论:

仅将CUDA表面用于书写而不用于读取纹理!