OpenCL雙邊濾波實現美顏功能

来源:http://www.cnblogs.com/betterwgo/archive/2017/02/22/6431522.html
-Advertisement-
Play Games

OpenCL是一個並行異構計算的框架，包括intel，AMD，英偉達等等許多廠家都有對它的支持，不過英偉達只到1.2版本，主要發展自己的CUDA去了。雖然沒有用過CUDA，但個人感覺CUDA比OpenCL更好一點，但OpenCL支持面更管，CPU,GPU,DSP,FPGA等多種晶元都能支持OpenC ...


    OpenCL是一個並行異構計算的框架，包括intel，AMD，英偉達等等許多廠家都有對它的支持，不過英偉達只到1.2版本，主要發展自己的CUDA去了。雖然沒有用過CUDA，但個人感覺CUDA比OpenCL更好一點，但OpenCL支持面更管，CPU,GPU,DSP,FPGA等多種晶元都能支持OpenCL。OpenCL與D3D中的像素著色器非常相似。

1.雙邊濾波原理

    雙邊濾波器的原理參考女神Rachel-Zhang的博客 雙邊濾波器的原理及實現. 引自Rachel-Zhang的博客，原理如下：

雙邊濾波（Bilateral filter）是一種可以保邊去噪的濾波器。之所以可以達到此去噪效果，是因為濾波器是由兩個函數構成。一個函數是由幾何空間距離決定濾波器繫數。另一個由像素差值決定濾波器繫數。可以與其相比較的兩個filter：高斯低通濾波器(http://en.wikipedia.org/wiki/Gaussian_filter)和α-截尾均值濾波器（去掉百分率為α的最小值和最大之後剩下像素的均值作為濾波器）。

雙邊濾波器中，輸出像素的值依賴於鄰域像素的值的加權組合，

，

          權重繫數w(i,j,k,l)取決於定義域核和值域核的乘積。同時考慮了空間域與值域的差別，而Gaussian Filter和α均值濾波分別隻考慮了空間域和值域差別。

本文基於這個公式用OpenCL實現雙邊濾波來做美顏。

2.核函數

    磨皮演算法原理參考自http://www.zealfilter.com/portal.php?mod=view&aid=138，其中的膚色檢測演算法不好，我給去掉了，本來還要做個銳化處理的，但發現不做銳化效果也蠻好，所以就先沒做，學下一步的OpenCL時在做銳化。

const sampler_t sampler = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;

kernel void bilateralBlur(read_only image2d_t src,write_only image2d_t dst)  
{
    int x = (int)get_global_id(0);  
    int y = (int)get_global_id(1);  
    if (x >= get_image_width(src) || y >= get_image_height(src))  
        return;  

    int ksize = 11;
    float sigma_d = 3.0;
    float sigma_r = 0.1;

    float4 fij = read_imagef(src, sampler, (int2)(x, y));
    float alpha = 0.2;

    float4 fkl;
    float dkl;
    float4 rkl;
    float4 wkl;

    float4 numerator = (float4)(0.0f,0.0f,0.0f,0.0f);
    float4 denominator = (float4)(1.0f, 1.0f, 1.0f, 1.0f);
    for (int K = -ksize / 2; K <= ksize / 2; K++)
    {
        for (int L = -ksize / 2; L <= ksize / 2; L++)
        {
            fkl = read_imagef(src, sampler, (int2)(x + K, y + L));

            dkl = -(K*K + L*L) / (2 * sigma_d*sigma_d);
            rkl.x = -(fij.x - fkl.x)*(fij.x - fkl.x) / (2 * sigma_r*sigma_r);
            rkl.y = -(fij.y - fkl.y)*(fij.y - fkl.y) / (2 * sigma_r*sigma_r);
            rkl.z = -(fij.z - fkl.z)*(fij.z - fkl.z) / (2 * sigma_r*sigma_r);

            wkl.x = exp(dkl + rkl.x);
            wkl.y = exp(dkl + rkl.y);
            wkl.z = exp(dkl + rkl.z);

            numerator.x += fkl.x * wkl.x;
            numerator.y += fkl.y * wkl.y;
            numerator.z += fkl.z * wkl.z;

            denominator.x += wkl.x;
            denominator.y += wkl.y;
            denominator.z += wkl.z;
        }
    }
    
    float4 gij = (float4)(0.0f, 0.0f, 0.0f, 1.0f);
    if (denominator.x > 0 && denominator.y > 0 && denominator.z)
    {
        gij.x = numerator.x / denominator.x;
        gij.y = numerator.y / denominator.y;
        gij.z = numerator.z / denominator.z;

        //雙邊濾波後再做一個融合
         gij.x = fij.x*alpha + gij.x*(1.0 - alpha);
        gij.y = fij.y*alpha + gij.y*(1.0 - alpha);
        gij.z = fij.z*alpha + gij.z*(1.0 - alpha);
    }

    write_imagef(dst, (int2)(x, y), gij);
}

kernel函數裡面基本就是把數學公式寫出來，可以說是非常簡單的。

3.host端代碼

    OpenCL代碼分為host端的代碼和device端的代碼，kernel是跑在並行設備device上的,host一般適合跑串列的邏輯性強的代碼，device則比較適合用來做計算，如捲積運算。電腦中，通常把CPU當host，把GPU當device。不過實際上CPU也可以作為device，因為intel也是支持OpenCL的。本文以CPU為host，GPU為device。

#include "stdafx.h"

#include <iostream>  
#include <fstream>  
#include <sstream>  
#include <malloc.h> 
#include <string.h>  
#include <opencv2/opencv.hpp>  

#include <CL/cl.h>  
 
 
 //----------獲取OpenCL平臺設備信息---------

void DisplayPlatformInfo(
    cl_platform_id id,
    cl_platform_info name,
    std::string str)
{
    cl_int errNum;
    std::size_t paramValueSize;

    errNum = clGetPlatformInfo(
        id,
        name,
        0,
        NULL,
        &paramValueSize);
    if (errNum != CL_SUCCESS)
    {
        std::cerr << "Failed to find OpenCL platform " << str << "." << std::endl;
        return;
    }

    char * info = (char *)alloca(sizeof(char) * paramValueSize);
    errNum = clGetPlatformInfo(
        id,
        name,
        paramValueSize,
        info,
        NULL);
    if (errNum != CL_SUCCESS)
    {
        std::cerr << "Failed to find OpenCL platform " << str << "." << std::endl;
        return;
    }

    std::cout << "\t" << str << ":\t" << info << std::endl;
}

template<typename T>
void appendBitfield(T info, T value, std::string name, std::string & str)
{
    if (info & value)
    {
        if (str.length() > 0)
        {
            str.append(" | ");
        }
        str.append(name);
    }
}

///
// Display information for a particular device.
// As different calls to clGetDeviceInfo may return
// values of different types a template is used. 
// As some values returned are arrays of values, a templated class is
// used so it can be specialized for this case, see below.
//
template <typename T>
class InfoDevice
{
public:
    static void display(
        cl_device_id id,
        cl_device_info name,
        std::string str)
    {
        cl_int errNum;
        std::size_t paramValueSize;

        errNum = clGetDeviceInfo(
            id,
            name,
            0,
            NULL,
            &paramValueSize);
        if (errNum != CL_SUCCESS)
        {
            std::cerr << "Failed to find OpenCL device info " << str << "." << std::endl;
            return;
        }

        T * info = (T *)alloca(sizeof(T) * paramValueSize);
        errNum = clGetDeviceInfo(
            id,
            name,
            paramValueSize,
            info,
            NULL);
        if (errNum != CL_SUCCESS)
        {
            std::cerr << "Failed to find OpenCL device info " << str << "." << std::endl;
            return;
        }

        // Handle a few special cases
        switch (name)
        {
        case CL_DEVICE_TYPE:
        {
            std::string deviceType;

            appendBitfield<cl_device_type>(
                *(reinterpret_cast<cl_device_type*>(info)),
                CL_DEVICE_TYPE_CPU,
                "CL_DEVICE_TYPE_CPU",
                deviceType);

            appendBitfield<cl_device_type>(
                *(reinterpret_cast<cl_device_type*>(info)),
                CL_DEVICE_TYPE_GPU,
                "CL_DEVICE_TYPE_GPU",
                deviceType);

            appendBitfield<cl_device_type>(
                *(reinterpret_cast<cl_device_type*>(info)),
                CL_DEVICE_TYPE_ACCELERATOR,
                "CL_DEVICE_TYPE_ACCELERATOR",
                deviceType);

            appendBitfield<cl_device_type>(
                *(reinterpret_cast<cl_device_type*>(info)),
                CL_DEVICE_TYPE_DEFAULT,
                "CL_DEVICE_TYPE_DEFAULT",
                deviceType);

            std::cout << "\t\t" << str << ":\t" << deviceType << std::endl;
        }
            break;
        case CL_DEVICE_SINGLE_FP_CONFIG:
        {
            std::string fpType;

            appendBitfield<cl_device_fp_config>(
                *(reinterpret_cast<cl_device_fp_config*>(info)),
                CL_FP_DENORM,
                "CL_FP_DENORM",
                fpType);

            appendBitfield<cl_device_fp_config>(
                *(reinterpret_cast<cl_device_fp_config*>(info)),
                CL_FP_INF_NAN,
                "CL_FP_INF_NAN",
                fpType);

            appendBitfield<cl_device_fp_config>(
                *(reinterpret_cast<cl_device_fp_config*>(info)),
                CL_FP_ROUND_TO_NEAREST,
                "CL_FP_ROUND_TO_NEAREST",
                fpType);

            appendBitfield<cl_device_fp_config>(
                *(reinterpret_cast<cl_device_fp_config*>(info)),
                CL_FP_ROUND_TO_ZERO,
                "CL_FP_ROUND_TO_ZERO",
                fpType);

            appendBitfield<cl_device_fp_config>(
                *(reinterpret_cast<cl_device_fp_config*>(info)),
                CL_FP_ROUND_TO_INF,
                "CL_FP_ROUND_TO_INF",
                fpType);

            appendBitfield<cl_device_fp_config>(
                *(reinterpret_cast<cl_device_fp_config*>(info)),
                CL_FP_FMA,
                "CL_FP_FMA",
                fpType);

#ifdef CL_FP_SOFT_FLOAT
            appendBitfield<cl_device_fp_config>(
                *(reinterpret_cast<cl_device_fp_config*>(info)),
                CL_FP_SOFT_FLOAT,
                "CL_FP_SOFT_FLOAT",
                fpType);
#endif

            std::cout << "\t\t" << str << ":\t" << fpType << std::endl;
        }
        case CL_DEVICE_GLOBAL_MEM_CACHE_TYPE:
        {
            std::string memType;

            appendBitfield<cl_device_mem_cache_type>(
                *(reinterpret_cast<cl_device_mem_cache_type*>(info)),
                CL_NONE,
                "CL_NONE",
                memType);
            appendBitfield<cl_device_mem_cache_type>(
                *(reinterpret_cast<cl_device_mem_cache_type*>(info)),
                CL_READ_ONLY_CACHE,
                "CL_READ_ONLY_CACHE",
                memType);

            appendBitfield<cl_device_mem_cache_type>(
                *(reinterpret_cast<cl_device_mem_cache_type*>(info)),
                CL_READ_WRITE_CACHE,
                "CL_READ_WRITE_CACHE",
                memType);

            std::cout << "\t\t" << str << ":\t" << memType << std::endl;
        }
            break;
        case CL_DEVICE_LOCAL_MEM_TYPE:
        {
            std::string memType;

            appendBitfield<cl_device_local_mem_type>(
                *(reinterpret_cast<cl_device_local_mem_type*>(info)),
                CL_GLOBAL,
                "CL_LOCAL",
                memType);

            appendBitfield<cl_device_local_mem_type>(
                *(reinterpret_cast<cl_device_local_mem_type*>(info)),
                CL_GLOBAL,
                "CL_GLOBAL",
                memType);

            std::cout << "\t\t" << str << ":\t" << memType << std::endl;
        }
            break;
        case CL_DEVICE_EXECUTION_CAPABILITIES:
        {
            std::string memType;

            appendBitfield<cl_device_exec_capabilities>(
                *(reinterpret_cast<cl_device_exec_capabilities*>(info)),
                CL_EXEC_KERNEL,
                "CL_EXEC_KERNEL",
                memType);

            appendBitfield<cl_device_exec_capabilities>(
                *(reinterpret_cast<cl_device_exec_capabilities*>(info)),
                CL_EXEC_NATIVE_KERNEL,
                "CL_EXEC_NATIVE_KERNEL",
                memType);

            std::cout << "\t\t" << str << ":\t" << memType << std::endl;
        }
            break;
        case CL_DEVICE_QUEUE_PROPERTIES:
        {
            std::string memType;

            appendBitfield<cl_device_exec_capabilities>(
                *(reinterpret_cast<cl_device_exec_capabilities*>(info)),
                CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE,
                "CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE",
                memType);

            appendBitfield<cl_device_exec_capabilities>(
                *(reinterpret_cast<cl_device_exec_capabilities*>(info)),
                CL_QUEUE_PROFILING_ENABLE,
                "CL_QUEUE_PROFILING_ENABLE",
                memType);

            std::cout << "\t\t" << str << ":\t" << memType << std::endl;
        }
            break;
        default:
            std::cout << "\t\t" << str << ":\t" << *info << std::endl;
            break;
        }
    }
};

///
// Simple trait class used to wrap base types.
//
template <typename T>
class ArrayType
{
public:
    static bool isChar() { return false; }
};

///
// Specialized for the char (i.e. null terminated string case).
//
template<>
class ArrayType<char>
{
public:
    static bool isChar() { return true; }
};

///
// Specialized instance of class InfoDevice for array types.
//
template <typename T>
class InfoDevice<ArrayType<T> >
{
public:
    static void display(
        cl_device_id id,
        cl_device_info name,
        std::string str)
    {
        cl_int errNum;
        std::size_t paramValueSize;

        errNum = clGetDeviceInfo(
            id,
            name,
            0,
            NULL,
            &paramValueSize);
        if (errNum != CL_SUCCESS)
        {
            std::cerr
                << "Failed to find OpenCL device info "
                << str
                << "."
                << std::endl;
            return;
        }

        T * info = (T *)alloca(sizeof(T) * paramValueSize);
        errNum = clGetDeviceInfo(
            id,
            name,
            paramValueSize,
            info,
            NULL);
        if (errNum != CL_SUCCESS)
        {
            std::cerr
                << "Failed to find OpenCL device info "
                << str
                << "."
                << std::endl;
            return;
        }

        if (ArrayType<T>::isChar())
        {
            std::cout << "\t" << str << ":\t" << info << std::endl;
        }
        else if (name == CL_DEVICE_MAX_WORK_ITEM_SIZES)
        {
            cl_uint maxWorkItemDimensions;

            errNum = clGetDeviceInfo(
                id,
                CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS,
                sizeof(cl_uint),
                &maxWorkItemDimensions,
                NULL);
            if (errNum != CL_SUCCESS)
            {
                std::cerr
                    << "Failed to find OpenCL device info "
                    << "CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS."
                    << std::endl;
                return;
            }

            std::cout << "\t" << str << ":\t";
            for (cl_uint i = 0; i < maxWorkItemDimensions; i++)
            {
                std::cout << info[i] << " ";
            }
            std::cout << std::endl;
        }
    }
};

///
//  Enumerate platforms and display information about them 
//  and their associated devices.
//
void displayInfo(void)
{
    cl_int errNum;
    cl_uint numPlatforms;
    cl_platform_id * platformIds;
    cl_context context = NULL;

    // First, query the total number of platforms
    errNum = clGetPlatformIDs(0, NULL, &numPlatforms);
    if (errNum != CL_SUCCESS || numPlatforms <= 0)
    {
        std::cerr << "Failed to find any OpenCL platform." << std::endl;
        return;
    }

    // Next, allocate memory for the installed plaforms, and qeury 
    // to get the list.
    platformIds = (cl_platform_id *)alloca(sizeof(cl_platform_id) * numPlatforms);
    // First, query the total number of platforms
    errNum = clGetPlatformIDs(numPlatforms, platformIds, NULL);
    if (errNum != CL_SUCCESS)
    {
        std::cerr << "Failed to find any OpenCL platforms." << std::endl;
        return;
    }

    std::cout << "Number of platforms: \t" << numPlatforms << std::endl;
    // Iterate through the list of platforms displaying associated information
    for (cl_uint i = 0; i < numPlatforms; i++) {
        // First we display information associated with the platform
        DisplayPlatformInfo(
            platformIds[i],
            CL_PLATFORM_PROFILE,
            "CL_PLATFORM_PROFILE");
        DisplayPlatformInfo(
            platformIds[i],
            CL_PLATFORM_VERSION,
            "CL_PLATFORM_VERSION");
        DisplayPlatformInfo(
            platformIds[i],
            CL_PLATFORM_VENDOR,
            "CL_PLATFORM_VENDOR");
        DisplayPlatformInfo(
            platformIds[i],
            CL_PLATFORM_EXTENSIONS,
            "CL_PLATFORM_EXTENSIONS");

        // Now query the set of devices associated with the platform
        cl_uint numDevices;
        errNum = clGetDeviceIDs(
            platformIds[i],
            CL_DEVICE_TYPE_ALL,
            0,
            NULL,
            &numDevices);
        if (errNum != CL_SUCCESS)
        {
            std::cerr << "Failed to find OpenCL devices." << std::endl;
            return;
        }

        cl_device_id * devices = (cl_device_id *)alloca(sizeof(cl_device_id) * numDevices);
        errNum = clGetDeviceIDs(
            platformIds[i],
            CL_DEVICE_TYPE_ALL,
            numDevices,
            devices,
            NULL);
        if (errNum != CL_SUCCESS)
        {
            std::cerr << "Failed to find OpenCL devices." << std::endl;
            return;
        }

        std::cout << "\tNumber of devices: \t" << numDevices << std::endl;
        // Iterate through each device, displaying associated information
        for (cl_uint j = 0; j < numDevices; j++)
        {
            InfoDevice<cl_device_type>::display(
                devices[j],
                CL_DEVICE_TYPE,
                "CL_DEVICE_TYPE");

            InfoDevice<cl_uint>::display(
                devices[j],
                CL_DEVICE_VENDOR_ID,
                "CL_DEVICE_VENDOR_ID");

            InfoDevice<cl_uint>::display(
                devices[j],
                CL_DEVICE_MAX_COMPUTE_UNITS,
                "CL_DEVICE_MAX_COMPUTE_UNITS");

            InfoDevice<cl_uint>::display(
                devices[j],
                CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS,
                "CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS");

            InfoDevice<ArrayType<size_t> >::display(
                devices[j],
                CL_DEVICE_MAX_WORK_ITEM_SIZES,
                "CL_DEVICE_MAX_WORK_ITEM_SIZES");

            InfoDevice<std::size_t>::display(
                devices[j],
                CL_DEVICE_MAX_WORK_GROUP_SIZE,
                "CL_DEVICE_MAX_WORK_GROUP_SIZE");

            InfoDevice<cl_uint>::display(
                devices[j],
                CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR,
                "CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR");

            InfoDevice<cl_uint>::display(
                devices[j],
                CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT,
                "CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT");

            InfoDevice<cl_uint>::display(
                devices[j],
                CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT,
                "CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT");

            InfoDevice<cl_uint>::display(
                devices[j],
                CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG,
                "CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG");

            InfoDevice<cl_uint>::display(
                devices[j],
                CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT,
                "CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT");

            InfoDevice<cl_uint>::display(
                devices[j],
                CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE,
                "CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE");

#ifdef CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF

            InfoDevice<cl_uint>::display(
                devices[j],
                CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF,
                "CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF");

            InfoDevice<cl_uint>::display(
                devices[j],
                CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR,
                "CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR");

            InfoDevice<cl_uint>::display(
                devices[j],
                CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT,
                "CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT");

            InfoDevice<cl_uint>::display(
                devices[j],
                CL_DEVICE_NATIVE_VECTOR_WIDTH_INT,
                "CL_DEVICE_NATIVE_VECTOR_WIDTH_INT");

            InfoDevice<cl_uint>::display(
                devices[j],
                CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG,
                "CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG");

            InfoDevice<cl_uint>::display(
                devices[j],
                CL_DEVICE_NATIVE_VECTOR_WIDTH_FLOAT,
                "CL_DEVICE_NATIVE_VECTOR_WIDTH_FLOAT");

            InfoDevice<cl_uint>::display(
                devices[j],
                CL_DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE,
                "CL_DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE");

            InfoDevice<cl_uint>::display(
                devices[j],
                CL_DEVICE_NATIVE_VECTOR_WIDTH_HALF,
                "CL_DEVICE_NATIVE_VECTOR_WIDTH_HALF");
#endif

            InfoDevice<cl_uint>::display(
                devices[j],
                CL_DEVICE_MAX_CLOCK_FREQUENCY,
                "CL_DEVICE_MAX_CLOCK_FREQUENCY");

            InfoDevice<cl_uint>::display(
                devices[j],
                CL_DEVICE_ADDRESS_BITS,
                "CL_DEVICE_ADDRESS_BITS");

            InfoDevice<cl_ulong>::display(
                devices[j],
                CL_DEVICE_MAX_MEM_ALLOC_SIZE,
                "CL_DEVICE_MAX_MEM_ALLOC_SIZE");

            InfoDevice<cl_bool>::display(
                devices[j],
                CL_DEVICE_IMAGE_SUPPORT,
                "CL_DEVICE_IMAGE_SUPPORT");

            InfoDevice<cl_uint>::display(
                devices[j],
                CL_DEVICE_MAX_READ_IMAGE_ARGS,
                "CL_DEVICE_MAX_READ_IMAGE_ARGS");

            InfoDevice<cl_uint>::display(
                devices[j],
                CL_DEVICE_MAX_WRITE_IMAGE_ARGS,
                "CL_DEVICE_MAX_WRITE_IMAGE_ARGS");

            InfoDevice<std::size_t>::display(
                devices[j],
                CL_DEVICE_IMAGE2D_MAX_WIDTH,
                "CL_DEVICE_IMAGE2D_MAX_WIDTH");

            InfoDevice<std::size_t>::display(
                devices[j],
                CL_DEVICE_IMAGE2D_MAX_WIDTH,
                "CL_DEVICE_IMAGE2D_MAX_WIDTH");

            InfoDevice<std::size_t>::display(
                devices[j],
                CL_DEVICE_IMAGE2D_MAX_HEIGHT,
                "CL_DEVICE_IMAGE2D_MAX_HEIGHT");

            InfoDevice<std::size_t>::display(
                devices[j],
                CL_DEVICE_IMAGE3D_MAX_WIDTH,
                "CL_DEVICE_IMAGE3D_MAX_WIDTH");

            InfoDevice<std::size_t>::display(
                devices[j],
                CL_DEVICE_IMAGE3D_MAX_HEIGHT,
                "CL_DEVICE_IMAGE3D_MAX_HEIGHT");

            InfoDevice<std::size_t>::display(
                devices[j],
                CL_DEVICE_IMAGE3D_MAX_DEPTH,
                "CL_DEVICE_IMAGE3D_MAX_DEPTH");

            InfoDevice<cl_uint>::display(
                devices[j],
                CL_DEVICE_MAX_SAMPLERS,
                "CL_DEVICE_MAX_SAMPLERS");

            InfoDevice<std::size_t>::display(
                devices[j],
                CL_DEVICE_MAX_PARAMETER_SIZE,
                "CL_DEVICE_MAX_PARAMETER_SIZE");

            InfoDevice<cl_uint>::display(
                devices[j],
                CL_DEVICE_MEM_BASE_ADDR_ALIGN,
                "CL_DEVICE_MEM_BASE_ADDR_ALIGN");

            InfoDevice<cl_uint>::display(
                devices[j],
                CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE,
                "CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE");

            InfoDevice<cl_device_fp_config>::display(
                devices[j],
                CL_DEVICE_SINGLE_FP_CONFIG,
                "CL_DEVICE_SINGLE_FP_CONFIG");

            InfoDevice<cl_device_mem_cache_type>::display(
                devices[j],
                CL_DEVICE_GLOBAL_MEM_CACHE_TYPE,
                "CL_DEVICE_GLOBAL_MEM_CACHE_TYPE");

            InfoDevice<cl_uint>::display(
                devices[j],
                CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE,
                "CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE");

            InfoDevice<cl_ulong>::display(
                devices[j],
                CL_DEVICE_GLOBAL_MEM_CACHE_SIZE,
                "CL_DEVICE_GLOBAL_MEM_CACHE_SIZE");

            InfoDevice<cl_ulong>::display(
                devices[j],
                CL_DEVICE_GLOBAL_MEM_SIZE,
                "CL_DEVICE_GLOBAL_MEM_SIZE");

            InfoDevice<cl_ulong>::display(
                devices[j],
                CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE,
                "CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE");

            InfoDevice<cl_uint>::display(
                devices[j],
                CL_DEVICE_MAX_CONSTANT_ARGS,
                "CL_DEVICE_MAX_CONSTANT_ARGS");

            InfoDevice<cl_device_local_mem_type>::display(
                devices[j],
                CL_DEVICE_LOCAL_MEM_TYPE,
                "CL_DEVICE_LOCAL_MEM_TYPE");

            InfoDevice<cl_ulong>::display(
                devices[j],
                CL_DEVICE_LOCAL_MEM_SIZE,
                "CL_DEVICE_LOCAL_MEM_SIZE");

            InfoDevice<cl_bool>::display(
                devices[j],
                CL_DEVICE_ERROR_CORRECTION_SUPPORT,
                "CL_DEVICE_ERROR_CORRECTION_SUPPORT");

#ifdef CL_DEVICE_HOST_UNIFIED_MEMORY
            InfoDevice<cl_bool>::display(
                devices[j],
                CL_DEVICE_HOST_UNIFIED_MEMORY,
                "CL_DEVICE_HOST_UNIFIED_MEMORY");
#endif

            InfoDevice<std::size_t>::display(
                devices[j],
                CL_DEVICE_PROFILING_TIMER_RESOLUTION,
                "CL_DEVICE_PROFILING_TIMER_RESOLUTION");

            InfoDevice<cl_bool>::display(
                devices[j],
                CL_DEVICE_ENDIAN_LITTLE,
                "CL_DEVICE_ENDIAN_LITTLE");

            InfoDevice<cl_bool>::display(
                devices[j],
                CL_DEVICE_AVAILABLE,
                "CL_DEVICE_AVAILABLE");

            InfoDevice<cl_bool>::display(
                devices[j],
                CL_DEVICE_COMPILER_AVAILABLE,
                "CL_DEVICE_COMPILER_AVAILABLE");

            InfoDevice<cl_device_exec_capabilities>::display(
                devices[j],
                CL_DEVICE_EXECUTION_CAPABILITIES,
                "CL_DEVICE_EXECUTION_CAPABILITIES");

            InfoDevice<cl_command_queue_properties>::display(
                devices[j],
                CL_DEVICE_QUEUE_PROPERTIES,
                "CL_DEVICE_QUEUE_PROPERTIES");

            InfoDevice<cl_platform_id>::display(
                devices[j],
                CL_DEVICE_PLATFORM,
                "CL_DEVICE_PLATFORM");

            InfoDevice<ArrayType<char> >::display(
                devices[j],
                CL_DEVICE_NAME,
                "CL_DEVICE_NAME");

            InfoDevice<ArrayType<char> >::display(
                devices[j],
                CL_DEVICE_VENDOR,
                "CL_DEVICE_VENDOR");

            InfoDevice<ArrayType<char> >::display(
                devices[j],
                CL_DRIVER_VERSION,
                "CL_DRIVER_VERSION");

            InfoDevice<ArrayType<char> >::display(
                devices[j],
                CL_DEVICE_PROFILE,
                "CL_DEVICE_PROFILE");

            InfoDevice<ArrayType<char> >::display(
                devices[j],
                CL_DEVICE_VERSION,
                "
您的分享是我們最大的動力!

-Advertisement-
Play Games
更多相關文章
  • ASP.NET Core Web API 開發-RESTful API實現
    ASP.NET Core Web API 開發-RESTful API實現 REST 介紹： 符合REST設計風格的Web API稱為RESTful API。 具象狀態傳輸（英文：Representational State Transfer，簡稱REST）是Roy Thomas Fielding博 ...
  • ASP.NET Core CORS 簡單使用
    CORS 全稱"跨域資源共用"（Cross-origin resource sharing）。 跨域就是不同域之間進行數據訪問，比如 a.sample.com 訪問 b.sample.com 中的數據，我們如果不做任何處理的話，就會出現下麵的錯誤： XMLHttpRequest cannot loa ...
  • HTTP Error 503. The service is unavailable
    今天發佈了MVC4網站，設置好IIS後老是提示“HTTP Error 503. The service is unavailable”， 網上找了好多資料都不行，最後改變一下應用池的ID就可以了。 ...
  • ASP.NET Web API（一）：使用初探，GET和POST數據
    概述 REST(Representational State Transfer表述性狀態轉移)而產生的REST API的討論越來越多，微軟在ASP.NET中也添加了Web API的功能。 我們看dudu的文章HttpClient + ASP.NET Web API, WCF之外的另一個選擇知道了博客 ...
  • python爬蟲利器Selenium使用詳解
    簡介： 用pyhon爬取動態頁面時普通的urllib2無法實現，例如下麵的京東首頁，隨著滾動條的下拉會載入新的內容，而urllib2就無法抓取這些內容，此時就需要今天的主角selenium。 Selenium是一個用於Web應用程式測試的工具。Selenium測試直接運行在瀏覽器中，就像真正的用戶在 ...
  • 網站安全的操作步聚
    概念介紹： 明文指加密前的內容，反之密文指加密後的內容。 密陰指加密或解密過程中，參與運算的關鍵數據。可以是一份映射表，也可以是一些演算法參數等。 演算法指加密或解密規則。 信息涉漏包括明文，密文，密陰或演算法涉漏。任何信息通過萬維網傳輸，就等同於該信息已經涉漏。演算法，密文和密陰（私陰）都涉漏的情況下，明 ...
  • Java 容器類練習題（一）
    題目1:創建一個名為Gerbil的類，該類擁有一個整數域gerbilNumber，通過構造器初始化gerbilNumber。創建方法hop()顯示該對象的gerbilNumber，以及“is hopping.” 創建一個ArrayList，並將Gerbil對象添加到該List中。用get()方法遍歷 ...
  • [Linux] PHP程式員玩轉Linux系列-搭建代碼開發環境
    1.PHP程式員玩轉Linux系列-怎麼安裝使用CentOS 2.PHP程式員玩轉Linux系列-lnmp環境的搭建 有些同學可能覺得我寫的都是啥yum安裝的，隨便配置一下而已，沒啥技術含量，我的目的是讓大家能夠以最簡單的方式實現目的，配置也是能不自定義的的統統不配置，不是為了炫耀會編譯安裝，我覺得 ...
一周排行
    -Advertisement-
    Play Games
  • 移動開發（一）：使用.NET MAUI開發第一個安卓APP
    移動開發（一）：使用.NET MAUI開發第一個安卓APP 對於工作多年的C#程式員來說，近來想嘗試開發一款安卓APP，考慮了很久最終選擇使用.NET MAUI這個微軟官方的框架來嘗試體驗開發安卓APP，畢竟是使用Visual Studio開發工具，使用起來也比較的順手，結合微軟官方的教程進行了安卓 ...
  • wpf ToggleButton選中效果和一個登錄界面
    前言 QuestPDF 是一個開源 .NET 庫，用於生成 PDF 文檔。使用了C# Fluent API方式可簡化開發、減少錯誤並提高工作效率。利用它可以輕鬆生成 PDF 報告、發票、導出文件等。 項目介紹 QuestPDF 是一個革命性的開源 .NET 庫，它徹底改變了我們生成 PDF 文檔的方 ...
  • 彈幕樹洞項目功能新增篇
    項目地址 項目後端地址： https://github.com/ZyPLJ/ZYTteeHole 項目前端頁面地址： ZyPLJ/TreeHoleVue (github.com) https://github.com/ZyPLJ/TreeHoleVue 目前項目測試訪問地址： http://tree ...
  • 第27篇 sqlserver2022詳細安裝步驟
    話不多說，直接開乾 一.下載 1.官方鏈接下載： https://www.microsoft.com/zh-cn/sql-server/sql-server-downloads 2.在下載目錄中找到下麵這個小的安裝包 SQL2022-SSEI-Dev.exe，運行開始下載SQL server； 二. ...
  • .NET 開源高性能 MQTT 類庫
    前言 隨著物聯網（IoT）技術的迅猛發展，MQTT（消息隊列遙測傳輸）協議憑藉其輕量級和高效性，已成為眾多物聯網應用的首選通信標準。 MQTTnet 作為一個高性能的 .NET 開源庫，為 .NET 平臺上的 MQTT 客戶端與伺服器開發提供了強大的支持。 本文將全面介紹 MQTTnet 的核心功能 ...
  • Serilog文檔翻譯系列（六） - 可用的接收器、增強器、格式化輸出
    Serilog支持多種接收器用於日誌存儲，增強器用於添加屬性，LogContext管理動態屬性，支持多種輸出格式包括純文本、JSON及ExpressionTemplate。還提供了自定義格式化選項，適用於不同需求。 ...
  • 警惕 Visual Studio 屬性求值副作用導致邏輯不符合預期
    目錄簡介獲取 HTML 文檔解析 HTML 文檔測試參考文章 簡介 動態內容網站使用 JavaScript 腳本動態檢索和渲染數據，爬取信息時需要模擬瀏覽器行為，否則獲取到的源碼基本是空的。 本文使用的爬取步驟如下： 使用 Selenium 獲取渲染後的 HTML 文檔 使用 HtmlAgility ...
  • [使用目前最新版]HybridCLR6.9.0+YooAsset2.2.4實現純C# Unity熱更新方案 (一)
    1.前言 什麼是熱更新 游戲或者軟體更新時，無需重新下載客戶端進行安裝，而是在應用程式啟動的情況下，在內部進行資源或者代碼更新 Unity目前常用熱更新解決方案 HybridCLR，Xlua，ILRuntime等 Unity目前常用資源管理解決方案 AssetBundles，Addressable， ...
  • 在 ASP.NET Core Web API 中使用操作篩選器統一處理通用操作
    本文章主要是在C# ASP.NET Core Web API框架實現向手機發送驗證碼簡訊功能。這裡我選擇是一個互億無線簡訊驗證碼平臺，其實像阿裡雲，騰訊雲上面也可以。 首先我們先去 互億無線 https://www.ihuyi.com/api/sms.html 去註冊一個賬號 註冊完成賬號後，它會送 ...
  • 第28篇 如何.net中實現高效可靠數據同步api
    通過以下方式可以高效，並保證數據同步的可靠性 1.API設計 使用RESTful設計，確保API端點明確，並使用適當的HTTP方法（如POST用於創建，PUT用於更新）。 設計清晰的請求和響應模型，以確保客戶端能夠理解預期格式。 2.數據驗證 在伺服器端進行嚴格的數據驗證，確保接收到的數據符合預期格 ...
所有分類