git.oblomov.eu Git - clinfo/blob - src/clinfo.c

   1 /* Collect all available information on all available devices
   2  * on all available OpenCL platforms present in the system
   3  */
   4
   5 #include <time.h>
   6 #include <string.h>
   7
   8 /* We will want to check for symbols in the OpenCL library.
   9  * On Windows, we must get the module handle for it, on Unix-like
  10  * systems we can just use RTLD_DEFAULT
  11  */
  12 #ifdef _MSC_VER
  13 # include <windows.h>
  14 # define dlsym GetProcAddress
  15 # define DL_MODULE GetModuleHandle("OpenCL")
  16 #else
  17 # include <dlfcn.h>
  18 #ifdef RTLD_DEFAULT
  19 # define DL_MODULE RTLD_DEFAULT
  20 #else
  21 # define DL_MODULE ((void*)0) /* This would be RTLD_DEFAULT */
  22 #endif
  23 #endif
  24
  25 /* Load STDC format macros (PRI*), or define them
  26  * for those crappy, non-standard compilers
  27  */
  28 #include "fmtmacros.h"
  29
  30 // More support for the horrible MS C compiler
  31 #ifdef _MSC_VER
  32 #include "ms_support.h"
  33 #endif
  34
  35 #include "error.h"
  36 #include "memory.h"
  37 #include "strbuf.h"
  38
  39 #include "ext.h"
  40 #include "ctx_prop.h"
  41 #include "info_loc.h"
  42 #include "info_ret.h"
  43 #include "opt_out.h"
  44
  45 #define ARRAY_SIZE(ar) (sizeof(ar)/sizeof(*ar))
  46
  47 #ifndef UNUSED
  48 #define UNUSED(x) x __attribute__((unused))
  49 #endif
  50
  51 struct platform_data {
  52         char *pname; /* CL_PLATFORM_NAME */
  53         char *sname; /* CL_PLATFORM_ICD_SUFFIX_KHR or surrogate */
  54         cl_uint ndevs; /* number of devices */
  55         cl_bool has_amd_offline; /* has cl_amd_offline_devices extension */
  56 };
  57
  58 struct platform_info_checks {
  59         cl_uint plat_version;
  60         cl_bool has_khr_icd;
  61         cl_bool has_amd_object_metadata;
  62         cl_bool has_extended_versioning;
  63         cl_bool has_external_memory;
  64         cl_bool has_semaphore;
  65         cl_bool has_external_semaphore;
  66 };
  67
  68 struct platform_list {
  69         /* Number of platforms in the system */
  70         cl_uint num_platforms;
  71         /* Total number of devices across all platforms */
  72         cl_uint ndevs_total;
  73         /* Number of devices allocated in all_devs array */
  74         cl_uint alloc_devs;
  75         /* Highest OpenCL version supported by any platform.
  76          * If the OpenCL library / ICD loader only supports
  77          * a lower version, problems may arise (such as
  78          * API calls causing segfaults or any other unexpected
  79          * behavior
  80          */
  81         cl_uint max_plat_version;
  82         /* Largest number of devices on any platform */
  83         cl_uint max_devs;
  84         /* Length of the longest platform sname */
  85         size_t max_sname_len;
  86         /* Array of platform IDs */
  87         cl_platform_id *platform;
  88         /* Array of device IDs (across all platforms) */
  89         cl_device_id *all_devs;
  90         /* Array of offsets in all_devs where the devices
  91          * of each platform begin */
  92         cl_uint *dev_offset;
  93         /* Array of clinfo-specific platform data */
  94         struct platform_data *pdata;
  95         /* Array of clinfo-specific platform checks */
  96         struct platform_info_checks *platform_checks;
  97 };
  98
  99 void
 100 init_plist(struct platform_list *plist)
 101 {
 102         plist->num_platforms = 0;
 103         plist->ndevs_total = 0;
 104         plist->alloc_devs = 0;
 105         plist->max_plat_version = 0;
 106         plist->max_devs = 0;
 107         plist->max_sname_len = 0;
 108         plist->platform = NULL;
 109         plist->all_devs = NULL;
 110         plist->dev_offset = NULL;
 111         plist->pdata = NULL;
 112         plist->platform_checks = NULL;
 113 }
 114
 115 void plist_devs_reserve(struct platform_list *plist, cl_uint amount)
 116 {
 117         if (amount > plist->alloc_devs) {
 118                 REALLOC(plist->all_devs, amount, "all devices");
 119                 plist->alloc_devs = amount;
 120         }
 121 }
 122
 123
 124 cl_uint
 125 alloc_plist(struct platform_list *plist, const struct opt_out *output)
 126 {
 127         cl_uint num_platforms = plist->num_platforms;
 128         if (output->null_platform)
 129                 num_platforms += 1;
 130         ALLOC(plist->platform, num_platforms, "platform IDs");
 131         ALLOC(plist->dev_offset, num_platforms, "platform device list offset");
 132         /* The actual sizing for this will change as we gather platform info,
 133          * but assume at least one device per platform
 134          */
 135         plist_devs_reserve(plist, num_platforms);
 136         ALLOC(plist->pdata, num_platforms, "platform data");
 137         ALLOC(plist->platform_checks, num_platforms, "platform checks data");
 138         return num_platforms;
 139 }
 140 void
 141 free_plist(struct platform_list *plist)
 142 {
 143         free(plist->platform);
 144         free(plist->all_devs);
 145         free(plist->dev_offset);
 146         for (cl_uint p = 0 ; p < plist->num_platforms; ++p) {
 147                 free(plist->pdata[p].sname);
 148                 free(plist->pdata[p].pname);
 149         }
 150         free(plist->pdata);
 151         free(plist->platform_checks);
 152         init_plist(plist);
 153 }
 154
 155 const cl_device_id *
 156 get_platform_devs(const struct platform_list *plist, cl_uint p)
 157 {
 158         return plist->all_devs + plist->dev_offset[p];
 159 }
 160
 161 cl_device_id
 162 get_platform_dev(const struct platform_list *plist, cl_uint p, cl_uint d)
 163 {
 164         return get_platform_devs(plist, p)[d];
 165 }
 166
 167 /* Data for the OpenCL library / ICD loader */
 168 struct icdl_data {
 169         /* auto-detected OpenCL version support for the ICD loader */
 170         cl_uint detected_version;
 171         /* OpenCL version support declared by the ICD loader */
 172         cl_uint reported_version;
 173 };
 174
 175 /* line prefix, used to identify the platform/device for each
 176  * device property in RAW output mode */
 177 char *line_pfx;
 178 int line_pfx_len;
 179
 180 #define CHECK_SIZE(ret, loc, val, cmd, ...) do { \
 181         /* check if the issue is with param size */ \
 182         if (output->check_size && ret->err == CL_INVALID_VALUE) { \
 183                 size_t _actual_sz; \
 184                 if (cmd(__VA_ARGS__, 0, NULL, &_actual_sz) == CL_SUCCESS) { \
 185                         REPORT_SIZE_MISMATCH(&(ret->err_str), loc, _actual_sz, sizeof(val)); \
 186                 } \
 187         } \
 188 } while (0)
 189
 190 static const char unk[] = "Unknown";
 191 static const char none[] = "None";
 192 static const char none_raw[] = "CL_NONE";
 193 static const char na[] = "n/a"; // not available
 194 static const char na_wrap[] = "(n/a)"; // not available
 195 static const char core[] = "core";
 196
 197 static const char bytes_str[] = " bytes";
 198 static const char pixels_str[] = " pixels";
 199 static const char images_str[] = " images";
 200
 201 static const char* bool_str[] = { "No", "Yes" };
 202 static const char* bool_raw_str[] = { "CL_FALSE", "CL_TRUE" };
 203 static const char* bool_json_str[] = { "false", "true" };
 204
 205 static const char* endian_str[] = { "Big-Endian", "Little-Endian" };
 206
 207 static const cl_device_type devtype[] = { 0,
 208         CL_DEVICE_TYPE_DEFAULT, CL_DEVICE_TYPE_CPU, CL_DEVICE_TYPE_GPU,
 209         CL_DEVICE_TYPE_ACCELERATOR, CL_DEVICE_TYPE_CUSTOM, CL_DEVICE_TYPE_ALL };
 210
 211 const size_t devtype_count = ARRAY_SIZE(devtype);
 212 /* number of actual device types, without ALL */
 213 const size_t actual_devtype_count = ARRAY_SIZE(devtype) - 1;
 214
 215 static const char* device_type_str[] = { unk, "Default", "CPU", "GPU", "Accelerator", "Custom", "All" };
 216 static const char* device_type_raw_str[] = { unk,
 217         "CL_DEVICE_TYPE_DEFAULT", "CL_DEVICE_TYPE_CPU", "CL_DEVICE_TYPE_GPU",
 218         "CL_DEVICE_TYPE_ACCELERATOR", "CL_DEVICE_TYPE_CUSTOM", "CL_DEVICE_TYPE_ALL"
 219 };
 220
 221 static const char* partition_type_str[] = {
 222         none, "equally", "by counts", "by affinity domain", "by names (Intel)"
 223 };
 224 static const char* partition_type_raw_str[] = {
 225         none_raw,
 226         "CL_DEVICE_PARTITION_EQUALLY_EXT",
 227         "CL_DEVICE_PARTITION_BY_COUNTS_EXT",
 228         "CL_DEVICE_PARTITION_BY_AFFINITY_DOMAIN_EXT",
 229         "CL_DEVICE_PARTITION_BY_NAMES_INTEL_EXT"
 230 };
 231
 232 static const char* atomic_cap_str[] = {
 233         "relaxed", "acquire/release", "sequentially-consistent",
 234         "work-item scope", "work-group scope", "device scope", "all-devices scope"
 235 };
 236 static const char* atomic_cap_raw_str[] = {
 237         "CL_DEVICE_ATOMIC_ORDER_RELAXED",
 238         "CL_DEVICE_ATOMIC_ORDER_ACQ_REL",
 239         "CL_DEVICE_ATOMIC_ORDER_SEQ_CST",
 240         "CL_DEVICE_ATOMIC_SCOPE_WORK_ITEM",
 241         "CL_DEVICE_ATOMIC_SCOPE_WORK_GROUP",
 242         "CL_DEVICE_ATOMIC_SCOPE_DEVICE",
 243         "CL_DEVICE_ATOMIC_SCOPE_ALL_DEVICES"
 244 };
 245 const size_t atomic_cap_count = ARRAY_SIZE(atomic_cap_str);
 246
 247 static const char *device_enqueue_cap_str[] = {
 248         "supported", "replaceable default queue"
 249 };
 250
 251 static const char *device_enqueue_cap_raw_str[] = {
 252         "CL_DEVICE_QUEUE_SUPPORTED",
 253         "CL_DEVICE_QUEUE_REPLACEABLE_DEFAULT"
 254 };
 255 const size_t device_enqueue_cap_count = ARRAY_SIZE(atomic_cap_str);
 256
 257 static const char *command_buffer_str[] = {
 258         "kernel printf", "device side enqueue", "simultaneous use", "out of order",
 259 };
 260
 261 static const char *command_buffer_raw_str[] = {
 262         "CL_COMMAND_BUFFER_CAPABILITY_KERNEL_PRINTF_KHR",
 263         "CL_COMMAND_BUFFER_CAPABILITY_DEVICE_SIDE_ENQUEUE_KHR",
 264         "CL_COMMAND_BUFFER_CAPABILITY_SIMULTANEOUS_USE_KHR",
 265         "CL_COMMAND_BUFFER_CAPABILITY_OUT_OF_ORDER_KHR",
 266 };
 267
 268 const size_t command_buffer_count = ARRAY_SIZE(command_buffer_str);
 269
 270 static const char *mutable_dispatch_str[] = {
 271         "Global Offset",
 272         "Local Offset",
 273         "Local Size",
 274         "Arguments",
 275         "Exec Info",
 276 };
 277
 278 static const char *mutable_dispatch_raw_str[] = {
 279         "CL_MUTABLE_DISPATCH_GLOBAL_OFFSET_KHR",
 280         "CL_MUTABLE_DISPATCH_GLOBAL_SIZE_KHR",
 281         "CL_MUTABLE_DISPATCH_LOCAL_SIZE_KHR",
 282         "CL_MUTABLE_DISPATCH_ARGUMENTS_KHR",
 283         "CL_MUTABLE_DISPATCH_EXEC_INFO_KHR",
 284 };
 285
 286 const size_t mutable_dispatch_count = ARRAY_SIZE(mutable_dispatch_str);
 287
 288 static const char numa[] = "NUMA";
 289 static const char l1cache[] = "L1 cache";
 290 static const char l2cache[] = "L2 cache";
 291 static const char l3cache[] = "L3 cache";
 292 static const char l4cache[] = "L4 cache";
 293
 294 static const char* affinity_domain_str[] = {
 295         numa, l4cache, l3cache, l2cache, l1cache, "next partitionable"
 296 };
 297
 298 static const char* affinity_domain_ext_str[] = {
 299         numa, l4cache, l3cache, l2cache, l1cache, "next fissionable"
 300 };
 301
 302 static const char* affinity_domain_raw_str[] = {
 303         "CL_DEVICE_AFFINITY_DOMAIN_NUMA",
 304         "CL_DEVICE_AFFINITY_DOMAIN_L4_CACHE",
 305         "CL_DEVICE_AFFINITY_DOMAIN_L3_CACHE",
 306         "CL_DEVICE_AFFINITY_DOMAIN_L2_CACHE",
 307         "CL_DEVICE_AFFINITY_DOMAIN_L1_CACHE",
 308         "CL_DEVICE_AFFINITY_DOMAIN_NEXT_PARTITIONABLE"
 309 };
 310
 311 static const char* affinity_domain_raw_ext_str[] = {
 312         "CL_AFFINITY_DOMAIN_NUMA_EXT",
 313         "CL_AFFINITY_DOMAIN_L4_CACHE_EXT",
 314         "CL_AFFINITY_DOMAIN_L3_CACHE_EXT",
 315         "CL_AFFINITY_DOMAIN_L2_CACHE_EXT",
 316         "CL_AFFINITY_DOMAIN_L1_CACHE_EXT",
 317         "CL_AFFINITY_DOMAIN_NEXT_FISSIONABLE_EXT"
 318 };
 319
 320 const size_t affinity_domain_count = ARRAY_SIZE(affinity_domain_str);
 321
 322 static const char *terminate_capability_str[] = {
 323         "Context"
 324 };
 325
 326 static const char *terminate_capability_raw_str[] = {
 327         "CL_DEVICE_TERMINATE_CAPABILITY_CONTEXT_KHR"
 328 };
 329
 330 const size_t terminate_capability_count = ARRAY_SIZE(terminate_capability_str);
 331
 332 static const char *terminate_capability_arm_str[] = {
 333         "Controlled Success",
 334         "Controlled Failurure",
 335         "Query"
 336 };
 337
 338 static const char * terminate_capability_arm_raw_str[] = {
 339         "CL_DEVICE_CONTROLLED_TERMINATION_SUCCESS_ARM",
 340         "CL_DEVICE_CONTROLLED_TERMINATION_FAILURE_ARM",
 341         "CL_DEVICE_CONTROLLED_TERMINATION_QUERY_ARM"
 342 };
 343
 344 const size_t terminate_capability_arm_count = ARRAY_SIZE(terminate_capability_arm_str);
 345
 346 static const char* fp_conf_str[] = {
 347         "Denormals", "Infinity and NANs", "Round to nearest", "Round to zero",
 348         "Round to infinity", "IEEE754-2008 fused multiply-add",
 349         "Support is emulated in software",
 350         "Correctly-rounded divide and sqrt operations"
 351 };
 352
 353 static const char* fp_conf_raw_str[] = {
 354         "CL_FP_DENORM",
 355         "CL_FP_INF_NAN",
 356         "CL_FP_ROUND_TO_NEAREST",
 357         "CL_FP_ROUND_TO_ZERO",
 358         "CL_FP_ROUND_TO_INF",
 359         "CL_FP_FMA",
 360         "CL_FP_SOFT_FLOAT",
 361         "CL_FP_CORRECTLY_ROUNDED_DIVIDE_SQRT"
 362 };
 363
 364 const size_t fp_conf_count = ARRAY_SIZE(fp_conf_str);
 365
 366 static const char* svm_cap_str[] = {
 367         "Coarse-grained buffer sharing",
 368         "Fine-grained buffer sharing",
 369         "Fine-grained system sharing",
 370         "Atomics"
 371 };
 372
 373 static const char* svm_cap_raw_str[] = {
 374         "CL_DEVICE_SVM_COARSE_GRAIN_BUFFER",
 375         "CL_DEVICE_SVM_FINE_GRAIN_BUFFER",
 376         "CL_DEVICE_SVM_FINE_GRAIN_SYSTEM",
 377         "CL_DEVICE_SVM_ATOMICS",
 378 };
 379
 380 const size_t svm_cap_count = ARRAY_SIZE(svm_cap_str);
 381
 382 static const char * intel_usm_cap_str[] = {
 383         "USM access",
 384         "USM atomic access",
 385         "USM concurrent access",
 386         "USM concurrent atomic access"
 387 };
 388
 389 static const char * intel_usm_cap_raw_str[] = {
 390         "CL_UNIFIED_SHARED_MEMORY_ACCESS_INTEL",
 391         "CL_UNIFIED_SHARED_MEMORY_ATOMIC_ACCESS_INTEL",
 392         "CL_UNIFIED_SHARED_MEMORY_CONCURRENT_ACCESS_INTEL",
 393         "CL_UNIFIED_SHARED_MEMORY_CONCURRENT_ATOMIC_ACCESS_INTEL",
 394 };
 395
 396 const size_t intel_usm_cap_count = ARRAY_SIZE(intel_usm_cap_str);
 397
 398 static const char* arm_scheduling_controls_str[] = {
 399         "Kernel batching",
 400         "Work-group batch size",
 401         "Work-group batch size modifier",
 402         "Deferred flush",
 403         "Register allocation",
 404         "Warp throttling",
 405         "Compute unit batch queue size",
 406 };
 407
 408 static const char* arm_scheduling_controls_raw_str[] = {
 409         "CL_DEVICE_SCHEDULING_KERNEL_BATCHING_ARM",
 410         "CL_DEVICE_SCHEDULING_WORKGROUP_BATCH_SIZE_ARM",
 411         "CL_DEVICE_SCHEDULING_WORKGROUP_BATCH_SIZE_MODIFIER_ARM",
 412         "CL_DEVICE_SCHEDULING_DEFERRED_FLUSH_ARM",
 413         "CL_DEVICE_SCHEDULING_REGISTER_ALLOCATION_ARM",
 414         "CL_DEVICE_SCHEDULING_WARP_THROTTLING_ARM",
 415         "CL_DEVICE_SCHEDULING_COMPUTE_UNIT_BATCH_QUEUE_SIZE_ARM",
 416 };
 417
 418 const size_t arm_scheduling_controls_count = ARRAY_SIZE(arm_scheduling_controls_str);
 419
 420 static const char* ext_mem_handle_str[] = {
 421         "Opaque FD",
 422         "Opaqe Win32",
 423         "Opaque Win32 KMT",
 424         "D3D11 Texture",
 425         "D3D11 Texture KMT",
 426         "D3D12 Heap",
 427         "D3D12 Resource",
 428         "DMA buffer"
 429 };
 430
 431 static const char* ext_mem_handle_raw_str[] = {
 432         "CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_FD_KHR",
 433         "CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KHR",
 434         "CL_EXTERNAL_MEMORY_HANDLE_OPAQUE_WIN32_KMT_KHR",
 435         "CL_EXTERNAL_MEMORY_HANDLE_D3D11_TEXTURE_KHR",
 436         "CL_EXTERNAL_MEMORY_HANDLE_D3D11_TEXTURE_KMT_KHR",
 437         "CL_EXTERNAL_MEMORY_HANDLE_D3D12_HEAP_KHR",
 438         "CL_EXTERNAL_MEMORY_HANDLE_D3D12_RESOURCE_KHR",
 439         "CL_EXTERNAL_MEMORY_HANDLE_DMA_BUF_KHR",
 440 };
 441
 442 const size_t ext_mem_handle_count = ARRAY_SIZE(ext_mem_handle_str);
 443 const size_t ext_mem_handle_offset = 0x2060;
 444
 445 static const char* semaphore_type_str[] = {
 446         "Binary"
 447 };
 448 static const char* semaphore_type_raw_str[] = {
 449         "CL_SEMAPHORE_TYPE_BINARY_KHR"
 450 };
 451 const size_t semaphore_type_count = ARRAY_SIZE(semaphore_type_str);
 452 const size_t semaphore_type_offset = 1;
 453
 454 static const char* semaphore_handle_str[] = {
 455         "Opaque FD",
 456         "Opaque Win32",
 457         "Opaque Win32 KMT",
 458         "Sync FD",
 459         "D3D12 Fence"
 460 };
 461 static const char* semaphore_handle_raw_str[] = {
 462         "CL_SEMAPHORE_HANDLE_OPAQUE_FD_KHR",
 463         "CL_SEMAPHORE_HANDLE_OPAQUE_WIN32_KHR",
 464         "CL_SEMAPHORE_HANDLE_OPAQUE_WIN32_KMT_KHR",
 465         "CL_SEMAPHORE_HANDLE_SYNC_FD_KHR",
 466         "CL_SEMAPHORE_HANDLE_D3D12_FENCE_KHR",
 467 };
 468 const size_t semaphore_handle_count = ARRAY_SIZE(semaphore_handle_str);
 469 const size_t semaphore_handle_offset = 0x2055;
 470
 471 /* SI suffixes for memory sizes. Note that in OpenCL most of them are
 472  * passed via a cl_ulong, which at most can mode 16 EiB, but hey,
 473  * let's be forward-thinking ;-)
 474  */
 475 static const char* memsfx[] = {
 476         "B", "KiB", "MiB", "GiB", "TiB", "PiB", "EiB", "ZiB", "YiB"
 477 };
 478
 479 const size_t memsfx_end = ARRAY_SIZE(memsfx) + 1;
 480
 481 static const char* lmem_type_str[] = { none, "Local", "Global" };
 482 static const char* lmem_type_raw_str[] = { none_raw, "CL_LOCAL", "CL_GLOBAL" };
 483 static const char* cache_type_str[] = { none, "Read-Only", "Read/Write" };
 484 static const char* cache_type_raw_str[] = { none_raw, "CL_READ_ONLY_CACHE", "CL_READ_WRITE_CACHE" };
 485
 486 static const char* queue_prop_str[] = { "Out-of-order execution", "Profiling" };
 487 static const char* queue_prop_raw_str[] = {
 488         "CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE",
 489         "CL_QUEUE_PROFILING_ENABLE"
 490 };
 491
 492 const size_t queue_prop_count = ARRAY_SIZE(queue_prop_str);
 493
 494 static const char* intel_queue_cap_str[] = {
 495         "create single-queue events",
 496         "create cross-queue events",
 497         "single-queue wait list",
 498         "cross-queue wait list",
 499         "unknown (bit 4)",
 500         "unknown (bit 5)",
 501         "unknown (bit 6)",
 502         "unknown (bit 7)",
 503         "transfer buffer",
 504         "transfer buffer rect",
 505         "map buffer",
 506         "fill buffer",
 507         "transfer image",
 508         "map image",
 509         "fill image",
 510         "transfer buffer to image",
 511         "transfer image to buffer",
 512         "unknown (bit 17)",
 513         "unknown (bit 18)",
 514         "unknown (bit 19)",
 515         "unknown (bit 20)",
 516         "unknown (bit 21)",
 517         "unknown (bit 22)",
 518         "unknown (bit 23)",
 519         "marker enqueue",
 520         "barrier enqueue",
 521         "kernel enqueue",
 522         "unknown (bit 27)",
 523         "unknown (bit 28)",
 524         "no sync operations",
 525 };
 526
 527 static const char* intel_queue_cap_raw_str[] = {
 528         "CL_QUEUE_CAPABILITY_CREATE_SINGLE_QUEUE_EVENTS_INTEL",
 529         "CL_QUEUE_CAPABILITY_CREATE_CROSS_QUEUE_EVENTS_INTEL",
 530         "CL_QUEUE_CAPABILITY_SINGLE_QUEUE_EVENT_WAIT_LIST_INTEL",
 531         "CL_QUEUE_CAPABILITY_CROSS_QUEUE_EVENT_WAIT_LIST_INTEL",
 532         "CL_QUEUE_CAPABILITY_UNKNOWN_4",
 533         "CL_QUEUE_CAPABILITY_UNKNOWN_5",
 534         "CL_QUEUE_CAPABILITY_UNKNOWN_6",
 535         "CL_QUEUE_CAPABILITY_UNKNOWN_7",
 536         "CL_QUEUE_CAPABILITY_TRANSFER_BUFFER_INTEL",
 537         "CL_QUEUE_CAPABILITY_TRANSFER_BUFFER_RECT_INTEL",
 538         "CL_QUEUE_CAPABILITY_MAP_BUFFER_INTEL",
 539         "CL_QUEUE_CAPABILITY_FILL_BUFFER_INTEL",
 540         "CL_QUEUE_CAPABILITY_TRANSFER_IMAGE_INTEL",
 541         "CL_QUEUE_CAPABILITY_MAP_IMAGE_INTEL",
 542         "CL_QUEUE_CAPABILITY_FILL_IMAGE_INTEL",
 543         "CL_QUEUE_CAPABILITY_TRANSFER_BUFFER_IMAGE_INTEL",
 544         "CL_QUEUE_CAPABILITY_TRANSFER_IMAGE_BUFFER_INTEL",
 545         "CL_QUEUE_CAPABILITY_UNKNOWN_17",
 546         "CL_QUEUE_CAPABILITY_UNKNOWN_18",
 547         "CL_QUEUE_CAPABILITY_UNKNOWN_19",
 548         "CL_QUEUE_CAPABILITY_UNKNOWN_20",
 549         "CL_QUEUE_CAPABILITY_UNKNOWN_21",
 550         "CL_QUEUE_CAPABILITY_UNKNOWN_22",
 551         "CL_QUEUE_CAPABILITY_UNKNOWN_23",
 552         "CL_QUEUE_CAPABILITY_MARKER_INTEL",
 553         "CL_QUEUE_CAPABILITY_BARRIER_INTEL",
 554         "CL_QUEUE_CAPABILITY_KERNEL_INTEL",
 555         "CL_QUEUE_CAPABILITY_UNKNOWN_27",
 556         "CL_QUEUE_CAPABILITY_UNKNOWN_28",
 557         "CL_QUEUE_NO_SYNC_OPERATIONS_INTEL",
 558 };
 559
 560 const size_t intel_queue_cap_count = ARRAY_SIZE(intel_queue_cap_str);
 561
 562 static const char* execap_str[] = { "Run OpenCL kernels", "Run native kernels" };
 563 static const char* execap_raw_str[] = {
 564         "CL_EXEC_KERNEL",
 565         "CL_EXEC_NATIVE_KERNEL"
 566 };
 567
 568 const size_t execap_count = ARRAY_SIZE(execap_str);
 569
 570 static const char* intel_features_str[] = { "DP4A", "DPAS" };
 571 static const char* intel_features_raw_str[] = { "CL_DEVICE_FEATURE_FLAG_DP4A_INTEL", "CL_DEVICE_FEATURE_FLAG_DPAS_INTEL" };
 572
 573 const size_t intel_features_count = ARRAY_SIZE(intel_features_str);
 574
 575 static const char* sources[] = {
 576         "#define GWO(type) global type* restrict\n",
 577         "#define GRO(type) global const type* restrict\n",
 578         "#define BODY int i = get_global_id(0); out[i] = in1[i] + in2[i]\n",
 579         "#define _KRN(T, N) kernel void sum##N(GWO(T##N) out, GRO(T##N) in1, GRO(T##N) in2) { BODY; }\n",
 580         "#define KRN(N) _KRN(float, N)\n",
 581         "KRN()\n/* KRN(2)\nKRN(4)\nKRN(8)\nKRN(16) */\n",
 582 };
 583
 584 const char *num_devs_header(const struct opt_out *output, cl_bool these_are_offline)
 585 {
 586         return output->mode == CLINFO_HUMAN ?
 587                 (these_are_offline ? "Number of offine devices (AMD)" : "Number of devices") :
 588                 (these_are_offline ? "#OFFDEVICES" : "#DEVICES");
 589 }
 590
 591 const char *not_specified(const struct opt_out *output)
 592 {
 593         return output->mode == CLINFO_HUMAN ?
 594                 na_wrap : "";
 595 }
 596
 597 const char *no_plat(const struct opt_out *output)
 598 {
 599         return output->mode == CLINFO_HUMAN ?
 600                 "No platform" :
 601                 "CL_INVALID_PLATFORM";
 602 }
 603
 604 const char *invalid_dev_type(const struct opt_out *output)
 605 {
 606         return output->mode == CLINFO_HUMAN ?
 607                 "Invalid device type for platform" :
 608                 "CL_INVALID_DEVICE_TYPE";
 609 }
 610
 611 const char *invalid_dev_value(const struct opt_out *output)
 612 {
 613         return output->mode == CLINFO_HUMAN ?
 614                 "Invalid device type value for platform" :
 615                 "CL_INVALID_VALUE";
 616 }
 617
 618 const char *no_dev_found(const struct opt_out *output)
 619 {
 620         return output->mode == CLINFO_HUMAN ?
 621                 "No devices found in platform" :
 622                 "CL_DEVICE_NOT_FOUND";
 623 }
 624
 625 const char *no_dev_avail(const struct opt_out *output)
 626 {
 627         return output->mode == CLINFO_HUMAN ?
 628                 "No devices available in platform" :
 629                 "CL_DEVICE_NOT_AVAILABLE";
 630 }
 631
 632 /* OpenCL context interop names */
 633
 634 typedef struct cl_interop_name {
 635         cl_uint from;
 636         cl_uint to;
 637         /* 5 because that's the largest we know of,
 638          * 2 because it's HUMAN, RAW */
 639         const char *value[5][2];
 640 } cl_interop_name;
 641
 642 static const cl_interop_name cl_interop_names[] = {
 643         { /* cl_khr_gl_sharing */
 644                  CL_GL_CONTEXT_KHR,
 645                  CL_CGL_SHAREGROUP_KHR,
 646                  {
 647                         { "GL", "CL_GL_CONTEXT_KHR" },
 648                         { "EGL", "CL_EGL_DISPALY_KHR" },
 649                         { "GLX", "CL_GLX_DISPLAY_KHR" },
 650                         { "WGL", "CL_WGL_HDC_KHR" },
 651                         { "CGL", "CL_CGL_SHAREGROUP_KHR" }
 652                 }
 653         },
 654         { /* cl_khr_dx9_media_sharing */
 655                 CL_CONTEXT_ADAPTER_D3D9_KHR,
 656                 CL_CONTEXT_ADAPTER_DXVA_KHR,
 657                 {
 658                         { "D3D9 (KHR)", "CL_CONTEXT_ADAPTER_D3D9_KHR" },
 659                         { "D3D9Ex (KHR)", "CL_CONTEXT_ADAPTER_D3D9EX_KHR" },
 660                         { "DXVA (KHR)", "CL_CONTEXT_ADAPTER_DXVA_KHR" }
 661                 }
 662         },
 663         { /* cl_khr_d3d10_sharing */
 664                 CL_CONTEXT_D3D10_DEVICE_KHR,
 665                 CL_CONTEXT_D3D10_DEVICE_KHR,
 666                 {
 667                         { "D3D10", "CL_CONTEXT_D3D10_DEVICE_KHR" }
 668                 }
 669         },
 670         { /* cl_khr_d3d11_sharing */
 671                 CL_CONTEXT_D3D11_DEVICE_KHR,
 672                 CL_CONTEXT_D3D11_DEVICE_KHR,
 673                 {
 674                         { "D3D11", "CL_CONTEXT_D3D11_DEVICE_KHR" }
 675                 }
 676         },
 677         /* cl_intel_dx9_media_sharing is split in two because the allowed values are not consecutive */
 678         { /* cl_intel_dx9_media_sharing part 1 */
 679                 CL_CONTEXT_D3D9_DEVICE_INTEL,
 680                 CL_CONTEXT_D3D9_DEVICE_INTEL,
 681                 {
 682                         { "D3D9 (INTEL)", "CL_CONTEXT_D3D9_DEVICE_INTEL" }
 683                 }
 684         },
 685         { /* cl_intel_dx9_media_sharing part 2 */
 686                 CL_CONTEXT_D3D9EX_DEVICE_INTEL,
 687                 CL_CONTEXT_DXVA_DEVICE_INTEL,
 688                 {
 689                         { "D3D9Ex (INTEL)", "CL_CONTEXT_D3D9EX_DEVICE_INTEL" },
 690                         { "DXVA (INTEL)", "CL_CONTEXT_DXVA_DEVICE_INTEL" }
 691                 }
 692         },
 693         { /* cl_intel_va_api_media_sharing */
 694                 CL_CONTEXT_VA_API_DISPLAY_INTEL,
 695                 CL_CONTEXT_VA_API_DISPLAY_INTEL,
 696                 {
 697                         { "VA-API", "CL_CONTEXT_VA_API_DISPLAY_INTEL" }
 698                 }
 699         }
 700 };
 701
 702 const size_t num_known_interops = ARRAY_SIZE(cl_interop_names);
 703
 704
 705 #define INDENT "  "
 706 #define I0_STR "%-48s  "
 707 #define I1_STR "  %-46s  "
 708 #define I2_STR "    %-44s  "
 709
 710 /* New line and a full padding */
 711 static const char full_padding[] = "\n"
 712 INDENT INDENT INDENT INDENT INDENT
 713 INDENT INDENT INDENT INDENT INDENT
 714 INDENT INDENT INDENT INDENT INDENT
 715 INDENT INDENT INDENT INDENT INDENT
 716 INDENT INDENT INDENT INDENT INDENT;
 717
 718 static const char empty_str[] = "";
 719 static const char spc_str[] = " ";
 720 static const char times_str[] = "x";
 721 static const char comma_str[] = ", ";
 722 static const char vbar_str[] = " | ";
 723
 724 const char *cur_sfx = empty_str;
 725
 726 /* parse a CL_DEVICE_VERSION or CL_PLATFORM_VERSION info to determine the OpenCL version.
 727  * Returns an unsigned integer in the form major*10 + minor
 728  */
 729 cl_uint
 730 getOpenCLVersion(const char *version)
 731 {
 732         cl_uint ret = 10;
 733         long parse = 0;
 734         const char *from = version;
 735         char *next = NULL;
 736         parse = strtol(from, &next, 10);
 737
 738         if (next != from) {
 739                 ret = parse*10;
 740                 // skip the dot TODO should we actually check for the dot?
 741                 from = ++next;
 742                 parse = strtol(from, &next, 10);
 743                 if (next != from)
 744                         ret += parse;
 745         }
 746         return ret;
 747 }
 748
 749 #define SPLIT_CL_VERSION(ver) ((ver)/10), ((ver)%10)
 750
 751 /* OpenCL 3.0 introduced “proper” versioning, in the form of a major.minor.patch struct
 752  * packed into a single cl_uint (type aliased to cl_version)
 753  */
 754 struct unpacked_cl_version {
 755         cl_uint major;
 756         cl_uint minor;
 757         cl_uint patch;
 758 };
 759
 760 struct unpacked_cl_version unpack_cl_version(cl_uint version)
 761 {
 762         struct unpacked_cl_version ret;
 763         ret.major = (version >> 22);
 764         ret.minor = (version >> 12) & ((1<<10)-1);
 765         ret.patch =  version & ((1<<12)-1);
 766         return ret;
 767 }
 768
 769 void strbuf_version(const char *what, struct _strbuf *str, const char *before, cl_uint version, const char *after)
 770 {
 771         struct unpacked_cl_version u = unpack_cl_version(version);
 772         strbuf_append(what, str, "%s%" PRIu32 ".%" PRIu32 ".%" PRIu32 "%s",
 773                                 before, u.major, u.minor, u.patch, after);
 774 }
 775
 776 void set_common_separator(const struct opt_out *output)
 777 {
 778         set_separator(output->json || output->mode == CLINFO_HUMAN ? comma_str : vbar_str);
 779 }
 780
 781 void strbuf_name_version(const char *what, struct _strbuf *str, const cl_name_version *ext, size_t num_exts,
 782         const struct opt_out *output)
 783 {
 784         realloc_strbuf(str, num_exts*(CL_NAME_VERSION_MAX_NAME_SIZE + 256), "extension versions");
 785         set_separator(output->mode == CLINFO_HUMAN ? full_padding : output->json ? comma_str : spc_str);
 786         if (output->json) {
 787                 strbuf_append_str(what, str, "{");
 788         }
 789         for (size_t i = 0; i < num_exts; ++i) {
 790                 const cl_name_version  *e = ext + i;
 791                 if (i > 0) strbuf_append_str(what, str, sep);
 792                 if (output->json || output->mode == CLINFO_HUMAN) {
 793                         struct unpacked_cl_version u = unpack_cl_version(e->version);
 794                         strbuf_append(what, str,
 795                                 output->json ?
 796                                 "\"%s\" : { \"raw\" : %" PRIu32 ", \"version\" : \"%d.%d.%d\" }" :
 797                                 "%-65s%#8" PRIx32 " (%d.%d.%d)",
 798                                 e->name, e->version, u.major, u.minor, u.patch);
 799                 } else {
 800                         strbuf_append(what, str, "%s:%#" PRIx32, e->name, e->version);
 801                 }
 802         }
 803         if (output->json)
 804                 strbuf_append_str(what, str, " }");
 805 }
 806
 807
 808 void strbuf_named_uint(const char *what, struct _strbuf *str, const cl_uint *ext, size_t num_exts, const struct opt_out *output,
 809         const char* const* human_str, const char* const* raw_str, const size_t count, const size_t offset)
 810 {
 811         const char *quote = output->json ? "\"" : "";
 812         const char * const * name_str = output->mode == CLINFO_HUMAN ? human_str : raw_str;
 813         set_common_separator(output);
 814         if (output->json)
 815                 strbuf_append_str_len(what, str, "[ ", 2);
 816
 817         for (size_t cursor = 0; cursor < num_exts; ++cursor) {
 818                 /* add separator for values past the first */
 819                 if (cursor > 0) strbuf_append_str(what, str, sep);
 820
 821                 cl_uint val = ext[cursor];
 822                 cl_bool known = (val >= offset && val < offset + count);
 823                 if (known)
 824                         strbuf_append(what, str, "%s%s%s", quote, name_str[val - offset], quote);
 825                 else
 826                         strbuf_append(what, str, "%s%#" PRIx32 "%s", quote, val, quote);
 827         }
 828         if (output->json)
 829                 strbuf_append_str_len(what, str, " ]", 2);
 830 }
 831
 832 void strbuf_ext_mem(const char *what, struct _strbuf *str, const cl_external_memory_handle_type_khr *ext, size_t num_exts,
 833         const struct opt_out *output)
 834 {
 835         strbuf_named_uint(what, str, ext, num_exts, output,
 836                 ext_mem_handle_str, ext_mem_handle_raw_str, ext_mem_handle_count, ext_mem_handle_offset);
 837 }
 838
 839 void strbuf_semaphore_type(const char *what, struct _strbuf *str, const cl_semaphore_type_khr *ext, size_t num_exts,
 840         const struct opt_out *output)
 841 {
 842         strbuf_named_uint(what, str, ext, num_exts, output,
 843                 semaphore_type_str, semaphore_type_raw_str, semaphore_type_count, semaphore_type_offset);
 844 }
 845
 846 void strbuf_ext_semaphore_handle(const char *what, struct _strbuf *str, const cl_external_semaphore_handle_type_khr *ext, size_t num_exts,
 847         const struct opt_out *output)
 848 {
 849         strbuf_named_uint(what, str, ext, num_exts, output,
 850                 semaphore_handle_str, semaphore_handle_raw_str, semaphore_handle_count, semaphore_handle_offset);
 851 }
 852
 853
 854 /* print strbuf, prefixed by pname, skipping leading whitespace if skip is nonzero,
 855  * affixing cur_sfx */
 856 static inline
 857 void show_strbuf(const struct _strbuf *strbuf, const char *pname, int skip, cl_int err)
 858 {
 859         printf("%s" I1_STR "%s%s\n",
 860                 line_pfx, pname,
 861                 (skip ? skip_leading_ws(strbuf->buf) : strbuf->buf),
 862                 err ? empty_str : cur_sfx);
 863 }
 864
 865 /* print a JSON string version of NULL-terminated string str, escaping \ and " and wrapping it all in "
 866  */
 867 static inline
 868 void json_stringify(const char *str)
 869 {
 870         putchar('"');
 871         while (*str) {
 872                 if (*str == '\\' || *str == '"')
 873                         putchar('\\');
 874                 putchar(*str);
 875                 ++str;
 876         }
 877         putchar('"');
 878 }
 879
 880 /* print JSON version of strbuf, prefixed by pname, skipping leading whitespace if skip is nonzero,
 881  * quoting and escaping as string if is_string is nonzero
 882  */
 883 static inline
 884 void json_strbuf(const struct _strbuf *strbuf, const char *pname, cl_uint n, cl_bool is_string)
 885 {
 886         printf("%s\"%s\" : ", (n > 0 ? comma_str : spc_str), pname);
 887         if (is_string)
 888                 json_stringify(strbuf->buf);
 889         else
 890                 fputs(strbuf->buf, stdout);
 891 }
 892
 893 void
 894 platform_info_str(struct platform_info_ret *ret,
 895         const struct info_loc *loc, const struct platform_info_checks* UNUSED(chk),
 896         const struct opt_out* UNUSED(output))
 897 {
 898         GET_STRING_LOC(ret, loc, clGetPlatformInfo, loc->plat, loc->param.plat);
 899         ret->needs_escaping = CL_TRUE;
 900 }
 901
 902 void
 903 platform_info_ulong(struct platform_info_ret *ret,
 904         const struct info_loc *loc, const struct platform_info_checks* UNUSED(chk),
 905         const struct opt_out *output)
 906 {
 907         ret->err = REPORT_ERROR_LOC(ret,
 908                 clGetPlatformInfo(loc->plat, loc->param.plat, sizeof(ret->value.u64), &ret->value.u64, NULL),
 909                 loc, "get %s");
 910         CHECK_SIZE(ret, loc, ret->value.u64, clGetPlatformInfo, loc->plat, loc->param.plat);
 911         strbuf_append(loc->pname, &ret->str, "%" PRIu64, ret->value.u64);
 912 }
 913
 914 void
 915 platform_info_sz(struct platform_info_ret *ret,
 916         const struct info_loc *loc, const struct platform_info_checks* UNUSED(chk),
 917         const struct opt_out *output)
 918 {
 919         ret->err = REPORT_ERROR_LOC(ret,
 920                 clGetPlatformInfo(loc->plat, loc->param.plat, sizeof(ret->value.s), &ret->value.s, NULL),
 921                 loc, "get %s");
 922         CHECK_SIZE(ret, loc, ret->value.s, clGetPlatformInfo, loc->plat, loc->param.plat);
 923         strbuf_append(loc->pname, &ret->str, "%" PRIuS, ret->value.s);
 924 }
 925
 926 void
 927 platform_info_version(struct platform_info_ret *ret,
 928         const struct info_loc *loc, const struct platform_info_checks* UNUSED(chk),
 929         const struct opt_out *output)
 930 {
 931         ret->err = REPORT_ERROR_LOC(ret,
 932                 clGetPlatformInfo(loc->plat, loc->param.plat, sizeof(ret->value.u32), &ret->value.u32, NULL),
 933                 loc, "get %s");
 934         CHECK_SIZE(ret, loc, ret->value.u32, clGetPlatformInfo, loc->plat, loc->param.plat);
 935         if (!ret->err) {
 936                 strbuf_append(loc->pname, &ret->str,
 937                         output->json ? "{ \"raw\" : %" PRIu32 ", \"version\" :" : "%#" PRIx32,
 938                         ret->value.u32);
 939                 if (output->json || output->mode == CLINFO_HUMAN) {
 940                         strbuf_version(loc->pname, &ret->str,
 941                                 output->json ? " \"" : " (",
 942                                 ret->value.u32,
 943                                 output->json ? "\" }" : ")");
 944                 }
 945         }
 946 }
 947
 948 void
 949 platform_info_ext_version(struct platform_info_ret *ret,
 950         const struct info_loc *loc, const struct platform_info_checks* UNUSED(chk),
 951         const struct opt_out *output)
 952 {
 953         cl_name_version *ext = NULL;
 954         size_t nusz = 0;
 955         ret->err = REPORT_ERROR_LOC(ret,
 956                 clGetPlatformInfo(loc->plat, loc->param.plat, 0, NULL, &nusz),
 957                 loc, "get %s size");
 958         if (!ret->err) {
 959                 REALLOC(ext, nusz, loc->sname);
 960                 ret->err = REPORT_ERROR_LOC(ret,
 961                         clGetPlatformInfo(loc->plat, loc->param.plat, nusz, ext, NULL),
 962                         loc, "get %s");
 963         }
 964         if (!ret->err) {
 965                 size_t num_exts = nusz / sizeof(*ext);
 966                 strbuf_name_version(loc->pname, &ret->str, ext, num_exts, output);
 967         }
 968         free(ext);
 969 }
 970
 971 void
 972 platform_info_ext_mem(struct platform_info_ret *ret,
 973         const struct info_loc *loc, const struct platform_info_checks* UNUSED(chk),
 974         const struct opt_out *output)
 975 {
 976         cl_external_memory_handle_type_khr *ext = NULL;
 977         size_t nusz = 0;
 978         ret->err = REPORT_ERROR_LOC(ret,
 979                 clGetPlatformInfo(loc->plat, loc->param.plat, 0, NULL, &nusz),
 980                 loc, "get %s size");
 981         if (!ret->err) {
 982                 REALLOC(ext, nusz, loc->sname);
 983                 ret->err = REPORT_ERROR_LOC(ret,
 984                         clGetPlatformInfo(loc->plat, loc->param.plat, nusz, ext, NULL),
 985                         loc, "get %s");
 986         }
 987         if (!ret->err) {
 988                 size_t num_exts = nusz / sizeof(*ext);
 989                 strbuf_ext_mem(loc->pname, &ret->str, ext, num_exts, output);
 990         }
 991         free(ext);
 992 }
 993
 994 void
 995 platform_info_semaphore_types(struct platform_info_ret *ret,
 996         const struct info_loc *loc, const struct platform_info_checks* UNUSED(chk),
 997         const struct opt_out *output)
 998 {
 999         cl_semaphore_type_khr *ext = NULL;
1000         size_t nusz = 0;
1001         ret->err = REPORT_ERROR_LOC(ret,
1002                 clGetPlatformInfo(loc->plat, loc->param.plat, 0, NULL, &nusz),
1003                 loc, "get %s size");
1004         if (!ret->err) {
1005                 REALLOC(ext, nusz, loc->sname);
1006                 ret->err = REPORT_ERROR_LOC(ret,
1007                         clGetPlatformInfo(loc->plat, loc->param.plat, nusz, ext, NULL),
1008                         loc, "get %s");
1009         }
1010         if (!ret->err) {
1011                 size_t num_exts = nusz / sizeof(*ext);
1012                 strbuf_semaphore_type(loc->pname, &ret->str, ext, num_exts, output);
1013         }
1014         free(ext);
1015 }
1016
1017 void
1018 platform_info_ext_semaphore_handles(struct platform_info_ret *ret,
1019         const struct info_loc *loc, const struct platform_info_checks* UNUSED(chk),
1020         const struct opt_out *output)
1021 {
1022         cl_external_semaphore_handle_type_khr *ext = NULL;
1023         size_t nusz = 0;
1024         ret->err = REPORT_ERROR_LOC(ret,
1025                 clGetPlatformInfo(loc->plat, loc->param.plat, 0, NULL, &nusz),
1026                 loc, "get %s size");
1027         if (!ret->err) {
1028                 REALLOC(ext, nusz, loc->sname);
1029                 ret->err = REPORT_ERROR_LOC(ret,
1030                         clGetPlatformInfo(loc->plat, loc->param.plat, nusz, ext, NULL),
1031                         loc, "get %s");
1032         }
1033         if (!ret->err) {
1034                 size_t num_exts = nusz / sizeof(*ext);
1035                 strbuf_ext_semaphore_handle(loc->pname, &ret->str, ext, num_exts, output);
1036         }
1037         free(ext);
1038 }
1039
1040 struct platform_info_traits {
1041         cl_platform_info param; // CL_PLATFORM_*
1042         const char *sname; // "CL_PLATFORM_*"
1043         const char *pname; // "Platform *"
1044         const char *sfx; // suffix for the output in non-raw mode
1045         /* pointer to function that retrieves the parameter */
1046         void (*show_func)(struct platform_info_ret *,
1047                 const struct info_loc *, const struct platform_info_checks *,
1048                 const struct opt_out *);
1049         /* pointer to function that checks if the parameter should be retrieved */
1050         cl_bool (*check_func)(const struct platform_info_checks *);
1051 };
1052
1053 cl_bool khr_icd_p(const struct platform_info_checks *chk)
1054 {
1055         return chk->has_khr_icd;
1056 }
1057
1058 cl_bool plat_is_12(const struct platform_info_checks *chk)
1059 {
1060         return !(chk->plat_version < 12);
1061 }
1062
1063 cl_bool plat_is_20(const struct platform_info_checks *chk)
1064 {
1065         return !(chk->plat_version < 20);
1066 }
1067
1068 cl_bool plat_is_21(const struct platform_info_checks *chk)
1069 {
1070         return !(chk->plat_version < 21);
1071 }
1072
1073 cl_bool plat_is_30(const struct platform_info_checks *chk)
1074 {
1075         return !(chk->plat_version < 30);
1076 }
1077
1078 cl_bool plat_has_amd_object_metadata(const struct platform_info_checks *chk)
1079 {
1080         return chk->has_amd_object_metadata;
1081 }
1082
1083 cl_bool plat_has_ext_ver(const struct platform_info_checks *chk)
1084 {
1085         return plat_is_30(chk) || chk->has_extended_versioning;
1086 }
1087
1088 cl_bool plat_has_ext_mem(const struct platform_info_checks *chk)
1089 {
1090         return chk->has_external_memory;
1091 }
1092
1093 cl_bool plat_has_semaphore(const struct platform_info_checks *chk)
1094 {
1095         return chk->has_semaphore;
1096 }
1097
1098 cl_bool plat_has_external_semaphore(const struct platform_info_checks *chk)
1099 {
1100         return chk->has_external_semaphore;
1101 }
1102
1103 #define PINFO_COND(symbol, name, sfx, typ, funcptr) { symbol, #symbol, "Platform " name, sfx, &platform_info_##typ, &funcptr }
1104 #define PINFO(symbol, name, sfx, typ) { symbol, #symbol, "Platform " name, sfx, &platform_info_##typ, NULL }
1105 struct platform_info_traits pinfo_traits[] = {
1106         PINFO(CL_PLATFORM_NAME, "Name", NULL, str),
1107         PINFO(CL_PLATFORM_VENDOR, "Vendor", NULL, str),
1108         PINFO(CL_PLATFORM_VERSION, "Version", NULL, str),
1109         PINFO(CL_PLATFORM_PROFILE, "Profile", NULL, str),
1110         PINFO(CL_PLATFORM_EXTENSIONS, "Extensions", NULL, str),
1111         PINFO_COND(CL_PLATFORM_EXTENSIONS_WITH_VERSION, "Extensions with Version", NULL, ext_version, plat_has_ext_ver),
1112         PINFO_COND(CL_PLATFORM_NUMERIC_VERSION, "Numeric Version", NULL, version, plat_has_ext_ver),
1113         PINFO_COND(CL_PLATFORM_ICD_SUFFIX_KHR, "Extensions function suffix", NULL, str, khr_icd_p),
1114         PINFO_COND(CL_PLATFORM_MAX_KEYS_AMD, "Max metadata object keys (AMD)", NULL, sz, plat_has_amd_object_metadata),
1115         PINFO_COND(CL_PLATFORM_HOST_TIMER_RESOLUTION, "Host timer resolution", "ns", ulong, plat_is_21),
1116         PINFO_COND(CL_PLATFORM_EXTERNAL_MEMORY_IMPORT_HANDLE_TYPES_KHR, "External memory handle types", NULL, ext_mem, plat_has_ext_mem),
1117         PINFO_COND(CL_PLATFORM_SEMAPHORE_TYPES_KHR, "Semaphore types", NULL, semaphore_types, plat_has_semaphore),
1118         PINFO_COND(CL_PLATFORM_SEMAPHORE_IMPORT_HANDLE_TYPES_KHR, "External semaphore import types", NULL, ext_semaphore_handles, plat_has_external_semaphore),
1119         PINFO_COND(CL_PLATFORM_SEMAPHORE_EXPORT_HANDLE_TYPES_KHR, "External semaphore export types", NULL, ext_semaphore_handles, plat_has_external_semaphore),
1120
1121 };
1122
1123 /* Collect (and optionally show) information on a specific platform,
1124  * initializing relevant arrays and optionally showing the collected
1125  * information
1126  */
1127 void
1128 gatherPlatformInfo(struct platform_list *plist, cl_uint p, const struct opt_out *output)
1129 {
1130         size_t len = 0;
1131         cl_uint n = 0; /* number of platform properties shown, for JSON */
1132
1133         struct platform_data *pdata = plist->pdata + p;
1134         struct platform_info_checks *pinfo_checks = plist->platform_checks + p;
1135         struct platform_info_ret ret;
1136         struct info_loc loc;
1137
1138         pinfo_checks->plat_version = 10;
1139
1140         INIT_RET(ret, "platform");
1141         reset_loc(&loc, __func__);
1142         loc.plat = plist->platform[p];
1143
1144         for (loc.line = 0; loc.line < ARRAY_SIZE(pinfo_traits); ++loc.line) {
1145                 const struct platform_info_traits *traits = pinfo_traits + loc.line;
1146                 cl_bool requested;
1147
1148                 /* checked is true if there was no condition to check for, or if the
1149                  * condition was satisfied
1150                  */
1151                 int checked = !(traits->check_func && !traits->check_func(pinfo_checks));
1152
1153                 if (output->cond == COND_PROP_CHECK && !checked)
1154                         continue;
1155
1156                 loc.sname = traits->sname;
1157                 loc.pname = (output->mode == CLINFO_HUMAN ?
1158                         traits->pname : traits->sname);
1159                 loc.param.plat = traits->param;
1160
1161                 cur_sfx = (output->mode == CLINFO_HUMAN && traits->sfx) ? traits->sfx : empty_str;
1162
1163                 reset_strbuf(&ret.str);
1164                 reset_strbuf(&ret.err_str);
1165                 ret.needs_escaping = CL_FALSE;
1166                 traits->show_func(&ret, &loc, pinfo_checks, output);
1167
1168                 /* The property is skipped if this was a conditional property,
1169                  * unsatisfied, there was an error retrieving it and cond_prop_mode is not
1170                  * COND_PROP_SHOW.
1171                  */
1172                 if (ret.err && !checked && output->cond != COND_PROP_SHOW)
1173                         continue;
1174
1175                 /* The property gets printed if we are not just listing,
1176                  * or if the user requested a property and this one matches.
1177                  * Otherwise, we're just gathering information */
1178                 requested = (output->prop && strstr(loc.sname, output->prop) != NULL);
1179                 if (output->detailed || requested) {
1180                         if (output->json) {
1181                                 json_strbuf(RET_BUF(ret), loc.pname, n++, ret.err || ret.needs_escaping);
1182                         } else {
1183                                 show_strbuf(RET_BUF(ret), loc.pname, CL_FALSE, ret.err);
1184                         }
1185                 }
1186
1187                 if (ret.err)
1188                         continue;
1189
1190                 /* post-processing */
1191
1192                 switch (traits->param) {
1193                 case CL_PLATFORM_NAME:
1194                         /* Store name for future reference */
1195                         len = strlen(ret.str.buf);
1196                         ALLOC(pdata->pname, len+1, "platform name copy");
1197                         /* memcpy instead of strncpy since we already have the len
1198                          * and memcpy is possibly more optimized */
1199                         memcpy(pdata->pname, ret.str.buf, len);
1200                         pdata->pname[len] = '\0';
1201                         break;
1202                 case CL_PLATFORM_VERSION:
1203                         /* compute numeric value for OpenCL version */
1204                         pinfo_checks->plat_version = getOpenCLVersion(ret.str.buf + 7);
1205                         break;
1206                 case CL_PLATFORM_EXTENSIONS:
1207                         pinfo_checks->has_khr_icd = !!strstr(ret.str.buf, "cl_khr_icd");
1208                         pinfo_checks->has_amd_object_metadata = !!strstr(ret.str.buf, "cl_amd_object_metadata");
1209                         pinfo_checks->has_external_memory = !!strstr(ret.str.buf, "cl_khr_external_memory");
1210                         pinfo_checks->has_semaphore = !!strstr(ret.str.buf, "cl_khr_semaphore");
1211                         pinfo_checks->has_external_semaphore = !!strstr(ret.str.buf, "cl_khr_external_semaphore");
1212                         pdata->has_amd_offline = !!strstr(ret.str.buf, "cl_amd_offline_devices");
1213                         break;
1214                 case CL_PLATFORM_ICD_SUFFIX_KHR:
1215                         /* Store ICD suffix for future reference */
1216                         len = strlen(ret.str.buf);
1217                         ALLOC(pdata->sname, len+1, "platform ICD suffix copy");
1218                         /* memcpy instead of strncpy since we already have the len
1219                          * and memcpy is possibly more optimized */
1220                         memcpy(pdata->sname, ret.str.buf, len);
1221                         pdata->sname[len] = '\0';
1222                 default:
1223                         /* do nothing */
1224                         break;
1225                 }
1226
1227         }
1228
1229         if (pinfo_checks->plat_version > plist->max_plat_version)
1230                 plist->max_plat_version = pinfo_checks->plat_version;
1231
1232         /* if no CL_PLATFORM_ICD_SUFFIX_KHR, use P### as short/symbolic name */
1233         if (!pdata->sname) {
1234 #define SNAME_MAX 32
1235                 ALLOC(pdata->sname, SNAME_MAX+1, "platform symbolic name");
1236                 snprintf(pdata->sname, SNAME_MAX, "P%" PRIu32 "", p);
1237         }
1238
1239         len = strlen(pdata->sname);
1240         if (len > plist->max_sname_len)
1241                 plist->max_sname_len = len;
1242
1243         ret.err = clGetDeviceIDs(loc.plat, CL_DEVICE_TYPE_ALL, 0, NULL, &pdata->ndevs);
1244         if (ret.err == CL_DEVICE_NOT_FOUND)
1245                 pdata->ndevs = 0;
1246         else
1247                 CHECK_ERROR(ret.err, "number of devices");
1248         plist->ndevs_total += pdata->ndevs;
1249         plist->dev_offset[p] = p ? plist->dev_offset[p-1] + (pdata-1)->ndevs : 0;
1250         plist_devs_reserve(plist, plist->ndevs_total);
1251
1252         if (pdata->ndevs > 0) {
1253                 ret.err = clGetDeviceIDs(loc.plat, CL_DEVICE_TYPE_ALL,
1254                         pdata->ndevs,
1255                         plist->all_devs + plist->dev_offset[p], NULL);
1256         }
1257
1258         if (pdata->ndevs > plist->max_devs)
1259                 plist->max_devs = pdata->ndevs;
1260
1261         UNINIT_RET(ret);
1262 }
1263
1264 /*
1265  * Device properties/extensions used in traits checks, and relevant functions
1266  * TODO add version control for 3.0+ platforms
1267  */
1268
1269 struct device_info_checks {
1270         const struct platform_info_checks *pinfo_checks;
1271         cl_device_type devtype;
1272         cl_device_mem_cache_type cachetype;
1273         cl_device_local_mem_type lmemtype;
1274         cl_bool image_support;
1275         cl_bool compiler_available;
1276         cl_bool arm_register_alloc_support;
1277         cl_bool arm_warp_count_support;
1278         char has_half[12];
1279         char has_double[24];
1280         char has_nv[29];
1281         char has_amd[30];
1282         char has_intel[32];
1283         char has_amd_svm[11];
1284         char has_arm_svm[29];
1285         char has_intel_usm[31];
1286         char has_external_memory[23];
1287         char has_semaphore[17];
1288         char has_external_semaphore[26];
1289         char has_arm_core_id[15];
1290         char has_arm_job_slots[26];
1291         char has_arm_scheduling_controls[27];
1292         char has_fission[22];
1293         char has_atomic_counters[26];
1294         char has_image2d_buffer[27];
1295         char has_il_program[18];
1296         char has_intel_queue_families[32];
1297         char has_intel_local_thread[30];
1298         char has_intel_AME[36];
1299         char has_intel_AVC_ME[43];
1300         char has_intel_planar_yuv[20];
1301         char has_intel_required_subgroup_size[32];
1302         char has_altera_dev_temp[29];
1303         char has_p2p[23];
1304         char has_pci_bus_info[20];
1305         char has_spir[12];
1306         char has_qcom_ext_host_ptr[21];
1307         char has_simultaneous_sharing[30];
1308         char has_subgroup_named_barrier[30];
1309         char has_command_buffer[25];
1310         char has_mutable_dispatch[27];
1311         char has_terminate_context[25];
1312         char has_terminate_arm[37];
1313         char has_extended_versioning[27];
1314         char has_cxx_for_opencl[22];
1315         char has_device_uuid[19];
1316         cl_uint dev_version;
1317         cl_uint p2p_num_devs;
1318 };
1319
1320 #define DEFINE_EXT_CHECK(ext) cl_bool dev_has_##ext(const struct device_info_checks *chk) \
1321 { \
1322         return !!(chk->has_##ext[0]); \
1323 }
1324
1325 DEFINE_EXT_CHECK(half)
1326 DEFINE_EXT_CHECK(double)
1327 DEFINE_EXT_CHECK(nv)
1328 DEFINE_EXT_CHECK(amd)
1329 DEFINE_EXT_CHECK(amd_svm)
1330 DEFINE_EXT_CHECK(arm_svm)
1331 DEFINE_EXT_CHECK(intel_usm)
1332 DEFINE_EXT_CHECK(external_memory)
1333 DEFINE_EXT_CHECK(semaphore)
1334 DEFINE_EXT_CHECK(external_semaphore)
1335 DEFINE_EXT_CHECK(arm_core_id)
1336 DEFINE_EXT_CHECK(arm_job_slots)
1337 DEFINE_EXT_CHECK(arm_scheduling_controls)
1338 DEFINE_EXT_CHECK(fission)
1339 DEFINE_EXT_CHECK(atomic_counters)
1340 DEFINE_EXT_CHECK(il_program)
1341 DEFINE_EXT_CHECK(intel)
1342 DEFINE_EXT_CHECK(intel_queue_families)
1343 DEFINE_EXT_CHECK(intel_local_thread)
1344 DEFINE_EXT_CHECK(intel_AME)
1345 DEFINE_EXT_CHECK(intel_AVC_ME)
1346 DEFINE_EXT_CHECK(intel_planar_yuv)
1347 DEFINE_EXT_CHECK(intel_required_subgroup_size)
1348 DEFINE_EXT_CHECK(altera_dev_temp)
1349 DEFINE_EXT_CHECK(p2p)
1350 DEFINE_EXT_CHECK(pci_bus_info)
1351 DEFINE_EXT_CHECK(spir)
1352 DEFINE_EXT_CHECK(qcom_ext_host_ptr)
1353 DEFINE_EXT_CHECK(simultaneous_sharing)
1354 DEFINE_EXT_CHECK(subgroup_named_barrier)
1355 DEFINE_EXT_CHECK(command_buffer)
1356 DEFINE_EXT_CHECK(mutable_dispatch)
1357 DEFINE_EXT_CHECK(terminate_context)
1358 DEFINE_EXT_CHECK(terminate_arm)
1359 DEFINE_EXT_CHECK(extended_versioning)
1360 DEFINE_EXT_CHECK(cxx_for_opencl)
1361 DEFINE_EXT_CHECK(device_uuid)
1362
1363 /* In the version checks we negate the opposite conditions
1364  * instead of double-negating the actual condition
1365  */
1366
1367 // device supports 1.1
1368 cl_bool dev_is_11(const struct device_info_checks *chk)
1369 {
1370         return !(chk->dev_version < 11);
1371 }
1372
1373
1374 // device supports 1.2
1375 cl_bool dev_is_12(const struct device_info_checks *chk)
1376 {
1377         return !(chk->dev_version < 12);
1378 }
1379
1380 // device supports 2.0
1381 cl_bool dev_is_20(const struct device_info_checks *chk)
1382 {
1383         return !(chk->dev_version < 20);
1384 }
1385
1386 // device supports 2.1
1387 cl_bool dev_is_21(const struct device_info_checks *chk)
1388 {
1389         return !(chk->dev_version < 21);
1390 }
1391
1392 // device does not support 2.0
1393 cl_bool dev_not_20(const struct device_info_checks *chk)
1394 {
1395         return !(chk->dev_version >= 20);
1396 }
1397
1398 // device supports 3.0
1399 cl_bool dev_is_30(const struct device_info_checks *chk)
1400 {
1401         return !(chk->dev_version < 30);
1402 }
1403
1404 // device has extended versioning: 3.0 or has_extended_versioning
1405 cl_bool dev_has_ext_ver(const struct device_info_checks *chk)
1406 {
1407         return dev_is_30(chk) || dev_has_extended_versioning(chk);
1408 }
1409
1410 cl_bool dev_is_gpu(const struct device_info_checks *chk)
1411 {
1412         return !!(chk->devtype & CL_DEVICE_TYPE_GPU);
1413 }
1414
1415 cl_bool dev_is_gpu_amd(const struct device_info_checks *chk)
1416 {
1417         return dev_is_gpu(chk) && dev_has_amd(chk);
1418 }
1419
1420 /* Device supports cl_amd_device_attribute_query v4 */
1421 cl_bool dev_has_amd_v4(const struct device_info_checks *chk)
1422 {
1423         /* We don't actually have a criterion to check if the device
1424          * supports a specific version of an extension, so for the time
1425          * being rely on them being GPU devices with cl_amd_device_attribute_query
1426          * and the platform supporting OpenCL 2.0 or later
1427          * TODO FIXME tune criteria
1428          */
1429         return dev_is_gpu(chk) && dev_has_amd(chk) && plat_is_20(chk->pinfo_checks);
1430 }
1431
1432 /* Device supports cl_intel_device_attribute_query and is a GPU */
1433 cl_bool dev_is_gpu_intel(const struct device_info_checks *chk)
1434 {
1435         return dev_is_gpu(chk) && dev_has_intel(chk);
1436 }
1437
1438 /* Device supports cl_arm_core_id v2 */
1439 cl_bool dev_has_arm_core_id_v2(const struct device_info_checks *chk)
1440 {
1441         /* We don't actually have a criterion to check if the device
1442          * supports a specific version of an extension, so for the time
1443          * being rely on them having cl_arm_core_id and the platform
1444          * supporting OpenCL 1.2 or later
1445          * TODO FIXME tune criteria
1446          */
1447         return dev_has_arm_core_id(chk) && plat_is_12(chk->pinfo_checks);
1448 }
1449
1450 /* Device supports register allocation queries */
1451 cl_bool dev_has_arm_register_alloc(const struct device_info_checks *chk)
1452 {
1453         return dev_has_arm_scheduling_controls(chk) && chk->arm_register_alloc_support;
1454 }
1455
1456 /* Device supports warp  */
1457 cl_bool dev_has_arm_warp_count_support(const struct device_info_checks *chk)
1458 {
1459         return dev_has_arm_scheduling_controls(chk) && chk->arm_warp_count_support;
1460 }
1461
1462 cl_bool dev_has_svm(const struct device_info_checks *chk)
1463 {
1464         return dev_is_20(chk) || dev_has_amd_svm(chk);
1465 }
1466
1467 cl_bool dev_has_partition(const struct device_info_checks *chk)
1468 {
1469         return dev_is_12(chk) || dev_has_fission(chk);
1470 }
1471
1472 cl_bool dev_has_cache(const struct device_info_checks *chk)
1473 {
1474         return chk->cachetype != CL_NONE;
1475 }
1476
1477 cl_bool dev_has_lmem(const struct device_info_checks *chk)
1478 {
1479         return chk->lmemtype != CL_NONE;
1480 }
1481
1482 cl_bool dev_has_il(const struct device_info_checks *chk)
1483 {
1484         return dev_is_21(chk) || dev_has_il_program(chk);
1485 }
1486
1487 cl_bool dev_has_images(const struct device_info_checks *chk)
1488 {
1489         return chk->image_support;
1490 }
1491
1492 cl_bool dev_has_images_12(const struct device_info_checks *chk)
1493 {
1494         return dev_has_images(chk) && dev_is_12(chk);
1495 }
1496
1497 cl_bool dev_has_images_20(const struct device_info_checks *chk)
1498 {
1499         return dev_has_images(chk) && dev_is_20(chk);
1500 }
1501
1502 cl_bool dev_has_image2d_buffer(const struct device_info_checks *chk)
1503 {
1504         return dev_has_images_20(chk) || !!(chk->has_image2d_buffer[0]);
1505 }
1506
1507 cl_bool dev_has_compiler(const struct device_info_checks *chk)
1508 {
1509         return chk->compiler_available;
1510 }
1511
1512 cl_bool dev_has_compiler_11(const struct device_info_checks *chk)
1513 {
1514         return dev_is_11(chk) && dev_has_compiler(chk);
1515 }
1516
1517 cl_bool dev_has_p2p_devs(const struct device_info_checks *chk)
1518 {
1519         return dev_has_p2p(chk) && chk->p2p_num_devs > 0;
1520 }
1521
1522
1523 void identify_device_extensions(const char *extensions, struct device_info_checks *chk)
1524 {
1525 #define _HAS_EXT(ext) (strstr(extensions, ext))
1526 #define CPY_EXT(what, ext) do { \
1527         strncpy(chk->has_##what, has+1, sizeof(ext)); \
1528         chk->has_##what[sizeof(ext)-1] = '\0'; \
1529 } while (0)
1530 #define CHECK_EXT(what, ext) do { \
1531         has = _HAS_EXT(" " #ext " "); \
1532         if (has) CPY_EXT(what, #ext); \
1533 } while(0)
1534
1535         char *has;
1536         CHECK_EXT(half, cl_khr_fp16);
1537         CHECK_EXT(spir, cl_khr_spir);
1538         CHECK_EXT(double, cl_khr_fp64);
1539         if (!dev_has_double(chk))
1540                 CHECK_EXT(double, cl_amd_fp64);
1541         if (!dev_has_double(chk))
1542                 CHECK_EXT(double, cl_APPLE_fp64_basic_ops);
1543         CHECK_EXT(nv, cl_nv_device_attribute_query);
1544         CHECK_EXT(amd, cl_amd_device_attribute_query);
1545         CHECK_EXT(intel, cl_intel_device_attribute_query);
1546         CHECK_EXT(amd_svm, cl_amd_svm);
1547         CHECK_EXT(arm_svm, cl_arm_shared_virtual_memory);
1548         CHECK_EXT(intel_usm, cl_intel_unified_shared_memory);
1549         CHECK_EXT(external_memory, cl_khr_external_memory);
1550         CHECK_EXT(semaphore, cl_khr_semaphore);
1551         CHECK_EXT(external_semaphore, cl_khr_external_semaphore);
1552         CHECK_EXT(arm_core_id, cl_arm_core_id);
1553         CHECK_EXT(arm_job_slots, cl_arm_job_slot_selection);
1554         CHECK_EXT(arm_scheduling_controls, cl_arm_scheduling_controls);
1555         CHECK_EXT(fission, cl_ext_device_fission);
1556         CHECK_EXT(atomic_counters, cl_ext_atomic_counters_64);
1557         if (dev_has_atomic_counters(chk))
1558                 CHECK_EXT(atomic_counters, cl_ext_atomic_counters_32);
1559         CHECK_EXT(image2d_buffer, cl_khr_image2d_from_buffer);
1560         CHECK_EXT(il_program, cl_khr_il_program);
1561         CHECK_EXT(intel_queue_families, cl_intel_command_queue_families);
1562         CHECK_EXT(intel_local_thread, cl_intel_exec_by_local_thread);
1563         CHECK_EXT(intel_AME, cl_intel_advanced_motion_estimation);
1564         CHECK_EXT(intel_AVC_ME, cl_intel_device_side_avc_motion_estimation);
1565         CHECK_EXT(intel_planar_yuv, cl_intel_planar_yuv);
1566         CHECK_EXT(intel_required_subgroup_size, cl_intel_required_subgroup_size);
1567         CHECK_EXT(altera_dev_temp, cl_altera_device_temperature);
1568         CHECK_EXT(p2p, cl_amd_copy_buffer_p2p);
1569         CHECK_EXT(pci_bus_info, cl_khr_pci_bus_info);
1570         CHECK_EXT(qcom_ext_host_ptr, cl_qcom_ext_host_ptr);
1571         CHECK_EXT(simultaneous_sharing, cl_intel_simultaneous_sharing);
1572         CHECK_EXT(subgroup_named_barrier, cl_khr_subgroup_named_barrier);
1573         CHECK_EXT(command_buffer, cl_khr_command_buffer);
1574         CHECK_EXT(mutable_dispatch, cl_khr_mutable_dispatch);
1575         CHECK_EXT(terminate_context, cl_khr_terminate_context);
1576         CHECK_EXT(terminate_arm, cl_arm_controlled_kernel_termination);
1577         CHECK_EXT(extended_versioning, cl_khr_extended_versioning);
1578         CHECK_EXT(cxx_for_opencl, cl_ext_cxx_for_opencl);
1579         CHECK_EXT(device_uuid, cl_khr_device_uuid);
1580 }
1581
1582
1583 /*
1584  * Device info print functions
1585  */
1586
1587 #define _GET_VAL(ret, loc, val) \
1588         ret->err = REPORT_ERROR_LOC(ret, \
1589                 clGetDeviceInfo((loc)->dev, (loc)->param.dev, sizeof(val), &(val), NULL), \
1590                 loc, "get %s"); \
1591         CHECK_SIZE(ret, loc, val, clGetDeviceInfo, (loc)->dev, (loc)->param.dev);
1592
1593 #define _GET_VAL_VALUES(ret, loc) \
1594         REALLOC(val, numval, loc->sname); \
1595         ret->err = REPORT_ERROR_LOC(ret, \
1596                 clGetDeviceInfo(loc->dev, loc->param.dev, szval, val, NULL), \
1597                 loc, "get %s"); \
1598         if (ret->err) { free(val); val = NULL; } \
1599
1600 #define _GET_VAL_ARRAY(ret, loc) \
1601         ret->err = REPORT_ERROR_LOC(ret, \
1602                 clGetDeviceInfo(loc->dev, loc->param.dev, 0, NULL, &szval), \
1603                 loc, "get number of %s"); \
1604         numval = szval/sizeof(*val); \
1605         if (!ret->err && numval > 0) { \
1606                 _GET_VAL_VALUES(ret, loc) \
1607         }
1608
1609 #define GET_VAL(ret, loc, field) do { \
1610         _GET_VAL(ret, (loc), ret->value.field) \
1611 } while (0)
1612
1613 #define GET_VAL_ARRAY(ret, loc) do { \
1614         _GET_VAL_ARRAY(ret, (loc)) \
1615 } while (0)
1616
1617 #define DEFINE_DEVINFO_FETCH(type, field) \
1618 type \
1619 device_fetch_##type(struct device_info_ret *ret, \
1620         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk), \
1621         const struct opt_out *output) \
1622 { \
1623         GET_VAL(ret, loc, field); \
1624         return ret->value.field; \
1625 }
1626
1627 DEFINE_DEVINFO_FETCH(size_t, s)
1628 DEFINE_DEVINFO_FETCH(cl_bool, b)
1629 DEFINE_DEVINFO_FETCH(cl_uint, u32)
1630 DEFINE_DEVINFO_FETCH(cl_version, u32)
1631 DEFINE_DEVINFO_FETCH(cl_ulong, u64)
1632 DEFINE_DEVINFO_FETCH(cl_bitfield, u64)
1633 DEFINE_DEVINFO_FETCH(cl_device_type, devtype)
1634 DEFINE_DEVINFO_FETCH(cl_device_mem_cache_type, cachetype)
1635 DEFINE_DEVINFO_FETCH(cl_device_local_mem_type, lmemtype)
1636 DEFINE_DEVINFO_FETCH(cl_device_topology_amd, devtopo_amd)
1637 DEFINE_DEVINFO_FETCH(cl_device_pci_bus_info_khr, devtopo_khr)
1638 DEFINE_DEVINFO_FETCH(cl_device_affinity_domain, affinity_domain)
1639 DEFINE_DEVINFO_FETCH(cl_device_fp_config, fpconfig)
1640 DEFINE_DEVINFO_FETCH(cl_command_queue_properties, qprop)
1641 DEFINE_DEVINFO_FETCH(cl_device_exec_capabilities, execap)
1642 DEFINE_DEVINFO_FETCH(cl_device_svm_capabilities, svmcap)
1643 DEFINE_DEVINFO_FETCH(cl_device_terminate_capability_khr, termcap)
1644
1645 #define DEV_FETCH_LOC(type, var, loc) \
1646         type var = device_fetch_##type(ret, loc, chk, output)
1647 #define DEV_FETCH(type, var) DEV_FETCH_LOC(type, var, loc)
1648
1649 #define FMT_VAL(loc, ret, fmt, val) if (!ret->err) strbuf_append(loc->pname, &ret->str, fmt, val)
1650
1651 #define DEFINE_DEVINFO_SHOW(how, type, field, fmt) \
1652 void \
1653 device_info_##how(struct device_info_ret *ret, \
1654         const struct info_loc *loc, const struct device_info_checks* chk, \
1655         const struct opt_out *output) \
1656 { \
1657         DEV_FETCH(type, val); \
1658         if (!ret->err) FMT_VAL(loc, ret, fmt, val); \
1659 }
1660
1661 DEFINE_DEVINFO_SHOW(int, cl_uint, u32, "%" PRIu32)
1662 DEFINE_DEVINFO_SHOW(hex, cl_uint, u32, output->json ? "%" PRIu32 : "%#" PRIx32)
1663 DEFINE_DEVINFO_SHOW(long, cl_ulong, u64, "%" PRIu64)
1664 DEFINE_DEVINFO_SHOW(sz, size_t, s, "%" PRIuS)
1665
1666 void
1667 device_info_str(struct device_info_ret *ret,
1668         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
1669         const struct opt_out* UNUSED(output))
1670 {
1671         GET_STRING_LOC(ret, loc, clGetDeviceInfo, loc->dev, loc->param.dev);
1672         ret->needs_escaping = CL_TRUE;
1673 }
1674
1675 void
1676 device_info_bool(struct device_info_ret *ret,
1677         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
1678         const struct opt_out *output)
1679 {
1680         DEV_FETCH(cl_bool, val);
1681         if (!ret->err) {
1682                 const char * const * str = (output->mode == CLINFO_HUMAN ?
1683                         bool_str : output->json ? bool_json_str : bool_raw_str);
1684                 strbuf_append(loc->pname, &ret->str, "%s", str[val]);
1685         }
1686 }
1687
1688 void
1689 device_info_bits(struct device_info_ret *ret,
1690         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
1691         const struct opt_out *output)
1692 {
1693         DEV_FETCH(cl_uint, val);
1694         if (!ret->err)
1695                 strbuf_append(loc->pname, &ret->str, "%" PRIu32 " bits (%" PRIu32 " bytes)", val, val/8);
1696 }
1697
1698 void
1699 device_info_version(struct device_info_ret *ret,
1700         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
1701         const struct opt_out *output)
1702 {
1703         GET_VAL(ret, loc, u32);
1704         if (!ret->err) {
1705                 strbuf_append(loc->pname, &ret->str,
1706                         output->json ? "{ \" raw \" : %" PRIu32 ", \"version\" :" : "%#" PRIx32,
1707                         ret->value.u32);
1708                 if (output->json || output->mode == CLINFO_HUMAN) {
1709                         strbuf_version(loc->pname, &ret->str,
1710                                 output->json ? " \"" : " (",
1711                                 ret->value.u32,
1712                                 output->json ? "\" }" : ")");
1713                 }
1714         }
1715 }
1716
1717 void
1718 device_info_ext_version(struct device_info_ret *ret,
1719         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
1720         const struct opt_out *output)
1721 {
1722         cl_name_version *val = NULL;
1723         size_t szval = 0, numval = 0;
1724         GET_VAL_ARRAY(ret, loc);
1725         if (!ret->err) {
1726                 strbuf_name_version(loc->pname, &ret->str, val, numval, output);
1727         }
1728         free(val);
1729 }
1730
1731 void
1732 device_info_ext_mem(struct device_info_ret *ret,
1733         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
1734         const struct opt_out *output)
1735 {
1736         cl_external_memory_handle_type_khr *val = NULL;
1737         size_t szval = 0, numval = 0;
1738         GET_VAL_ARRAY(ret, loc);
1739         if (!ret->err) {
1740                 strbuf_ext_mem(loc->pname, &ret->str, val, numval, output);
1741         }
1742         free(val);
1743 }
1744
1745 void
1746 device_info_semaphore_types(struct device_info_ret *ret,
1747         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
1748         const struct opt_out *output)
1749 {
1750         cl_semaphore_type_khr *val = NULL;
1751         size_t szval = 0, numval = 0;
1752         GET_VAL_ARRAY(ret, loc);
1753         if (!ret->err) {
1754                 strbuf_semaphore_type(loc->pname, &ret->str, val, numval, output);
1755         }
1756         free(val);
1757 }
1758
1759 void
1760 device_info_ext_semaphore_handles(struct device_info_ret *ret,
1761         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
1762         const struct opt_out *output)
1763 {
1764         cl_external_semaphore_handle_type_khr *val = NULL;
1765         size_t szval = 0, numval = 0;
1766         GET_VAL_ARRAY(ret, loc);
1767         if (!ret->err) {
1768                 strbuf_ext_semaphore_handle(loc->pname, &ret->str, val, numval, output);
1769         }
1770         free(val);
1771 }
1772
1773 void strbuf_mem(const char *what, struct _strbuf *str, cl_ulong val)
1774 {
1775         double dbl = (double)val;
1776         size_t sfx = 0;
1777         while (dbl > 1024 && sfx < memsfx_end) {
1778                 dbl /= 1024;
1779                 ++sfx;
1780         }
1781         strbuf_append(what, str, " (%.4lg%s)", dbl, memsfx[sfx]);
1782 }
1783
1784 void
1785 device_info_mem(struct device_info_ret *ret,
1786         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
1787         const struct opt_out *output)
1788 {
1789         GET_VAL(ret, loc, u64);
1790         if (!ret->err) {
1791                 strbuf_append(loc->pname, &ret->str, "%" PRIu64, ret->value.u64);
1792                 if (output->mode == CLINFO_HUMAN && ret->value.u64 > 1024)
1793                         strbuf_mem(loc->pname, &ret->str, ret->value.u64);
1794         }
1795 }
1796
1797 void
1798 device_info_mem_int(struct device_info_ret *ret,
1799         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
1800         const struct opt_out *output)
1801 {
1802         GET_VAL(ret, loc, u32);
1803         if (!ret->err) {
1804                 strbuf_append(loc->pname, &ret->str, "%" PRIu32, ret->value.u32);
1805                 if (output->mode == CLINFO_HUMAN && ret->value.u32 > 1024)
1806                         strbuf_mem(loc->pname, &ret->str, ret->value.u32);
1807         }
1808 }
1809
1810 void
1811 device_info_mem_sz(struct device_info_ret *ret,
1812         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
1813         const struct opt_out *output)
1814 {
1815         GET_VAL(ret, loc, s);
1816         if (!ret->err) {
1817                 strbuf_append(loc->pname, &ret->str, "%" PRIuS, ret->value.s);
1818                 if (output->mode == CLINFO_HUMAN && ret->value.s > 1024)
1819                         strbuf_mem(loc->pname, &ret->str, ret->value.s);
1820         }
1821 }
1822
1823 void
1824 device_info_free_mem_amd(struct device_info_ret *ret,
1825         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
1826         const struct opt_out *output)
1827 {
1828         // Apparently, with the introduction of ROCm, CL_DEVICE_GLOBAL_FREE_MEMORY_AMD
1829         // returns 1 or 2 values depending on how it's called: if it's called with a
1830         // szval < 2*sizeof(size_t), it will only return 1 value, otherwise it will return 2.
1831         // At least now these are documented in the ROCm source code: the first value
1832         // is the total amount of free memory, and the second is the size of the largest
1833         // free block. So let's just manually ask for both values
1834         GET_VAL(ret, loc, u64v2);
1835         if (!ret->err) {
1836                 size_t cursor = 0;
1837                 if (output->json)
1838                         strbuf_append_str_len(loc->pname, &ret->str, " [", 2);
1839                 for (cursor = 0; cursor < 2; ++cursor) {
1840                         cl_ulong v = ret->value.u64v2.s[cursor];
1841                         if (cursor > 0)
1842                                 strbuf_append_str(loc->pname, &ret->str,
1843                                         output->json ? comma_str : spc_str);
1844                         strbuf_append(loc->pname, &ret->str, "%" PRIuS, v);
1845                         if (output->mode == CLINFO_HUMAN)
1846                                 strbuf_mem(loc->pname, &ret->str, v*UINT64_C(1024));
1847                 }
1848                 if (output->json)
1849                         strbuf_append_str_len(loc->pname, &ret->str, " ]", 2);
1850         }
1851 }
1852
1853 void
1854 device_info_time_offset(struct device_info_ret *ret,
1855         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
1856         const struct opt_out *output)
1857 {
1858         GET_VAL(ret, loc, u64);
1859         if (!ret->err) {
1860                 time_t time = ret->value.u64/UINT64_C(1000000000);
1861                 strbuf_append(loc->pname, &ret->str, "%" PRIu64 "ns (", ret->value.u64);
1862                 strbuf_append_str(loc->pname, &ret->str, ctime(&time));
1863                 /* overwrite ctime's newline with the closing parenthesis */
1864                 ret->str.buf[ret->str.end - 1] = ')';
1865         }
1866 }
1867
1868 void
1869 device_info_intptr(struct device_info_ret *ret,
1870         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
1871         const struct opt_out *output)
1872 {
1873         cl_int *val = NULL;
1874         size_t szval = 0, numval = 0;
1875         GET_VAL_ARRAY(ret, loc);
1876         if (!ret->err) {
1877                 size_t counter = 0;
1878                 set_separator(output->mode == CLINFO_HUMAN ? comma_str : output->json ? comma_str : spc_str);
1879                 if (output->json)
1880                         strbuf_append_str_len(loc->pname, &ret->str, " [", 2);
1881                 for (counter = 0; counter < numval; ++counter) {
1882                         if (counter > 0) strbuf_append_str(loc->pname, &ret->str, sep);
1883                         strbuf_append(loc->pname, &ret->str, "%" PRId32, val[counter]);
1884                 }
1885                 if (output->json)
1886                         strbuf_append_str_len(loc->pname, &ret->str, " ]", 2);
1887                 // TODO: ret->value.??? = val;
1888         }
1889         free(val);
1890 }
1891
1892 void
1893 device_info_szptr_sep(struct device_info_ret *ret, const char *human_sep,
1894         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
1895         const struct opt_out *output)
1896 {
1897         size_t *val = NULL;
1898         size_t szval = 0, numval = 0;
1899         GET_VAL_ARRAY(ret, loc);
1900         if (!ret->err) {
1901                 size_t counter = 0;
1902                 set_separator(output->mode == CLINFO_HUMAN ? human_sep : output->json ? comma_str : spc_str);
1903                 if (output->json)
1904                         strbuf_append_str_len(loc->pname, &ret->str, " [", 2);
1905                 for (counter = 0; counter < numval; ++counter) {
1906                         if (counter > 0) strbuf_append_str(loc->pname, &ret->str, sep);
1907                         strbuf_append(loc->pname, &ret->str, "%" PRIuS, val[counter]);
1908                 }
1909                 if (output->json)
1910                         strbuf_append_str_len(loc->pname, &ret->str, " ]", 2);
1911                 // TODO: ret->value.??? = val;
1912         }
1913         free(val);
1914 }
1915
1916
1917 void
1918 device_info_szptr_times(struct device_info_ret *ret,
1919         const struct info_loc *loc, const struct device_info_checks* chk,
1920         const struct opt_out *output)
1921 {
1922         device_info_szptr_sep(ret, times_str, loc, chk, output);
1923 }
1924
1925 void
1926 device_info_szptr_comma(struct device_info_ret *ret,
1927         const struct info_loc *loc, const struct device_info_checks* chk,
1928         const struct opt_out *output)
1929 {
1930         device_info_szptr_sep(ret, comma_str, loc, chk, output);
1931 }
1932
1933 void
1934 getWGsizes(struct device_info_ret *ret, const struct info_loc *loc, size_t *wgm, size_t wgm_sz,
1935         const struct opt_out* UNUSED(output))
1936 {
1937         cl_int log_err;
1938
1939         cl_context_properties ctxpft[] = {
1940                 CL_CONTEXT_PLATFORM, (cl_context_properties)loc->plat,
1941                 0, 0 };
1942         cl_uint cursor = 0;
1943         cl_context ctx = NULL;
1944         cl_program prg = NULL;
1945         cl_kernel krn = NULL;
1946
1947         ret->err = CL_SUCCESS;
1948
1949         ctx = clCreateContext(ctxpft, 1, &loc->dev, NULL, NULL, &ret->err);
1950         if (REPORT_ERROR(&ret->err_str, ret->err, "create context")) goto out;
1951         prg = clCreateProgramWithSource(ctx, ARRAY_SIZE(sources), sources, NULL, &ret->err);
1952         if (REPORT_ERROR(&ret->err_str, ret->err, "create program")) goto out;
1953         ret->err = clBuildProgram(prg, 1, &loc->dev, NULL, NULL, NULL);
1954         log_err = REPORT_ERROR(&ret->err_str, ret->err, "build program");
1955
1956         /* for a program build failure, dump the log to stderr before bailing */
1957         if (log_err == CL_BUILD_PROGRAM_FAILURE) {
1958                 struct _strbuf logbuf;
1959                 init_strbuf(&logbuf, "program build log");
1960                 GET_STRING(&logbuf, ret->err,
1961                         clGetProgramBuildInfo, CL_PROGRAM_BUILD_LOG, "CL_PROGRAM_BUILD_LOG", prg, loc->dev);
1962                 if (ret->err == CL_SUCCESS) {
1963                         fflush(stdout);
1964                         fflush(stderr);
1965                         fputs("=== CL_PROGRAM_BUILD_LOG ===\n", stderr);
1966                         fputs(logbuf.buf, stderr);
1967                         fflush(stderr);
1968                 }
1969                 free_strbuf(&logbuf);
1970         }
1971         if (ret->err)
1972                 goto out;
1973
1974         for (cursor = 0; cursor < wgm_sz; ++cursor) {
1975                 strbuf_append(__func__, &ret->str, "sum%u", 1<<cursor);
1976                 if (cursor == 0)
1977                         ret->str.buf[3] = 0; // scalar kernel is called 'sum'
1978                 krn = clCreateKernel(prg, ret->str.buf, &ret->err);
1979                 reset_strbuf(&ret->str);
1980                 if (REPORT_ERROR(&ret->err_str, ret->err, "create kernel")) goto out;
1981                 ret->err = clGetKernelWorkGroupInfo(krn, loc->dev, CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE,
1982                         sizeof(*wgm), wgm + cursor, NULL);
1983                 if (REPORT_ERROR(&ret->err_str, ret->err, "get kernel info")) goto out;
1984                 clReleaseKernel(krn);
1985                 krn = NULL;
1986         }
1987
1988 out:
1989         if (krn)
1990                 clReleaseKernel(krn);
1991         if (prg)
1992                 clReleaseProgram(prg);
1993         if (ctx)
1994                 clReleaseContext(ctx);
1995 }
1996
1997
1998 void
1999 device_info_wg(struct device_info_ret *ret,
2000         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
2001         const struct opt_out *output)
2002 {
2003         /* preferred workgroup size multiple for each kernel
2004          * have not found a platform where the WG multiple changes,
2005          * but keep this flexible (this can grow up to 5)
2006          */
2007 #define NUM_KERNELS 1
2008         size_t wgm[NUM_KERNELS] = {0};
2009
2010         getWGsizes(ret, loc, wgm, NUM_KERNELS, output);
2011         if (!ret->err) {
2012                 strbuf_append("get WG sizes", &ret->str, "%" PRIuS, wgm[0]);
2013         }
2014         ret->value.s = wgm[0];
2015 }
2016
2017 void
2018 device_info_img_sz_2d(struct device_info_ret *ret,
2019         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
2020         const struct opt_out *output)
2021 {
2022         struct info_loc loc2 = *loc;
2023         size_t width = 0, height = 0;
2024         _GET_VAL(ret, loc, height); /* HEIGHT */
2025         if (!ret->err) {
2026                 RESET_LOC_PARAM(loc2, dev, CL_DEVICE_IMAGE2D_MAX_WIDTH);
2027                 _GET_VAL(ret, &loc2, width);
2028                 if (!ret->err) {
2029                         strbuf_append("image size 2D", &ret->str, "%" PRIuS "x%" PRIuS, width, height);
2030                 }
2031         }
2032         ret->value.u64v.s[0] = width;
2033         ret->value.u64v.s[1] = height;
2034 }
2035
2036 void
2037 device_info_img_sz_intel_planar_yuv(struct device_info_ret *ret,
2038         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
2039         const struct opt_out *output)
2040 {
2041         struct info_loc loc2 = *loc;
2042         size_t width = 0, height = 0;
2043         _GET_VAL(ret, loc, height); /* HEIGHT */
2044         if (!ret->err) {
2045                 RESET_LOC_PARAM(loc2, dev, CL_DEVICE_PLANAR_YUV_MAX_WIDTH_INTEL);
2046                 _GET_VAL(ret, &loc2, width);
2047                 if (!ret->err) {
2048                          strbuf_append("image size planar YUV", &ret->str, "%" PRIuS "x%" PRIuS, width, height);
2049                 }
2050         }
2051         ret->value.u64v.s[0] = width;
2052         ret->value.u64v.s[1] = height;
2053 }
2054
2055
2056 void
2057 device_info_img_sz_3d(struct device_info_ret *ret,
2058         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
2059         const struct opt_out *output)
2060 {
2061         struct info_loc loc2 = *loc;
2062         size_t width = 0, height = 0, depth = 0;
2063         _GET_VAL(ret, loc, height); /* HEIGHT */
2064         if (!ret->err) {
2065                 RESET_LOC_PARAM(loc2, dev, CL_DEVICE_IMAGE3D_MAX_WIDTH);
2066                 _GET_VAL(ret, &loc2, width);
2067                 if (!ret->err) {
2068                         RESET_LOC_PARAM(loc2, dev, CL_DEVICE_IMAGE3D_MAX_DEPTH);
2069                         _GET_VAL(ret, &loc2, depth);
2070                         if (!ret->err) {
2071                                 strbuf_append("image size 3D", &ret->str,
2072                                         "%" PRIuS "x%" PRIuS "x%" PRIuS,
2073                                         width, height, depth);
2074                         }
2075                 }
2076         }
2077         ret->value.u64v.s[0] = width;
2078         ret->value.u64v.s[1] = height;
2079         ret->value.u64v.s[2] = depth;
2080 }
2081
2082 void strbuf_bitfield(const char *what, struct _strbuf *str,
2083         cl_bitfield bits, const char *bits_name,
2084         const char * const *bit_str, size_t bit_str_count,
2085         const struct opt_out *output)
2086 {
2087         const char *quote = output->json ? "\"" : "";
2088         /* number of matches so far, for separator placement */
2089         cl_uint count = 0;
2090         /* iterator */
2091         cl_uint i = 0;
2092         /* leftovers bits */
2093         cl_bitfield known_mask, extra;
2094
2095         set_common_separator(output);
2096
2097         if (output->json)
2098                 strbuf_append(what, str,
2099                         "{ \"raw\" : %" PRIu64 ", \"%s\" : [ ",
2100                         bits, bits_name);
2101
2102         if (bits) {
2103                 for (i = 0; i < bit_str_count; ++i) {
2104                         if (bits & (1UL << i)) {
2105                                 strbuf_append(what, str, "%s%s%s%s",
2106                                         (count > 0 ? sep : ""),
2107                                         quote, bit_str[i], quote);
2108                                 ++count;
2109                         }
2110                 }
2111
2112                 /* check for extra bits */
2113                 known_mask = ((cl_bitfield)(1) << bit_str_count) - 1;
2114                 extra = bits & ~known_mask;
2115                 if (extra) {
2116                         strbuf_append(what, str, "%s%s%#" PRIx64 "%s",
2117                                 (count > 0 ? sep : ""), quote, extra, quote);
2118                 }
2119         }
2120
2121         if (output->json)
2122                 strbuf_append_str(what, str, " ] }");
2123 }
2124
2125
2126 void
2127 device_info_bitfield(struct device_info_ret *ret,
2128         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
2129         const struct opt_out *output,
2130         const cl_bitfield bits,
2131         const size_t bit_str_count, /* number of entries in bit_str */
2132         const char * const * bit_str, /* array of strings describing the bits */
2133         const char * bits_name) /* JSON name for this bitfield */
2134 {
2135         strbuf_bitfield(loc->pname, &ret->str, bits, bits_name, bit_str, bit_str_count, output);
2136 }
2137
2138
2139 /* This could use device_info_bitfield, but we prefer to go through fields in reverse,
2140  * so we just dup the code
2141  */
2142 void
2143 device_info_devtype(struct device_info_ret *ret,
2144         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
2145         const struct opt_out *output)
2146 {
2147         GET_VAL(ret, loc, devtype);
2148         if (!ret->err) {
2149                 const char *quote = output->json ? "\"" : "";
2150                 const char * const *devstr = (output->mode == CLINFO_HUMAN ?
2151                         device_type_str : device_type_raw_str);
2152                 cl_uint i = (cl_uint)actual_devtype_count;
2153                 /* number of matches so far, for separator placement */
2154                 cl_uint count = 0;
2155                 /* leftovers bits */
2156                 cl_device_type known_mask, extra;
2157
2158                 set_common_separator(output);
2159
2160                 if (output->json)
2161                         strbuf_append(loc->pname, &ret->str,
2162                                 "{ \"raw\" : %" PRIu64 ", \"type\" : [ ",
2163                                 ret->value.devtype);
2164
2165                 /* iterate over device type strings, appending their textual form
2166                  * to ret->str */
2167                 for (; i > 0; --i) {
2168                         /* assemble CL_DEVICE_TYPE_* from index i */
2169                         cl_device_type cur = (cl_device_type)(1) << (i-1);
2170                         if (ret->value.devtype & cur) {
2171                                 /* match: add separator if not first match */
2172                                 strbuf_append(loc->pname, &ret->str, "%s%s%s%s",
2173                                         (count > 0 ? sep : ""),
2174                                         quote, devstr[i], quote);
2175                                 ++count;
2176                         }
2177                 }
2178
2179                 /* check for extra bits */
2180                 known_mask = ((cl_device_type)(1) << actual_devtype_count) - 1;
2181                 extra = ret->value.devtype & ~known_mask;
2182                 if (extra) {
2183                         strbuf_append(loc->pname, &ret->str, "%s%s%#" PRIx64 "%s",
2184                                 (count > 0 ? sep : ""), quote, extra, quote);
2185                 }
2186
2187                 if (output->json)
2188                         strbuf_append_str(loc->pname, &ret->str, " ] }");
2189         }
2190 }
2191
2192 void
2193 device_info_cachetype(struct device_info_ret *ret,
2194         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
2195         const struct opt_out *output)
2196 {
2197         GET_VAL(ret, loc, cachetype);
2198         if (!ret->err) {
2199                 const char * const *ar = (output->mode == CLINFO_HUMAN ?
2200                         cache_type_str : cache_type_raw_str);
2201                 strbuf_append_str(loc->pname, &ret->str, ar[ret->value.cachetype]);
2202                 ret->needs_escaping = CL_TRUE;
2203         }
2204 }
2205
2206 void
2207 device_info_lmemtype(struct device_info_ret *ret,
2208         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
2209         const struct opt_out *output)
2210 {
2211         GET_VAL(ret, loc, lmemtype);
2212         if (!ret->err) {
2213                 const char * const *ar = (output->mode == CLINFO_HUMAN ?
2214                         lmem_type_str : lmem_type_raw_str);
2215                 strbuf_append_str(loc->pname, &ret->str, ar[ret->value.lmemtype]);
2216                 ret->needs_escaping = CL_TRUE;
2217         }
2218 }
2219
2220 void
2221 device_info_atomic_caps(struct device_info_ret *ret,
2222         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
2223         const struct opt_out *output)
2224 {
2225         GET_VAL(ret, loc, bits);
2226         if (!ret->err) {
2227                 device_info_bitfield(ret, loc, chk, output, ret->value.bits,
2228                         atomic_cap_count, (output->mode == CLINFO_HUMAN ?
2229                                 atomic_cap_str : atomic_cap_raw_str),
2230                         "capabilities");
2231         }
2232 }
2233
2234 void
2235 device_info_device_enqueue_caps(struct device_info_ret *ret,
2236         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
2237         const struct opt_out *output)
2238 {
2239         GET_VAL(ret, loc, bits);
2240         if (!ret->err) {
2241                 device_info_bitfield(ret, loc, chk, output, ret->value.bits,
2242                         device_enqueue_cap_count, (output->mode == CLINFO_HUMAN ?
2243                                 device_enqueue_cap_str : device_enqueue_cap_raw_str),
2244                         "capabilities");
2245         }
2246 }
2247
2248 /* cl_arm_core_id */
2249 void
2250 device_info_core_ids(struct device_info_ret *ret,
2251         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
2252         const struct opt_out *output)
2253 {
2254         cl_ulong val;
2255         GET_VAL(ret, loc, u64);
2256         val = ret->value.u64;
2257
2258         if (!ret->err) {
2259                 const char *quote = output->json ? "\"" : "";
2260                 /* The value is a bitfield where each set bit corresponds to a core ID
2261                  * value that can be returned by the device-side function. We print them
2262                  * here as ranges, such as 0-4, 8-12 */
2263                 int range_start = -1;
2264                 int cur_bit = 0;
2265
2266                 if (output->json)
2267                         strbuf_append(loc->pname, &ret->str,
2268                                 "{ \"raw\" : %" PRIu64 ", \"core_ids\" : [ ",
2269                                 ret->value.u64);
2270
2271                 set_separator(empty_str);
2272 #define CORE_ID_END 64
2273                 do {
2274                         /* Find the start of the range */
2275                         while ((cur_bit < CORE_ID_END) && !((val >> cur_bit) & 1))
2276                                 ++cur_bit;
2277                         range_start = cur_bit++;
2278
2279                         /* Find the end of the range */
2280                         while ((cur_bit < CORE_ID_END) && ((val >> cur_bit) & 1))
2281                                 ++cur_bit;
2282
2283                         /* print the range [range_start, cur_bit[ */
2284                         if (range_start >= 0 && range_start < CORE_ID_END) {
2285                                 strbuf_append(loc->pname, &ret->str, "%s%s%d", sep, quote, range_start);
2286                                 if (cur_bit - range_start > 1)
2287                                         strbuf_append(loc->pname, &ret->str, "-%d", cur_bit - 1);
2288                                 set_separator(comma_str);
2289                                 if (output->json)
2290                                         strbuf_append_str(loc->pname, &ret->str, quote);
2291                         }
2292                 } while (cur_bit < CORE_ID_END);
2293
2294                 if (output->json)
2295                         strbuf_append_str(loc->pname, &ret->str, " ] }");
2296         }
2297 }
2298
2299 /* cl_arm_job_slot_selection */
2300 void
2301 device_info_job_slots(struct device_info_ret *ret,
2302         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
2303         const struct opt_out *output)
2304 {
2305         cl_uint val;
2306         GET_VAL(ret, loc, u32);
2307         val = ret->value.u32;
2308
2309         if (!ret->err) {
2310                 const char *quote = output->json ? "\"" : "";
2311                 /* The value is a bitfield where each set bit corresponds to an available job slot.
2312                  * We print them here as ranges, such as 0-4, 8-12 */
2313                 int range_start = -1;
2314                 int cur_bit = 0;
2315
2316                 if (output->json)
2317                         strbuf_append(loc->pname, &ret->str,
2318                                 "{ \"raw\" : %" PRIu32 ", \"slots\" : [ ",
2319                                 ret->value.u32);
2320
2321                 set_separator(empty_str);
2322 #define JOB_SLOT_END 32
2323                 do {
2324                         /* Find the start of the range */
2325                         while ((cur_bit < JOB_SLOT_END) && !((val >> cur_bit) & 1))
2326                                 ++cur_bit;
2327                         range_start = cur_bit++;
2328
2329                         /* Find the end of the range */
2330                         while ((cur_bit < JOB_SLOT_END) && ((val >> cur_bit) & 1))
2331                                 ++cur_bit;
2332
2333                         /* print the range [range_start, cur_bit[ */
2334                         if (range_start >= 0 && range_start < JOB_SLOT_END) {
2335                                 strbuf_append(loc->pname, &ret->str, "%s%s%d", sep, quote, range_start);
2336                                 if (cur_bit - range_start > 1)
2337                                         strbuf_append(loc->pname, &ret->str, "-%d", cur_bit - 1);
2338                                 set_separator(comma_str);
2339                                 if (output->json)
2340                                         strbuf_append_str(loc->pname, &ret->str, quote);
2341                         }
2342                 } while (cur_bit < JOB_SLOT_END);
2343
2344                 if (output->json)
2345                         strbuf_append_str(loc->pname, &ret->str, " ] }");
2346         }
2347 }
2348
2349 void devtopo_pci_str(struct device_info_ret *ret, const cl_device_pci_bus_info_khr *devtopo)
2350 {
2351         strbuf_append("devtopo", &ret->str, "PCI-E, %04x:%02x:%02x.%u",
2352                 devtopo->pci_domain,
2353                 devtopo->pci_bus,
2354                 devtopo->pci_device, devtopo->pci_function);
2355         ret->value.devtopo_khr = *devtopo;
2356 }
2357
2358 void
2359 device_info_devtopo_khr(struct device_info_ret *ret,
2360         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
2361         const struct opt_out *output)
2362 {
2363         GET_VAL(ret, loc, devtopo_khr);
2364         /* TODO how to do this in CLINFO_RAW mode */
2365         if (!ret->err) {
2366                 devtopo_pci_str(ret, &ret->value.devtopo_khr);
2367                 /* TODO JSONify */
2368                 ret->needs_escaping = CL_TRUE;
2369         }
2370 }
2371
2372
2373 /* stringify a cl_device_topology_amd */
2374 void devtopo_amd_str(struct device_info_ret *ret, const cl_device_topology_amd *devtopo)
2375 {
2376         cl_device_pci_bus_info_khr devtopo_info;
2377
2378         switch (devtopo->raw.type) {
2379         case 0:
2380                 /* leave empty */
2381                 break;
2382         case CL_DEVICE_TOPOLOGY_TYPE_PCIE_AMD:
2383                 devtopo_info.pci_domain = 0;
2384                 devtopo_info.pci_bus = devtopo->pcie.bus;
2385                 devtopo_info.pci_device = devtopo->pcie.device;
2386                 devtopo_info.pci_function = devtopo->pcie.function;
2387                 devtopo_pci_str(ret, &devtopo_info);
2388                 break;
2389         default:
2390                 strbuf_append("devtopo", &ret->str, "<unknown (%u): %u %u %u %u %u>",
2391                         devtopo->raw.type,
2392                         devtopo->raw.data[0], devtopo->raw.data[1],
2393                         devtopo->raw.data[2],
2394                         devtopo->raw.data[3], devtopo->raw.data[4]);
2395         }
2396 }
2397
2398 void
2399 device_info_devtopo_amd(struct device_info_ret *ret,
2400         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
2401         const struct opt_out *output)
2402 {
2403         GET_VAL(ret, loc, devtopo_amd);
2404         /* TODO how to do this in CLINFO_RAW mode */
2405         if (!ret->err) {
2406                 devtopo_amd_str(ret, &ret->value.devtopo_amd);
2407                 /* TODO JSONify */
2408                 ret->needs_escaping = CL_TRUE;
2409         }
2410 }
2411
2412 /* we assemble a clinfo_device_topology_pci struct from the NVIDIA info */
2413 void
2414 device_info_devtopo_nv(struct device_info_ret *ret,
2415         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
2416         const struct opt_out *output)
2417 {
2418         struct info_loc loc2 = *loc;
2419         cl_device_pci_bus_info_khr devtopo;
2420         DEV_FETCH(cl_uint, val); /* CL_DEVICE_PCI_BUS_ID_NV */
2421         if (!ret->err) {
2422                 devtopo.pci_bus = val & 0xff;
2423                 RESET_LOC_PARAM(loc2, dev, CL_DEVICE_PCI_SLOT_ID_NV);
2424                 _GET_VAL(ret, &loc2, val);
2425
2426                 if (!ret->err) {
2427                         cl_int safe_err;
2428                         devtopo.pci_device = (val >> 3) & 0xff;
2429                         devtopo.pci_function = val & 7;
2430
2431                         /* CL_DEVICE_PCI_DOMAIN_ID_NV is not supported in older drivers,
2432                          * but we have no way to check other than querying, and recovering
2433                          * in the CL_INVALID_VALUE case */
2434                         RESET_LOC_PARAM(loc2, dev, CL_DEVICE_PCI_DOMAIN_ID_NV);
2435                         safe_err = clGetDeviceInfo(loc2.dev, CL_DEVICE_PCI_DOMAIN_ID_NV,
2436                                 sizeof(val), &val, NULL);
2437                         if (safe_err == CL_SUCCESS) {
2438                                 devtopo.pci_domain = val;
2439                         } else if (safe_err == CL_INVALID_VALUE) {
2440                                 devtopo.pci_domain = 0;
2441                         } else {
2442                                 REPORT_ERROR_LOC(ret, safe_err, &loc2, "get CL_DEVICE_PCI_DOMAIN_ID_NV");
2443                         }
2444                         if (!ret->err)
2445                                 devtopo_pci_str(ret, &devtopo);
2446                 }
2447         }
2448 }
2449
2450 /* NVIDIA Compute Capability */
2451 void
2452 device_info_cc_nv(struct device_info_ret *ret,
2453         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
2454         const struct opt_out *output)
2455 {
2456         struct info_loc loc2 = *loc;
2457         cl_uint major = 0, minor = 0;
2458         _GET_VAL(ret, loc, major); /* MAJOR */
2459         if (!ret->err) {
2460                 RESET_LOC_PARAM(loc2, dev, CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV);
2461                 _GET_VAL(ret, &loc2, minor);
2462                 if (!ret->err) {
2463                         strbuf_append("NV CC", &ret->str, "%" PRIu32 ".%" PRIu32 "", major, minor);
2464                 }
2465         }
2466         ret->value.u32v.s[0] = major;
2467         ret->value.u32v.s[1] = minor;
2468 }
2469
2470 /* AMD GFXIP */
2471 void
2472 device_info_gfxip_amd(struct device_info_ret *ret,
2473         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
2474         const struct opt_out *output)
2475 {
2476         struct info_loc loc2 = *loc;
2477         cl_uint major = 0, minor = 0;
2478         _GET_VAL(ret, loc, major); /* MAJOR */
2479         if (!ret->err) {
2480                 RESET_LOC_PARAM(loc2, dev, CL_DEVICE_GFXIP_MINOR_AMD);
2481                 _GET_VAL(ret, &loc2, minor);
2482                 if (!ret->err) {
2483                         strbuf_append("AMD GFXIP", &ret->str, "%" PRIu32 ".%" PRIu32 "", major, minor);
2484                 }
2485         }
2486         ret->value.u32v.s[0] = major;
2487         ret->value.u32v.s[1] = minor;
2488 }
2489
2490 /* Intel feature capabilities */
2491 void
2492 device_info_intel_features(struct device_info_ret *ret,
2493         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
2494         const struct opt_out *output)
2495 {
2496         GET_VAL(ret, loc, bits);
2497         device_info_bitfield(ret, loc, chk, output, ret->value.bits, intel_features_count,
2498                 (output->mode == CLINFO_HUMAN ? intel_features_str : intel_features_raw_str),
2499                 "features_intel");
2500 }
2501
2502
2503
2504 /* Device Partition, CLINFO_HUMAN header */
2505 void
2506 device_info_partition_header(struct device_info_ret *ret,
2507         const struct info_loc *UNUSED(loc), const struct device_info_checks *chk,
2508         const struct opt_out* UNUSED(output))
2509 {
2510         cl_bool is_12 = dev_is_12(chk);
2511         cl_bool has_fission = dev_has_fission(chk);
2512         strbuf_append("dev partition", &ret->str, "(%s%s%s%s)",
2513                 (is_12 ? core : empty_str),
2514                 (is_12 && has_fission ? comma_str : empty_str),
2515                 chk->has_fission,
2516                 (!(is_12 || has_fission) ? na : empty_str));
2517
2518         ret->err = CL_SUCCESS;
2519 }
2520
2521 /* Device partition properties */
2522 void
2523 device_info_partition_types(struct device_info_ret *ret,
2524         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
2525         const struct opt_out *output)
2526 {
2527         size_t numval = 0, szval = 0, cursor = 0;
2528         cl_device_partition_property *val = NULL;
2529         const char * const *ptstr = (output->mode == CLINFO_HUMAN ?
2530                 partition_type_str : partition_type_raw_str);
2531
2532         GET_VAL_ARRAY(ret, loc);
2533
2534         if (!ret->err) {
2535                 const char *quote = output->json ? "\"" : "";
2536                 set_common_separator(output);
2537                 if (output->json)
2538                         strbuf_append_str_len(loc->pname, &ret->str, "[ ", 2);
2539
2540                 for (cursor = 0; cursor < numval; ++cursor) {
2541                         int str_idx = -1;
2542
2543                         /* add separator for values past the first */
2544                         if (cursor > 0) strbuf_append_str(loc->pname, &ret->str, sep);
2545
2546                         switch (val[cursor]) {
2547                         case 0: str_idx = 0; break;
2548                         case CL_DEVICE_PARTITION_EQUALLY: str_idx = 1; break;
2549                         case CL_DEVICE_PARTITION_BY_COUNTS: str_idx = 2; break;
2550                         case CL_DEVICE_PARTITION_BY_AFFINITY_DOMAIN: str_idx = 3; break;
2551                         case CL_DEVICE_PARTITION_BY_NAMES_INTEL: str_idx = 4; break;
2552                         default:
2553                                 strbuf_append(loc->pname, &ret->str,
2554                                         "%sby <unknown> (%#" PRIxPTR ")%s",
2555                                         quote, val[cursor], quote);
2556                                 break;
2557                         }
2558                         if (str_idx >= 0) {
2559                                 /* string length, minus _EXT */
2560                                 size_t slen = strlen(ptstr[str_idx]);
2561                                 if (output->mode == CLINFO_RAW && str_idx > 0)
2562                                         slen -= 4;
2563                                 strbuf_append_str(loc->pname, &ret->str, quote);
2564                                 strbuf_append_str_len(loc->pname, &ret->str, ptstr[str_idx], slen);
2565                                 strbuf_append_str(loc->pname, &ret->str, quote);
2566                         }
2567                 }
2568                 if (output->json)
2569                         strbuf_append_str_len(loc->pname, &ret->str, " ]", 2);
2570                 // TODO ret->value.??? = val
2571         }
2572         free(val);
2573 }
2574
2575 void
2576 device_info_partition_types_ext(struct device_info_ret *ret,
2577         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
2578         const struct opt_out *output)
2579 {
2580         size_t numval = 0, szval = 0, cursor = 0;
2581         cl_device_partition_property_ext *val = NULL;
2582         const char * const *ptstr = (output->mode == CLINFO_HUMAN ?
2583                 partition_type_str : partition_type_raw_str);
2584
2585         GET_VAL_ARRAY(ret, loc);
2586
2587         if (!ret->err) {
2588                 const char *quote = output->json ? "\"" : "";
2589                 set_common_separator(output);
2590                 if (output->json)
2591                         strbuf_append_str_len(loc->pname, &ret->str, "[ ", 1);
2592
2593                 for (cursor = 0; cursor < numval; ++cursor) {
2594                         int str_idx = -1;
2595
2596                         /* add separator for values past the first */
2597                         if (cursor > 0) strbuf_append_str(loc->pname, &ret->str, sep);
2598
2599                         switch (val[cursor]) {
2600                         case 0: str_idx = 0; break;
2601                         case CL_DEVICE_PARTITION_EQUALLY_EXT: str_idx = 1; break;
2602                         case CL_DEVICE_PARTITION_BY_COUNTS_EXT: str_idx = 2; break;
2603                         case CL_DEVICE_PARTITION_BY_AFFINITY_DOMAIN_EXT: str_idx = 3; break;
2604                         case CL_DEVICE_PARTITION_BY_NAMES_EXT: str_idx = 4; break;
2605                         default:
2606                                 strbuf_append(loc->pname, &ret->str,
2607                                         "%sby <unknown> (%#" PRIx64 ")%s",
2608                                         quote, val[cursor], quote);
2609                                 break;
2610                         }
2611                         if (str_idx >= 0) {
2612                                 strbuf_append(loc->pname, &ret->str, "%s%s%s",
2613                                         quote, ptstr[str_idx], quote);
2614                         }
2615                 }
2616                 if (output->json)
2617                         strbuf_append_str_len(loc->pname, &ret->str, " ]", 2);
2618                 // TODO ret->value.??? = val
2619         }
2620         free(val);
2621 }
2622
2623
2624 /* Device partition affinity domains */
2625 void
2626 device_info_partition_affinities(struct device_info_ret *ret,
2627         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
2628         const struct opt_out *output)
2629 {
2630         GET_VAL(ret, loc, affinity_domain);
2631
2632         if (!ret->err) {
2633                 device_info_bitfield(ret, loc, chk, output, ret->value.affinity_domain,
2634                         affinity_domain_count, (output->mode == CLINFO_HUMAN ?
2635                                 affinity_domain_str : affinity_domain_raw_str),
2636                         "domain");
2637         }
2638 }
2639
2640 void
2641 device_info_partition_affinities_ext(struct device_info_ret *ret,
2642         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
2643         const struct opt_out *output)
2644 {
2645         size_t numval = 0, szval = 0, cursor = 0;
2646         cl_device_partition_property_ext *val = NULL;
2647         const char * const *ptstr = (output->mode == CLINFO_HUMAN ?
2648                 affinity_domain_ext_str : affinity_domain_raw_ext_str);
2649
2650         GET_VAL_ARRAY(ret, loc);
2651
2652         if (!ret->err) {
2653                 const char *quote = output->json ? "\"" : "";
2654                 set_common_separator(output);
2655                 if (output->json)
2656                         strbuf_append_str_len(loc->pname, &ret->str, "[ ", 2);
2657
2658                 for (cursor = 0; cursor < numval; ++cursor) {
2659                         int str_idx = -1;
2660
2661                         /* add separator for values past the first */
2662                         if (cursor > 0) strbuf_append_str(loc->pname, &ret->str, sep);
2663
2664                         switch (val[cursor]) {
2665                         case CL_AFFINITY_DOMAIN_NUMA_EXT: str_idx = 0; break;
2666                         case CL_AFFINITY_DOMAIN_L4_CACHE_EXT: str_idx = 1; break;
2667                         case CL_AFFINITY_DOMAIN_L3_CACHE_EXT: str_idx = 2; break;
2668                         case CL_AFFINITY_DOMAIN_L2_CACHE_EXT: str_idx = 3; break;
2669                         case CL_AFFINITY_DOMAIN_L1_CACHE_EXT: str_idx = 4; break;
2670                         case CL_AFFINITY_DOMAIN_NEXT_FISSIONABLE_EXT: str_idx = 5; break;
2671                         default:
2672                                 strbuf_append(loc->pname, &ret->str,
2673                                         "%s<unknown> (%#" PRIx64 ")%s",
2674                                         quote, val[cursor], quote);
2675                                 break;
2676                         }
2677                         if (str_idx >= 0) {
2678                                 strbuf_append(loc->pname, &ret->str, "%s%s%s",
2679                                         quote, ptstr[str_idx], quote);
2680                         }
2681                 }
2682                 if (output->json)
2683                         strbuf_append_str_len(loc->pname, &ret->str, " ]", 2);
2684                 // TODO: ret->value.??? = val
2685         }
2686         free(val);
2687 }
2688
2689 /* Preferred / native vector widths */
2690 void
2691 device_info_vecwidth(struct device_info_ret *ret,
2692         const struct info_loc *loc, const struct device_info_checks *chk,
2693         const struct opt_out *output)
2694 {
2695         struct info_loc loc2 = *loc;
2696         cl_uint preferred = 0, native = 0;
2697         _GET_VAL(ret, loc, preferred);
2698         if (!ret->err) {
2699                 /* we get called with PREFERRED, NATIVE is at +0x30 offset, except for HALF,
2700                  * which is at +0x08 */
2701                 loc2.param.dev +=
2702                         (loc2.param.dev == CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF ? 0x08 : 0x30);
2703                 /* TODO update loc2.sname */
2704                 _GET_VAL(ret, &loc2, native);
2705
2706                 if (!ret->err) {
2707                         const char *ext = (loc2.param.dev == CL_DEVICE_NATIVE_VECTOR_WIDTH_HALF ?
2708                                 chk->has_half : (loc2.param.dev == CL_DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE ?
2709                                 chk->has_double : NULL));
2710                         strbuf_append(loc->pname, &ret->str, "%8u / %-8u", preferred, native);
2711                         if (ext)
2712                                 strbuf_append(loc->pname, &ret->str, " (%s)", *ext ? ext : na);
2713                 }
2714         }
2715         ret->value.u32v.s[0] = preferred;
2716         ret->value.u32v.s[1] = native;
2717 }
2718
2719 /* Floating-point configurations */
2720 void
2721 device_info_fpconf(struct device_info_ret *ret,
2722         const struct info_loc *loc, const struct device_info_checks *chk,
2723         const struct opt_out *output)
2724 {
2725         /* When in HUMAN output, we are called unconditionally,
2726          * so we have to do some manual checks ourselves */
2727         const cl_bool get_it = (output->mode != CLINFO_HUMAN) ||
2728                 (loc->param.dev == CL_DEVICE_SINGLE_FP_CONFIG) ||
2729                 (loc->param.dev == CL_DEVICE_HALF_FP_CONFIG && dev_has_half(chk)) ||
2730                 (loc->param.dev == CL_DEVICE_DOUBLE_FP_CONFIG && dev_has_double(chk));
2731
2732         GET_VAL(ret, loc, fpconfig);
2733         /* Sanitize! */
2734         if (ret->err && !get_it) {
2735                 ret->err = CL_SUCCESS;
2736                 ret->value.fpconfig = 0;
2737         }
2738
2739         if (output->json)
2740                 strbuf_append(loc->pname, &ret->str,
2741                         "{ \"raw\" : %" PRIu64 ", \"config\" : [ ",
2742                         ret->value.fpconfig);
2743
2744         if (!ret->err) {
2745                 cl_uint i = 0;
2746                 cl_uint count = 0;
2747                 const char * const *fpstr = (output->mode == CLINFO_HUMAN ?
2748                         fp_conf_str : fp_conf_raw_str);
2749                 set_common_separator(output);
2750                 if (output->mode == CLINFO_HUMAN) {
2751                         const char *why = na;
2752                         switch (loc->param.dev) {
2753                         case CL_DEVICE_HALF_FP_CONFIG:
2754                                 if (get_it)
2755                                         why = chk->has_half;
2756                                 break;
2757                         case CL_DEVICE_SINGLE_FP_CONFIG:
2758                                 why = core;
2759                                 break;
2760                         case CL_DEVICE_DOUBLE_FP_CONFIG:
2761                                 if (get_it)
2762                                         why = chk->has_double;
2763                                 break;
2764                         default:
2765                                 /* "this can't happen" (unless OpenCL starts supporting _other_ floating-point formats, maybe) */
2766                                 fprintf(stderr, "unsupported floating-point configuration parameter %s\n", loc->pname);
2767                         }
2768                         /* show 'why' it's being shown */
2769                         strbuf_append(loc->pname, &ret->str, "(%s)", why);
2770                 }
2771                 if (get_it) {
2772                         const char *quote = output->json ? "\"" : "";
2773                         size_t num_flags = fp_conf_count;
2774                         /* The last flag, CL_FP_CORRECTLY_ROUNDED_DIVIDE_SQRT is only considered
2775                          * in the single-precision case. half and double don't consider it,
2776                          * so we skip it altogether */
2777                         if (loc->param.dev != CL_DEVICE_SINGLE_FP_CONFIG)
2778                                 num_flags -= 1;
2779
2780                         for (i = 0; i < num_flags; ++i) {
2781                                 cl_device_fp_config cur = (cl_device_fp_config)(1) << i;
2782                                 cl_bool present = !!(ret->value.fpconfig & cur);
2783                                 if (output->mode == CLINFO_HUMAN) {
2784                                         strbuf_append(loc->pname, &ret->str, "\n%s" I2_STR "%s",
2785                                                 line_pfx, fpstr[i], bool_str[present]);
2786                                 } else if (present) {
2787                                         strbuf_append(loc->pname, &ret->str, "%s%s%s%s",
2788                                                 (count > 0 ? sep : ""), quote, fpstr[i], quote);
2789                                         ++count;
2790                                 }
2791                         }
2792                 }
2793         }
2794         if (output->json)
2795                 strbuf_append_str(loc->pname, &ret->str, " ] }");
2796 }
2797
2798 /* Queue properties */
2799 void
2800 device_info_qprop(struct device_info_ret *ret,
2801         const struct info_loc *loc, const struct device_info_checks *chk,
2802         const struct opt_out *output)
2803 {
2804         GET_VAL(ret, loc, qprop);
2805         if (!ret->err) {
2806                 const char * const *qpstr = (output->mode == CLINFO_HUMAN ?
2807                         queue_prop_str : queue_prop_raw_str);
2808
2809                 if (output->mode != CLINFO_HUMAN) {
2810                         device_info_bitfield(ret, loc, chk, output, ret->value.qprop,
2811                                 queue_prop_count, qpstr, "queue_prop");
2812                 } else { /* output->mode == CLINFO_HUMAN */
2813                         for (cl_uint i = 0; i < queue_prop_count; ++i) {
2814                                 cl_command_queue_properties cur = (cl_command_queue_properties)(1) << i;
2815                                 cl_bool present =!!(ret->value.qprop & cur);
2816                                 strbuf_append(loc->pname, &ret->str, "\n%s" I2_STR "%s",
2817                                         line_pfx, qpstr[i], bool_str[present]);
2818                         }
2819                         /* TODO FIXME extra bits? */
2820                         if (loc->param.dev == CL_DEVICE_QUEUE_PROPERTIES && dev_has_intel_local_thread(chk))
2821                                 strbuf_append(loc->pname, &ret->str, "\n%s" I2_STR "%s",
2822                                         line_pfx, "Local thread execution (Intel)", bool_str[CL_TRUE]);
2823                 }
2824         }
2825 }
2826
2827 void
2828 device_info_command_buffer_caps(struct device_info_ret *ret,
2829         const struct info_loc *loc, const struct device_info_checks *chk,
2830         const struct opt_out *output)
2831 {
2832         GET_VAL(ret, loc, cmdbufcap);
2833         if (!ret->err) {
2834                 device_info_bitfield(ret, loc, chk, output, ret->value.cmdbufcap,
2835                         command_buffer_count,
2836                         (output->mode == CLINFO_RAW ? command_buffer_raw_str : command_buffer_str),
2837                         "capabilities");
2838         }
2839 }
2840
2841 void
2842 device_info_mutable_dispatch_caps(struct device_info_ret *ret,
2843         const struct info_loc *loc, const struct device_info_checks *chk,
2844         const struct opt_out *output)
2845 {
2846         GET_VAL(ret, loc, cmdbufcap);
2847         if (!ret->err) {
2848                 device_info_bitfield(ret, loc, chk, output, ret->value.cmdbufcap,
2849                         mutable_dispatch_count,
2850                         (output->mode == CLINFO_RAW ? mutable_dispatch_raw_str : mutable_dispatch_str),
2851                         "capabilities");
2852         }
2853 }
2854
2855 void
2856 device_info_intel_usm_cap(struct device_info_ret *ret,
2857         const struct info_loc *loc, const struct device_info_checks *chk,
2858         const struct opt_out *output)
2859 {
2860         GET_VAL(ret, loc, svmcap);
2861         if (!ret->err) {
2862                 device_info_bitfield(ret, loc, chk, output, ret->value.svmcap,
2863                         intel_usm_cap_count,
2864                         (output->mode == CLINFO_RAW ? intel_usm_cap_raw_str : intel_usm_cap_str),
2865                         "capabilities");
2866         }
2867 }
2868
2869 /* Device queue family properties */
2870 void
2871 strbuf_intel_queue_family(const char *what, struct _strbuf *str, const cl_queue_family_properties_intel *fams, size_t num_fams,
2872         const struct opt_out *output)
2873 {
2874         realloc_strbuf(str, num_fams*(CL_QUEUE_FAMILY_MAX_NAME_SIZE_INTEL + 512), "queue families");
2875         if (output->json) {
2876                 strbuf_append_str(what, str, "{");
2877         }
2878         for (size_t i = 0; i < num_fams; ++i) {
2879                 const cl_queue_family_properties_intel  *fam = fams + i;
2880                 set_separator(output->mode == CLINFO_HUMAN ? full_padding : output->json ? comma_str : spc_str);
2881                 if (i > 0) strbuf_append_str(what, str, sep);
2882                 if (output->json || output->mode == CLINFO_HUMAN) {
2883                         strbuf_append(what, str,
2884                                 output->json ?
2885                                 "\"%s\" : { \"count\" : %u" :
2886                                 "%-65s(%u)",
2887                                 fam->name, fam->count);
2888                 } else {
2889                         strbuf_append(what, str, "%s:%u:", fam->name, fam->count);
2890                 }
2891
2892                 if (output->json)
2893                         strbuf_append(what, str, ", \"proprerties\" : ");
2894                 else if (output->mode == CLINFO_HUMAN)
2895                         strbuf_append(what, str, "\n%115s", "Queue properties" INDENT);
2896                 strbuf_bitfield(what, str, fam->properties, "properties",
2897                         output->mode == CLINFO_RAW ? queue_prop_raw_str : queue_prop_str,
2898                         queue_prop_count, output);
2899
2900                 if (output->json)
2901                         strbuf_append(what, str, ", \"capabilities\" : ");
2902                 else if (output->mode == CLINFO_HUMAN)
2903                         strbuf_append(what, str, "\n%115s", "Capabilities" INDENT);
2904                 else strbuf_append(what, str, ":");
2905                 strbuf_bitfield(what, str, fam->properties, "capabilities",
2906                         output->mode == CLINFO_RAW ? intel_queue_cap_raw_str : intel_queue_cap_str,
2907                         intel_queue_cap_count, output);
2908                 if (output->json)
2909                         strbuf_append(what, str, "}");
2910         }
2911         if (output->json)
2912                 strbuf_append_str(what, str, " }");
2913 }
2914
2915 void
2916 device_info_qfamily_prop(struct device_info_ret *ret,
2917         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
2918         const struct opt_out *output)
2919 {
2920         cl_queue_family_properties_intel *val = NULL;
2921         size_t szval = 0, numval = 0;
2922         GET_VAL_ARRAY(ret, loc);
2923         if (!ret->err) {
2924                 strbuf_intel_queue_family(loc->pname, &ret->str, val, numval, output);
2925                 // TODO: ret->value.??? = val;
2926         }
2927         free(val);
2928 }
2929
2930
2931 /* Execution capabilities */
2932 void
2933 device_info_execap(struct device_info_ret *ret,
2934         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
2935         const struct opt_out *output)
2936 {
2937         GET_VAL(ret, loc, execap);
2938         if (!ret->err) {
2939                 const char * const *qpstr = (output->mode == CLINFO_HUMAN ?
2940                         execap_str : execap_raw_str);
2941
2942                 if (output->mode != CLINFO_HUMAN) {
2943                         device_info_bitfield(ret, loc, chk, output, ret->value.execap,
2944                                 execap_count, qpstr, "type");
2945                 } else { /* output->mode == CLINFO_HUMAN */
2946                         for (cl_uint i = 0; i < execap_count; ++i) {
2947                                 cl_device_exec_capabilities cur = (cl_device_exec_capabilities)(1) << i;
2948                                 cl_bool present =!!(ret->value.execap & cur);
2949                                 strbuf_append(loc->pname, &ret->str, "\n%s" I2_STR "%s",
2950                                         line_pfx, qpstr[i], bool_str[present]);
2951                         }
2952                 }
2953         }
2954 }
2955
2956 /* Arch bits and endianness (HUMAN) */
2957 void
2958 device_info_arch(struct device_info_ret *ret,
2959         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
2960         const struct opt_out *output)
2961 {
2962         struct info_loc loc2 = *loc;
2963         DEV_FETCH(cl_uint, bits);
2964         RESET_LOC_PARAM(loc2, dev, CL_DEVICE_ENDIAN_LITTLE);
2965         if (!ret->err) {
2966                 DEV_FETCH_LOC(cl_bool, val, &loc2);
2967                 if (!ret->err) {
2968                         strbuf_append(loc->pname, &ret->str, "%" PRIu32 ", %s", bits, endian_str[val]);
2969                 }
2970         }
2971 }
2972
2973 /* SVM capabilities */
2974 void
2975 device_info_svm_cap(struct device_info_ret *ret,
2976         const struct info_loc *loc, const struct device_info_checks *chk,
2977         const struct opt_out *output)
2978 {
2979         const cl_bool is_20 = dev_is_20(chk);
2980         const cl_bool checking_core = (loc->param.dev == CL_DEVICE_SVM_CAPABILITIES);
2981         const cl_bool has_amd_svm = (checking_core && dev_has_amd_svm(chk));
2982         GET_VAL(ret, loc, svmcap);
2983
2984         if (!ret->err) {
2985                 const char * const *scstr = (output->mode == CLINFO_HUMAN ?
2986                         svm_cap_str : svm_cap_raw_str);
2987
2988                 if (output->mode != CLINFO_HUMAN) {
2989                         device_info_bitfield(ret, loc, chk, output, ret->value.svmcap,
2990                                 svm_cap_count, scstr, "capabilities");
2991                 } else { /* output->mode == CLINFO_HUMAN */
2992                         if (checking_core) {
2993                                 /* show 'why' it's being shown */
2994                                 strbuf_append(loc->pname, &ret->str, "(%s%s%s)",
2995                                         (is_20 ? core : empty_str),
2996                                         (is_20 && has_amd_svm ? comma_str : empty_str),
2997                                         chk->has_amd_svm);
2998                         }
2999                         for (cl_uint i = 0; i < svm_cap_count; ++i) {
3000                                 cl_device_svm_capabilities cur = (cl_device_svm_capabilities)(1) << i;
3001                                 cl_bool present = !!(ret->value.svmcap & cur);
3002                                 strbuf_append(loc->pname, &ret->str, "\n%s" I2_STR "%s",
3003                                         line_pfx, scstr[i], bool_str[present]);
3004                         }
3005                 }
3006         }
3007 }
3008
3009 /* Device terminate capability */
3010 void
3011 device_info_terminate_capability(struct device_info_ret *ret,
3012         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
3013         const struct opt_out *output)
3014 {
3015         GET_VAL(ret, loc, termcap);
3016
3017         if (!ret->err) {
3018                 device_info_bitfield(ret, loc, chk, output, ret->value.termcap,
3019                         terminate_capability_count, (output->mode == CLINFO_HUMAN ?
3020                                 terminate_capability_str : terminate_capability_raw_str),
3021                         "terminate");
3022         }
3023 }
3024
3025 /* Device terminate capability */
3026 void
3027 device_info_terminate_arm(struct device_info_ret *ret,
3028         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
3029         const struct opt_out *output)
3030 {
3031         GET_VAL(ret, loc, termcap);
3032
3033         if (!ret->err) {
3034                 device_info_bitfield(ret, loc, chk, output, ret->value.termcap,
3035                         terminate_capability_arm_count, (output->mode == CLINFO_HUMAN ?
3036                                 terminate_capability_arm_str : terminate_capability_arm_raw_str),
3037                         "terminate");
3038         }
3039 }
3040
3041
3042 /* ARM scheduling controls */
3043 void
3044 device_info_arm_scheduling_controls(struct device_info_ret *ret,
3045         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
3046         const struct opt_out *output)
3047 {
3048         GET_VAL(ret, loc, sched_controls);
3049
3050         if (!ret->err) {
3051                 device_info_bitfield(ret, loc, chk, output, ret->value.sched_controls,
3052                         arm_scheduling_controls_count, (output->mode == CLINFO_HUMAN ?
3053                                 arm_scheduling_controls_str : arm_scheduling_controls_raw_str),
3054                         "scheduling controls");
3055         }
3056 }
3057
3058 void
3059 device_info_p2p_dev_list(struct device_info_ret *ret,
3060         const struct info_loc *loc, const struct device_info_checks *chk,
3061         const struct opt_out* UNUSED(output))
3062 {
3063         // Contrary to most array values in OpenCL, the AMD platform does not support querying
3064         // CL_DEVICE_P2P_DEVICES_AMD with a NULL ptr to get the number of results.
3065         // The user is assumed to have queried for the CL_DEVICE_NUM_P2P_DEVICES_AMD first,
3066         // and to have allocated the return array beforehand.
3067         cl_device_id *val = NULL;
3068         size_t numval = chk->p2p_num_devs, szval = numval*sizeof(*val);
3069         _GET_VAL_VALUES(ret, loc);
3070         if (!ret->err) {
3071                 size_t cursor = 0;
3072                 strbuf_append_str_len(loc->pname, &ret->str, "[ ", 2);
3073                 set_common_separator(output);
3074                 for (cursor = 0; cursor < numval; ++cursor) {
3075                         strbuf_append(loc->pname, &ret->str, "%s%p",
3076                                 (cursor > 0 ? sep : ""), (void*)val[cursor]);
3077                 }
3078                 strbuf_append_str_len(loc->pname, &ret->str, " ]", 2);
3079                 // TODO: ret->value.??? = val;
3080         }
3081         free(val);
3082 }
3083
3084 void
3085 device_info_interop_list(struct device_info_ret *ret,
3086         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
3087         const struct opt_out *output)
3088 {
3089         cl_uint *val = NULL;
3090         size_t szval = 0, numval = 0;
3091         GET_VAL_ARRAY(ret, loc);
3092         if (!ret->err) {
3093                 size_t cursor = 0;
3094                 const cl_interop_name *interop_name_end = cl_interop_names + num_known_interops;
3095                 cl_uint human_raw = output->mode - CLINFO_HUMAN;
3096                 const char *groupsep = (output->mode == CLINFO_HUMAN ? comma_str : vbar_str);
3097                 cl_bool first = CL_TRUE;
3098                 szval = 0;
3099                 for (cursor = 0; cursor < numval; ++cursor) {
3100                         cl_uint current = val[cursor];
3101                         if (!current && cursor < numval - 1) {
3102                                 /* A null value is used as group terminator, but we only print it
3103                                  * if it's not the final one
3104                                  */
3105                                 strbuf_append_str(loc->pname, &ret->str, groupsep);
3106                                 first = CL_TRUE;
3107                         }
3108                         if (current) {
3109                                 cl_bool found = CL_FALSE;
3110                                 const cl_interop_name *n = cl_interop_names;
3111
3112                                 if (!first) {
3113                                         strbuf_append_str(loc->pname, &ret->str, " ");
3114                                 }
3115
3116                                 while (n < interop_name_end) {
3117                                         if (current >= n->from && current <= n->to) {
3118                                                 found = CL_TRUE;
3119                                                 break;
3120                                         }
3121                                         ++n;
3122                                 }
3123                                 if (found) {
3124                                         cl_uint i = current - n->from;
3125                                         strbuf_append(loc->pname, &ret->str, "%s", n->value[i][human_raw]);
3126                                 } else {
3127                                         strbuf_append(loc->pname, &ret->str, "%#" PRIx32, val[cursor]);
3128                                 }
3129                                 first = CL_FALSE;
3130                         }
3131                 }
3132                 // TODO: ret->value.??? = val;
3133         }
3134         // TODO JSONify
3135         ret->needs_escaping = CL_TRUE;
3136         free(val);
3137 }
3138
3139 void device_info_uuid(struct device_info_ret *ret,
3140         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
3141         const struct opt_out *output)
3142 {
3143         cl_uchar uuid[CL_UUID_SIZE_KHR];
3144         _GET_VAL(ret, loc, uuid);
3145         if (!ret->err) {
3146                 strbuf_append(loc->pname, &ret->str,
3147                         "%02x%02x%02x%02x-"
3148                         "%02x%02x-"
3149                         "%02x%02x-"
3150                         "%02x%02x-"
3151                         "%02x%02x%02x%02x%02x%02x",
3152                         uuid[0],  uuid[1],  uuid[2],  uuid[3],  uuid[4],
3153                         uuid[5],  uuid[6],
3154                         uuid[7],  uuid[8],
3155                         uuid[9],  uuid[10],
3156                         uuid[11], uuid[12], uuid[13], uuid[14], uuid[15]);
3157         }
3158         ret->needs_escaping = CL_TRUE;
3159 }
3160
3161 void device_info_luid(struct device_info_ret *ret,
3162         const struct info_loc *loc, const struct device_info_checks* UNUSED(chk),
3163         const struct opt_out *output)
3164 {
3165         cl_uchar uuid[CL_LUID_SIZE_KHR];
3166         _GET_VAL(ret, loc, uuid);
3167         if (!ret->err) {
3168                 /* TODO not sure this is the correct representation for LUIDs? */
3169                 strbuf_append(loc->pname, &ret->str, "%02x%02x-%02x%02x%02x%02x%02x%02x",
3170                         uuid[0], uuid[1],
3171                         uuid[2], uuid[3], uuid[4], uuid[5], uuid[6], uuid[7]);
3172         }
3173         ret->needs_escaping = CL_TRUE;
3174 }
3175
3176
3177 /*
3178  * Device info traits
3179  */
3180
3181 /* A CL_FALSE param means "just print pname" */
3182
3183 struct device_info_traits {
3184         enum output_modes output_mode;
3185         cl_device_info param; // CL_DEVICE_*
3186         const char *sname; // "CL_DEVICE_*"
3187         const char *pname; // "Device *"
3188         const char *sfx; // suffix for the output in non-raw mode
3189         /* pointer to function that retrieves the parameter */
3190         void (*show_func)(struct device_info_ret *,
3191                 const struct info_loc *, const struct device_info_checks *,
3192                 const struct opt_out *);
3193         /* pointer to function that checks if the parameter should be retrieved */
3194         cl_bool (*check_func)(const struct device_info_checks *);
3195 };
3196
3197 #define DINFO_SFX(symbol, name, sfx, typ) symbol, #symbol, name, sfx, device_info_##typ
3198 #define DINFO(symbol, name, typ) symbol, #symbol, name, NULL, device_info_##typ
3199
3200 struct device_info_traits dinfo_traits[] = {
3201         { CLINFO_BOTH, DINFO(CL_DEVICE_NAME, "Device Name", str), NULL },
3202         { CLINFO_BOTH, DINFO(CL_DEVICE_VENDOR, "Device Vendor", str), NULL },
3203         { CLINFO_BOTH, DINFO(CL_DEVICE_VENDOR_ID, "Device Vendor ID", hex), NULL },
3204         { CLINFO_BOTH, DINFO(CL_DEVICE_VERSION, "Device Version", str), NULL },
3205
3206         /* This has to be made before calling NUMERIC_VERSION , since to know if it's supported
3207          * we need to know about the extensions */
3208         { CLINFO_BOTH, DINFO(CL_DEVICE_EXTENSIONS, "Device Extensions", str), NULL },
3209         { CLINFO_BOTH, DINFO(CL_DEVICE_EXTENSIONS_WITH_VERSION, "Device Extensions with Version", ext_version), dev_has_ext_ver },
3210
3211         { CLINFO_BOTH, DINFO(CL_DEVICE_UUID_KHR, "Device UUID", uuid), dev_has_device_uuid },
3212         { CLINFO_BOTH, DINFO(CL_DRIVER_UUID_KHR, "Driver UUID", uuid), dev_has_device_uuid },
3213         { CLINFO_BOTH, DINFO(CL_DEVICE_LUID_VALID_KHR, "Valid Device LUID", bool), dev_has_device_uuid },
3214         { CLINFO_BOTH, DINFO(CL_DEVICE_LUID_KHR, "Device LUID", luid), dev_has_device_uuid },
3215         { CLINFO_BOTH, DINFO(CL_DEVICE_NODE_MASK_KHR, "Device Node Mask", hex), dev_has_device_uuid },
3216
3217         { CLINFO_BOTH, DINFO(CL_DEVICE_NUMERIC_VERSION, "Device Numeric Version", version), dev_has_ext_ver },
3218         { CLINFO_BOTH, DINFO(CL_DRIVER_VERSION, "Driver Version", str), NULL },
3219         { CLINFO_BOTH, DINFO(CL_DEVICE_OPENCL_C_VERSION, "Device OpenCL C Version", str), dev_is_11 },
3220         { CLINFO_BOTH, DINFO(CL_DEVICE_OPENCL_C_NUMERIC_VERSION_KHR, "Device OpenCL C Numeric Version", version), dev_has_extended_versioning },
3221         { CLINFO_BOTH, DINFO(CL_DEVICE_OPENCL_C_ALL_VERSIONS, "Device OpenCL C all versions", ext_version), dev_is_30 },
3222         { CLINFO_BOTH, DINFO(CL_DEVICE_OPENCL_C_FEATURES, "Device OpenCL C features", ext_version), dev_is_30 },
3223
3224         { CLINFO_BOTH, DINFO(CL_DEVICE_CXX_FOR_OPENCL_NUMERIC_VERSION_EXT, "Device C++ for OpenCL Numeric Version", version), dev_has_cxx_for_opencl },
3225
3226         { CLINFO_BOTH, DINFO(CL_DEVICE_LATEST_CONFORMANCE_VERSION_PASSED, "Latest conformance test passed", str), dev_is_30 },
3227         { CLINFO_BOTH, DINFO(CL_DEVICE_TYPE, "Device Type", devtype), NULL },
3228
3229         { CLINFO_BOTH, DINFO(CL_DEVICE_BOARD_NAME_AMD, "Device Board Name (AMD)", str), dev_has_amd },
3230         { CLINFO_BOTH, DINFO(CL_DEVICE_PCIE_ID_AMD, "Device PCI-e ID (AMD)", hex), dev_is_gpu_amd },
3231         { CLINFO_BOTH, DINFO(CL_DEVICE_TOPOLOGY_AMD, "Device Topology (AMD)", devtopo_amd), dev_has_amd },
3232
3233         /* Device Topology (NV) is multipart, so different for HUMAN and RAW */
3234         { CLINFO_HUMAN, DINFO(CL_DEVICE_PCI_BUS_ID_NV, "Device Topology (NV)", devtopo_nv), dev_has_nv },
3235         { CLINFO_RAW, DINFO(CL_DEVICE_PCI_BUS_ID_NV, "Device PCI bus (NV)", int), dev_has_nv },
3236         { CLINFO_RAW, DINFO(CL_DEVICE_PCI_SLOT_ID_NV, "Device PCI slot (NV)", int), dev_has_nv },
3237         { CLINFO_RAW, DINFO(CL_DEVICE_PCI_DOMAIN_ID_NV, "Device PCI domain (NV)", int), dev_has_nv },
3238
3239         /* Device Topology / PCI bus info (KHR) */
3240         { CLINFO_BOTH, DINFO(CL_DEVICE_PCI_BUS_INFO_KHR, "Device PCI bus info (KHR)", devtopo_khr), dev_has_pci_bus_info },
3241
3242         { CLINFO_BOTH, DINFO(CL_DEVICE_PROFILE, "Device Profile", str), NULL },
3243         { CLINFO_BOTH, DINFO(CL_DEVICE_AVAILABLE, "Device Available", bool), NULL },
3244         { CLINFO_BOTH, DINFO(CL_DEVICE_COMPILER_AVAILABLE, "Compiler Available", bool), NULL },
3245         { CLINFO_BOTH, DINFO(CL_DEVICE_LINKER_AVAILABLE, "Linker Available", bool), dev_is_12 },
3246
3247         { CLINFO_BOTH, DINFO(CL_DEVICE_MAX_COMPUTE_UNITS, "Max compute units", int), NULL },
3248         { CLINFO_HUMAN, DINFO(CL_DEVICE_COMPUTE_UNITS_BITFIELD_ARM, "Available core IDs (ARM)", core_ids), dev_has_arm_core_id_v2 },
3249         { CLINFO_RAW, DINFO(CL_DEVICE_COMPUTE_UNITS_BITFIELD_ARM, "Available core IDs (ARM)", long), dev_has_arm_core_id_v2 },
3250         { CLINFO_HUMAN, DINFO(CL_DEVICE_JOB_SLOTS_ARM, "Available job slots (ARM)", job_slots), dev_has_arm_job_slots },
3251         { CLINFO_RAW, DINFO(CL_DEVICE_JOB_SLOTS_ARM, "Available job slots (ARM)", int), dev_has_arm_job_slots },
3252         { CLINFO_BOTH, DINFO(CL_DEVICE_SIMD_PER_COMPUTE_UNIT_AMD, "SIMD per compute unit (AMD)", int), dev_is_gpu_amd },
3253         { CLINFO_BOTH, DINFO(CL_DEVICE_SIMD_WIDTH_AMD, "SIMD width (AMD)", int), dev_is_gpu_amd },
3254         { CLINFO_BOTH, DINFO(CL_DEVICE_SIMD_INSTRUCTION_WIDTH_AMD, "SIMD instruction width (AMD)", int), dev_is_gpu_amd },
3255         { CLINFO_BOTH, DINFO_SFX(CL_DEVICE_MAX_CLOCK_FREQUENCY, "Max clock frequency", "MHz", int), NULL },
3256
3257         /* Device Compute Capability (NV) is multipart, so different for HUMAN and RAW */
3258         { CLINFO_HUMAN, DINFO(CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV, "Compute Capability (NV)", cc_nv), dev_has_nv },
3259         { CLINFO_RAW, DINFO(CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV, INDENT "Compute Capability Major (NV)", int), dev_has_nv },
3260         { CLINFO_RAW, DINFO(CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV, INDENT "Compute Capability Minor (NV)", int), dev_has_nv },
3261
3262         /* GFXIP (AMD) is multipart, so different for HUMAN and RAW */
3263         /* TODO: find a better human-friendly name than GFXIP; v3 of the cl_amd_device_attribute_query
3264          * extension specification calls it “core engine GFXIP”, which honestly is not better than
3265          * our name choice. */
3266         { CLINFO_HUMAN, DINFO(CL_DEVICE_GFXIP_MAJOR_AMD, "Graphics IP (AMD)", gfxip_amd), dev_is_gpu_amd },
3267         { CLINFO_RAW, DINFO(CL_DEVICE_GFXIP_MAJOR_AMD, INDENT "Graphics IP MAJOR (AMD)", int), dev_is_gpu_amd },
3268         { CLINFO_RAW, DINFO(CL_DEVICE_GFXIP_MINOR_AMD, INDENT "Graphics IP MINOR (AMD)", int), dev_is_gpu_amd },
3269
3270         /* Device IP version (Intel) */
3271         { CLINFO_BOTH, DINFO(CL_DEVICE_IP_VERSION_INTEL, "Device IP (Intel)", version), dev_is_gpu_intel },
3272         { CLINFO_BOTH, DINFO(CL_DEVICE_ID_INTEL, "Device ID (Intel)", int), dev_is_gpu_intel },
3273         { CLINFO_BOTH, DINFO(CL_DEVICE_NUM_SLICES_INTEL, "Slices (Intel)", int), dev_is_gpu_intel },
3274         { CLINFO_BOTH, DINFO(CL_DEVICE_NUM_SUB_SLICES_PER_SLICE_INTEL, "Sub-slices per slice (Intel)", int), dev_is_gpu_intel },
3275         { CLINFO_BOTH, DINFO(CL_DEVICE_NUM_EUS_PER_SUB_SLICE_INTEL, "EUs per sub-slice (Intel)", int), dev_is_gpu_intel },
3276         { CLINFO_BOTH, DINFO(CL_DEVICE_NUM_THREADS_PER_EU_INTEL, "Threads per EU (Intel)", int), dev_is_gpu_intel },
3277         { CLINFO_BOTH, DINFO(CL_DEVICE_FEATURE_CAPABILITIES_INTEL, "Feature capabilities (Intel)", intel_features), dev_is_gpu_intel },
3278
3279         { CLINFO_BOTH, DINFO_SFX(CL_DEVICE_CORE_TEMPERATURE_ALTERA, "Core Temperature (Altera)", " C", int), dev_has_altera_dev_temp },
3280
3281         /* Device partition support: summary is only presented in HUMAN case */
3282         { CLINFO_HUMAN, DINFO(CL_DEVICE_PARTITION_MAX_SUB_DEVICES, "Device Partition", partition_header), dev_has_partition },
3283         { CLINFO_BOTH, DINFO(CL_DEVICE_PARTITION_MAX_SUB_DEVICES, INDENT "Max number of sub-devices", int), dev_is_12 },
3284         { CLINFO_BOTH, DINFO(CL_DEVICE_PARTITION_PROPERTIES, INDENT "Supported partition types", partition_types), dev_is_12 },
3285         { CLINFO_BOTH, DINFO(CL_DEVICE_PARTITION_AFFINITY_DOMAIN, INDENT "Supported affinity domains", partition_affinities), dev_is_12 },
3286         { CLINFO_BOTH, DINFO(CL_DEVICE_PARTITION_TYPES_EXT, INDENT "Supported partition types (ext)", partition_types_ext), dev_has_fission },
3287         { CLINFO_BOTH, DINFO(CL_DEVICE_AFFINITY_DOMAINS_EXT, INDENT "Supported affinity domains (ext)", partition_affinities_ext), dev_has_fission },
3288
3289         { CLINFO_BOTH, DINFO(CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS, "Max work item dimensions", int), NULL },
3290         { CLINFO_BOTH, DINFO(CL_DEVICE_MAX_WORK_ITEM_SIZES, "Max work item sizes", szptr_times), NULL },
3291         { CLINFO_BOTH, DINFO(CL_DEVICE_MAX_WORK_GROUP_SIZE, "Max work group size", sz), NULL },
3292
3293         /* cl_amd_device_attribute_query v4 */
3294         { CLINFO_BOTH, DINFO(CL_DEVICE_PREFERRED_WORK_GROUP_SIZE_AMD, "Preferred work group size (AMD)", sz), dev_has_amd_v4 },
3295         { CLINFO_BOTH, DINFO(CL_DEVICE_MAX_WORK_GROUP_SIZE_AMD, "Max work group size (AMD)", sz), dev_has_amd_v4 },
3296
3297         { CLINFO_BOTH, DINFO(CL_DEVICE_PREFERRED_WORK_GROUP_SIZE_MULTIPLE, "Preferred work group size multiple (device)", sz), dev_is_30 },
3298         { CLINFO_BOTH, DINFO(CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE, "Preferred work group size multiple (kernel)", wg), dev_has_compiler_11 },
3299         { CLINFO_BOTH, DINFO(CL_DEVICE_WARP_SIZE_NV, "Warp size (NV)", int), dev_has_nv },
3300         { CLINFO_BOTH, DINFO(CL_DEVICE_WAVEFRONT_WIDTH_AMD, "Wavefront width (AMD)", int), dev_is_gpu_amd },
3301         { CLINFO_BOTH, DINFO(CL_DEVICE_MAX_NUM_SUB_GROUPS, "Max sub-groups per work group", int), dev_is_21 },
3302         { CLINFO_BOTH, DINFO(CL_DEVICE_MAX_NAMED_BARRIER_COUNT_KHR, "Max named sub-group barriers", int), dev_has_subgroup_named_barrier },
3303         { CLINFO_BOTH, DINFO(CL_DEVICE_SUB_GROUP_SIZES_INTEL, "Sub-group sizes (Intel)", szptr_comma), dev_has_intel_required_subgroup_size },
3304
3305         /* Preferred/native vector widths: header is only presented in HUMAN case, that also pairs
3306          * PREFERRED and NATIVE in a single line */
3307 #define DINFO_VECWIDTH(Type, type) \
3308         { CLINFO_HUMAN, DINFO(CL_DEVICE_PREFERRED_VECTOR_WIDTH_##Type, INDENT #type, vecwidth), NULL }, \
3309         { CLINFO_RAW, DINFO(CL_DEVICE_PREFERRED_VECTOR_WIDTH_##Type, INDENT #type, int), NULL }, \
3310         { CLINFO_RAW, DINFO(CL_DEVICE_NATIVE_VECTOR_WIDTH_##Type, INDENT #type, int), dev_is_11 }
3311
3312         { CLINFO_HUMAN, DINFO(CL_FALSE, "Preferred / native vector sizes", str), NULL },
3313         DINFO_VECWIDTH(CHAR, char),
3314         DINFO_VECWIDTH(SHORT, short),
3315         DINFO_VECWIDTH(INT, int),
3316         DINFO_VECWIDTH(LONG, long),
3317         DINFO_VECWIDTH(HALF, half), /* this should be excluded for 1.0 */
3318         DINFO_VECWIDTH(FLOAT, float),
3319         DINFO_VECWIDTH(DOUBLE, double),
3320
3321         /* Floating point configurations */
3322 #define DINFO_FPCONF(Type, type, cond) \
3323         { CLINFO_HUMAN, DINFO(CL_DEVICE_##Type##_FP_CONFIG, #type "-precision Floating-point support", fpconf), NULL }, \
3324         { CLINFO_RAW, DINFO(CL_DEVICE_##Type##_FP_CONFIG, #type "-precision Floating-point support", fpconf), cond }
3325
3326         DINFO_FPCONF(HALF, Half, dev_has_half),
3327         DINFO_FPCONF(SINGLE, Single, NULL),
3328         DINFO_FPCONF(DOUBLE, Double, dev_has_double),
3329
3330         /* Address bits and endianness are written together for HUMAN, separate for RAW */
3331         { CLINFO_HUMAN, DINFO(CL_DEVICE_ADDRESS_BITS, "Address bits", arch), NULL },
3332         { CLINFO_RAW, DINFO(CL_DEVICE_ADDRESS_BITS, "Address bits", int), NULL },
3333         { CLINFO_RAW, DINFO(CL_DEVICE_ENDIAN_LITTLE, "Little Endian", bool), NULL },
3334
3335         /* External memory */
3336         { CLINFO_BOTH, DINFO(CL_DEVICE_EXTERNAL_MEMORY_IMPORT_HANDLE_TYPES_KHR, "External memory handle types", ext_mem), dev_has_external_memory },
3337
3338         /* Semaphores */
3339         { CLINFO_BOTH, DINFO(CL_DEVICE_SEMAPHORE_TYPES_KHR, "Semaphore types", semaphore_types), dev_has_semaphore },
3340         { CLINFO_BOTH, DINFO(CL_DEVICE_SEMAPHORE_IMPORT_HANDLE_TYPES_KHR, "External semaphore import types", ext_semaphore_handles), dev_has_external_semaphore },
3341         { CLINFO_BOTH, DINFO(CL_DEVICE_SEMAPHORE_EXPORT_HANDLE_TYPES_KHR, "External semaphore export types", ext_semaphore_handles), dev_has_external_semaphore },
3342
3343         /* Global memory */
3344         { CLINFO_BOTH, DINFO(CL_DEVICE_GLOBAL_MEM_SIZE, "Global memory size", mem), NULL },
3345         { CLINFO_BOTH, DINFO(CL_DEVICE_GLOBAL_FREE_MEMORY_AMD, "Global free memory (AMD)", free_mem_amd), dev_is_gpu_amd },
3346         { CLINFO_BOTH, DINFO(CL_DEVICE_GLOBAL_MEM_CHANNELS_AMD, "Global memory channels (AMD)", int), dev_is_gpu_amd },
3347         { CLINFO_BOTH, DINFO(CL_DEVICE_GLOBAL_MEM_CHANNEL_BANKS_AMD, "Global memory banks per channel (AMD)", int), dev_is_gpu_amd },
3348         { CLINFO_BOTH, DINFO_SFX(CL_DEVICE_GLOBAL_MEM_CHANNEL_BANK_WIDTH_AMD, "Global memory bank width (AMD)", bytes_str, int), dev_is_gpu_amd },
3349         { CLINFO_BOTH, DINFO(CL_DEVICE_ERROR_CORRECTION_SUPPORT, "Error Correction support", bool), NULL },
3350         { CLINFO_BOTH, DINFO(CL_DEVICE_MAX_MEM_ALLOC_SIZE, "Max memory allocation", mem), NULL },
3351         { CLINFO_BOTH, DINFO(CL_DEVICE_HOST_UNIFIED_MEMORY, "Unified memory for Host and Device", bool), dev_is_11 },
3352         { CLINFO_BOTH, DINFO(CL_DEVICE_INTEGRATED_MEMORY_NV, "Integrated memory (NV)", bool), dev_has_nv },
3353
3354         { CLINFO_BOTH, DINFO(CL_DEVICE_SVM_CAPABILITIES, "Shared Virtual Memory (SVM) capabilities", svm_cap), dev_has_svm },
3355         { CLINFO_BOTH, DINFO(CL_DEVICE_SVM_CAPABILITIES_ARM, "Shared Virtual Memory (SVM) capabilities (ARM)", svm_cap), dev_has_arm_svm },
3356
3357         { CLINFO_HUMAN, DINFO_SFX(CL_FALSE, "Unified Shared Memory (USM)", "(cl_intel_unified_shared_memory)", str), dev_has_intel_usm },
3358         { CLINFO_BOTH, DINFO(CL_DEVICE_HOST_MEM_CAPABILITIES_INTEL, "Host USM capabilities (Intel)", intel_usm_cap), dev_has_intel_usm },
3359         { CLINFO_BOTH, DINFO(CL_DEVICE_DEVICE_MEM_CAPABILITIES_INTEL, "Device USM capabilities (Intel)", intel_usm_cap), dev_has_intel_usm },
3360         { CLINFO_BOTH, DINFO(CL_DEVICE_SINGLE_DEVICE_SHARED_MEM_CAPABILITIES_INTEL, "Single-Device USM caps (Intel)", intel_usm_cap), dev_has_intel_usm },
3361         { CLINFO_BOTH, DINFO(CL_DEVICE_CROSS_DEVICE_SHARED_MEM_CAPABILITIES_INTEL, "Cross-Device USM caps (Intel)", intel_usm_cap), dev_has_intel_usm },
3362         { CLINFO_BOTH, DINFO(CL_DEVICE_SHARED_SYSTEM_MEM_CAPABILITIES_INTEL, "Shared System USM caps (Intel)", intel_usm_cap), dev_has_intel_usm },
3363
3364         /* Alignment */
3365         { CLINFO_BOTH, DINFO_SFX(CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE, "Minimum alignment for any data type", bytes_str, int), NULL },
3366         { CLINFO_HUMAN, DINFO(CL_DEVICE_MEM_BASE_ADDR_ALIGN, "Alignment of base address", bits), NULL },
3367         { CLINFO_RAW, DINFO(CL_DEVICE_MEM_BASE_ADDR_ALIGN, "Alignment of base address", int), NULL },
3368
3369         { CLINFO_BOTH, DINFO_SFX(CL_DEVICE_PAGE_SIZE_QCOM, "Page size (QCOM)", bytes_str, sz), dev_has_qcom_ext_host_ptr },
3370         { CLINFO_BOTH, DINFO_SFX(CL_DEVICE_EXT_MEM_PADDING_IN_BYTES_QCOM, "External memory padding (QCOM)", bytes_str, sz), dev_has_qcom_ext_host_ptr },
3371
3372         /* Atomics alignment, with HUMAN-only header */
3373         { CLINFO_HUMAN, DINFO(CL_FALSE, "Preferred alignment for atomics", str), dev_is_20 },
3374         { CLINFO_BOTH, DINFO_SFX(CL_DEVICE_PREFERRED_PLATFORM_ATOMIC_ALIGNMENT, INDENT "SVM", bytes_str, int), dev_is_20 },
3375         { CLINFO_BOTH, DINFO_SFX(CL_DEVICE_PREFERRED_GLOBAL_ATOMIC_ALIGNMENT, INDENT "Global", bytes_str, int), dev_is_20 },
3376         { CLINFO_BOTH, DINFO_SFX(CL_DEVICE_PREFERRED_LOCAL_ATOMIC_ALIGNMENT, INDENT "Local", bytes_str, int), dev_is_20 },
3377
3378         /* 3.0+ Atomic memory and fence capabilities */
3379         { CLINFO_BOTH, DINFO(CL_DEVICE_ATOMIC_MEMORY_CAPABILITIES, "Atomic memory capabilities", atomic_caps), dev_is_30 },
3380         { CLINFO_BOTH, DINFO(CL_DEVICE_ATOMIC_FENCE_CAPABILITIES, "Atomic fence capabilities", atomic_caps), dev_is_30 },
3381
3382         /* Global variables. TODO some 1.2 devices respond to this too */
3383         { CLINFO_BOTH, DINFO(CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE, "Max size for global variable", mem), dev_is_20 },
3384         { CLINFO_BOTH, DINFO(CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE, "Preferred total size of global vars", mem), dev_is_20 },
3385
3386         /* Global memory cache */
3387         { CLINFO_BOTH, DINFO(CL_DEVICE_GLOBAL_MEM_CACHE_TYPE, "Global Memory cache type", cachetype), NULL },
3388         { CLINFO_BOTH, DINFO(CL_DEVICE_GLOBAL_MEM_CACHE_SIZE, "Global Memory cache size", mem), dev_has_cache },
3389         { CLINFO_BOTH, DINFO_SFX(CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE, "Global Memory cache line size", " bytes", int), dev_has_cache },
3390
3391         /* Image support */
3392         { CLINFO_BOTH, DINFO(CL_DEVICE_IMAGE_SUPPORT, "Image support", bool), NULL },
3393         { CLINFO_BOTH, DINFO(CL_DEVICE_MAX_SAMPLERS, INDENT "Max number of samplers per kernel", int), dev_has_images },
3394         { CLINFO_BOTH, DINFO_SFX(CL_DEVICE_IMAGE_MAX_BUFFER_SIZE, INDENT "Max size for 1D images from buffer", pixels_str, sz), dev_has_images_12 },
3395         { CLINFO_BOTH, DINFO_SFX(CL_DEVICE_IMAGE_MAX_ARRAY_SIZE, INDENT "Max 1D or 2D image array size", images_str, sz), dev_has_images_12 },
3396         { CLINFO_BOTH, DINFO_SFX(CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT, INDENT "Base address alignment for 2D image buffers", bytes_str, sz), dev_has_image2d_buffer },
3397         { CLINFO_BOTH, DINFO_SFX(CL_DEVICE_IMAGE_PITCH_ALIGNMENT, INDENT "Pitch alignment for 2D image buffers", pixels_str, sz), dev_has_image2d_buffer },
3398
3399         /* Image dimensions are split for RAW, combined for HUMAN */
3400         { CLINFO_HUMAN, DINFO_SFX(CL_DEVICE_IMAGE2D_MAX_HEIGHT, INDENT "Max 2D image size",  pixels_str, img_sz_2d), dev_has_images },
3401         { CLINFO_RAW, DINFO(CL_DEVICE_IMAGE2D_MAX_HEIGHT, INDENT "Max 2D image height",  sz), dev_has_images },
3402         { CLINFO_RAW, DINFO(CL_DEVICE_IMAGE2D_MAX_WIDTH, INDENT "Max 2D image width",  sz), dev_has_images },
3403         { CLINFO_HUMAN, DINFO_SFX(CL_DEVICE_PLANAR_YUV_MAX_HEIGHT_INTEL, INDENT "Max planar YUV image size",  pixels_str, img_sz_2d), dev_has_intel_planar_yuv },
3404         { CLINFO_RAW, DINFO(CL_DEVICE_PLANAR_YUV_MAX_HEIGHT_INTEL, INDENT "Max planar YUV image height",  sz), dev_has_intel_planar_yuv },
3405         { CLINFO_RAW, DINFO(CL_DEVICE_PLANAR_YUV_MAX_WIDTH_INTEL, INDENT "Max planar YUV image width",  sz), dev_has_intel_planar_yuv },
3406         { CLINFO_HUMAN, DINFO_SFX(CL_DEVICE_IMAGE3D_MAX_HEIGHT, INDENT "Max 3D image size",  pixels_str, img_sz_3d), dev_has_images },
3407         { CLINFO_RAW, DINFO(CL_DEVICE_IMAGE3D_MAX_HEIGHT, INDENT "Max 3D image height",  sz), dev_has_images },
3408         { CLINFO_RAW, DINFO(CL_DEVICE_IMAGE3D_MAX_WIDTH, INDENT "Max 3D image width",  sz), dev_has_images },
3409         { CLINFO_RAW, DINFO(CL_DEVICE_IMAGE3D_MAX_DEPTH, INDENT "Max 3D image depth",  sz), dev_has_images },
3410
3411         { CLINFO_BOTH, DINFO(CL_DEVICE_MAX_READ_IMAGE_ARGS, INDENT "Max number of read image args", int), dev_has_images },
3412         { CLINFO_BOTH, DINFO(CL_DEVICE_MAX_WRITE_IMAGE_ARGS, INDENT "Max number of write image args", int), dev_has_images },
3413         { CLINFO_BOTH, DINFO(CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS, INDENT "Max number of read/write image args", int), dev_has_images_20 },
3414
3415         /* Pipes */
3416         { CLINFO_BOTH, DINFO(CL_DEVICE_PIPE_SUPPORT, "Pipe support", bool), dev_is_30 },
3417         /* TODO FIXME: the above should be true if dev is [2.0, 3.0[, and the next properties should be nested */
3418         { CLINFO_BOTH, DINFO(CL_DEVICE_MAX_PIPE_ARGS, "Max number of pipe args", int), dev_is_20 },
3419         { CLINFO_BOTH, DINFO(CL_DEVICE_PIPE_MAX_ACTIVE_RESERVATIONS, "Max active pipe reservations", int), dev_is_20 },
3420         { CLINFO_BOTH, DINFO(CL_DEVICE_PIPE_MAX_PACKET_SIZE, "Max pipe packet size", mem_int), dev_is_20 },
3421
3422         /* Local memory */
3423         { CLINFO_BOTH, DINFO(CL_DEVICE_LOCAL_MEM_TYPE, "Local memory type", lmemtype), NULL },
3424         { CLINFO_BOTH, DINFO(CL_DEVICE_LOCAL_MEM_SIZE, "Local memory size", mem), dev_has_lmem },
3425         { CLINFO_BOTH, DINFO(CL_DEVICE_LOCAL_MEM_SIZE_PER_COMPUTE_UNIT_AMD, "Local memory size per CU (AMD)", mem), dev_is_gpu_amd },
3426         { CLINFO_BOTH, DINFO(CL_DEVICE_LOCAL_MEM_BANKS_AMD, "Local memory banks (AMD)", int), dev_is_gpu_amd },
3427         { CLINFO_BOTH, DINFO(CL_DEVICE_REGISTERS_PER_BLOCK_NV, "Registers per block (NV)", int), dev_has_nv },
3428
3429         /* Constant memory */
3430         { CLINFO_BOTH, DINFO(CL_DEVICE_MAX_CONSTANT_ARGS, "Max number of constant args", int), NULL },
3431         { CLINFO_BOTH, DINFO(CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE, "Max constant buffer size", mem), NULL },
3432         { CLINFO_BOTH, DINFO(CL_DEVICE_PREFERRED_CONSTANT_BUFFER_SIZE_AMD, "Preferred constant buffer size (AMD)", mem_sz), dev_has_amd_v4 },
3433
3434         /* Generic address space support */
3435         { CLINFO_BOTH, DINFO(CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT, "Generic address space support", bool), dev_is_30},
3436
3437         { CLINFO_BOTH, DINFO(CL_DEVICE_MAX_PARAMETER_SIZE, "Max size of kernel argument", mem), NULL },
3438         { CLINFO_BOTH, DINFO(CL_DEVICE_MAX_ATOMIC_COUNTERS_EXT, "Max number of atomic counters", sz), dev_has_atomic_counters },
3439
3440         /* Queue properties */
3441         { CLINFO_BOTH, DINFO(CL_DEVICE_QUEUE_PROPERTIES, "Queue properties", qprop), dev_not_20 },
3442         { CLINFO_BOTH, DINFO(CL_DEVICE_QUEUE_ON_HOST_PROPERTIES, "Queue properties (on host)", qprop), dev_is_20 },
3443         { CLINFO_BOTH, DINFO(CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES, "Device enqueue capabilities", device_enqueue_caps), dev_is_30 },
3444         /* TODO FIXME: the above should be true if dev is [2.0, 3.0[, and the next properties should be nested */
3445         { CLINFO_BOTH, DINFO(CL_DEVICE_QUEUE_ON_DEVICE_PROPERTIES, "Queue properties (on device)", qprop), dev_is_20 },
3446         { CLINFO_BOTH, DINFO(CL_DEVICE_QUEUE_ON_DEVICE_PREFERRED_SIZE, INDENT "Preferred size", mem), dev_is_20 },
3447         { CLINFO_BOTH, DINFO(CL_DEVICE_QUEUE_ON_DEVICE_MAX_SIZE, INDENT "Max size", mem), dev_is_20 },
3448         { CLINFO_BOTH, DINFO(CL_DEVICE_MAX_ON_DEVICE_QUEUES, "Max queues on device", int), dev_is_20 },
3449         { CLINFO_BOTH, DINFO(CL_DEVICE_MAX_ON_DEVICE_EVENTS, "Max events on device", int), dev_is_20 },
3450         { CLINFO_BOTH, DINFO(CL_DEVICE_QUEUE_FAMILY_PROPERTIES_INTEL, "Device queue families", qfamily_prop), dev_has_intel_queue_families },
3451
3452         /* Command buffers */
3453         { CLINFO_BOTH, DINFO(CL_DEVICE_COMMAND_BUFFER_CAPABILITIES_KHR, "Command buffer capabilities", command_buffer_caps), dev_has_command_buffer },
3454         { CLINFO_BOTH, DINFO(CL_DEVICE_COMMAND_BUFFER_REQUIRED_QUEUE_PROPERTIES_KHR, INDENT "Required queue properties for command buffer", qprop), dev_has_command_buffer },
3455         { CLINFO_BOTH, DINFO(CL_DEVICE_MUTABLE_DISPATCH_CAPABILITIES_KHR, "Mutable dispatch capabilities", mutable_dispatch_caps), dev_has_mutable_dispatch },
3456
3457         /* Terminate context */
3458         { CLINFO_BOTH, DINFO(CL_DEVICE_TERMINATE_CAPABILITY_KHR_1x, "Terminate capability (1.2 define)", terminate_capability), dev_has_terminate_context },
3459         { CLINFO_BOTH, DINFO(CL_DEVICE_TERMINATE_CAPABILITY_KHR, "Terminate capability (2.x and later)", terminate_capability), dev_has_terminate_context },
3460
3461         { CLINFO_BOTH, DINFO(CL_DEVICE_CONTROLLED_TERMINATION_CAPABILITIES_ARM, "Controlled termination caps. (ARM)", terminate_arm), dev_has_terminate_arm },
3462
3463         /* Interop */
3464         { CLINFO_BOTH, DINFO(CL_DEVICE_PREFERRED_INTEROP_USER_SYNC, "Prefer user sync for interop", bool), dev_is_12 },
3465         { CLINFO_BOTH, DINFO(CL_DEVICE_NUM_SIMULTANEOUS_INTEROPS_INTEL, "Number of simultaneous interops (Intel)", int), dev_has_simultaneous_sharing },
3466         { CLINFO_BOTH, DINFO(CL_DEVICE_SIMULTANEOUS_INTEROPS_INTEL, "Simultaneous interops", interop_list), dev_has_simultaneous_sharing },
3467
3468         /* P2P buffer copy */
3469         { CLINFO_BOTH, DINFO(CL_DEVICE_NUM_P2P_DEVICES_AMD, "Number of P2P devices (AMD)", int), dev_has_p2p },
3470         { CLINFO_BOTH, DINFO(CL_DEVICE_P2P_DEVICES_AMD, "P2P devices (AMD)", p2p_dev_list), dev_has_p2p_devs },
3471
3472         /* Profiling resolution */
3473         { CLINFO_BOTH, DINFO_SFX(CL_DEVICE_PROFILING_TIMER_RESOLUTION, "Profiling timer resolution", "ns", sz), NULL },
3474         { CLINFO_HUMAN, DINFO(CL_DEVICE_PROFILING_TIMER_OFFSET_AMD, "Profiling timer offset since Epoch (AMD)", time_offset), dev_has_amd },
3475         { CLINFO_RAW, DINFO(CL_DEVICE_PROFILING_TIMER_OFFSET_AMD, "Profiling timer offset since Epoch (AMD)", long), dev_has_amd },
3476
3477         /* Kernel execution capabilities */
3478         { CLINFO_BOTH, DINFO(CL_DEVICE_EXECUTION_CAPABILITIES, "Execution capabilities", execap), NULL },
3479         { CLINFO_BOTH, DINFO(CL_DEVICE_NON_UNIFORM_WORK_GROUP_SUPPORT, INDENT "Non-uniform work-groups",  bool), dev_is_30 },
3480         { CLINFO_BOTH, DINFO(CL_DEVICE_WORK_GROUP_COLLECTIVE_FUNCTIONS_SUPPORT, INDENT "Work-group collective functions",  bool), dev_is_30 },
3481         { CLINFO_BOTH, DINFO(CL_DEVICE_SUB_GROUP_INDEPENDENT_FORWARD_PROGRESS, INDENT "Sub-group independent forward progress", bool), dev_is_21 },
3482         { CLINFO_BOTH, DINFO(CL_DEVICE_THREAD_TRACE_SUPPORTED_AMD, INDENT "Thread trace supported (AMD)", bool), dev_is_gpu_amd },
3483         { CLINFO_BOTH, DINFO(CL_DEVICE_KERNEL_EXEC_TIMEOUT_NV, INDENT "Kernel execution timeout (NV)", bool), dev_has_nv },
3484         { CLINFO_BOTH, DINFO(CL_DEVICE_GPU_OVERLAP_NV, INDENT "Concurrent copy and kernel execution (NV)", bool), dev_has_nv },
3485         { CLINFO_BOTH, DINFO(CL_DEVICE_ATTRIBUTE_ASYNC_ENGINE_COUNT_NV, INDENT INDENT "Number of async copy engines", int), dev_has_nv },
3486         { CLINFO_BOTH, DINFO(CL_DEVICE_AVAILABLE_ASYNC_QUEUES_AMD, INDENT "Number of async queues (AMD)", int), dev_has_amd_v4 },
3487         /* TODO FIXME undocumented, experimental */
3488         { CLINFO_BOTH, DINFO(CL_DEVICE_MAX_REAL_TIME_COMPUTE_QUEUES_AMD, INDENT "Max real-time compute queues (AMD)", int), dev_has_amd_v4 },
3489         { CLINFO_BOTH, DINFO(CL_DEVICE_MAX_REAL_TIME_COMPUTE_UNITS_AMD, INDENT "Max real-time compute units (AMD)", int), dev_has_amd_v4 },
3490
3491         { CLINFO_BOTH, DINFO(CL_DEVICE_SCHEDULING_CONTROLS_CAPABILITIES_ARM, INDENT "Scheduling controls (ARM)", arm_scheduling_controls), dev_has_arm_scheduling_controls },
3492         { CLINFO_BOTH, DINFO(CL_DEVICE_SUPPORTED_REGISTER_ALLOCATIONS_ARM, INDENT "Supported reg allocs (ARM)", intptr), dev_has_arm_register_alloc },
3493         { CLINFO_BOTH, DINFO(CL_DEVICE_MAX_WARP_COUNT_ARM, INDENT "Max warps/CU (ARM)", int), dev_has_arm_warp_count_support },
3494
3495         /* TODO: this should tell if it's being done due to the device being 2.1 or due to it having the extension */
3496         { CLINFO_BOTH, DINFO(CL_DEVICE_IL_VERSION, INDENT "IL version", str), dev_has_il },
3497         { CLINFO_BOTH, DINFO(CL_DEVICE_ILS_WITH_VERSION, INDENT "ILs with version", ext_version), dev_has_ext_ver },
3498         { CLINFO_BOTH, DINFO(CL_DEVICE_SPIR_VERSIONS, INDENT "SPIR versions", str), dev_has_spir },
3499         { CLINFO_BOTH, DINFO(CL_DEVICE_PRINTF_BUFFER_SIZE, "printf() buffer size", mem_sz), dev_is_12 },
3500         { CLINFO_BOTH, DINFO(CL_DEVICE_BUILT_IN_KERNELS, "Built-in kernels", str), dev_is_12 },
3501         { CLINFO_BOTH, DINFO(CL_DEVICE_BUILT_IN_KERNELS_WITH_VERSION, "Built-in kernels with version", ext_version), dev_has_ext_ver },
3502         { CLINFO_BOTH, DINFO(CL_DEVICE_ME_VERSION_INTEL, "Motion Estimation accelerator version (Intel)", int), dev_has_intel_AME },
3503         { CLINFO_BOTH, DINFO(CL_DEVICE_AVC_ME_VERSION_INTEL, INDENT "Device-side AVC Motion Estimation version", int), dev_has_intel_AVC_ME },
3504         { CLINFO_BOTH, DINFO(CL_DEVICE_AVC_ME_SUPPORTS_TEXTURE_SAMPLER_USE_INTEL, INDENT INDENT "Supports texture sampler use", bool), dev_has_intel_AVC_ME },
3505         { CLINFO_BOTH, DINFO(CL_DEVICE_AVC_ME_SUPPORTS_PREEMPTION_INTEL, INDENT INDENT "Supports preemption", bool), dev_has_intel_AVC_ME },
3506 };
3507
3508 /* Process all the device info in the traits, except if param_whitelist is not NULL,
3509  * in which case only those in the whitelist will be processed.
3510  * If present, the whitelist should be sorted in the order of appearance of the parameters
3511  * in the traits table, and terminated by the value CL_FALSE
3512  */
3513
3514 void
3515 printDeviceInfo(cl_device_id dev, const struct platform_list *plist, cl_uint p,
3516         const cl_device_info *param_whitelist, /* list of device info to process, or NULL */
3517         const struct opt_out *output)
3518 {
3519         char *extensions = NULL;
3520         size_t ext_len = 0;
3521         char *versioned_extensions = NULL;
3522
3523         /* pointers to the traits for CL_DEVICE_EXTENSIONS and CL_DEVICE_EXTENSIONS_WITH_VERSION */
3524         const struct device_info_traits *extensions_traits = NULL;
3525         const struct device_info_traits *versioned_extensions_traits = NULL;
3526
3527         struct device_info_checks chk;
3528         struct device_info_ret ret;
3529         struct info_loc loc;
3530
3531         cl_uint n = 0; /* number of device properties shown, for JSON */
3532
3533         memset(&chk, 0, sizeof(chk));
3534         chk.pinfo_checks = plist->platform_checks + p;
3535         chk.dev_version = 10;
3536
3537         INIT_RET(ret, "device");
3538
3539         reset_loc(&loc, __func__);
3540         loc.plat = plist->platform[p];
3541         loc.dev = dev;
3542
3543         for (loc.line = 0; loc.line < ARRAY_SIZE(dinfo_traits); ++loc.line) {
3544
3545                 const struct device_info_traits *traits = dinfo_traits + loc.line;
3546                 cl_bool requested;
3547
3548                 /* checked is true if there was no condition to check for, or if the
3549                  * condition was satisfied
3550                  */
3551                 int checked = !(traits->check_func && !traits->check_func(&chk));
3552
3553                 loc.sname = traits->sname;
3554                 loc.pname = (output->mode == CLINFO_HUMAN ?
3555                         traits->pname : traits->sname);
3556                 loc.param.dev = traits->param;
3557
3558                 /* Whitelist check: finish if done traversing the list,
3559                  * skip current param if it's not the right one
3560                  */
3561                 if ((output->cond == COND_PROP_CHECK || output->brief) && param_whitelist) {
3562                         if (*param_whitelist == CL_FALSE)
3563                                 break;
3564                         if (traits->param != *param_whitelist)
3565                                 continue;
3566                         ++param_whitelist;
3567                 }
3568
3569                 /* skip if it's not for this output mode */
3570                 if (!(output->mode & traits->output_mode))
3571                         continue;
3572
3573                 if (output->cond == COND_PROP_CHECK && !checked)
3574                         continue;
3575
3576                 cur_sfx = (output->mode == CLINFO_HUMAN && traits->sfx) ? traits->sfx : empty_str;
3577
3578                 reset_strbuf(&ret.str);
3579                 reset_strbuf(&ret.err_str);
3580                 ret.needs_escaping = CL_FALSE;
3581
3582                 /* Handle headers */
3583                 if (traits->param == CL_FALSE) {
3584                         ret.err = CL_SUCCESS;
3585                         show_strbuf(&ret.str, loc.pname, 0, ret.err);
3586                         continue;
3587                 }
3588
3589                 traits->show_func(&ret, &loc, &chk, output);
3590
3591                 /* Do not print this property if the user requested one and this does not match */
3592                 requested = !(output->prop && strstr(loc.sname, output->prop) == NULL);
3593                 if (traits->param == CL_DEVICE_EXTENSIONS) {
3594                         /* make a backup of the extensions string, regardless of
3595                          * errors and requested, because we need the information
3596                          * to fetch further information */
3597                         const char *msg = RET_BUF(ret)->buf;
3598                         ext_len = strlen(msg);
3599                         extensions_traits = traits;
3600                         /* pad with spaces: this will make it easier to check for extension presence
3601                          * without erroneously matching substrings by simply padding the extension name
3602                          * with spaces.
3603                          */
3604                         ALLOC(extensions, ext_len+3, "extensions");
3605                         memcpy(extensions + 1, msg, ext_len);
3606                         extensions[0] = ' ';
3607                         extensions[ext_len+1] = ' ';
3608                         extensions[ext_len+2] = '\0';
3609                 } else if (traits->param == CL_DEVICE_EXTENSIONS_WITH_VERSION) {
3610                         if (ret.err && !checked && output->cond != COND_PROP_SHOW)
3611                                 continue;
3612                         /* This will be displayed at the end, after we display the output of CL_DEVICE_EXTENSIONS */
3613                         const char *msg = RET_BUF(ret)->buf;
3614                         const size_t len = RET_BUF(ret)->sz;
3615                         if (!requested)
3616                                 continue;
3617                         versioned_extensions_traits = traits;
3618                         ALLOC(versioned_extensions, len, "versioned extensions");
3619                         memcpy(versioned_extensions, msg, len);
3620                 } else if (requested) {
3621                         if (ret.err) {
3622                                 /* if there was an error retrieving the property,
3623                                  * skip if it wasn't expected to work and we
3624                                  * weren't asked to show everything regardless of
3625                                  * error */
3626                                 if (!checked && output->cond != COND_PROP_SHOW)
3627                                         continue;
3628
3629                         } else {
3630                                 /* on success, but empty result, show (n/a) */
3631                                 if (ret.str.buf[0] == '\0') {
3632                                         reset_strbuf(&ret.str);
3633                                         strbuf_append_str(loc.pname, &ret.str, not_specified(output));
3634                                 }
3635                         }
3636                         if (output->brief)
3637                                 printf("%s%s\n", line_pfx, RET_BUF(ret)->buf);
3638                         else if (output->json)
3639                                 json_strbuf(RET_BUF(ret), loc.pname, n++, ret.err || ret.needs_escaping);
3640                         else
3641                                 show_strbuf(RET_BUF(ret), loc.pname, 0, ret.err);
3642                 }
3643
3644                 if (ret.err)
3645                         continue;
3646
3647                 switch (traits->param) {
3648                 case CL_DEVICE_VERSION:
3649                         /* compute numeric value for OpenCL version */
3650                         chk.dev_version = getOpenCLVersion(ret.str.buf + 7);
3651                         break;
3652                 case CL_DEVICE_EXTENSIONS:
3653                         identify_device_extensions(extensions, &chk);
3654                         if (!requested) {
3655                                 free(extensions);
3656                                 extensions = NULL;
3657                         }
3658                         break;
3659                 case CL_DEVICE_TYPE:
3660                         chk.devtype = ret.value.devtype;
3661                         break;
3662                 case CL_DEVICE_GLOBAL_MEM_CACHE_TYPE:
3663                         chk.cachetype = ret.value.cachetype;
3664                         break;
3665                 case CL_DEVICE_LOCAL_MEM_TYPE:
3666                         chk.lmemtype = ret.value.lmemtype;
3667                         break;
3668                 case CL_DEVICE_IMAGE_SUPPORT:
3669                         chk.image_support = ret.value.b;
3670                         break;
3671                 case CL_DEVICE_COMPILER_AVAILABLE:
3672                         chk.compiler_available = ret.value.b;
3673                         break;
3674                 case CL_DEVICE_NUM_P2P_DEVICES_AMD:
3675                         chk.p2p_num_devs = ret.value.u32;
3676                         break;
3677                 case CL_DEVICE_SCHEDULING_CONTROLS_CAPABILITIES_ARM:
3678                         chk.arm_register_alloc_support = !!(ret.value.sched_controls & CL_DEVICE_SCHEDULING_REGISTER_ALLOCATION_ARM);
3679                         // TODO warp count support should check for extension version >= 0.4
3680                         chk.arm_warp_count_support = !!(ret.value.sched_controls);
3681                         break;
3682                 default:
3683                         /* do nothing */
3684                         break;
3685                 }
3686         }
3687
3688         // and finally the extensions, if we retrieved them
3689         if (extensions) {
3690                 // undo the padding
3691                 extensions[ext_len + 1] = '\0';
3692                 if (output->json) {
3693                         printf("%s\"%s\" : ", (n > 0 ? comma_str : spc_str),
3694                                 (output->mode == CLINFO_HUMAN ?
3695                                  extensions_traits->pname : extensions_traits->sname));
3696                         json_stringify(extensions + 1);
3697                         ++n;
3698                 } else
3699                         printf("%s" I1_STR "%s\n", line_pfx, (output->mode == CLINFO_HUMAN ?
3700                                         extensions_traits->pname : extensions_traits->sname),
3701                                 extensions + 1);
3702         }
3703         if (versioned_extensions) {
3704                 if (output->json) {
3705                         printf("%s\"%s\" : ", (n > 0 ? comma_str : spc_str),
3706                                 (output->mode == CLINFO_HUMAN ?
3707                                  versioned_extensions_traits->pname : versioned_extensions_traits->sname));
3708                         fputs(versioned_extensions, stdout);
3709                         ++n;
3710                 } else {
3711                         printf("%s" I1_STR "%s\n", line_pfx, (output->mode == CLINFO_HUMAN ?
3712                                         versioned_extensions_traits->pname :
3713                                         versioned_extensions_traits->sname),
3714                                 versioned_extensions);
3715                 }
3716         }
3717         free(extensions);
3718         free(versioned_extensions);
3719         extensions = NULL;
3720         UNINIT_RET(ret);
3721 }
3722
3723 /* list of allowed properties for AMD offline devices */
3724 /* everything else seems to be set to 0, and all the other string properties
3725  * actually segfault the driver */
3726
3727 static const cl_device_info amd_offline_info_whitelist[] = {
3728         CL_DEVICE_NAME,
3729         /* These are present, but all the same, so just skip them:
3730         CL_DEVICE_VENDOR,
3731         CL_DEVICE_VENDOR_ID,
3732         CL_DEVICE_VERSION,
3733         CL_DRIVER_VERSION,
3734         CL_DEVICE_OPENCL_C_VERSION,
3735         */
3736         CL_DEVICE_EXTENSIONS,
3737         CL_DEVICE_TYPE,
3738         CL_DEVICE_GFXIP_MAJOR_AMD,
3739         CL_DEVICE_GFXIP_MINOR_AMD,
3740         CL_DEVICE_MAX_WORK_GROUP_SIZE,
3741         CL_FALSE
3742 };
3743
3744 static const cl_device_info list_info_whitelist[] = {
3745         CL_DEVICE_NAME,
3746         CL_FALSE
3747 };
3748
3749 /* return a list of offline devices from the AMD extension */
3750 cl_device_id *
3751 fetchOfflineDevicesAMD(const struct platform_list *plist, cl_uint p,
3752         /* the number of devices will be returned in ret->value.u32,
3753          * the associated context in ret->base.ctx;
3754          */
3755         struct device_info_ret *ret)
3756 {
3757         cl_platform_id pid = plist->platform[p];
3758         cl_device_id *device = NULL;
3759         cl_uint num_devs = 0;
3760         cl_context ctx = NULL;
3761
3762         cl_context_properties ctxpft[] = {
3763                 CL_CONTEXT_PLATFORM, (cl_context_properties)pid,
3764                 CL_CONTEXT_OFFLINE_DEVICES_AMD, (cl_context_properties)CL_TRUE,
3765                 0
3766         };
3767
3768         ctx = clCreateContextFromType(ctxpft, CL_DEVICE_TYPE_ALL,
3769                 NULL, NULL, &ret->err);
3770         REPORT_ERROR(&ret->err_str, ret->err, "create context");
3771
3772         if (!ret->err) {
3773                 ret->err = REPORT_ERROR(&ret->err_str,
3774                         clGetContextInfo(ctx, CL_CONTEXT_NUM_DEVICES,
3775                                 sizeof(num_devs), &num_devs, NULL),
3776                         "get num devs");
3777         }
3778
3779         if (!ret->err) {
3780                 ALLOC(device, num_devs, "offline devices");
3781
3782                 ret->err = REPORT_ERROR(&ret->err_str,
3783                         clGetContextInfo(ctx, CL_CONTEXT_DEVICES,
3784                                 num_devs*sizeof(*device), device, NULL),
3785                         "get devs");
3786         }
3787
3788         if (ret->err) {
3789                 if (ctx) clReleaseContext(ctx);
3790                 free(device);
3791                 device = NULL;
3792         } else {
3793                 ret->value.u32 = num_devs;
3794                 ret->base.ctx = ctx;
3795         }
3796         return device;
3797 }
3798
3799 void printPlatformName(const struct platform_list *plist, cl_uint p, struct _strbuf *str,
3800         const struct opt_out *output)
3801 {
3802         const struct platform_data *pdata = plist->pdata + p;
3803         const char *brief_prefix = (output->mode == CLINFO_HUMAN ? "Platform #" : "");
3804         const char *title = (output->mode == CLINFO_HUMAN  ? pinfo_traits[0].pname :
3805                 pinfo_traits[0].sname);
3806         const int prefix_width = -line_pfx_len*(!output->brief);
3807         if (output->brief) {
3808                 strbuf_append(__func__, str, "%s%" PRIu32 ": ", brief_prefix, p);
3809         } else if (output->mode == CLINFO_RAW) {
3810                 strbuf_append(__func__, str, "[%s/*]", pdata->sname);
3811         }
3812         sprintf(line_pfx, "%*s", prefix_width, str->buf);
3813         reset_strbuf(str);
3814
3815         if (output->brief)
3816                 printf("%s%s\n", line_pfx, pdata->pname);
3817         else
3818                 printf("%s" I1_STR "%s\n", line_pfx, title, pdata->pname);
3819 }
3820
3821 void printPlatformDevices(const struct platform_list *plist, cl_uint p,
3822         const cl_device_id *device, cl_uint ndevs,
3823         struct _strbuf *str, const struct opt_out *output, cl_bool these_are_offline)
3824 {
3825         const struct platform_data *pdata = plist->pdata + p;
3826         const cl_device_info *param_whitelist = output->brief ? list_info_whitelist :
3827                 these_are_offline ? amd_offline_info_whitelist : NULL;
3828         cl_uint d;
3829
3830         if (output->json)
3831                 printf("%s\"%s\" : [", (these_are_offline ? comma_str : spc_str),
3832                         (these_are_offline ? "offline" : "online"));
3833         else if (output->detailed)
3834                 printf("%s" I0_STR "%" PRIu32 "\n",
3835                         line_pfx,
3836                         num_devs_header(output, these_are_offline),
3837                         ndevs);
3838
3839         for (d = 0; d < ndevs; ++d) {
3840                 const cl_device_id dev = device[d];
3841                 if (output->selected && output->device != d) continue;
3842                 if (output->brief) {
3843                         const cl_bool last_device = (d == ndevs - 1 &&
3844                                 output->mode != CLINFO_RAW &&
3845                                 (!output->offline ||
3846                                  !pdata->has_amd_offline ||
3847                                  these_are_offline));
3848                         if (output->mode == CLINFO_RAW)
3849                                 sprintf(line_pfx, "%" PRIu32 "%c%" PRIu32 ": ",
3850                                         p,
3851                                         these_are_offline ? '*' : '.',
3852                                         d);
3853                         else
3854                                 sprintf(line_pfx, " +-- %sDevice #%" PRIu32 ": ",
3855                                         these_are_offline ? "Offline " : "",
3856                                         d);
3857                         if (last_device)
3858                                 line_pfx[1] = '`';
3859                 } else if (line_pfx_len > 0) {
3860                         cl_int sd = (these_are_offline ? -1 : 1)*(cl_int)d;
3861                         strbuf_append(__func__, str, "[%s/%" PRId32 "]", pdata->sname, sd);
3862                         sprintf(line_pfx, "%*s", -line_pfx_len, str->buf);
3863                         reset_strbuf(str);
3864                 }
3865
3866                 if (output->json)
3867                         printf("%s{", d > 0 ? comma_str : spc_str);
3868
3869                 printDeviceInfo(dev, plist, p, param_whitelist, output);
3870
3871                 if (output->json)
3872                         printf(" }");
3873                 else if (output->detailed && d < pdata[p].ndevs - 1)
3874                         puts("");
3875
3876                 fflush(stdout);
3877                 fflush(stderr);
3878         }
3879         if (output->json)
3880                 fputs(" ]", stdout);
3881 }
3882
3883
3884 void showDevices(const struct platform_list *plist, const struct opt_out *output)
3885 {
3886         const cl_uint num_platforms = plist->num_platforms + (output->null_platform ? 1 : 0);
3887         const cl_uint maxdevs = plist->max_devs;
3888         const struct platform_data *pdata = plist->pdata;
3889
3890         cl_uint p;
3891         struct _strbuf str;
3892         init_strbuf(&str, __func__);
3893
3894         if (output->mode == CLINFO_RAW) {
3895                 if (output->brief)
3896                         strbuf_append(__func__, &str, "%" PRIu32 ".%" PRIu32 ": ", num_platforms, maxdevs);
3897                 else
3898                         strbuf_append(__func__, &str, "[%*s/%" PRIu32 "] ",
3899                                 plist->max_sname_len, "", maxdevs);
3900         } else {
3901                 if (output->brief)
3902                         strbuf_append(__func__, &str, " +-- %sDevice #%" PRIu32 ": ",
3903                                 (output->offline ? "Offline " : ""), maxdevs);
3904                 /* TODO we have no prefix in HUMAN detailed output mode,
3905                  * consider adding one
3906                  */
3907         }
3908
3909         if (str.buf[0]) {
3910                 line_pfx_len = (int)(strlen(str.buf) + 1);
3911                 REALLOC(line_pfx, line_pfx_len, "line prefix");
3912                 reset_strbuf(&str);
3913         }
3914
3915         for (p = 0; p < num_platforms; ++p) {
3916                 /* skip non-selected platforms altogether */
3917                 if (output->selected && output->platform != p) continue;
3918
3919                 /* Open the JSON devices list for this platform */
3920                 if (output->json)
3921                         printf("%s{", p > 0 ? comma_str : spc_str);
3922                 /* skip platform header if only printing specfic properties, */
3923                 else if (!output->prop)
3924                         printPlatformName(plist, p, &str, output);
3925
3926                 printPlatformDevices(plist, p,
3927                         get_platform_devs(plist, p), pdata[p].ndevs,
3928                         &str, output, CL_FALSE);
3929
3930                 if (output->offline && pdata[p].has_amd_offline) {
3931                         struct device_info_ret ret;
3932                         cl_device_id *devs = NULL;
3933
3934                         INIT_RET(ret, "offline device");
3935                         if (output->detailed)
3936                                 puts("");
3937
3938                         devs = fetchOfflineDevicesAMD(plist, p, &ret);
3939                         if (ret.err) {
3940                                 puts(ret.err_str.buf);
3941                         } else {
3942                                 printPlatformDevices(plist, p, devs, ret.value.u32,
3943                                         &str, output, CL_TRUE);
3944                                 clReleaseContext(ret.base.ctx);
3945                                 free(devs);
3946                         }
3947                         UNINIT_RET(ret);
3948                 }
3949
3950                 /* Close JSON object for this platform */
3951                 if (output->json)
3952                         fputs(" }", stdout);
3953                 else if (output->detailed)
3954                         puts("");
3955         }
3956         free_strbuf(&str);
3957 }
3958
3959 /* check the behavior of clGetPlatformInfo() when given a NULL platform ID */
3960 void checkNullGetPlatformName(const struct opt_out *output)
3961 {
3962         struct device_info_ret ret;
3963         struct info_loc loc;
3964
3965         INIT_RET(ret, "null ctx");
3966         reset_loc(&loc, __func__);
3967         RESET_LOC_PARAM(loc, plat, CL_PLATFORM_NAME);
3968
3969         ret.err = clGetPlatformInfo(NULL, CL_PLATFORM_NAME, ret.str.sz, ret.str.buf, NULL);
3970         if (ret.err == CL_INVALID_PLATFORM) {
3971                 strbuf_append(__func__, &ret.err_str, no_plat(output));
3972         } else {
3973                 loc.line = __LINE__ + 1;
3974                 REPORT_ERROR_LOC(&ret, ret.err, &loc, "get %s");
3975         }
3976         printf(I1_STR "%s\n",
3977                 "clGetPlatformInfo(NULL, CL_PLATFORM_NAME, ...)", RET_BUF(ret)->buf);
3978         UNINIT_RET(ret);
3979 }
3980
3981 /* check the behavior of clGetDeviceIDs() when given a NULL platform ID;
3982  * return the index of the default platform in our array of platform IDs,
3983  * or num_platforms (which is an invalid platform index) in case of errors
3984  * or no platform or device found.
3985  */
3986 cl_uint checkNullGetDevices(const struct platform_list *plist, const struct opt_out *output)
3987 {
3988         const cl_uint num_platforms = plist->num_platforms;
3989         const struct platform_data *pdata = plist->pdata;
3990         const cl_platform_id *platform = plist->platform;
3991
3992         struct device_info_ret ret;
3993         struct info_loc loc;
3994
3995         cl_uint i = 0; /* generic iterator */
3996         cl_device_id dev = NULL; /* sample device */
3997         cl_platform_id plat = NULL; /* detected platform */
3998
3999         cl_uint found = 0; /* number of platforms found */
4000         cl_uint pidx = num_platforms; /* index of the platform found */
4001         cl_uint numdevs = 0;
4002
4003         INIT_RET(ret, "null get devices");
4004
4005         reset_loc(&loc, __func__);
4006         loc.sname = "device IDs";
4007
4008         ret.err = clGetDeviceIDs(NULL, CL_DEVICE_TYPE_ALL, 0, NULL, &numdevs);
4009         /* TODO we should check other CL_DEVICE_TYPE_* combinations, since a smart
4010          * implementation might give you a different default platform for GPUs
4011          * and for CPUs.
4012          * Of course the “no devices” case would then need to be handled differently.
4013          * The logic might be maintained similarly, provided we also gather
4014          * the number of devices of each type for each platform, although it's
4015          * obviously more likely to have multiple platforms with no devices
4016          * of a given type.
4017          */
4018
4019         switch (ret.err) {
4020         case CL_INVALID_PLATFORM:
4021                 strbuf_append_str(__func__, &ret.err_str, no_plat(output));
4022                 break;
4023         case CL_DEVICE_NOT_FOUND:
4024                  /* No devices were found, see if there are platforms with
4025                   * no devices, and if there's only one, assume this is the
4026                   * one being used as default by the ICD loader */
4027                 for (i = 0; i < num_platforms; ++i) {
4028                         if (pdata[i].ndevs == 0) {
4029                                 ++found;
4030                                 if (found > 1)
4031                                         break;
4032                                 else {
4033                                         plat = platform[i];
4034                                         pidx = i;
4035                                 }
4036                         }
4037                 }
4038
4039                 switch (found) {
4040                 case 0:
4041                         strbuf_append_str(__func__, &ret.err_str, (output->mode == CLINFO_HUMAN ?
4042                                 "<error: 0 devices, no matching platform!>" :
4043                                 "CL_DEVICE_NOT_FOUND | CL_INVALID_PLATFORM"));
4044                         break;
4045                 case 1:
4046                         strbuf_append(__func__, &ret.err_str, "%s%s%s%s",
4047                                 no_dev_found(output),
4048                                 (output->mode == CLINFO_HUMAN ? " [" : " | "),
4049                                 (output->mode == CLINFO_HUMAN ? pdata[pidx].pname : pdata[pidx].sname),
4050                                 (output->mode == CLINFO_HUMAN ? "?]" : "?"));
4051                         break;
4052                 default: /* found > 1 */
4053                         strbuf_append_str(__func__, &ret.err_str, (output->mode == CLINFO_HUMAN ?
4054                                 "<error: 0 devices, multiple matching platforms!>" :
4055                                 "CL_DEVICE_NOT_FOUND | ????"));
4056                         break;
4057                 }
4058                 break;
4059         default:
4060                 loc.line = __LINE__+1;
4061                 if (REPORT_ERROR_LOC(&ret, ret.err, &loc, "get number of %s")) break;
4062
4063                 /* Determine platform by looking at the CL_DEVICE_PLATFORM of
4064                  * one of the devices */
4065                 ret.err = clGetDeviceIDs(NULL, CL_DEVICE_TYPE_ALL, 1, &dev, NULL);
4066                 loc.line = __LINE__+1;
4067                 if (REPORT_ERROR_LOC(&ret, ret.err, &loc, "get %s")) break;
4068
4069                 RESET_LOC_PARAM(loc, dev, CL_DEVICE_PLATFORM);
4070                 ret.err = clGetDeviceInfo(dev, CL_DEVICE_PLATFORM,
4071                         sizeof(plat), &plat, NULL);
4072                 loc.line = __LINE__+1;
4073                 if (REPORT_ERROR_LOC(&ret, ret.err, &loc, "get %s")) break;
4074
4075                 for (i = 0; i < num_platforms; ++i) {
4076                         if (platform[i] == plat) {
4077                                 pidx = i;
4078                                 strbuf_append(__func__, &ret.str, "%s [%s]",
4079                                         (output->mode == CLINFO_HUMAN ? "Success" : "CL_SUCCESS"),
4080                                         pdata[i].sname);
4081                                 break;
4082                         }
4083                 }
4084                 if (i == num_platforms) {
4085                         ret.err = CL_INVALID_PLATFORM;
4086                         strbuf_append(__func__, &ret.err_str, "<error: platform %p not found>", (void*)plat);
4087                 }
4088         }
4089         printf(I1_STR "%s\n",
4090                 "clGetDeviceIDs(NULL, CL_DEVICE_TYPE_ALL, ...)", RET_BUF(ret)->buf);
4091
4092         UNINIT_RET(ret);
4093         return pidx;
4094 }
4095
4096 void checkNullCtx(struct device_info_ret *ret,
4097         const struct platform_list *plist, cl_uint pidx, const char *which,
4098         const struct opt_out *output)
4099 {
4100         const cl_device_id *dev = plist->all_devs + plist->dev_offset[pidx];
4101         struct info_loc loc;
4102         cl_context ctx = clCreateContext(NULL, 1, dev, NULL, NULL, &ret->err);
4103
4104         reset_loc(&loc, __func__);
4105         loc.sname = which;
4106         loc.line = __LINE__+2;
4107
4108         if (!REPORT_ERROR_LOC(ret, ret->err, &loc, "create context with device from %s platform"))
4109                 strbuf_append(__func__, &ret->str, "%s [%s]",
4110                         (output->mode == CLINFO_HUMAN ? "Success" : "CL_SUCCESS"),
4111                         plist->pdata[pidx].sname);
4112         if (ctx) {
4113                 clReleaseContext(ctx);
4114                 ctx = NULL;
4115         }
4116 }
4117
4118 /* check behavior of clCreateContextFromType() with NULL cl_context_properties */
4119 void checkNullCtxFromType(const struct platform_list *plist, const struct opt_out *output)
4120 {
4121         const cl_uint num_platforms = plist->num_platforms;
4122         const struct platform_data *pdata = plist->pdata;
4123         const cl_platform_id *platform = plist->platform;
4124
4125         size_t t; /* type iterator */
4126         size_t i; /* generic iterator */
4127         char def[1024];
4128         cl_context ctx = NULL;
4129
4130         size_t ndevs = 8;
4131         size_t szval = 0;
4132         size_t cursz = ndevs*sizeof(cl_device_id);
4133         cl_platform_id plat = NULL;
4134         cl_device_id *devs = NULL;
4135
4136         struct device_info_ret ret;
4137         struct info_loc loc;
4138
4139         const char *platname_prop = (output->mode == CLINFO_HUMAN ?
4140                 pinfo_traits[0].pname :
4141                 pinfo_traits[0].sname);
4142
4143         const char *devname_prop = (output->mode == CLINFO_HUMAN ?
4144                 dinfo_traits[0].pname :
4145                 dinfo_traits[0].sname);
4146
4147         reset_loc(&loc, __func__);
4148         INIT_RET(ret, "null ctx from type");
4149
4150         ALLOC(devs, ndevs, "context devices");
4151
4152         for (t = 1; t < devtype_count; ++t) { /* we skip 0 */
4153                 loc.sname = device_type_raw_str[t];
4154
4155                 strbuf_append(__func__, &ret.str, "clCreateContextFromType(NULL, %s)", loc.sname);
4156                 sprintf(def, I1_STR, ret.str.buf);
4157                 reset_strbuf(&ret.str);
4158
4159                 loc.line = __LINE__+1;
4160                 ctx = clCreateContextFromType(NULL, devtype[t], NULL, NULL, &ret.err);
4161
4162                 switch (ret.err) {
4163                 case CL_INVALID_PLATFORM:
4164                         strbuf_append_str(__func__, &ret.err_str, no_plat(output)); break;
4165                 case CL_DEVICE_NOT_FOUND:
4166                         strbuf_append_str(__func__, &ret.err_str, no_dev_found(output)); break;
4167                 case CL_INVALID_DEVICE_TYPE: /* e.g. _CUSTOM device on 1.1 platform */
4168                         strbuf_append_str(__func__, &ret.err_str, invalid_dev_type(output)); break;
4169                 case CL_INVALID_VALUE: /* This is what apple returns for the case above */
4170                         strbuf_append_str(__func__, &ret.err_str, invalid_dev_type(output)); break;
4171                 case CL_DEVICE_NOT_AVAILABLE:
4172                         strbuf_append_str(__func__, &ret.err_str, no_dev_avail(output)); break;
4173                 default:
4174                         if (REPORT_ERROR_LOC(&ret, ret.err, &loc, "create context from type %s")) break;
4175
4176                         /* get the devices */
4177                         loc.sname = "CL_CONTEXT_DEVICES";
4178                         loc.line = __LINE__+2;
4179
4180                         ret.err = clGetContextInfo(ctx, CL_CONTEXT_DEVICES, 0, NULL, &szval);
4181                         if (REPORT_ERROR_LOC(&ret, ret.err, &loc, "get %s size")) break;
4182                         if (szval > cursz) {
4183                                 REALLOC(devs, szval, "context devices");
4184                                 cursz = szval;
4185                         }
4186
4187                         loc.line = __LINE__+1;
4188                         ret.err = clGetContextInfo(ctx, CL_CONTEXT_DEVICES, cursz, devs, NULL);
4189                         if (REPORT_ERROR_LOC(&ret, ret.err, &loc, "get %s")) break;
4190                         ndevs = szval/sizeof(cl_device_id);
4191                         if (ndevs < 1) {
4192                                 ret.err = CL_DEVICE_NOT_FOUND;
4193                                 strbuf_append_str(__func__, &ret.err_str, "<error: context created with no devices>");
4194                         }
4195
4196                         /* get the platform from the first device */
4197                         RESET_LOC_PARAM(loc, dev, CL_DEVICE_PLATFORM);
4198                         loc.line = __LINE__+1;
4199                         ret.err = clGetDeviceInfo(*devs, CL_DEVICE_PLATFORM, sizeof(plat), &plat, NULL);
4200                         if (REPORT_ERROR_LOC(&ret, ret.err, &loc, "get %s")) break;
4201                         loc.plat = plat;
4202
4203                         for (i = 0; i < num_platforms; ++i) {
4204                                 if (platform[i] == plat)
4205                                         break;
4206                         }
4207                         if (i == num_platforms) {
4208                                 ret.err = CL_INVALID_PLATFORM;
4209                                 strbuf_append(__func__, &ret.err_str, "<error: platform %p not found>", (void*)plat);
4210                                 break;
4211                         } else {
4212                                 strbuf_append(__func__, &ret.str, "%s (%" PRIuS ")",
4213                                         (output->mode == CLINFO_HUMAN ? "Success" : "CL_SUCCESS"),
4214                                         ndevs);
4215                                 strbuf_append(__func__, &ret.str, "\n" I2_STR "%s",
4216                                         platname_prop, pdata[i].pname);
4217                         }
4218                         for (i = 0; i < ndevs; ++i) {
4219                                 size_t szname = 0;
4220                                 /* for each device, show the device name */
4221                                 /* TODO some other unique ID too, e.g. PCI address, if available? */
4222
4223                                 strbuf_append(__func__, &ret.str, "\n" I2_STR, devname_prop);
4224
4225                                 RESET_LOC_PARAM(loc, dev, CL_DEVICE_NAME);
4226                                 loc.dev = devs[i];
4227                                 loc.line = __LINE__+1;
4228                                 ret.err = clGetDeviceInfo(devs[i], CL_DEVICE_NAME, ret.str.sz - ret.str.end, ret.str.buf + ret.str.end, &szname);
4229                                 if (REPORT_ERROR_LOC(&ret, ret.err, &loc, "get %s")) break;
4230                                 ret.str.end += szname - 1;
4231                         }
4232                         if (i != ndevs)
4233                                 break; /* had an error earlier, bail */
4234                 }
4235
4236                 if (ctx) {
4237                         clReleaseContext(ctx);
4238                         ctx = NULL;
4239                 }
4240                 printf("%s%s\n", def, RET_BUF(ret)->buf);
4241                 reset_strbuf(&ret.str);
4242                 reset_strbuf(&ret.err_str);
4243         }
4244         free(devs);
4245         UNINIT_RET(ret);
4246 }
4247
4248 /* check the behavior of NULL platform in clGetDeviceIDs (see checkNullGetDevices)
4249  * and in clCreateContext() */
4250 void checkNullBehavior(const struct platform_list *plist, const struct opt_out *output)
4251 {
4252         const cl_uint num_platforms = plist->num_platforms;
4253         const struct platform_data *pdata = plist->pdata;
4254
4255         cl_uint p = 0;
4256         struct device_info_ret ret;
4257
4258         INIT_RET(ret, "null behavior");
4259
4260         printf("NULL platform behavior\n");
4261
4262         checkNullGetPlatformName(output);
4263
4264         p = checkNullGetDevices(plist, output);
4265
4266         /* If there's a default platform, and it has devices, try
4267          * creating a context with its first device and see if it works */
4268
4269         if (p == num_platforms) {
4270                 ret.err = CL_INVALID_PLATFORM;
4271                 strbuf_append(__func__, &ret.err_str, no_plat(output));
4272         } else if (pdata[p].ndevs == 0) {
4273                 ret.err = CL_DEVICE_NOT_FOUND;
4274                 strbuf_append(__func__, &ret.err_str, no_dev_found(output));
4275         } else {
4276                 if (p < num_platforms) {
4277                         checkNullCtx(&ret, plist, p, "default", output);
4278                 } else {
4279                         /* this shouldn't happen, but still ... */
4280                         ret.err = CL_OUT_OF_HOST_MEMORY;
4281                         strbuf_append_str(__func__, &ret.err_str, "<error: overflow in default platform scan>");
4282                 }
4283         }
4284         printf(I1_STR "%s\n", "clCreateContext(NULL, ...) [default]", RET_BUF(ret)->buf);
4285
4286         /* Look for a device from a non-default platform, if there are any */
4287         if (p == num_platforms || num_platforms > 1) {
4288                 cl_uint p2 = 0;
4289                 reset_strbuf(&ret.str);
4290                 reset_strbuf(&ret.err_str);
4291                 while (p2 < num_platforms && (p2 == p || pdata[p2].ndevs == 0)) {
4292                         p2++;
4293                 }
4294                 if (p2 < num_platforms) {
4295                         checkNullCtx(&ret, plist, p2, "non-default", output);
4296                 } else {
4297                         ret.err = CL_DEVICE_NOT_FOUND;
4298                         strbuf_append(__func__, &ret.err_str, "<error: no devices in non-default plaforms>");
4299                 }
4300                 printf(I1_STR "%s\n", "clCreateContext(NULL, ...) [other]", RET_BUF(ret)->buf);
4301         }
4302
4303         checkNullCtxFromType(plist, output);
4304
4305         UNINIT_RET(ret);
4306 }
4307
4308
4309 /* Get properties of the ocl-icd loader, if available */
4310 /* All properties are currently char[] */
4311
4312 /* Function pointer to the ICD loader info function */
4313
4314 typedef cl_int (*icdl_info_fn_ptr)(cl_icdl_info, size_t, void*, size_t*);
4315 icdl_info_fn_ptr clGetICDLoaderInfoOCLICD;
4316
4317 /* We want to auto-detect the OpenCL version supported by the ICD loader.
4318  * To do this, we will progressively find symbols introduced in new APIs,
4319  * until a NULL symbol is found.
4320  */
4321
4322 struct icd_loader_test {
4323         cl_uint version;
4324         const char *symbol;
4325 } icd_loader_tests[] = {
4326         { 11, "clCreateSubBuffer" },
4327         { 12, "clCreateImage" },
4328         { 20, "clSVMAlloc" },
4329         { 21, "clGetHostTimer" },
4330         { 22, "clSetProgramSpecializationConstant" },
4331         { 30, "clSetContextDestructorCallback" },
4332         { 0, NULL }
4333 };
4334
4335 void
4336 icdl_info_str(struct icdl_info_ret *ret, const struct info_loc *loc)
4337 {
4338         GET_STRING_LOC(ret, loc, clGetICDLoaderInfoOCLICD, loc->param.icdl);
4339         return;
4340 }
4341
4342 struct icdl_info_traits {
4343         cl_icdl_info param; // CL_ICDL_*
4344         const char *sname; // "CL_ICDL_*"
4345         const char *pname; // "ICD loader *"
4346 };
4347
4348 static const char * const oclicdl_pfx = "OCLICD";
4349
4350 #define LINFO(symbol, name) { symbol, #symbol, "ICD loader " name }
4351 struct icdl_info_traits linfo_traits[] = {
4352         LINFO(CL_ICDL_NAME, "Name"),
4353         LINFO(CL_ICDL_VENDOR, "Vendor"),
4354         LINFO(CL_ICDL_VERSION, "Version"),
4355         LINFO(CL_ICDL_OCL_VERSION, "Profile")
4356 };
4357
4358 /* The ICD loader info function must be retrieved via clGetExtensionFunctionAddress,
4359  * which returns a void pointer.
4360  * ISO C forbids assignments between function pointers and void pointers,
4361  * but POSIX allows it. To compile without warnings even in -pedantic mode,
4362  * we take advantage of the fact that we _can_ do the conversion via
4363  * pointers-to-pointers. This is supported on most compilers, except
4364  * for some rather old GCC versions whose strict aliasing rules are
4365  * too strict. Disable strict aliasing warnings for these compilers.
4366  */
4367 #if defined __GNUC__ && ((__GNUC__*10 + __GNUC_MINOR__) < 46)
4368 #pragma GCC diagnostic ignored "-Wstrict-aliasing"
4369 #endif
4370
4371 struct icdl_data oclIcdProps(const struct platform_list *plist, const struct opt_out *output)
4372 {
4373         const cl_uint max_plat_version = plist->max_plat_version;
4374
4375         struct icdl_data icdl;
4376
4377         /* clinfo may lag behind the OpenCL standard or loader version,
4378          * and we don't want to give a warning if we can't tell if the loader
4379          * correctly supports a version unknown to us
4380          */
4381         cl_uint clinfo_highest_known_version = 0;
4382
4383         /* Counter that'll be used to walk the icd_loader_tests */
4384         int i = 0;
4385
4386         /* We find the clGetICDLoaderInfoOCLICD extension address, which will be used
4387          * to query the ICD loader properties.
4388          * It should be noted that in this specific case we cannot replace the
4389          * call to clGetExtensionFunctionAddress with a call to the superseding function
4390          * clGetExtensionFunctionAddressForPlatform because the extension is in the
4391          * loader itself, not in a specific platform.
4392          */
4393         void *ptrHack = clGetExtensionFunctionAddress("clGetICDLoaderInfoOCLICD");
4394         clGetICDLoaderInfoOCLICD = *(icdl_info_fn_ptr*)(&ptrHack);
4395
4396         /* Initialize icdl_data ret versions */
4397         icdl.detected_version = 10;
4398         icdl.reported_version = 0;
4399
4400         /* Step #1: try to auto-detect the supported ICD loader version */
4401         do {
4402                 struct icd_loader_test check = icd_loader_tests[i];
4403                 if (check.symbol == NULL)
4404                         break;
4405                 if (dlsym(DL_MODULE, check.symbol) == NULL)
4406                         break;
4407                 clinfo_highest_known_version = icdl.detected_version = check.version;
4408                 ++i;
4409         } while (1);
4410
4411         /* Step #2: query properties from extension, if available */
4412         if (clGetICDLoaderInfoOCLICD != NULL) {
4413                 cl_uint n = 0; /* number of ICD loader properties shown, for JSON */
4414                 struct info_loc loc;
4415                 struct icdl_info_ret ret;
4416                 reset_loc(&loc, __func__);
4417                 INIT_RET(ret, "ICD loader");
4418
4419                 /* TODO think of a sensible header in CLINFO_RAW */
4420                 if (output->mode != CLINFO_RAW)
4421                         puts("\nICD loader properties");
4422
4423                 if (output->json) {
4424                         fputs(", \"icd_loader\" : {", stdout);
4425                 } else if (output->mode == CLINFO_RAW) {
4426                         line_pfx_len = (int)(strlen(oclicdl_pfx) + 5);
4427                         REALLOC(line_pfx, line_pfx_len, "line prefix OCL ICD");
4428                         strbuf_append(loc.pname, &ret.str, "[%s/*]", oclicdl_pfx);
4429                         sprintf(line_pfx, "%*s", -line_pfx_len, ret.str.buf);
4430                         reset_strbuf(&ret.str);
4431                 }
4432
4433                 for (loc.line = 0; loc.line < ARRAY_SIZE(linfo_traits); ++loc.line) {
4434                         const struct icdl_info_traits *traits = linfo_traits + loc.line;
4435                         cl_bool requested;
4436                         loc.sname = traits->sname;
4437                         loc.pname = (output->mode == CLINFO_HUMAN ?
4438                                 traits->pname : traits->sname);
4439                         loc.param.icdl = traits->param;
4440
4441                         reset_strbuf(&ret.str);
4442                         reset_strbuf(&ret.err_str);
4443                         icdl_info_str(&ret, &loc);
4444
4445                         /* Do not print this property if the user requested one and this does not match */
4446                         requested = !(output->prop && strstr(loc.sname, output->prop) == NULL);
4447                         if (requested) {
4448                                 if (output->json)
4449                                         json_strbuf(RET_BUF(ret), loc.pname, n++, CL_TRUE);
4450                                 else
4451                                         show_strbuf(RET_BUF(ret), loc.pname, 1, ret.err);
4452                         }
4453
4454                         if (!ret.err && traits->param == CL_ICDL_OCL_VERSION) {
4455                                 icdl.reported_version = getOpenCLVersion(ret.str.buf + 7);
4456                         }
4457                 }
4458
4459                 if (output->json)
4460                         printf("%s\"_detected_version\" : \"%" PRIu32 ".%" PRIu32 "\" }",
4461                                 (n > 0 ? comma_str : spc_str),
4462                                 SPLIT_CL_VERSION(icdl.detected_version));
4463                 UNINIT_RET(ret);
4464         }
4465
4466         /* Step #3: show it */
4467         if (output->mode == CLINFO_HUMAN) {
4468                 // for the loader vs platform max version check we use the version we detected
4469                 // if the reported version is known to us, and the reported version if it's higher
4470                 // than the standard versions we know about
4471                 cl_uint max_version_check = icdl.reported_version > clinfo_highest_known_version ?
4472                         icdl.reported_version : icdl.detected_version;
4473
4474                 if (icdl.reported_version &&
4475                         icdl.reported_version <= clinfo_highest_known_version &&
4476                         icdl.reported_version != icdl.detected_version) {
4477                         printf( "\tNOTE:\tyour OpenCL library declares to support OpenCL %" PRIu32 ".%" PRIu32 ",\n"
4478                                 "\t\tbut it seems to support up to OpenCL %" PRIu32 ".%" PRIu32 " %s.\n",
4479                                 SPLIT_CL_VERSION(icdl.reported_version),
4480                                 SPLIT_CL_VERSION(icdl.detected_version),
4481                                 icdl.detected_version < icdl.reported_version  ?
4482                                 "only" : "too");
4483                 }
4484
4485                 if (max_version_check < max_plat_version) {
4486                         printf( "\tNOTE:\tyour OpenCL library only supports OpenCL %" PRIu32 ".%" PRIu32 ",\n"
4487                                 "\t\tbut some installed platforms support OpenCL %" PRIu32 ".%" PRIu32 ".\n"
4488                                 "\t\tPrograms using %" PRIu32 ".%" PRIu32 " features may crash\n"
4489                                 "\t\tor behave unexpectedly\n",
4490                                 SPLIT_CL_VERSION(icdl.detected_version),
4491                                 SPLIT_CL_VERSION(max_plat_version),
4492                                 SPLIT_CL_VERSION(max_plat_version));
4493                 }
4494         }
4495         return icdl;
4496 }
4497
4498 #if defined __GNUC__ && ((__GNUC__*10 + __GNUC_MINOR__) < 46)
4499 #pragma GCC diagnostic warning "-Wstrict-aliasing"
4500 #endif
4501
4502 void version(void)
4503 {
4504         puts("clinfo version 3.0.23.01.25");
4505 }
4506
4507 void parse_device_spec(const char *str, struct opt_out *output)
4508 {
4509         int p, d, n;
4510         if (!str) {
4511                 fprintf(stderr, "please specify a device in the form P:D where P is the platform number and D the device number\n");
4512                 exit(1);
4513         }
4514         n = sscanf(str, "%d:%d", &p, &d);
4515         if (n != 2 || p < 0 || d < 0) {
4516                 fprintf(stderr, "invalid device specification '%s'\n", str);
4517                 exit(1);
4518         }
4519         output->platform = p;
4520         output->device = d;
4521 }
4522
4523 void free_output(struct opt_out *output)
4524 {
4525         free((char*)output->prop);
4526         output->prop = NULL;
4527 }
4528
4529 void parse_prop(const char *input, struct opt_out *output)
4530 {
4531         /* We normalize the property name by upcasing it and replacing the minus sign (-)
4532          * with the underscore (_). If any other character is found, we consider it an error
4533          */
4534
4535         size_t len = strlen(input);
4536         char *normalized;
4537         ALLOC(normalized, len+1, "normalized property name");
4538         for (size_t i = 0; i < len; ++i)
4539         {
4540                 char c = input[i];
4541                 if ( (c == '_') || ( c >= 'A' && c <= 'Z'))
4542                         normalized[i] = c;
4543                 else if (c >= 'a' && c <= 'z')
4544                         normalized[i] = 'A' + (c - 'a');
4545                 else if (c == '-')
4546                         normalized[i] = '_';
4547                 else {
4548                         fprintf(stderr, "invalid property name substring '%s'\n", input);
4549                         exit(1);
4550                 }
4551         }
4552
4553         if (output->prop) {
4554                 fprintf(stderr, "WARNING: only one property name substring supported, discarding %s in favor of %s\n",
4555                         output->prop, normalized);
4556                 free_output(output);
4557         }
4558         output->prop = normalized;
4559 }
4560
4561 void usage(void)
4562 {
4563         version();
4564         puts("Display properties of all available OpenCL platforms and devices");
4565         puts("Usage: clinfo [options ...]\n");
4566         puts("Options:");
4567         puts("\t--all-props, -a\t\ttry all properties, only show valid ones");
4568         puts("\t--always-all-props, -At\tshow all properties, even if invalid");
4569         puts("\t--human\t\thuman-friendly output (default)");
4570         puts("\t--raw\t\traw output");
4571         puts("\t--offline\talso show offline devices");
4572         puts("\t--null-platform\talso show the NULL platform devices");
4573         puts("\t--list, -l\tonly list the platforms and devices by name");
4574         puts("\t--prop prop-name\tonly list properties matching the given name");
4575         puts("\t--device p:d,");
4576         puts("\t-d p:d\t\tonly show information about device number d from platform number p");
4577         puts("\t-h, -?\t\tshow usage");
4578         puts("\t--version, -v\tshow version\n");
4579         puts("Defaults to raw mode if invoked with");
4580         puts("a name that contains the string \"raw\"");
4581 }
4582
4583 int main(int argc, char *argv[])
4584 {
4585         cl_uint p;
4586         cl_int err;
4587         int a = 0;
4588
4589         struct opt_out output;
4590
4591         struct platform_list plist;
4592         init_plist(&plist);
4593
4594         output.platform = CL_UINT_MAX;
4595         output.device = CL_UINT_MAX;
4596         output.prop = NULL;
4597         output.mode = CLINFO_HUMAN;
4598         output.cond = COND_PROP_CHECK;
4599         output.brief = CL_FALSE;
4600         output.offline = CL_FALSE;
4601         output.null_platform = CL_FALSE;
4602         output.json = CL_FALSE;
4603         output.check_size = CL_FALSE;
4604
4605         /* if there's a 'raw' in the program name, switch to raw output mode */
4606         if (strstr(argv[0], "raw"))
4607                 output.mode = CLINFO_RAW;
4608
4609         /* process command-line arguments */
4610         while (++a < argc) {
4611                 if (!strcmp(argv[a], "-a") || !strcmp(argv[a], "--all-props"))
4612                         output.cond = COND_PROP_TRY;
4613                 else if (!strcmp(argv[a], "-A") || !strcmp(argv[a], "--always-all-props"))
4614                         output.cond = COND_PROP_SHOW;
4615                 else if (!strcmp(argv[a], "--raw"))
4616                         output.mode = CLINFO_RAW;
4617                 else if (!strcmp(argv[a], "--human"))
4618                         output.mode = CLINFO_HUMAN;
4619                 else if (!strcmp(argv[a], "--offline"))
4620                         output.offline = CL_TRUE;
4621                 else if (!strcmp(argv[a], "--null-platform"))
4622                         output.null_platform = CL_TRUE;
4623                 else if (!strcmp(argv[a], "--json"))
4624                         output.json = CL_TRUE;
4625                 else if (!strcmp(argv[a], "-l") || !strcmp(argv[a], "--list"))
4626                         output.brief = CL_TRUE;
4627                 else if (!strcmp(argv[a], "-d") || !strcmp(argv[a], "--device")) {
4628                         ++a;
4629                         parse_device_spec(argv[a], &output);
4630                 } else if (!strncmp(argv[a], "-d", 2)) {
4631                         parse_device_spec(argv[a] + 2, &output);
4632                 } else if (!strcmp(argv[a], "--prop")) {
4633                         ++a;
4634                         parse_prop(argv[a], &output);
4635                 } else if (!strcmp(argv[a], "-?") || !strcmp(argv[a], "-h")) {
4636                         usage();
4637                         free_output(&output);
4638                         return 0;
4639                 } else if (!strcmp(argv[a], "--version") || !strcmp(argv[a], "-v")) {
4640                         version();
4641                         free_output(&output);
4642                         return 0;
4643                 } else {
4644                         fprintf(stderr, "ignoring unknown command-line parameter %s\n", argv[a]);
4645                 }
4646         }
4647         /* If a property was specified, we only print in RAW mode.
4648          * Likewise, JSON format assumes RAW
4649          */
4650         if (output.prop || output.json)
4651                 output.mode = CLINFO_RAW;
4652         output.selected = (output.device != CL_UINT_MAX);
4653         output.detailed = !output.brief && !output.selected && !output.prop;
4654
4655         err = clGetPlatformIDs(0, NULL, &plist.num_platforms);
4656         if (err != CL_PLATFORM_NOT_FOUND_KHR)
4657                 CHECK_ERROR(err, "number of platforms");
4658
4659         if (output.detailed && !output.json)
4660                 printf(I0_STR "%" PRIu32 "\n",
4661                         (output.mode == CLINFO_HUMAN ?
4662                          "Number of platforms" : "#PLATFORMS"),
4663                         plist.num_platforms);
4664
4665         cl_uint alloced_platforms = 0;
4666         if (plist.num_platforms) {
4667                 alloced_platforms = alloc_plist(&plist, &output);
4668                 err = clGetPlatformIDs(plist.num_platforms, plist.platform, NULL);
4669                 CHECK_ERROR(err, "platform IDs");
4670         }
4671
4672         ALLOC(line_pfx, 1, "line prefix");
4673
4674         /* Open the JSON object and the JSON platforms list */
4675         if (output.json)
4676                 fputs("{ \"platforms\" : [", stdout);
4677
4678         for (p = 0; p < alloced_platforms; ++p) {
4679                 // skip non-selected platforms altogether
4680                 if (output.selected && output.platform != p) continue;
4681
4682                 /* Open a JSON object for this platform */
4683                 if (output.json)
4684                         printf("%s{", p > 0 ? comma_str : spc_str);
4685
4686                 gatherPlatformInfo(&plist, p, &output);
4687
4688                 /* Close JSON object for this platform */
4689                 if (output.json)
4690                         fputs(" }", stdout);
4691                 else if (output.detailed)
4692                         puts("");
4693         }
4694
4695         /* Close JSON platforms list, open JSON devices list */
4696         if (alloced_platforms) {
4697                 if (output.json)
4698                         fputs(" ], \"devices\" : [", stdout);
4699
4700                 showDevices(&plist, &output);
4701         }
4702
4703         /* Close JSON devices list */
4704         if (output.json)
4705                 fputs(" ]", stdout);
4706
4707         if (output.prop || (output.detailed && !output.selected)) {
4708                 if (output.mode != CLINFO_RAW && plist.num_platforms)
4709                         checkNullBehavior(&plist, &output);
4710                 oclIcdProps(&plist, &output);
4711         }
4712
4713         /* Close the JSON object */
4714         if (output.json)
4715                 fputs(" }", stdout);
4716
4717
4718         free_plist(&plist);
4719         free(line_pfx);
4720         free_output(&output);
4721         return 0;
4722 }