zoukankan      html  css  js  c++  java
  • 《OpenCL编程指南》之 与Direct3D互操作

        介绍OpenCL与D3D 10之间的互操作。

    1.初始化OpenCL上下文实现Direct3D互操作

        OpenCL共享由pragma cl_khr_d3d10_sharing启用:

        #pragma OPENCL EXTENSION cl_khr_d3d10_sharing: enable

        启用D3D共享时,很多OpenCL函数会有所扩展,将接受一些处理D3D10共享的参数类型和值。

        可以用D3D互操作属性来创建OpenCL上下文:

        ·CL_CONTEXT_D3D10_DEVICE_KHR   在clCreateContext和clCreateContextFromtype的属性参数中作为一个属性名。

        函数可以查询D3D互操作特定的对象参数:

        ·CL_CONTEXT_D3D10_PREFER_SHARED_RESOURCES_KHR  作为clGetContextInfo的param_name参数值。

        ·CL_MEM_D3D10_RESOURCE_KHR 作为clGetMemObjectInfo的param_name参数值。

        ·CL_IMAGE_D3D10_SUBRESOURCE_KHR 作为clGetImageInfo的param_name参数值。

        ·CL_COMMAND_ACQUIRE_D3D10_OBJECTS_KHR CL_COMMAND_RELEASE_D3D10_OBJECTS_KHR 当param_name为CL_ENCENT_COMMAND_TYPE时,在clGetEventInfo的参数param_value中返回。

        OpenCL D3D10互操作函数在头文件cl_d3d10.h中。D3D10的Khronos扩展可以从Khronos网站得到。对于某些发布版本,可能需要下载这个扩展。

        初始化OpenCL的过程与平常基本相同,只有几点细小差别。首先平台可以使用clGetPlatformIDs函数列出。由于我们在搜索一个支持D3D共享的平台,要在各个平台上使用clGetPlatformInfo()调用来查询它支持的扩展。如果扩展串中包含cl_khr_d3d10_sharing,说明可以选用这个平台来实现D3D共享。

        给定一个支持D3D共享的cl_platform_id,可以在这个平台上使用clGetDeviceIDsFromD3D10KHR()查询相应的OpenCL设备ID:

    cl_int clGetDeviceIDsFromD3D10KHR(
        cl_platform_id             platform,
        cl_d3d10_device_source_khr d3d_device_source,
        void *                     d3d_object,
        cl_d3d10_device_set_khr    d3d_device_set,
        cl_uint                    num_entries,
        cl_device_id *             devices,
        cl_uint *                  num_devices)

    例如:

    errNum = clGetDeviceIDsFromD3D10KHR(
        platformIds[index_platform],
        CL_D3D10_DEVICE_KHR,
        g_pD3DDevice,
        CL_PREFERRED_DEVICES_FOR_D3D10_KHR,
        1,
        &cdDevice,
        &num_devices);
    
    if (errNum == CL_INVALID_PLATFORM) {
        printf("Invalid Platform: Specified platform is not valid
    ");
    } else if( errNum == CL_INVALID_VALUE) {
        printf("Invalid Value: d3d_device_source, d3d_device_set is not valid or num_entries = 0 and devices != NULL or num_devices == devices == NULL
    ");
    } else if( errNum == CL_DEVICE_NOT_FOUND) {
        printf("No OpenCL devices corresponding to the d3d_object were found
    ");
    }

    代码为选择的OpenCL平台(platformIds[index_platform])获取一个OpenCL设备ID(cdDevice)。常量CL_D3D10_DEVICE_KHR指示发送的D3D10对象(g_pD3DDevice)是一个D3D10设备,通过CL_PREFERRED_DEVICES_FOR_D3D10_KHR来选择该平台的期望设备。这会返回与平台和D3D10设备关联的期望OpenCL设备。

        这个函数返回的设备ID可以用来创建一个支持D3D共享的上下文。创建OpenCL上下文时,clCreateContext*()调用中的cl_context_properties域应当包括要共享的D3D10设备的指针。例如:

    cl_context_properties contextProperties[] =
    {
        CL_CONTEXT_D3D10_DEVICE_KHR, 
        (cl_context_properties)g_pD3DDevice,
        CL_CONTEXT_PLATFORM,
        (cl_context_properties)platformIds[index_platform],
        0
    };
    
    
    context = clCreateContextFromType( contextProperties, CL_DEVICE_TYPE_GPU, NULL, NULL, &errNum ) ;

    这个示例代码中,会从D3D10CreateDeviceAndSwapChain()调用返回D3D10设备g_pD3DDevice的指针。

    2.从D3D缓冲区和纹理创建OpenCL内存对象

        可以使用clCreateFromD3D10*KHR() OpenCL函数由现有的D3D缓冲区对象和纹理创建OpenCL缓冲区和图像对象。

        可以使用clCreateFromD3D10BufferKHR()由现有的D3D缓冲区创建OpenCL内存对象:

    cl_mem clCreateFromD3D10BufferKHR(
        cl_context     context,
        cl_mem_flags   flags,
        ID3D10Buffer * resource,
        cl_int *       errcode_ret)

        所返回的OpenCL缓冲区对象的大小与resource的大小相同。这个调用将使resource上的内部Direct3D引用计数增1.所返回OpenCL内存对象上的OpenCL引用计数减至0时,resource上的内部Direct3D引用计数会减1.

        缓冲区与纹理都可以与OpenCL共享。

         在D3D10中,纹理可以如下创建:

    // 2D texture
    D3D10_TEXTURE2D_DESC desc;
    ZeroMemory( &desc, sizeof(D3D10_TEXTURE2D_DESC) );
    desc.Width = g_WindowWidth;
    desc.Height = g_WindowHeight;
    desc.MipLevels = 1;
    desc.ArraySize = 1;
    desc.Format = DXGI_FORMAT_R8G8B8A8_UNORM;
    desc.SampleDesc.Count = 1;
    desc.Usage = D3D10_USAGE_DEFAULT;
    desc.BindFlags = D3D10_BIND_SHADER_RESOURCE;
    if (FAILED(g_pD3DDevice->CreateTexture2D( &desc, NULL, &g_pTexture2D)))
        return E_FAIL;

    这个共享的纹理格式为DXGI_FORMAT_R8G8B8A8_UNORM。然后可以使用

    cl_mem clCreateFromD3D10Texture2DKHR(
        cl_context        context,
        cl_mem_flags      flags,
        ID3D10Texture2D * resource,
        UINT              subresource,
        cl_int *          errcode_ret)

    创建一个OpenCL图像对象。所返回的OpenCL图像对象的宽度、高度和深度由resource得子资源subresource的宽度、高度、深度决定。所返回的OpenCL图像对象的通道类型和次序由resource的格式确定。

        这个调用将使resource上的内部Direct3D引用计数增1.所返回的OpenCL内存对象上的OpenCL引用计数减至0时,resource上的内部Direct3D引用计数减1.

        类似有3D的,

    cl_mem clCreateFromD3D10Texture3DKHR(
        cl_context        context,
        cl_mem_flags      flags,
        ID3D10Texture3D * resource,
        UINT              subresource,
        cl_int *          errcode_ret)

    微信图片_20170606080720

    微信图片_20170606080731

    3.OpenCL中获取和释放Direct3D对象

            在opencl中处理之前必须先获取direct3d对象,在由direct3d使用之前必须先释放direct3d对象。

    cl_int clEnqueueAcquireD3D10ObjectsKHR(
        cl_command_queue command_queue,
        cl_uint          num_objects,
        const cl_mem *   mem_objects,
        cl_uint          num_events_in_wait_list,
        const cl_event * event_wait_list,
        cl_event *       event)

    这会获得由D3D10资源创建的的OpenCL内存对象。

    cl_int clEnqueueAcquireD3D10ObjectsKHR(
        cl_command_queue command_queue,
        cl_uint          num_objects,
        const cl_mem *   mem_objects,
        cl_uint          num_events_in_wait_list,
        const cl_event * event_wait_list,
        cl_event *       event)

    这会获得由Direct3D 10资源创建OpenCL内存对象。clEnqueueAcquireD3D10ObjectsKHR()提供了同步保证,在调用clEnqueueAcquireD3D10ObjectsKHR()之前做出的所有D3D 10调用都必须先完全执行,之后event才能报告完成,command_queue中的所有后续OpenCL工作才能开始执行。

           释放函数为:

    cl_int clEnqueueReleaseD3D10ObjectsKHR(
        cl_command_queue command_queue,
        cl_uint          num_objects,
        const cl_mem *   mem_objects,
        cl_uint          num_events_in_wait_list,
        const cl_event * event_wait_list,
        cl_event *       event)

    这会获得由Direct3D 10资源创建OpenCL内存对象。clEnqueueReleaseD3D10ObjectsKHR()提供了同步保证,在调用clEnqueueReleaseD3D10ObjectsKHR()之后做出的所有D3D 10调用不会立即开始执行,直到event_wait_list中所有事件都已经完成,而且提交到command_queue中的所有工作都已经完成执行之后这些D3D 10调用才会开始。

        另外,与D3D10不同,OpenGL获取函数不会提供同步保证。另外,获取和释放纹理时,最高效的做法是同时获取和释放所有共享的纹理和资源。另外,最好在切换回D3D处理之前处理完所有opencl内核。采用这种方式,获取和释放调用可以用来构成opencl和D3D处理的边界。

    4.OpenCL中处理D3D纹理

          opencl修改纹理内容:

    cl_int computeTexture()
    {
        cl_int errNum;
    
        static cl_int seq =0;
        seq = (seq+1)%(g_WindowWidth*2);
    
        errNum = clSetKernelArg(tex_kernel, 0, sizeof(cl_mem), &g_clTexture2D);
        errNum = clSetKernelArg(tex_kernel, 1, sizeof(cl_int), &g_WindowWidth);
        errNum = clSetKernelArg(tex_kernel, 2, sizeof(cl_int), &g_WindowHeight);
        errNum = clSetKernelArg(tex_kernel, 3, sizeof(cl_int), &seq);
        
        size_t tex_globalWorkSize[2] = { g_WindowWidth, g_WindowHeight };
        size_t tex_localWorkSize[2] = { 32, 4 } ;
    
        errNum = clEnqueueAcquireD3D10ObjectsKHR(commandQueue, 1, &g_clTexture2D, 0, NULL, NULL );
    
        errNum = clEnqueueNDRangeKernel(commandQueue, tex_kernel, 2, NULL,
                                        tex_globalWorkSize, tex_localWorkSize,
                                        0, NULL, NULL);
        if (errNum != CL_SUCCESS)
        {
            std::cerr << "Error queuing kernel for execution." << std::endl;
        }
        errNum = clEnqueueReleaseD3D10ObjectsKHR(commandQueue, 1, &g_clTexture2D, 0, NULL, NULL );
        clFinish(commandQueue);
        return 0;
    }

    用opencl内核计算生成一个D3D纹理对象的内容:

    __kernel void xyz_init_texture_kernel(__write_only image2d_t im, int w, int h, int seq )
    {
        int2 coord = { get_global_id(0), get_global_id(1) };
        float4 color =  { 
                          (float)coord.x/(float)w,
                          (float)coord.y/(float)h,
                          (float)abs(seq-w)/(float)w,
                          1.0f};
        write_imagef( im, coord, color );
    }

    这个纹理使用write_imagef()函数写至内核。这里seq是一个序列号变量,在宿主机上每一帧会循环递增,并发送至内核。在内核中,seq变量用于生成纹理颜色值。seq递增时,颜色会改变来实现纹理动画。

            另外,代码中使用了一种渲染技术g_pTechnique。这是一个基本处理管线,会用到一个简单的顶点着色器,将顶点和纹理坐标传递到一个像素着色器:

    //
    // Vertex Shader
    //
    PS_INPUT VS( VS_INPUT input )
    {
        PS_INPUT output = (PS_INPUT)0;
        output.Pos = input.Pos;
        output.Tex = input.Tex;
            
        return output;
    }
    
    technique10 Render
    {
        pass P0
        {
            SetVertexShader( CompileShader( vs_4_0, VS() ) );
            SetGeometryShader( NULL );
            SetPixelShader( CompileShader( ps_4_0, PS() ) );
        }
    }

            这个技术使用常规的D3D10调用加载。像素着色器再对OpenCL内核修改的纹理完成纹理查找,比提供显示:

    SamplerState samLinear
    {
        Filter = MIN_MAG_MIP_LINEAR;
        AddressU = Wrap;
        AddressV = Wrap;
    };
    
    float4 PS( PS_INPUT input) : SV_Target
    {
        return txDiffuse.Sample( samLinear, input.Tex );
    }

    在像素着色器中,samLinear是输入纹理的一个线性采样器。对于渲染循环的每次迭代,OpenCL在computeTexture()中更新纹理内容,有D3D10显示更新的纹理。

    5.OpenCL中处理D3D顶点数据

         现考虑 使用一个包含顶点数据的D3D缓冲区在屏幕上绘制一个正弦曲线。首先为D3D中的顶点缓冲区定义一个简单的结构:

    struct SimpleSineVertex
    {
        D3DXVECTOR4 Pos;
    };

    可以为这个结构创建一个D3D10缓冲区,这里缓冲区中包含256个元素:

    bd.Usage = D3D10_USAGE_DEFAULT;
    bd.ByteWidth = sizeof( SimpleSineVertex ) * 256;
    bd.BindFlags = D3D10_BIND_VERTEX_BUFFER;
    bd.CPUAccessFlags = 0;
    bd.MiscFlags = 0;
    
    hr = g_pD3DDevice->CreateBuffer( &bd, NULL, &g_pSineVertexBuffer );

    因为要使用OpenCL设置缓冲区中的数据,所以为第二个参数pInitialData传入NULL,只分配空间。一旦创建了D3D缓冲区 g_pSineVertexBuffer,可以使用clCreateFromD3D10BufferKHR()函数从g_pSineVertexBuffer创建一个OpenCL缓冲区:

    g_clBuffer = clCreateFromD3D10BufferKHR( context, CL_MEM_READ_WRITE, g_pSineVertexBuffer, &errNum );
    if( errNum != CL_SUCCESS)
    {
    
        std::cerr << "Error creating buffer from D3D10" << std::endl;
        return E_FAIL;
    }

    与前类似,g_clBuffer可以作为一个内核参数发送到一个生产数据的OpenCL内核。 在示例代码中,正弦曲线的顶点位置在内核中生成:

    __kernel void init_vbo_kernel(__global float4 *vbo, int w, int h, int seq)
    {
        int gid = get_global_id(0);
        float4 linepts;
        float f = 1.0f;
        float a = 0.4f;
        float b = 0.0f;
    
        linepts.x = gid/(w/2.0f)-1.0f;
        linepts.y = b + a*sin(3.14*2.0*((float)gid/(float)w*f + (float)seq/(float)w));
        linepts.z = 0.5f;
        linepts.w = 0.0f;
    
        vbo[gid] = linepts;
    }

    渲染时,设置布局和缓冲区,并指定一个线条带。接下来,computeBuffer()调用前面的内核更新缓冲区。激活一个简单的渲染管线,并绘制256个数据点:

    // Set the input layout
    g_pD3DDevice->IASetInputLayout( g_pSineVertexLayout );
    // Set vertex buffer
    stride = sizeof( SimpleSineVertex );
    offset = 0;
    g_pD3DDevice->IASetVertexBuffers( 0, 1, &g_pSineVertexBuffer, &stride, &offset );
    
    // Set primitive topology
    g_pD3DDevice->IASetPrimitiveTopology( D3D10_PRIMITIVE_TOPOLOGY_LINESTRIP );
    
    computeBuffer();
    g_pTechnique->GetPassByIndex( 1 )->Apply( 0 );
    g_pD3DDevice->Draw( 256, 0 );

            运行时,程序会应用这个内核生成纹理内容,然后运行D3D管线对纹理采样,并在屏幕上显示。然后还会绘制顶点缓冲区,在屏幕上得到一个正弦曲线。

    image

    示例工程源码:http://download.csdn.net/download/qq_33892166/9867159

  • 相关阅读:
    【TouchGFX】使用CubeMX创建touchgfx项目 -- 初始篇
    opencv haarcascades 下载
    更换 Anaconda 源(贼快)
    yii2 允许跨域
    sublime text 3 安装 pyv8 失败的解决办法
    win10 anaconda cuda11.1 安装 tensorlfow-gpu 环境
    nginx 403转404
    python requests 全部异常
    win10 anaconda 安装 tensorflow-gpu 及 jupyter notebook
    【PHP】自有图片服务器,图片动态裁剪缩放
  • 原文地址:https://www.cnblogs.com/betterwgo/p/6947043.html
Copyright © 2011-2022 走看看