以编程方式选择 OpenCL 中最佳 GPU 的最佳方法是什么?
What is the best way to programmatically choose the best GPU in OpenCL?
在我的笔记本电脑上，我有两个显卡 - Intel Iris 和 Nvidia GeForce GT 750M。我正在尝试使用 OpenCL 进行简单的矢量添加。我知道，Nvidia卡要快得多，可以更好地完成工作。原则上，我可以在代码中放置一个 if 语句，该语句将在 VENDOR 属性中查找NVIDIA。但我想有一些优雅的东西。在 OpenCL C/C++ 中以编程方式选择更好(更快(GPU 的最佳方法是什么？
我开发了一个实时光线追踪器(不仅仅是光线投射器(，它以编程方式选择两个GPU和一个CPU，并实时渲染和平衡所有三个GPU的负载。 这是我是如何做到的。
假设有三种设备，d1、d2和d3。为每个设备分配一个权重：w1、w2 和 w3 。调用要渲染的像素数n 。假设有一个名为 alpha 的自由参数。
- 为每个设备分配 1/3 的权重。
- 让
alpha = 0.5. - 在
d1上渲染前n1=w1*n像素，在d2上渲染接下来的n2=w2*n像素，在d3上渲染最后n3=w3*n像素，并记录每个 deivcet1、t2和t3的渲染时间。 - 计算一个值
vsum = n1/t1 + n2/t2 + n3/t3。 - 重新计算
w_i = alpha*w_i + (1-alpha)*n_i/t_i/vsum权重。 - 返回到步骤 3。
值alpha的要点是允许平滑过渡。而不是根据混合一些旧权重的时间重新分配所有权重。如果不使用alpha我就会不稳定。可以调整alpha的值。在实践中，它可能可以设置为 1% 左右，但不能设置为 0%。
让我们举个例子。
我有一个GTX 590，这是一个双GPU卡，有两个低于频率的GTX580。我还有一个Sandy Bridge 2600K处理器。 GPU比CPU快得多。 让我们假设它们的速度快了大约 10 倍。假设有 900 像素。
使用 GPU1 渲染前 300 像素、使用 GPU2 渲染接下来的 300 像素和使用 CPU1 渲染后 300 像素，并分别记录10 s, 10 s, and 100 s的时间。因此，一个GPU用于整个图像需要30秒，而仅CPU就需要300秒。两个 GPU 一起需要15 s .
计算vsum = 30 + 30 + 3 = 63 .再次重新计算权重： w1,w2 = 0.5*(1/3) + 0.5*300/10/63 = 0.4和w3 = 0.5*(1/3) + 0.5*300/100/63 = 0.2 .
渲染下一帧：使用 GPU1 渲染 360 像素，使用 GPU2 渲染 360 像素，使用 CPU1 渲染 180 像素，时间变得更加平衡，比如说11 s, 11 s, and 55 s。
经过多个帧后，(1-alpha)项占主导地位，直到最终权重都基于该项。在这种情况下，权重分别变为 47%(427 像素(、47%、6%(46 像素(，时间分别变为 14 s, 14 s, 14 s。 在这种情况下，CPU 仅将仅使用 GPU 的结果提高了一秒钟。
我在这个计算中假设了一个统一的负载。在真实的光线追踪器中，负载因扫描线和像素而异，但算法在确定权重方面保持不变。
在实践中，一旦发现权重，它们就不会有太大变化，除非场景的负载发生显着变化，例如，如果场景的一个区域具有高折射和反射，而其余区域是漫射的，但即使在这种情况下，我也会限制树的深度，因此这不会产生显着效果。
通过循环将此方法轻松扩展到多个设备。我曾经在四台设备上测试了我的光线追踪器。两个 12 核至强 CPU 和两个 GPU。在这种情况下，CPU具有更大的影响力，但GPU仍然占主导地位。
万一有人想知道。我为每个设备创建了一个上下文，并在单独的线程中使用了每个上下文(使用 pthreads(。对于三个设备，我使用了三个线程。
实际上，您可以使用它在不同供应商的同一设备上运行。例如，我在 2600K 上同时使用了 AMD 和 Intel CPU 驱动程序(每个驱动程序生成大约一半的帧(，以查看哪个供应商更好。当我第一次这样做时(2012 年(，如果我没记错的话，具有讽刺意味的是，AMD 在英特尔 CPU 上击败了英特尔。
如果有人对我如何提出权重公式感兴趣，我使用了物理学的想法(我的背景是物理学而不是编程(。
速度(v(=距离/时间。在这种情况下，距离(d(是要处理的像素数。然后总距离为
d = v1*t1 + v2*t2 + v3*t3
我们希望他们每个人都在同一时间内完成，所以
d = (v1 + v2 + v3)*t
然后得到权重定义
v_i*t = w_i*d
这给了
w_i = v_i*t/d
并将 ( t/d ( 从 ( d = (v1 + v2 + v3)*t ( 替换得到：
w_i = v_i /(v1 + v2 + v3)
很容易看出，这可以推广到任意数量的设备k
w_i = v_i/(v1 + v2 + ...v_k)
所以vsum在我的算法中代表"速度的总和"。最后，由于v_i是随着时间的推移而像素，因此n_i/t_i最终给出了
w_i = n_i/t_i/(n1/t1 + n2/t2 + ...n_k/t_k)
这是我公式中计算权重的第二项。
好：只需选择第一个兼容的设备。在大多数系统上，只有一个。
更好：通过将CL_DEVICE_MAX_COMPUTE_UNITS设备信息结果乘以CL_DEVICE_MAX_CLOCK_FREQUENCY设备信息结果，可以非常粗略地估计设备性能。根据您的工作负载，您可能希望包含其他指标，例如内存大小。您可以根据工作负载混合这些内容。
最佳：在每台设备上使用您的确切工作流程进行基准测试。这确实是唯一确定的方法，因为其他任何事情都只是猜测。
最后，用户可能关心您正在使用哪个 GPU，因此无论您选择哪种方法，都应该有某种方法来覆盖您的编程选择。
如果它只是一个矢量添加，并且您的应用程序驻留在主机端，则 cpu 将获胜。甚至更好的是，集成CPU将更快。整体性能取决于算法、opencl 缓冲区类型(use_host_ptr、read_write 等(和计算与数据比率。即使不复制而是固定阵列和访问，CPU 的延迟也会小于 pci-e 延迟。
如果要使用 opengl + opencl 互操作，则需要知道计算设备是否与渲染输出设备相同。(如果您的屏幕从 igpu 获取数据，那么它是虹膜，如果不是，那么它是 NVIDIA(
如果您只需要在c ++数组(主机端(上执行一些操作并以最快的方式获得结果，那么我建议您使用"负载平衡"。
在配备 Iris pro 和两个 gt750m 的酷睿 i7-5775C 上矢量添加 4k 元素的示例(一个超频 10%(
首先，为所有设备提供相同数量的 ndrange 愤怒。在每个计算阶段结束时，检查计时。
CPU iGPU dGPU-1 dGPU-2 oc
Intel Intel Nvidia Nvidia
1024 1024 1024 1024
34 ms 5ms 10ms 9ms
然后计算加权(取决于最后一个 ndrange 范围(但放宽(不精确但接近(计算带宽的近似值，并相应地更改 ndrange 范围：
Intel Intel Nvidia Nvidia
512 1536 1024 1024
16 ms 8ms 10ms 9ms
然后继续计算，直到它真正变得稳定。
Intel Intel Nvidia Nvidia
256 1792 1024 1024
9ms 10ms 10ms 9ms
或者直到您可以启用更细的谷物。
Intel Intel Nvidia Nvidia
320 1728 1024 1024
10ms 10ms 10ms 9ms
Intel Intel Nvidia Nvidia
320 1728 960 1088
10ms 10ms 10ms 10ms
^ ^
| |
| PCI-E bandwidth not more than 16 GB/s per device
closer to RAM, better bandwidth (20-40 GB/s) and less kernel overhead
您可以获取最近 10 个结果的平均值(或 PID(，而不是仅获取最新的平衡迭代，以消除误导平衡的峰值。此外，缓冲区副本可能比计算花费更多时间，如果将其包含在平衡中，则可以关闭不必要/无益的设备。
如果你创建了一个库，那么你就不必为你的每一个新项目尝试基准测试。当您加速矩阵乘法、流体运动、sql 表连接和财务近似时，它们将在设备之间自动平衡。
对于平衡的解决方案：
如果可以将线性系统求解为 n 个未知数(每个设备的负载(和 n 个方程(所有设备的基准结果(，则可以一步找到目标负载。如果选择迭代，则需要更多步骤才能收敛。后者并不比编写基准更难。前者对我来说更难，但随着时间的推移应该更有效率。
虽然只向量添加内核不是现实世界的场景，但这是我系统的真实基准：
__kernel void bench(__global float * a, __global float *b, __global float *c)
{
int i=get_global_id(0);
c[i]=a[i]+b[i];
}
2560 768 768
AMD FX(tm)-8150 Eight-Core Processor Oland Pitcairn
这是在几次迭代之后(即使使用额外的缓冲区副本，FX 也更快，不使用任何主机指针(。甚至 oland GPU 也在赶上 pitcairn，因为它们的 pci-e 带宽是相同的。
现在有一些三角函数：
__kernel void bench(__global float * a, __global float *b, __global float *c)
{
int i=get_global_id(0);
c[i]=sin(a[i])+cos(b[i])+sin(cos((float)i));
}
1792 1024 1280
测试 GDDR3-128 位与 GDDR5-256 位(超频(和缓存。
__kernel void bench(__global float * a, __global float *b, __global float *c)
{
int i=get_global_id(0);
c[i]=a[i]+b[i]-a[i]-b[i]+a[i]+b[i]-b[i]-a[i]+b[i];
for(int j=0;j<12000;j++)
c[i]+=a[i]+b[i]-a[i]-b[i]+a[i]+b[i]-b[i]-a[i]+b[i];
}
256 256 3584
高计算数据比：
__kernel void bench(__global float * a, __global float *b, __global float *c)
{
int i=get_global_id(0);
c[i]=0.0f; float c0=c[i];float a0=a[i];float b0=b[i];
for(int j=0;j<12000;j++)
c0+=sin(a0)+cos(b0*a0)+cos(sin(b0)*19.95f);
c[i]=c0;
}
256 2048 1792
现在，Oland GPU 再次值得一提，即使只有 320 个内核，也赢得了胜利。因为 4k 元素很容易缠绕所有 320 个内核超过 10 次，但 pitcairn gpu(1280 个内核(没有完全填充折叠阵列(波前(，这导致执行单元占用率较低--->无法隐藏延迟。我认为低负载的低端设备更好。也许当 directx-12 推出一些负载均衡器时，我可以使用它，这个 Oland 可以从游戏中计算 5000 - 10000 个粒子的物理特性，而 pitcairn 可以计算烟雾密度。
看看这个用于 GPU 识别的代码：
#include <iostream>
#ifdef __APPLE__
#include <OpenCL/cl.h>
#else
#include <CL/cl.h>
#pragma comment (lib, "x86_64/opencl.lib")
#endif
//OpenCL saxpy kernel used for benchmarking
const char* saxpy_kernel =
"__kernel n"
"void saxpy_kernel(float alpha, n"
" __global float *A, n"
" __global float *B, n"
" __global float *C) n"
"{ n"
" int idx = get_global_id(0); n"
" C[idx] = alpha * A[idx] + B[idx]; n"
"} ";
const char* clErrName[] = {
"CL_SUCCESS", //0
"CL_DEVICE_NOT_FOUND", //-1
"CL_DEVICE_NOT_AVAILABLE", //-2
"CL_COMPILER_NOT_AVAILABLE", //-3
"CL_MEM_OBJECT_ALLOCATION_FAILURE", //-4
"CL_OUT_OF_RESOURCES", //-5
"CL_OUT_OF_HOST_MEMORY", //-6
"CL_PROFILING_INFO_NOT_AVAILABLE", //-7
"CL_MEM_COPY_OVERLAP", //-8
"CL_IMAGE_FORMAT_MISMATCH", //-9
"CL_IMAGE_FORMAT_NOT_SUPPORTED", //-10
"CL_BUILD_PROGRAM_FAILURE", //-11
"CL_MAP_FAILURE", //-12
"CL_MISALIGNED_SUB_BUFFER_OFFSET", //-13
"CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST", //-14
"CL_COMPILE_PROGRAM_FAILURE", //-15
"CL_LINKER_NOT_AVAILABLE", //-16
"CL_LINK_PROGRAM_FAILURE", //-17
"CL_DEVICE_PARTITION_FAILED", //-18
"CL_KERNEL_ARG_INFO_NOT_AVAILABLE", //-19
"CL_UNDEFINED_ERROR_20", //-20
"CL_UNDEFINED_ERROR_21", //-21
"CL_UNDEFINED_ERROR_22", //-22
"CL_UNDEFINED_ERROR_23", //-23
"CL_UNDEFINED_ERROR_24", //-24
"CL_UNDEFINED_ERROR_25", //-25
"CL_UNDEFINED_ERROR_26", //-26
"CL_UNDEFINED_ERROR_27", //-27
"CL_UNDEFINED_ERROR_28", //-28
"CL_UNDEFINED_ERROR_29", //-29
"CL_INVALID_VALUE", //-30
"CL_INVALID_DEVICE_TYPE", //-31
"CL_INVALID_PLATFORM", //-32
"CL_INVALID_DEVICE", //-33
"CL_INVALID_CONTEXT", //-34
"CL_INVALID_QUEUE_PROPERTIES", //-35
"CL_INVALID_COMMAND_QUEUE", //-36
"CL_INVALID_HOST_PTR", //-37
"CL_INVALID_MEM_OBJECT", //-38
"CL_INVALID_IMAGE_FORMAT_DESCRIPTOR", //-39
"CL_INVALID_IMAGE_SIZE", //-40
"CL_INVALID_SAMPLER", //-41
"CL_INVALID_BINARY", //-42
"CL_INVALID_BUILD_OPTIONS", //-43
"CL_INVALID_PROGRAM", //-44
"CL_INVALID_PROGRAM_EXECUTABLE", //-45
"CL_INVALID_KERNEL_NAME", //-46
"CL_INVALID_KERNEL_DEFINITION", //-47
"CL_INVALID_KERNEL", //-48
"CL_INVALID_ARG_INDEX", //-49
"CL_INVALID_ARG_VALUE", //-50
"CL_INVALID_ARG_SIZE", //-51
"CL_INVALID_KERNEL_ARGS", //-52
"CL_INVALID_WORK_DIMENSION", //-53
"CL_INVALID_WORK_GROUP_SIZE", //-54
"CL_INVALID_WORK_ITEM_SIZE", //-55
"CL_INVALID_GLOBAL_OFFSET", //-56
"CL_INVALID_EVENT_WAIT_LIST", //-57
"CL_INVALID_EVENT", //-58
"CL_INVALID_OPERATION", //-59
"CL_INVALID_GL_OBJECT", //-60
"CL_INVALID_BUFFER_SIZE", //-61
"CL_INVALID_MIP_LEVEL", //-62
"CL_INVALID_GLOBAL_WORK_SIZE", //-63
"CL_INVALID_PROPERTY", //-64
"CL_INVALID_IMAGE_DESCRIPTOR", //-65
"CL_INVALID_COMPILER_OPTIONS", //-66
"CL_INVALID_LINKER_OPTIONS", //-67
"CL_INVALID_DEVICE_PARTITION_COUNT", //-68
"CL_INVALID_PIPE_SIZE", //-69
"CL_INVALID_DEVICE_QUEUE", //-70
};
const int MAX_ERR_CODE = 70;
inline bool __clCallSuccess(cl_int err_code, const char* source_file, const int source_line)
{
if (err_code == CL_SUCCESS)
return true;
if ((err_code > 0) || (err_code < -MAX_ERR_CODE))
std::clog << "t - unknown CL error: " << err_code;
else
std::clog << "t - CL call error: " << clErrName[-err_code];
std::clog << " [" << source_file << " : " << source_line << "]" << std::endl;
return false;
}
#define clCallSuccess(err_code) __clCallSuccess(err_code, __FILE__, __LINE__)
float cl_BenchmarkDevice(cl_context context, cl_command_queue command_queue, cl_device_id device_id)
{
float microSeconds = -1.;
int i;
cl_int clStatus;
const int VECTOR_SIZE = 512 * 1024;
// Allocate space for vectors A, B and C
float* A = (float*)malloc(sizeof(float) * VECTOR_SIZE); if(A) {
float* B = (float*)malloc(sizeof(float) * VECTOR_SIZE); if(B) {
float* C = (float*)malloc(sizeof(float) * VECTOR_SIZE); if(C) {
for (i = 0; i < VECTOR_SIZE; i++)
{
A[i] = (float)i;
B[i] = (float)(VECTOR_SIZE - i);
C[i] = 0;
}
// Create memory buffers on the device for each vector
cl_mem A_clmem = clCreateBuffer(context, CL_MEM_READ_ONLY, VECTOR_SIZE * sizeof(float), NULL, &clStatus); if (clCallSuccess(clStatus)) {
cl_mem B_clmem = clCreateBuffer(context, CL_MEM_READ_ONLY, VECTOR_SIZE * sizeof(float), NULL, &clStatus); if (clCallSuccess(clStatus)) {
cl_mem C_clmem = clCreateBuffer(context, CL_MEM_WRITE_ONLY, VECTOR_SIZE * sizeof(float), NULL, &clStatus); if (clCallSuccess(clStatus)) {
// Copy the Buffer A and B to the device
clStatus = clEnqueueWriteBuffer(command_queue, A_clmem, CL_TRUE, 0, VECTOR_SIZE * sizeof(float), A, 0, NULL, NULL); if (clCallSuccess(clStatus)) {
clStatus = clEnqueueWriteBuffer(command_queue, B_clmem, CL_TRUE, 0, VECTOR_SIZE * sizeof(float), B, 0, NULL, NULL); if (clCallSuccess(clStatus)) {
// Create a program from the kernel source and build it
cl_program program = clCreateProgramWithSource(context, 1, (const char**)&saxpy_kernel, NULL, &clStatus); if (clCallSuccess(clStatus) && program) {
clStatus = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL); if (clCallSuccess(clStatus)) {
// Create the OpenCL kernel
cl_kernel kernel = clCreateKernel(program, "saxpy_kernel", &clStatus); if (clCallSuccess(clStatus) && kernel) {
float alpha = 2.5;
// Set the arguments of the kernel
clStatus = clSetKernelArg(kernel, 0, sizeof(float), (void*)&alpha); if (clCallSuccess(clStatus)) {
clStatus = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void*)&A_clmem); if (clCallSuccess(clStatus)) {
clStatus = clSetKernelArg(kernel, 2, sizeof(cl_mem), (void*)&B_clmem); if (clCallSuccess(clStatus)) {
clStatus = clSetKernelArg(kernel, 3, sizeof(cl_mem), (void*)&C_clmem); if (clCallSuccess(clStatus)) {
// Execute the OpenCL kernel on the list
cl_event event;
size_t global_size = VECTOR_SIZE; // Process the entire lists
size_t local_size = 512; // Process one item at a time
//clStatus = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, &global_size, &local_size, 0, NULL, &event);
clStatus = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, &global_size, NULL, 0, NULL, &event); if (clCallSuccess(clStatus)) {
clStatus = clWaitForEvents(1, &event); if (clCallSuccess(clStatus)) {
//measure duration
cl_ulong time_start;
cl_ulong time_end;
clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_START, sizeof(time_start), &time_start, NULL);
clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_END, sizeof(time_end), &time_end, NULL);
microSeconds = (float)(time_end - time_start) / 1000.0f;
std::clog << "nOpenCl benchmarking time: " << microSeconds << " microseconds n";
std::clog << "nt*****************************nn";
}
// Read the cl memory C_clmem on device to the host variable C
clCallSuccess(clEnqueueReadBuffer(command_queue, C_clmem, CL_TRUE, 0, VECTOR_SIZE * sizeof(float), C, 0, NULL, NULL));
// Clean up and wait for all the comands to complete.
clCallSuccess(clFlush(command_queue));
clCallSuccess(clFinish(command_queue));
} //Kernel
}}}} //SetKErnelArg
// Finally release all OpenCL allocated objects and host buffers.
clCallSuccess(clReleaseKernel(kernel)); }
} //BuildProgram
clCallSuccess(clReleaseProgram(program)); }
} } //EnqueueWriteBuffer
clCallSuccess(clReleaseMemObject(C_clmem)); }
clCallSuccess(clReleaseMemObject(B_clmem)); }
clCallSuccess(clReleaseMemObject(A_clmem)); }
free(C); }
free(B); }
free(A); }
return microSeconds;
}
/*
struct _dev_info {
cl_platform_id platfID;
cl_device_id devID;
};
typedef struct _dev_info dev_info;
*/
cl_device_id cl_GetBestDevice(void)
{
cl_int err;
cl_uint numPlatforms, numDevices;
cl_platform_id platfIDs[10];
cl_device_id devIDsAll[10];
int countGPUs = 0;
cl_device_id best_device = NULL;
float best_perf = 100000000.;
if (clCallSuccess(clGetPlatformIDs(10, platfIDs, &numPlatforms)))
{
std::clog << "OpenCL platforms detected: " << numPlatforms << std::endl;
for (unsigned int i = 0; i < numPlatforms; i++)
{
std::clog << "PlatformInfo for platform no." << (i + 1) << std::endl;
const int SZ_INFO = 1024;
char info[SZ_INFO];
size_t sz;
if (clCallSuccess(clGetPlatformInfo(platfIDs[i], CL_PLATFORM_NAME, SZ_INFO, info, &sz)))
std::clog << " - - Name: " << info << std::endl;
if (clCallSuccess(clGetPlatformInfo(platfIDs[i], CL_PLATFORM_VENDOR, SZ_INFO, info, &sz)))
std::clog << " - - Vendor: " << info << std::endl;
if (clCallSuccess(clGetPlatformInfo(platfIDs[i], CL_PLATFORM_PROFILE, SZ_INFO, info, &sz)))
std::clog << " - - Profile: " << info << std::endl;
if (clCallSuccess(clGetPlatformInfo(platfIDs[i], CL_PLATFORM_VERSION, SZ_INFO, info, &sz)))
std::clog << " - - Version: " << info << std::endl;
if (clCallSuccess(clGetPlatformInfo(platfIDs[i], CL_PLATFORM_EXTENSIONS, SZ_INFO, info, &sz)))
std::clog << " - - Extensions: " << info << std::endl;
if (clCallSuccess(clGetDeviceIDs(platfIDs[i], CL_DEVICE_TYPE_ALL, 10, devIDsAll, &numDevices)))
{
cl_context_properties cProperties[] = { CL_CONTEXT_PLATFORM, (cl_context_properties)(platfIDs[i]), 0 };
cl_command_queue_properties qProperties[] = { CL_QUEUE_PROPERTIES, CL_QUEUE_PROFILING_ENABLE, 0 };
for (unsigned int ii = 0; ii < numDevices; ii++)
{
cl_uint val;
cl_ulong memsz;
cl_device_type dt;
size_t mws;
std::clog << " >> DeviceInfo for device no." << (ii + 1) << std::endl;
if (clCallSuccess(clGetDeviceInfo(devIDsAll[ii], CL_DEVICE_NAME, SZ_INFO, info, &sz)))
std::clog << "t - Name: " << info << std::endl;
if (clCallSuccess(clGetDeviceInfo(devIDsAll[ii], CL_DEVICE_VENDOR, SZ_INFO, info, &sz)))
std::clog << "t - Vendor: " << info << std::endl;
if (clCallSuccess(clGetDeviceInfo(devIDsAll[ii], CL_DEVICE_VERSION, SZ_INFO, info, &sz)))
std::clog << "t - Version: " << info << std::endl;
if (clCallSuccess(clGetDeviceInfo(devIDsAll[ii], CL_DEVICE_TYPE, sizeof(dt), &dt, &sz)))
{
std::clog << "t - Type: ";
switch (dt)
{
case CL_DEVICE_TYPE_CPU: std::clog << "CPU"; break;
case CL_DEVICE_TYPE_GPU: std::clog << "GPU"; break;
case CL_DEVICE_TYPE_ACCELERATOR: std::clog << "Accelerator"; break;
case CL_DEVICE_TYPE_DEFAULT: std::clog << "Default"; break;
default: std::clog << "ERROR";
}
std::clog << std::endl;
}
if (clCallSuccess(clGetDeviceInfo(devIDsAll[ii], CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(memsz), &memsz, &sz)))
std::clog << "t - Memory: " << (memsz / 1024 / 1024) << " MB" << std::endl;
if (clCallSuccess(clGetDeviceInfo(devIDsAll[ii], CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof(val), &val, &sz)))
std::clog << "t - Max Frequency: " << val << " MHz" << std::endl;
if (clCallSuccess(clGetDeviceInfo(devIDsAll[ii], CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(val), &val, &sz)))
std::clog << "t - Compute units: " << val << std::endl;
if (clCallSuccess(clGetDeviceInfo(devIDsAll[ii], CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(mws), &mws, &sz)))
std::clog << "t - Max workgroup size: " << mws << std::endl;
// Create an OpenCL context
cl_context context = clCreateContext(NULL, 1, devIDsAll+ii, NULL, NULL, &err);
if (clCallSuccess(err) && context)
{
// Create a command queue
cl_command_queue command_queue = clCreateCommandQueueWithProperties(context, devIDsAll[ii], qProperties, &err);
if (clCallSuccess(err) && command_queue)
{
float perf = cl_BenchmarkDevice(context, command_queue, devIDsAll[ii]);
if ((perf > 0) && (perf < best_perf))
{
best_perf = perf;
best_device = devIDsAll[ii];
}
clCallSuccess(clReleaseCommandQueue(command_queue));
}
clCallSuccess(clReleaseContext(context));
}
}
}
}
}
return best_device;
}
这是我电脑上的输出
相关文章：
- 在C#中处理C++指针而不使用unsafe的最佳方法
- 在C++中，将大的无符号浮点数四舍五入为整数的最佳方法是什么
- 实现无开销push_back的最佳方法是什么
- 在c代码之间共享数据的最佳方式
- 使用std::source_location报告错误的最佳实践
- 派生类销毁的最佳实践是什么
- 将寄存器设计成可由C和C++访问的外设的最佳实践
- 在两台机器之间进行时间戳的最佳c++chrono函数是什么
- 在使用GPU支持编译Tensorflow时，会遇到CUDA_TOOLKIT_PATH未绑定变量
- 使用QQuickFramebufferObject时同步数据的最佳方式是什么
- 在C++中向零方向近似的最佳方法
- 使用不同的CRT将新的C++代码与旧的(二进制）组件隔离开来的最佳方法是什么
- 从嵌套在std::映射中的std::列表中删除元素的最佳方式
- 如果条件为TRUE(最佳方式?），则在do while循环中后置增量
- 检测win32服务创建和删除的最佳方法
- 在reactor中存储eventHandlers的最佳方式是什么
- 利用 GPU 的最佳方式
- 在C++-GPU中解决稀疏线性系统的最佳方法可能吗
- 什么是使应用程序具有CPU和GPU计算后端的最佳方法
- 以编程方式选择 OpenCL 中最佳 GPU 的最佳方法是什么?
