首页 > > 程序设计 > C/C++ >

《OpenCL编程指南》之 与Direct3D互操作

2018-06-17 22:27:49来源：未知 阅读 ()

新老客户大回馈,云服务器低至5折

    介绍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\n");
} 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\n");
} else if( errNum == CL_DEVICE_NOT_FOUND) {
    printf("No OpenCL devices corresponding to the d3d_object were found\n");
}

代码为选择的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

标签：

版权申明：本站文章部分自网络，如有侵权，请联系：west999com@outlook.com
特别注意：本站所有转载文章言论不代表本站观点，本站所提供的摄影照片，插画，设计作品，如需使用，请与原作者联系，版权归原作者所有

上一篇：C语言进制转换的一个小错误

下一篇：Leetcode: 85. Maximal Rectangle

相关文章

IDC资讯： 主机资讯 注册资讯 托管资讯 vps资讯 网站建设

网站运营： 建站经验 策划盈利 搜索优化 网站推广 免费资源

网站联盟： 联盟新闻 联盟介绍 联盟点评 网赚技巧

行业资讯： 搜索引擎 网络游戏 电子商务 广告传媒

网络编程： Asp.Net编程 Asp编程 Php编程 Xml编程 Access Mssql Mysql 其它

服务器技术： Web服务器 Ftp服务器 Mail服务器 Dns服务器 安全防护

软件技巧： 其它软件 Word Excel Powerpoint Ghost Vista QQ空间 QQ FlashGet 迅雷

网页制作： FrontPages Dreamweaver Javascript css photoshop fireworks Flash

程序设计： Java技术 C/C++ VB delphi

网络知识： 网络协议 网络安全 网络管理 组网方案 Cisco技术

操作系统： Win2000 WinXP Win2003 Mac OS Linux FreeBSD

热门词条
最新资讯
热门关注
热门标签