ze_bindings_base.rb 75.4 KB
Newer Older
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
using ZE::ZERefinements
module ZE
  class Pointer < FFI::Pointer
    def initialize(*args)
      if args.length == 2 then
        super(ZE::find_type(args[0]), args[1])
      else
        super(*args)
      end
    end
  end

  class Function < FFI::Function
    def initialize(ret, args, *opts, &block)
      super(ZE::find_type(ret), args.collect { |a| ZE::find_type(a) }, *opts, &block)
    end
  end

  class MemoryPointer < FFI::MemoryPointer

    def initialize(size, count = 1, clear = true)
      if size.is_a?(Symbol)
        size = ZE::find_type(size)
      end
      super(size, count, clear)
    end

  end

Brice Videau's avatar
Brice Videau committed
30
  UINT32_MAX = 0xffffffff
31
32
33
  UINT64_MAX = 0xffffffffffffffff
  ZES_EVENT_WAIT_NONE = 0x0
  ZES_EVENT_WAIT_INFINITE = UINT32_MAX
Brice Videau's avatar
Brice Videau committed
34

35
36
37
38
  ZES_DIAG_FIRST_TEST_INDEX =  0x0
  ZES_DIAG_LAST_TEST_INDEX = 0xffffffff
  ZES_FAN_TEMP_SPEED_PAIR_COUNT = 32
  ZES_MAX_RAS_ERROR_CATEGORY_COUNT = 7
Brice Videau's avatar
Brice Videau committed
39

Brice Videau's avatar
Brice Videau committed
40
41
42
43
  ZET_DEBUG_TIMEOUT_INFINITE = 0xffffffffffffffff
  ZET_DEBUG_THREAD_NONE = 0xffffffffffffffff
  ZET_DEBUG_THREAD_ALL = 0xfffffffffffffffe

Brice Videau's avatar
Brice Videau committed
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
  ZE_OBJECTS = {}
  ZE_OBJECTS_MUTEX = Mutex.new

  def self.register(handle, destructor)
    if handle
      ZE_OBJECTS_MUTEX.synchronize {
        ZE_OBJECTS[handle] = destructor
      }
    else
      raise "invalid handle"
    end
  end

  def self.unregister(handle)
    if handle
      ZE_OBJECTS_MUTEX.synchronize {
        d = ZE_OBJECTS.delete(handle) { raise "non-existing handle" }
        result = method(d).call(handle)
        ZE.error_check(result)
      }
    else
      raise "invalid handle"
    end
  end

69
70
  ENV["ZE_ENABLE_VALIDATION_LAYER"] = "1"
  ENV["ZE_ENABLE_LOADER_INTERCEPT"] = "1"
71
72
73
74
  ENV["ZE_ENABLE_PARAMETER_VALIDATION"] = "1"
  ENV["ZET_ENABLE_API_TRACING_EXP"] = "1"
  ENV["ZET_ENABLE_METRICS"] = "1"
#  ENV["ZET_ENABLE_PROGRAM_INSTRUMENTATION"] = "1"
Brice Videau's avatar
Bugfix.    
Brice Videau committed
75
#  ENV["ZET_ENABLE_PROGRAM_DEBUGGING"] = "1"
76
  ENV["ZES_ENABLE_SYSMAN"] = "1"
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
  if ENV["LIBZE_LOADER_SO"]
    ffi_lib ENV["LIBZE_LOADER_SO"]
  else
    begin
      ffi_lib "ze_loader"
    rescue FFI::LoadError
      ffi_lib "libze_loader.so.0.91"
    end
  end

  class Error < StandardError
  end

  def self.error_check(result)
    if result != :ZE_RESULT_SUCCESS &&
Brice Videau's avatar
Brice Videau committed
92
       result != :ZE_RESULT_NOT_READY
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
      raise Error, result
    end
  end

  class Object
    attr_reader :handle
    def initialize(handle)
      @handle = handle
    end

    def to_ptr
      @handle
    end

    def self.process_property_name(pname)
      res = pname.to_s.split("_").collect(&:capitalize).join
109
      res.gsub("Uuid","UUID").gsub("Dditable", "DDITable").gsub(/\AFp/, "FP").gsub("P2p", "P2P")
110
111
    end

Brice Videau's avatar
WIP    
Brice Videau committed
112
113
    def self.process_ffi_name(name)
      res = name.to_s.gsub(/_t\z/, "").split("_").collect(&:capitalize).join
114
      res.gsub(/\AZet/,"ZET").gsub(/\AZes/,"ZES").gsub(/\AZe/, "ZE").gsub("Uuid","UUID").gsub("Dditable", "DDITable").gsub(/\AFp/, "FP").gsub("P2p", "P2P")
115
116
    end

117
    def self.add_object_array(aname, oname, fname, memoize: false)
118
119
      src = <<EOF
    def #{aname}
Brice Videau's avatar
WIP    
Brice Videau committed
120
121
122
EOF
      if memoize
        src << <<EOF
123
      return @#{aname} if @#{aname}
Brice Videau's avatar
WIP    
Brice Videau committed
124
125
126
EOF
      end
      src << <<EOF
127
128
129
130
131
132
133
134
      pCount = MemoryPointer::new(:uint32_t)
      result = ZE.#{fname}(@handle, pCount, nil)
      ZE.error_check(result)
      count = pCount.read(:uint32)
      return [] if count == 0
      pArr = MemoryPointer::new(:pointer, count)
      result = ZE.#{fname}(@handle, pCount, pArr)
      ZE.error_check(result)
Brice Videau's avatar
WIP    
Brice Videau committed
135
136
137
EOF
      if memoize
        src << <<EOF
138
139
140
      @#{aname} = pArr.read_array_of_pointer(count).collect { |h| #{oname}::new(h) }
    end
EOF
Brice Videau's avatar
WIP    
Brice Videau committed
141
142
143
144
145
146
      else
        src << <<EOF
      pArr.read_array_of_pointer(count).collect { |h| #{oname}::new(h) }
    end
EOF
      end
Brice Videau's avatar
Brice Videau committed
147
      class_eval src
148
149
    end

150
    def self.add_array_property(aname, sname, fname, memoize: false)
151
152
      src = <<EOF
    def #{aname}
Brice Videau's avatar
WIP    
Brice Videau committed
153
154
155
EOF
      if memoize
        src << <<EOF
156
      return @#{aname} if @#{aname}
Brice Videau's avatar
WIP    
Brice Videau committed
157
158
159
EOF
      end
      src << <<EOF
160
161
162
163
164
165
166
167
168
169
      pCount = MemoryPointer::new(:uint32_t)
      result = ZE.#{fname}(@handle, pCount, nil)
      ZE.error_check(result)
      count = pCount.read(:uint32)
      return [] if count == 0
      pArr = MemoryPointer::new(#{sname}, count)
      sz = #{sname}.size
      #{aname} = count.times.collect { |i| #{sname}::new(pArr.slice(sz * i, sz)) }
      result = ZE.#{fname}(@handle, pCount, #{aname}.first)
      ZE.error_check(result)
Brice Videau's avatar
WIP    
Brice Videau committed
170
171
172
EOF
      if memoize
        src << <<EOF
173
174
175
      @#{aname} = #{aname}
    end
EOF
Brice Videau's avatar
WIP    
Brice Videau committed
176
177
178
179
180
181
182
183
184
      else
        src << <<EOF
      #{aname}
    end
EOF
      end
      class_eval src
    end

185
    def self.add_enum_array_property(aname, ename, fname, memoize: false)
Brice Videau's avatar
WIP    
Brice Videau committed
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
      pname = process_ffi_name(ename)
      src = <<EOF
    def #{aname}
EOF
      if memoize
        src << <<EOF
      return @#{aname} if @#{aname}
EOF
      end
      src << <<EOF
      p_count = MemoryPointer::new(:uint32_t)
      result = ZE.#{fname}(@handle, p_count, nil)
      ZE.error_check(result)
      count = p_count.read(:uint32)
      return [] if count == 0
      p_arr = MemoryPointer::new(:#{ename}, count)
      result = ZE.#{fname}(@handle, p_count, p_arr)
      ZE.error_check(result)
EOF
      if memoize
        src << <<EOF
Brice Videau's avatar
Brice Videau committed
207
      @#{aname} = p_arr.read_array_of_int32.collect { |i| #{pname}.from_native(i, nil) }
Brice Videau's avatar
WIP    
Brice Videau committed
208
209
210
211
    end
EOF
      else
        src << <<EOF
Brice Videau's avatar
Brice Videau committed
212
      p_arr.read_array_of_int32.collect { |i| #{pname}.from_native(i, nil) }
Brice Videau's avatar
WIP    
Brice Videau committed
213
214
215
216
217
218
    end
EOF
      end
      class_eval src
    end

219
    def self.add_property(pname, sname, fname, memoize: false)
Brice Videau's avatar
WIP    
Brice Videau committed
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
      src = <<EOF
    def #{pname}
EOF
      if memoize
        src << <<EOF
      return @#{pname} if @#{pname}
EOF
      end
      src << <<EOF
      #{pname} = #{sname}::new
      result = ZE.#{fname}(@handle, #{pname})
      ZE.error_check(result)
EOF
      if memoize
        src << <<EOF
      @#{pname} = #{pname}
    end
EOF
      else
        src << <<EOF
      #{pname}
    end
EOF
      end
      class_eval src
    end
246
    def self.add_enum_property(pname, ename, fname, memoize: false)
Brice Videau's avatar
WIP    
Brice Videau committed
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
      pename = process_ffi_name(ename)
      src = <<EOF
    def #{pname}
EOF
      if memoize
        src << <<EOF
      return @#{pname} if @#{pname}
EOF
      end
      src << <<EOF
      p_prop = MemoryPointer::new(:#{ename})
      result = ZE.#{fname}(@handle, p_prop)
      ZE.error_check(result)
EOF
      if memoize
        src << <<EOF
Brice Videau's avatar
Brice Videau committed
263
      @#{pname} = #{pename}.from_native(p_prop.read_int32, nil)
Brice Videau's avatar
WIP    
Brice Videau committed
264
265
266
267
    end
EOF
      else
        src << <<EOF
Brice Videau's avatar
Brice Videau committed
268
      #{pename}.from_native(p_prop.read_int32, nil)
Brice Videau's avatar
WIP    
Brice Videau committed
269
270
271
    end
EOF
      end
Brice Videau's avatar
Brice Videau committed
272
      class_eval src
273
274
275
    end
  end

Brice Videau's avatar
WIP    
Brice Videau committed
276
  class ZEObject < Object
277
    def self.add_property(pname, sname: nil, fname: nil, memoize: false)
Brice Videau's avatar
WIP    
Brice Videau committed
278
279
280
281
282
      n = name.split("::").last
      ppn = process_property_name(pname) unless sname && fname
      sname = "ZE" << n << ppn unless sname
      fname = "ze#{n}Get" << ppn unless fname
      super(pname, sname, fname, memoize: memoize)
Brice Videau's avatar
Brice Videau committed
283
    end
Brice Videau's avatar
WIP    
Brice Videau committed
284
  end
Brice Videau's avatar
Brice Videau committed
285

Brice Videau's avatar
WIP    
Brice Videau committed
286
  class ZETObject < Object
287
    def self.add_property(pname, sname: nil, fname: nil, memoize: false)
Brice Videau's avatar
WIP    
Brice Videau committed
288
289
290
291
292
293
294
295
      n = name.split("::").last
      ppn = process_property_name(pname) unless sname && fname
      sname = "ZET" << n << ppn unless sname
      fname = "zet#{n}Get" << ppn unless fname
      super(pname, sname, fname, memoize: memoize)
    end
  end

296
297
298
299
300
301
302
303
304
305
  class ZESObject < Object
    def self.add_property(pname, sname: nil, fname: nil, memoize: false)
      n = name.split("::").last
      ppn = process_property_name(pname) unless sname && fname
      sname = "ZES" << n << ppn unless sname
      fname = "zes#{n}Get" << ppn unless fname
      super(pname, sname, fname, memoize: memoize)
    end
  end

Brice Videau's avatar
WIP    
Brice Videau committed
306
  module Lifetime
Brice Videau's avatar
Brice Videau committed
307
308
309
310
311
312
313
314
315
316
317
318
319
    def initialize(handle)
      super
      ZE.register(handle, self.class.destructor)
    end

    def destroy
      handle = @handle
      @handle = nil
      ZE.unregister(handle)
      nil
    end
  end

Brice Videau's avatar
WIP    
Brice Videau committed
320
321
322
323
324
325
326
327
328
329
330
331
332
333
  module Destructor
    attr_reader :destructor
  end

  class ZEManagedObject < ZEObject
    extend Destructor
    include Lifetime
  end

  class ZETManagedObject < ZETObject
    extend Destructor
    include Lifetime
  end

334
335
336
337
338
  class ZESManagedObject < ZESObject
    extend Destructor
    include Lifetime
  end

Brice Videau's avatar
WIP    
Brice Videau committed
339
  class Driver < ZEObject
340
341
342
    add_property :api_version, sname: :ZEApiVersion, memoize: true
    add_property :properties, memoize: true
    add_property :ipc_properties, memoize: true
Brice Videau's avatar
Brice Videau committed
343

344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
359
    def extension_properties
      @extension_properties ||= begin
        pCount = MemoryPointer::new(:uint32_t)
        result = ZE.zeDriverGetExtensionProperties(@handle, pCount, nil);
        ZE.error_check(result)
        count = pCount.read(:uint32)
        if count == 0
          []
        else
          pArr = MemoryPointer::new(:ze_driver_extension_properties_t, count)
          result = ZE.zeDriverGetExtensionProperties(@handle, pCount, pArr)
          ZE.error_check(result)
          sz = ZEDriverExtensionProperties.size
          count.times.collect { |i| ZEDriverExtensionProperties::new(pArr.slice(sz * i, sz)) } 
        end
      end
360
361
    end

Brice Videau's avatar
Brice Videau committed
362
    def devices
363
364
365
366
367
368
369
370
371
372
373
374
375
376
377
378
379
380
381
382
383
      @devices ||= begin
        pCount = MemoryPointer::new(:uint32_t)
        result = ZE.zeDeviceGet(@handle, pCount, nil)
        ZE.error_check(result)
        count = pCount.read(:uint32)
        if count == 0
          []
        else
          pArr = MemoryPointer::new(:pointer, count)
          result = ZE.zeDeviceGet(@handle, pCount, pArr)
          ZE.error_check(result)
          pArr.read_array_of_pointer(count).collect { |h| Device::new(h, self) }
        end
      end
    end

    def context_create(flags: 0)
      desc = ZEContextDesc::new()
      desc[:flags] = flags
      ph_context = MemoryPointer::new(:ze_context_handle_t)
      result = ZE.zeContextCreate(@handle, desc, ph_context)
Brice Videau's avatar
Brice Videau committed
384
      ZE.error_check(result)
385
386
387
388
389
390
391
392
393
394
395
396
397
398
399
400
401
402
403
404
405
406
407
408
409
410
411
412
413
414
415
416
417
418
419
420
421
422
423
424
425
426
427
428
429
430
431
432
433
434
435
436
437
438
439
440
441
442
443
444
445
446
447
448
449
450
451
452
453
      Context::new(ph_context.read_ze_context_handle_t, self)
    end

    def event_listen(*devices, timeout: ZET_EVENT_WAIT_NONE)
      count = devices.length
      ph_devices = MemoryPointer::new(:zes_device_handle_t, count)
      ph_devices.write_array_of_zes_device_handle_t(devices.collect(&:handle).collect(&:address))
      p_num_device_events = MemoryPointer::new(:uint32_t)
      p_events = MemoryPointer::new(:zes_event_type_flags_t, count)
      result = ZE.zesDriverEventListen(@handle, timeout, count, ph_devices, p_num_device_events, p_events)
      ZE.error_check(result)
      p_events.read_array_of_uint32(count).collect { |e| ZESEventTypeFlag.from_native(e, nil) }
    end

  end

  class Context < ZEManagedObject
    @destructor = :zeContextDestroy

    attr_reader :driver
    def initialize(handle, driver)
      super(handle)
      @driver = driver
    end

    def status
      result = ZE.zeContextGetStatus(@handle)
      ZE.error_check(result)
      result
    end

    def command_queue_create(device, ordinal: 0, index: 0, flags: 0, mode: 0, priority: 0)
      desc = ZECommandQueueDesc::new
      desc[:ordinal] = ordinal
      desc[:index] = index
      desc[:flags] = flags
      desc[:mode] = mode
      desc[:priority] = priority
      phCommandQueue = MemoryPointer::new(:ze_command_queue_handle_t)
      result = ZE.zeCommandQueueCreate(@handle, device, desc, phCommandQueue)
      ZE.error_check(result)
      CommandQueue::new(phCommandQueue.read_ze_command_queue_handle_t)
    end

    def command_list_create(device, command_queue_group_ordinal: 0, flags: 0)
      desc = ZECommandListDesc::new
      desc[:commandQueueGroupOrdinal] = command_queue_group_ordinal
      desc[:flags] = flags
      phCommandList = MemoryPointer::new(:ze_command_list_handle_t)
      result = ZE.zeCommandListCreate(@handle, device, desc, phCommandList)
      ZE.error_check(result)
      CommandList::new(phCommandList.read_ze_command_list_handle_t)
    end

    def command_list_create_immediate(device, ordinal: 0, index: 0, flags: 0, mode: 0, priority: 0)
      desc = ZECommandQueueDesc::new
      desc[:ordinal] = ordinal
      desc[:index] = index
      desc[:flags] = flags
      desc[:mode] = mode
      desc[:priority] = priority
      phCommandList = MemoryPointer::new(:ze_command_list_handle_t)
      result = ZE.zeCommandListCreateImmediate(@handle, device, desc, phCommandList)
      ZE.error_check(result)
      CommandList::new(phCommandList.read_ze_command_list_handle_t)
    end

    def system_barrier(device)
      result = ZE.zeContextSystemBarrier(@handle, device)
Brice Videau's avatar
Brice Videau committed
454
      ZE.error_check(result)
455
      self
Brice Videau's avatar
Brice Videau committed
456
457
    end

Brice Videau's avatar
Brice Videau committed
458
    def event_pool_create(count, flags: 0, devices: nil)
459
460
461
462
463
464
465
466
467
468
469
470
471
472
473
474
475
476
477
478
479
480
481
482
      desc = ZEEventPoolDesc::new()
      desc[:flags] = flags
      desc[:count] = count
      ph_devices = nil
      num_devices = 0
      if devices
        devices = [devices].flatten
        num_devices = devices.length
        ph_devices = MemoryPointer::new(:ze_device_handle_t, num_devices)
        ph_devices.write_array_of_ze_device_handle_t(devices.collect(&:to_ptr))
      end
      ph_event_pool = MemoryPointer::new(:ze_event_pool_handle_t)
      result = ZE.zeEventPoolCreate(@handle, desc, num_devices, ph_devices, ph_event_pool)
      ZE.error_check(result)
      EventPool::new(ph_event_pool.read_ze_event_pool_handle_t)
    end

    def event_pool_open_ipc_handle(h_ipc)
      ph_event_pool = MemoryPointer::new(:ze_event_pool_handle_t)
      result = ZE.zeEventPoolOpenIpcHandle(@handle, h_ipc, ph_event_pool)
      ZE.error_check(result)
      IPCEventPool::new(ph_event_pool.read_ze_event_pool_handle_t)
    end

483
484
485
486
487
488
489
490
491
492
493
494
495
496
497
    def image_create(device, width, height = 0, depth = 0,
                     arraylevels: 0,
                     miplevels: 0,
                     flags: :ZE_IMAGE_FLAG_PROGRAM_READ,
                     type: nil,
                     format_layout: :ZE_IMAGE_FORMAT_LAYOUT_8_8_8_8,
                     format_type: :ZE_IMAGE_FORMAT_TYPE_UINT,
                     format_x: :ZE_IMAGE_FORMAT_SWIZZLE_R,
                     format_y: :ZE_IMAGE_FORMAT_SWIZZLE_G,
                     format_z: :ZE_IMAGE_FORMAT_SWIZZLE_B,
                     format_w: :ZE_IMAGE_FORMAT_SWIZZLE_A)
      if !type
        type = _get_image_type(width, height, depth, arraylevels)
      end
      desc = ZEImageDesc::new
Brice Videau's avatar
Brice Videau committed
498
      desc[:flags] = flags
499
500
501
502
503
504
505
506
507
508
509
510
511
512
      desc[:type] = type
      desc[:format][:layout] = format_layout
      desc[:format][:type] = format_type
      desc[:format][:x] = format_x
      desc[:format][:y] = format_y
      desc[:format][:z] = format_z
      desc[:format][:w] = format_w
      desc[:width] = width
      desc[:height] = height
      desc[:depth] = depth
      desc[:arraylevels] = arraylevels
      desc[:miplevels] = miplevels
      phImage = MemoryPointer::new(:ze_image_handle_t)
      result = ZE.zeImageCreate(@handle, device, desc, phImage)
Brice Videau's avatar
Brice Videau committed
513
      ZE.error_check(result)
514
515
516
517
518
519
520
521
522
523
524
525
      Image::new(phImage.read_ze_image_handle_t)
    end

    # missing support for raytracing
    def mem_alloc_shared(size, device: nil, alignment: 0, ordinal: 0, host_flags: 0, device_flags: 0)
      pptr = MemoryPointer::new(:pointer)
      host_desc = ZEHostMemAllocDesc::new
      device_desc = ZEDeviceMemAllocDesc::new
      host_desc[:flags] = host_flags
      device_desc[:flags] = device_flags
      device_desc[:ordinal] = ordinal
      result = ZE.zeMemAllocShared(@handle, device_desc, host_desc, size, alignment, device, pptr)
Brice Videau's avatar
Brice Videau committed
526
527
528
      pptr.read_pointer.slice(0, size)
    end

529
530
    # missing support for file descriptor
    def mem_alloc_device(size, device, alignment: 0, ordinal: 0, flags: 0)
Brice Videau's avatar
Brice Videau committed
531
532
533
534
      pptr = MemoryPointer::new(:pointer)
      desc = ZEDeviceMemAllocDesc::new
      desc[:flags] = flags
      desc[:ordinal] = ordinal
535
      result = ZE.zeMemAllocDevice(@handle, desc, size, alignment, device, pptr)
Brice Videau's avatar
Brice Videau committed
536
537
538
539
      ZE.error_check(result)
      pptr.read_pointer.slice(0, size)
    end

540
    def mem_alloc_host(size, alignment: 0, flags: 0)
Brice Videau's avatar
Brice Videau committed
541
      pptr = MemoryPointer::new(:pointer)
542
543
544
545
      desc = ZEHostMemAllocDesc::new
      desc[:flags] = flags
      result = ZE.zeMemAllocHost(@handle, desc, size, alignment, pptr)
      ZE.error_check(result)
Brice Videau's avatar
Brice Videau committed
546
547
548
      pptr.read_pointer.slice(0, size)
    end

549
550
    def mem_free(ptr)
      result = ZE.zeMemFree(@handle, ptr)
Brice Videau's avatar
Brice Videau committed
551
552
553
      ZE.error_check(result)
      return self
    end
Brice Videau's avatar
Brice Videau committed
554

555
556
    # missing support for file descriptor
    def mem_get_alloc_properties(ptr)
557
558
      props = ZEMemoryAllocationProperties::new
      ph_device = MemoryPointer::new(:ze_device_handle_t)
559
      result = ZE.zeMemGetAllocProperties(@handle, ptr, props, ph_device)
Brice Videau's avatar
Brice Videau committed
560
      ZE.error_check(result)
561
562
563
      h_device = ph_device.read_ze_device_handle_t
      device = h_device.null? ? nil : Device::new(h_device, self)
      return [props, device]
Brice Videau's avatar
Brice Videau committed
564
    end
565
    alias mem_alloc_properties mem_get_alloc_properties
Brice Videau's avatar
Brice Videau committed
566

567
    def mem_get_address_range(ptr)
568
569
      p_base = MemoryPointer::new(:pointer)
      p_size = MemoryPointer::new(:size_t)
570
      result = ZE.zeMemGetAddressRange(@handle, ptr, p_base, p_size)
Brice Videau's avatar
Brice Videau committed
571
      ZE.error_check(result)
572
573
      return p_base.read_pointer.slice(0, p_size.read_size_t)
    end
574
    alias mem_address_range mem_get_address_range
575

576
577
578
    def mem_get_ipc_handle(ptr)
      ipc_handle = ZEIpcMemHandle.new
      result = ZE.zeMemGetIpcHandle(@handle, ptr, ipc_handle)
579
580
581
      ZE.error_check(result)
      return ipc_handle
    end
582
    alias mem_ipc_handle mem_get_ipc_handle
583

584
    def mem_open_ipc_handle(device, ipc_handle, flags: 0)
585
      pptr = MemoryPointer::new(:pointer)
586
      result = ZE.zeMemOpenIpcHandle(@handle, device, ipc_handle, flags, pptr)
587
588
      ZE.error_check(result)
      return pptr.read_pointer
Brice Videau's avatar
Brice Videau committed
589
    end
590

591
592
    def mem_close_ipc_handle(ptr)
      result = ZE.zeMemCloseIpcHandle(@handle, ptr)
593
594
595
596
      ZE.error_check(result)
      return self
    end

597
598
599
600
601
602
603
604
605
606
607
608
609
610
611
612
613
614
615
616
617
    def module_create(device, input_module,
                      format: :ZE_MODULE_FORMAT_IL_SPIRV,
                      build_flags: "",
                      profile_flags: nil,
                      constants: nil)
      desc = ZEModuleDesc::new
      input_size = input_module.bytesize
      p_input_module = MemoryPointer::new(input_size)
      p_input_module.write_bytes(input_module)
      if profile_flags
       build_flags += " -zet-profile-flags 0x#{ZETProfileFlag.to_native(profile_flags, nil).to_s(16)}"
      end
      p_build_flags = MemoryPointer.from_string(build_flags)
      desc[:format] = format
      desc[:inputSize] = input_size
      desc[:pInputModule] = p_input_module
      desc[:pBuildFlags] = p_build_flags
      desc[:pConstants] = constants
      ph_module = MemoryPointer::new(:ze_module_handle_t)
      ph_build_log = MemoryPointer::new(:ze_module_build_log_handle_t)
      result = ZE.zeModuleCreate(@handle, device, desc, ph_module, ph_build_log)
618
619
620
621
622
623
624
      begin
        ZE.error_check(result)
      rescue
        l = Module::BuildLog::new(ph_build_log.read_ze_module_build_log_handle_t)
        p l.string
        raise
      end
625
626
627
628
629
630
631
632
633
634
635
636
637
638
639
640
641
642
643
644
645
646
647
648
649
650
651
652
653
654
655
656
657
658
659
660
661
662
663
664
665
666
667
668
669
670
671
672
673
674
675
676
677
678
679
680
681
682
683
684
685
686
687
688
689
690
691
692
693
694
695
696
697
698
699
700
701
702
703
704
705
706
707
708
709
710
711
712
713
714
715
716
717
718
719
720
721
722
723
724
725
726
727
728
729
730
731
732
733
734
735
736
737
738
739
740
741
742
743
744
745
746
747
748
749
750
751
      [Module::new(ph_module.read_ze_module_handle_t, device),
       Module::BuildLog::new(ph_build_log.read_ze_module_build_log_handle_t)]
    end

    def make_memory_resident(device, ptr, size: nil)
      size = ptr.to_ptr.size unless size
      result = ZE.zeContextMakeMemoryResident(@handle, device, ptr, size)
      ZE.error_check(result)
      ptr.to_ptr.slice(0, size)
    end

    def evict_memory(device, ptr, size: nil)
      size = ptr.to_ptr.size unless size
      result = ZE.zeContextEvictMemory(@handle, device, ptr, size)
      ZE.error_check(result)
      ptr.to_ptr.slice(0, size)
    end

    def make_image_resident(device, image)
      result = ZE.zeContextMakeImageResident(@handle, device, image)
      ZE.error_check(result)
      self
    end

    def evict_image(device, image)
      result = ZE.zeContextEvictImage(@handle, device, image)
      ZE.error_check(result)
      self
    end

    def sampler_create(device, address_mode, filter_mode, normalized)
      desc = ZESamplerDesc::new
      desc[:addressMode] = address_mode
      desc[:filterMode] = filter_mode
      desc[:isNormalized] = normalized
      ptr = MemoryPointer::new(:ze_sampler_handle_t)
      result = ZE.zeSamplerCreate(@handle, device, desc, ptr)
      ZE.error_check(result)
      Sampler::new(ptr.read_ze_sampler_handle_t)
    end

    def virtual_mem_reserve(size, start: nil)
      pptr = MemoryPointer::new(:pointer)
      result = ZE.zeVirtualMemReserve(@handle, start, size, pptr)
      ZE.error_check(result)
      pptr.read_pointer.slice(0, size)
    end

    def virtual_mem_free(ptr, size: nil)
      size = ptr.to_ptr.size unless size
      result = ZE.zeVirtualMemFree(@handle, ptr, size)
      ZE.error_check(result)
      self
    end

    def virtual_mem_query_page_size(device, size)
      pagesize = MemoryPointer::new(:size_t)
      result = ZE.zeVirtualMemQueryPageSize(@handle, device, size, pagesize)
      ZE.error_check(result)
      pagesize.read_size_t
    end

    def physical_mem_create(device, size, flags: 0)
      desc = ZEPhysicalMemDesc::new
      desc[:flags] = flags
      desc[:size] = size
      ph = MemoryPointer::new(:ze_physical_mem_handle_t)
      result = ZE.zePhysicalMemCreate(@handle, device, desc, ph)
      ZE.error_check(result)
      PhysicalMem::new(ph.read_ze_physical_mem_handle_t)
    end

    def virtual_mem_map(ptr, physical_memory, access, size: nil, offset: 0)
      size = ptr.to_ptr.size unless size
      result = ZE.zeVirtualMemMap(@handle, ptr, size, physical_memory, offset, access)
      ZE.error_check(result)
      ptr.to_ptr.slice(0, size)
    end

    def virtual_mem_unmap(ptr, size: nil)
      size = ptr.to_ptr.size unless size
      result = ZE.zeVirtualMemUnmap(@handle, ptr, size)
      ZE.error_check(result)
      ptr.to_ptr.slice(0, size)
    end

    def virtual_mem_set_access_attribute(ptr, access, size: nil)
      size = ptr.to_ptr.size unless size
      result = ZE.zeVirtualMemSetAccessAttribute(@handle, ptr, size, access)
      ZE.error_check(result)
      ptr.to_ptr.slice(0, size)
    end

    def virtual_mem_get_access_attribute(ptr, size: nil)
      size = ptr.to_ptr.size unless size
      p_access = MemoryPointer::new(:ze_memory_access_attribute_t)
      out_size = MemoryPointer::new(:size_t)
      result = ZE.zeVirtualMemGetAccessAttribute(@handle, ptr, size, p_access, out_size)
      ZE.error_check(result)
      [ptr.to_ptr.slice(0, out_size.read_size_t), p_access.read_ze_memory_access_attribute_t]
    end

    def activate_metric_groups(device, *metric_groups)
      count = metric_groups.length
      ph_metric_groups = MemoryPointer::new(:zet_metric_group_handle_t, count)
      ph_metric_groups.write_array_of_zet_metric_group_handle_t(metric_group.collect(&:to_ptr))
      result = ZE.zetContextActivateMetricGroups(@handle, device, count, ph_metric_groups)
      ZE.error_check(result)
      self
    end

    def metric_streamer_open(device, metric_group, sampling_period: 1000,  notify_every_n_reports: 0, signal_event: nil)
      desc = ZETMetricStreamerDesc::new
      desc[:notifyEveryNReports] = notify_every_n_reports
      desc[:samplingPeriod] = sampling_period
      ph_metric_tracer = MemoryPointer::new(:zet_metric_tracer_handle_t)
      result = ZE.zetMetricStreamerOpen(@handle, device, metric_group, desc, signal_event, ph_metric_tracer)
      ZE.error_check(result)
      MetricStreamer::new(ph_metric_tracer.read_zet_metric_tracer_handle_t)
    end

    def metric_query_pool_create(device, metric_group, count, type: 0)
      desc = ZETMetricQueryPoolDesc::new
      desc[:type] = type
      desc[:count] = count
      ph_metric_query_pool = MemoryPointer::new(:zet_metric_query_pool_handle_t)
      result = ZE.zetMetricQueryPoolCreate(@handle, device, metric_group, desc, ph_metric_query_pool)
Brice Videau's avatar
Brice Videau committed
752
      ZE.error_check(result)
753
      MetricQueryPool::new(ph_metric_query_pool.read_zet_metric_query_pool_handle_t)
Brice Videau's avatar
Brice Videau committed
754
755
    end

756
    def tracer_exp_create(user_data: nil)
Brice Videau's avatar
Brice Videau committed
757
      user_data = @handle unless user_data
758
759
      ph_tracer = MemoryPointer::new(:zet_tracer_exp_handle_t)
      desc = ZETTracerExpDesc::new
Brice Videau's avatar
Brice Videau committed
760
      desc[:pUserData] = user_data
761
      result = ZE.zetTracerExpCreate(@handle, desc, ph_tracer)
Brice Videau's avatar
Brice Videau committed
762
      ZE.error_check(result)
763
      TracerExp::new(ph_tracer.read_zet_tracer_exp_handle_t)
Brice Videau's avatar
Brice Videau committed
764
765
    end

766
767
  end

768
769
770
771
772
773
774
775
776
777
  def self.module_dynamic_link(*modules)
    count = modules.size
    ph_modules = MemoryPointer::new(:ze_module_handle_t, count)
    ph_modules.write_array_of_ze_module_handle_t(events.collect(&:handle).collect(&:address))
    ph_build_log = MemoryPointer::new(:ze_module_build_log_handle_t)
    result = zeModuleDynamicLink(count, ph_modules, ph_build_log)
    ZE.error_check(result)
    Module::BuildLog::new(ph_build_log.read_ze_module_build_log_handle_t)
  end

Brice Videau's avatar
WIP    
Brice Videau committed
778
  class Device < ZEObject
Brice Videau's avatar
Brice Videau committed
779
780
781
782
783
784
785
    attr_reader :driver

    def initialize(handle, driver)
      super(handle)
      @driver = driver
    end

786
787
    add_property :properties, memoize: true
    add_property :compute_properties, memoize: true
788
789
    add_property :module_properties, memoize: true
    add_array_property :command_queue_group_properties, :ZECommandQueueGroupProperties, :zeDeviceGetCommandQueueGroupProperties, memoize: true
790
791
    add_array_property :memory_properties, :ZEDeviceMemoryProperties, :zeDeviceGetMemoryProperties, memoize: true
    add_property :memory_access_properties, memoize: true
792
    add_array_property :cache_properties, :ZEDeviceCacheProperties, :zeDeviceGetCacheProperties, memoize: true
793
    add_property :image_properties, memoize: true
794
795
    add_property :external_memory_properties, memoize: true
    add_property :debug_properties, sname: :ZETDeviceDebugProperties, fname: :zetDeviceGetDebugProperties, memoize: true
Brice Videau's avatar
Brice Videau committed
796
    add_object_array :metric_groups, :MetricGroup, :zetMetricGroupGet, memoize: true
797
798
799
800
801
802
803
804
805
806
807
808
809
810
811
812
813
814
815
816
    add_array_property :processes_state, :ZESProcessState, :zesDeviceProcessesGetState
    add_property :pci_properties, sname: :ZESPciProperties, fname: :zesDevicePciGetProperties
    add_property :pci_state, sname: :ZESPciState, fname: :zesDevicePciGetState
    add_array_property :pci_bars, :ZESPciBarProperties, :zesDevicePciGetBars
    add_property :pci_stats, sname: :ZESPciStats, fname: :zesDevicePciGetStats
    add_object_array :diagnostics_test_suites, :Diagnostics, :zesDeviceEnumDiagnosticTestSuites, memoize: true
    add_object_array :engine_groups, :Engine, :zesDeviceEnumEngineGroups, memoize: true
    add_object_array :fabric_ports, :FabricPort, :zesDeviceEnumFabricPorts, memoize: true
    add_object_array :fans, :Fan, :zesDeviceEnumFans, memoize: true
    add_object_array :firmwares, :Firmware, :zesDeviceEnumFirmwares, memoize: true
    add_object_array :frequency_domains, :Frequency, :zesDeviceEnumFrequencyDomains, memoize: true
    add_object_array :leds, :Led, :zesDeviceEnumLeds, memoize: true
    add_object_array :memory_modules, :Memory, :zesDeviceEnumMemoryModules, memoize: true
    add_object_array :performance_factor_domains, :PerformanceFactor, :zesDeviceEnumPerformanceFactorDomains, memoize: true
    add_object_array :power_domains, :Power, :zesDeviceEnumPowerDomains, memoize: true
    add_object_array :psus, :Psu, :zesDeviceEnumPsus, memoize: true
    add_object_array :ras_error_sets, :Ras, :zesDeviceEnumRasErrorSets, memoize: true
    add_object_array :schedulers, :Scheduler, :zesDeviceEnumSchedulers, memoize: true
    add_object_array :standby_domains, :Standby, :zesDeviceEnumStandbyDomains, memoize: true
    add_object_array :temperature_sensors, :Temperature, :zesDeviceEnumTemperatureSensors, memoize: true
817

Brice Videau's avatar
Brice Videau committed
818
819
820
821
822
823
824
825
826
827
828
829
830
    def sub_devices
      return @sub_devices if @sub_devices
      pCount = MemoryPointer::new(:uint32_t)
      result = ZE.zeDeviceGetSubDevices(@handle, pCount, nil)
      ZE.error_check(result)
      count = pCount.read(:uint32)
      return [] if count == 0
      pArr = MemoryPointer::new(:pointer, count)
      result = ZE.zeDeviceGetSubDevices(@handle, pCount, pArr)
      ZE.error_check(result)
      @sub_devices = pArr.read_array_of_pointer(count).collect { |h| Device::new(h, @driver) }
    end

831
832
833
834
835
836
837
838
839
840
841
842
843
844
    def p2p_properties(device)
      p2p_properties = ZEDeviceP2PProperties::new
      result = ZE.zeDeviceGetP2PProperties(@handle, device, p2p_properties)
      ZE.error_check(result)
      return p2p_properties
    end

    def can_access_peer?(device)
      can_access_peer = MemoryPointer::new(:ze_bool_t)
      result = ZE.zeDeviceCanAccessPeer(@handle, device, can_access_peer)
      ZE.error_check(result)
      return can_access_peer.read_ze_bool_t == 1
    end

845
846
    def status
      result = ZE.zeDeviceGetStatus(@handle)
847
      ZE.error_check(result)
848
      result
Brice Videau's avatar
Brice Videau committed
849
850
    end

851
852
853
854
855
856
857
858
859
860
861
862
863
864
865
    def image_get_properties(width, height = 0, depth = 0,
                             arraylevels: 0,
                             miplevels: 0,
                             flags: :ZE_IMAGE_FLAG_PROGRAM_READ,
                             type: nil,
                             format_layout: :ZE_IMAGE_FORMAT_LAYOUT_8_8_8_8,
                             format_type: :ZE_IMAGE_FORMAT_TYPE_UINT,
                             format_x: :ZE_IMAGE_FORMAT_SWIZZLE_R,
                             format_y: :ZE_IMAGE_FORMAT_SWIZZLE_G,
                             format_z: :ZE_IMAGE_FORMAT_SWIZZLE_B,
                             format_w: :ZE_IMAGE_FORMAT_SWIZZLE_A)
      if !type
        type = _get_image_type(width, height, depth, arraylevels)
      end
      desc = ZEImageDesc::new
Brice Videau's avatar
Brice Videau committed
866
      desc[:flags] = flags
867
868
869
870
871
872
873
874
875
876
877
878
879
880
      desc[:type] = type
      desc[:format][:layout] = format_layout
      desc[:format][:type] = format_type
      desc[:format][:x] = format_x
      desc[:format][:y] = format_y
      desc[:format][:z] = format_z
      desc[:format][:w] = format_w
      desc[:width] = width
      desc[:height] = height
      desc[:depth] = depth
      desc[:arraylevels] = arraylevels
      desc[:miplevels] = miplevels
      props = ZEImageProperties::new
      result = ZE.zeImageGetProperties(handle, desc, props)
Brice Videau's avatar
Brice Videau committed
881
      ZE.error_check(result)
882
      props
Brice Videau's avatar
Brice Videau committed
883
884
885
886
887
888
889
890
    end

    def system_barrier
      result = ZE.zeDeviceSystemBarrier(@handle)
      ZE.error_check(result)
      self
    end

891
892
893
894
895
    def debug_attach(pid)
      config = ZETDebugConfig::new
      config[:pid] = pid
      ph_debug = MemoryPointer::new(:zet_debug_session_handle_t)
      result = ZE.zetDebugAttach(@handle, config, ph_debug)
Brice Videau's avatar
Brice Videau committed
896
      ZE.error_check(result)
897
      Debug::new(ph_debug.read_zet_debug_session_handle_t)
Brice Videau's avatar
Brice Videau committed
898
899
    end

900
901
902
    def debug_get_register_set_properties
      p_count = MemoryPointer::new(:uint32)
      result = ZE.zetDebugGetRegisterSetProperties(@handle, p_count, nil)
Brice Videau's avatar
Brice Videau committed
903
      ZE.error_check(result)
904
905
906
907
      count = p_count.read_uint32
      return [] if count == 0
      p_regset_props = MemoryPointer::new(:zet_debug_regset_properties_t, count)
      result = ZE.zetDebugGetRegisterSetProperties(@handle, p_count, p_regset_props)
Brice Videau's avatar
Brice Videau committed
908
      ZE.error_check(result)
909
910
      sz = ZETDebugRegsetProperties.size
      count.times.collect { |i| ZETDebugRegsetProperties::new(p_regset_props.slice(i*sz, sz) ) }
Brice Videau's avatar
Brice Videau committed
911
912
    end

913
914
915
    def sysman_properties
      props = ZESDeviceProperties::new
      result = ZE.zesDeviceGetProperties(@handle, props)
916
      ZE.error_check(result)
917
      props
918
919
    end

920
921
922
    def state
      state = ZESDeviceState::new
      result = ZE.zesDeviceGetState(@handle, state)
Brice Videau's avatar
WIP    
Brice Videau committed
923
      ZE.error_check(result)
924
      state
Brice Videau's avatar
WIP    
Brice Videau committed
925
926
    end

927
928
    def reset(force: false)
      result = ZE.zesDeviceReset(@handle, force ? 1 : 0)
Brice Videau's avatar
Brice Videau committed
929
930
931
932
      ZE.error_check(result)
      self
    end

933
934
    def event_register(events)
      result = ZE.zesDeviceEventRegister(@handle, events)
Brice Videau's avatar
Brice Videau committed
935
      ZE.error_check(result)
936
      self
Brice Videau's avatar
Brice Videau committed
937
938
    end

Brice Videau's avatar
Brice Videau committed
939
940
941
942
943
944
945
946
947
948
949
950
951
952
953
954
955
956
957
958
959
960
961
962
    private
    def _get_image_type(width, height, depth, arraylevels)
      if arraylevels > 0
        if height > 0
          type = :ZE_IMAGE_TYPE_2DARRAY
        else
          type = :ZE_IMAGE_TYPE_1DARRAY
        end
      else
        if height > 0
          if depth > 0
            type = :ZE_IMAGE_TYPE_3D
          else
            type = :ZE_IMAGE_TYPE_2D
          end
        else
          type = :ZE_IMAGE_TYPE_1D
        end
      end
      type
    end

  end

Brice Videau's avatar
WIP    
Brice Videau committed
963
  class CommandQueue < ZEManagedObject
Brice Videau's avatar
Brice Videau committed
964
965
966
    @destructor = :zeCommandQueueDestroy

    def execute_command_lists(command_lists, fence: nil)
Brice Videau's avatar
WIP    
Brice Videau committed
967
968
969
970
971
      cl = [command_lists].flatten
      num_command_lists = cl.length
      ph_command_lists = MemoryPointer::new(:ze_command_list_handle_t, num_command_lists)
      ph_command_lists.write_array_of_ze_command_list_handle_t(cl.collect(&:handle))
      result = ZE.zeCommandQueueExecuteCommandLists(@handle, num_command_lists, ph_command_lists, fence)
Brice Videau's avatar
Brice Videau committed
972
973
974
975
      ZE.error_check(result)
      self
    end

976
    def synchronize(timeout: ZE::UINT64_MAX)
Brice Videau's avatar
Brice Videau committed
977
978
979
980
      result = ZE.zeCommandQueueSynchronize(@handle, timeout)
      ZE.error_check(result)
      result
    end
981
982
983

    def fence_create(flags: 0)
      desc = ZEFenceDesc::new
984
      desc[:flags] = flags
985
986
987
988
989
      ph_fence = MemoryPointer::new(:ze_fence_handle_t)
      result = ZE.zeFenceCreate(@handle, desc, ph_fence)
      ZE.error_check(result)
      return Fence::new(ph_fence.read_ze_fence_handle_t)
    end
Brice Videau's avatar
Brice Videau committed
990
991
  end

Brice Videau's avatar
WIP    
Brice Videau committed
992
  class CommandList < ZEManagedObject
Brice Videau's avatar
Brice Videau committed
993
994
995
996
997
998
999
1000
1001
1002
1003
1004
1005
1006
    @destructor = :zeCommandListDestroy

    def close
      result = ZE.zeCommandListClose(@handle)
      ZE.error_check(result)
      self
    end

    def reset
      result = ZE.zeCommandListReset(@handle)
      ZE.error_check(result)
      self
    end

1007
1008
1009
1010
1011
1012
1013
    def append_write_global_timestamp(dstptr, signal_event: nil, wait_events: nil)
      count, ph_wait_events = _create_event_list(wait_events)
      result = ZE.zeCommandListAppendWriteGlobalTimestamp(@handle, dstptr, count, ph_wait_events)
      ZE.error_check(result)
      self
    end

Brice Videau's avatar
Brice Videau committed
1014
    def append_barrier(signal_event: nil, wait_events: nil)
Brice Videau's avatar
Brice Videau committed
1015
1016
      count, ph_wait_events = _create_event_list(wait_events)
      result = ZE.zeCommandListAppendBarrier(@handle, signal_event, count, ph_wait_events)
Brice Videau's avatar
Brice Videau committed
1017
1018
1019
1020
1021
      ZE.error_check(result)
      self
    end

    def append_memory_ranges_barrier(*ranges, signal_event: nil, wait_events: nil)
Brice Videau's avatar
Brice Videau committed
1022
      count, ph_wait_events = _create_event_list(wait_events)
Brice Videau's avatar
Brice Videau committed
1023
1024
1025
1026
1027
1028
1029
      num_ranges = ranges.length
      if num_ranges > 0
        pRangesSizes = MemoryPointer::new(:size_t, num_ranges)
        pRanges = MemoryPointer::new(:pointer, num_ranges)
        rranges = ranges.collect { |r|
          case r
          when Array
1030
            [r.last, r.first]
Brice Videau's avatar
Brice Videau committed
1031
1032
          when Range
            [r.size, r.first]
1033
1034
          when Pointer
            [r.size, r.address]
Brice Videau's avatar
Brice Videau committed
1035
1036
1037
1038
1039
1040
1041
1042
1043
1044
1045
          else
            ZE.error_check(:ZE_RESULT_ERROR_INVALID_ARGUMENT)
          end
        }
        sizes, starts = rranges.transpose
        pRangesSizes.write_array_of_size_t(sizes)
        pRanges.write_array_of_pointer(starts)
      else
        pRanges = nil
        pRangesSizes = nil
      end
Brice Videau's avatar
Brice Videau committed
1046
      result = ZE.zeCommandListAppendMemoryRangesBarrier(@handle, num_ranges, pRangesSizes, pRanges, signal_event, count, ph_wait_events)
Brice Videau's avatar
Brice Videau committed
1047
1048
1049
1050
      ZE.error_check(result)
      self
    end

1051
    def append_memory_copy(dstptr, srcptr, size: nil, signal_event: nil, wait_events: nil, context_src: nil)
Brice Videau's avatar
Brice Videau committed
1052
      count, ph_wait_events = _create_event_list(wait_events)
1053
1054
1055
      unless size
        size = [dstptr.size, srcptr.size].min
      end
1056
1057
1058
1059
1060
      if context
        result = ZE.zeCommandListAppendMemoryCopyFromContext(@handle, dstptr, context_src, srcptr, size, signal_event, count, ph_wait_events)
      else
        result = ZE.zeCommandListAppendMemoryCopy(@handle, dstptr, srcptr, size, signal_event, count, ph_wait_events)
      end
1061
1062
1063
1064
      ZE.error_check(result)
      self
    end

1065
1066
    def append_memory_fill(ptr, pattern, size: nil, pattern_size: nil, signal_event: nil, wait_events: nil)
      count, ph_wait_events = _create_event_list(wait_events)
1067
1068
      size = ptr.size unless size
      pattern_size = pattern.size unless size
Brice Videau's avatar
Brice Videau committed
1069
      result = ZE.zeCommandListAppendMemoryFill(@handle, ptr, pattern, pattern_size, size, signal_event, count, ph_wait_events)
1070
1071
1072
1073
1074
1075
1076
1077
1078
      ZE.error_check(result)
      self
    end

    def append_memory_copy_region(dstptr, srcptr, width, height, depth = 0,
                                  dst_origin: [0, 0, 0], src_origin: [0, 0, 0],
                                  dst_pitch: width, src_pitch: width,
                                  dst_slice_pitch: depth == 0 ? 0 : dst_pitch * height,
                                  src_slice_pitch: depth == 0 ? 0 : src_pitch * height,
1079
1080
                                  signal_event: nil, wait_events: nil)
      count, ph_wait_events = _create_event_list(wait_events)
1081
1082
1083
1084
1085
1086
1087
1088
1089
1090
1091
1092
1093
1094
      dst_region = ZECopyRegion::new
      dst_region[:originX] = dst_origin[0]
      dst_region[:originY] = dst_origin[1]
      dst_region[:originZ] = dst_origin[2]
      dst_region[:width] = width
      dst_region[:height] = height
      dst_region[:depth] = depth
      src_region = ZECopyRegion::new
      src_region[:originX] = src_origin[0]
      src_region[:originY] = src_origin[1]
      src_region[:originZ] = src_origin[2]
      src_region[:width] = width
      src_region[:height] = height
      src_region[:depth] = depth
1095
      result = ZE.zeCommandListAppendMemoryCopyRegion(@handle, dstptr, dst_region, dst_pitch, dst_slice_pitch, srcptr, src_region, src_pitch, src_slice_pitch, signal_event, count, ph_wait_events)
1096
1097
1098
1099
      ZE.error_check(result)
      self
    end

1100
1101
1102
    def append_image_copy(dst_image, src_image, signal_event: nil, wait_events: nil)
      count, ph_wait_events = _create_event_list(wait_events)
      result = ZE.zeCommandListAppendImageCopy(@handle, dst_image, src_image, signal_event, count, ph_wait_events)
1103
1104
1105
1106
1107
1108
      ZE.error_check(result)
      self
    end

    def append_image_copy_region(dst_image, src_image, width, height = 1, depth = 1,
                                 dst_origin: [0, 0, 0], src_origin: [0, 0, 0],
1109
1110
                                 signal_event: nil, wait_events: nil)
      count, ph_wait_events = _create_event_list(wait_events)
1111
1112
1113
1114
1115
1116
1117
1118
1119
1120
1121
1122
1123
1124
      dst_region = ZEImageRegion::new
      dst_region[:originX] = dst_origin[0]
      dst_region[:originY] = dst_origin[1]
      dst_region[:originZ] = dst_origin[2]
      dst_region[:width] = width
      dst_region[:height] = height
      dst_region[:depth] = depth
      src_region = ZEImageRegion::new
      src_region[:originX] = src_origin[0]
      src_region[:originY] = src_origin[1]
      src_region[:originZ] = src_origin[2]
      src_region[:width] = width
      src_region[:height] = height
      src_region[:depth] = depth
1125
      result = ZE.zeCommandListAppendImageCopyRegion(@handle, dst_image, src_image, dst_region, src_region, signal_event, count, ph_wait_events)
1126
1127
1128
1129
      ZE.error_check(result)
      self
    end

1130
1131
    def append_image_copy_to_memory(dstptr, image, ranges: nil, signal_event: nil, wait_events: nil)
      count, ph_wait_events = _create_event_list(wait_events)
1132
1133
      src_region = nil
      src_region = _create_image_region_from_ranges(ranges) if ranges
1134
      result = ZE.zeCommandListAppendImageCopyToMemory(@handle, dstptr, image, src_region, signal_event, count, ph_wait_events)
1135
1136
1137
1138
      ZE.error_check(result)
      self
    end

1139
1140
    def append_image_copy_from_memory(image, srcptr, ranges: nil, signal_event: nil, wait_events: nil)
      count, ph_wait_events = _create_event_list(wait_events)
1141
1142
      dst_region = nil
      dst_region = _create_image_region_from_ranges(ranges) if ranges
1143
      result = ZE.zeCommandListAppendImageCopyFromMemory(@handle, image, srcptr, dst_region, signal_event, count, ph_wait_events)
1144
1145
1146
1147
1148
1149
1150
1151
1152
1153
1154
1155
1156
1157
1158
1159
1160
1161
      ZE.error_check(result)
      self
    end

    def append_memory_prefetch(ptr, size: nil)
      size = ptr.to_ptr.size unless size
      result = ZE.zeCommandListAppendMemoryPrefetch(@handle, ptr, size)
      ZE.error_check(result)
      self
    end

    def append_mem_advise(device, ptr, advice, size: nil)
      size = ptr.to_ptr.size unless size
      result = ZE.zeCommandListAppendMemAdvise(@handle, device, ptr, size, advice)
      ZE.error_check(result)
      self
    end

1162
1163
    def append_signal_event(event)
      result = ZE.zeCommandListAppendSignalEvent(@handle, event)
Brice Videau's avatar
Brice Videau committed
1164
1165
1166
1167
      ZE.error_check(result)
      self
    end

1168
1169
1170
    def append_wait_on_events(events)
      count, ph_events = _create_event_list(events)
      result = ZE.zeCommandListAppendWaitOnEvents(@handle, count, ph_events)
Brice Videau's avatar
Brice Videau committed
1171
1172
1173
1174
      ZE.error_check(result)
      self
    end

1175
1176
    def append_event_reset(event)
      result = ZE.zeCommandListAppendEventReset(@handle, event)
Brice Videau's avatar
Brice Videau committed
1177
1178
1179
1180
      ZE.error_check(result)
      self
    end

1181
1182
1183
1184
1185
1186
1187
1188
1189
1190
    def append_query_kernel_timestamps(*events, dstptr, offsets: nil, signal_event: nil, wait_events: nil)
      count, ph_wait_events = _create_event_list(wait_events)
      num, ph_events = _create_event_list(events)
      if offsets
        p_offsets = MemoryPointer(:size_t, count)
        p_offsets.write_array_of_size_t(offsets)
        offsets = p_offsets
      end
      result = ZE.zeCommandListAppendQueryKernelTimestamps(@handle, num, ph_events, dstptr, 
offsets, signal_event, count, ph_wait_events)
Brice Videau's avatar
Brice Videau committed
1191
1192
1193
1194
      ZE.error_check(result)
      self
    end

1195
1196
1197
1198
1199
1200
1201
1202
1203
1204
1205
1206
1207
1208
1209
1210
1211
1212
1213
1214
1215
1216
1217
1218
1219
1220
1221
1222
1223
1224
1225
1226
1227
1228
1229
1230
1231
1232
1233
1234
1235
1236
1237
1238
1239
1240
1241
1242
1243
1244
1245
1246
1247
1248
1249
1250
1251
1252
1253
1254
1255
1256
1257
1258
1259
1260
1261
1262
1263
1264
1265
1266
1267
1268
1269
1270
1271
1272
1273
1274
    def append_launch_kernel(kernel, x, y = 1, z = 1, signal_event: nil, wait_events: nil)
      count, ph_wait_events = _create_event_list(wait_events)
      group_count = ZEGroupCount::new
      group_count[:groupCountX] = x
      group_count[:groupCountY] = y
      group_count[:groupCountZ] = z
      result = ZE.zeCommandListAppendLaunchKernel(@handle, kernel, group_count, signal_event, count, ph_wait_events)
      ZE.error_check(result)
      self
    end

    def append_launch_cooperative_kernel(kernel, x, y = 1, z = 1, signal_event: nil, wait_events: nil)
      count, ph_wait_events = _create_event_list(wait_events)
      group_count = ZEGroupCount::new
      group_count[:groupCountX] = x
      group_count[:groupCountY] = y
      group_count[:groupCountZ] = z
      result = ZE.zeCommandListAppendLaunchCooperativeKernel(@handle, kernel, group_count, signal_event, count, ph_wait_events)
      ZE.error_check(result)
      self
    end

    def append_launch_kernel_indirect(kernel, group_count, signal_event: nil, wait_events: nil)
      count, ph_wait_events = _create_event_list(wait_events)
      group_count = ZEGroupCount::new(group_count.to_ptr)
      result = ZE.zeCommandListAppendLaunchKernelIndirect(@handle, kernel, group_count, signal_event, count, ph_wait_events)
      ZE.error_check(result)
      self
    end

    def append_launch_nultiple_kernels_indirect(kernels, count_buffer, group_counts, signal_event: nil, wait_events: nil)
      count, ph_wait_events = _create_event_list(wait_events)
      kernels = [kernels].flatten
      num_kernels = kernels.length
      ph_kernels = MemoryPointer::new(:ze_kernel_handle_t, num_kernels)
      ph_kernels.write_array_of_ze_kernel_handle_t(kernels.collect(&:to_ptr))
      group_counts = ZEGroupCount::new(group_counts.to_ptr)
      result = ZE.zeCommandListAppendLaunchMultipleKernelsIndirect(@handle, num_kernels, kernels, count_buffer, group_counts, signal_event, count, ph_wait_events)
      ZE.error_check(result)
      self
    end

    def append_metric_streamer_marker(metric_streamer, value)
      result = ZE.zetCommandListAppendMetricStreamerMarker(@handle, metric_streamer, value)
      ZE.error_check(result)
      self
    end

    def append_metric_query_begin(metric_query)
      result = ZE.zetCommandListAppendMetricQueryBegin(@handle, metric_query)
      ZE.error_check(result)
      self
    end

    def append_metric_query_end(metric_query, signal_event: nil, wait_events: nil)
      count, ph_wait_events = _create_event_list(wait_events)
      result = ZE.zetCommandListAppendMetricQueryEnd(@handle, metric_query, signal_event, count, ph_wait_events)
      ZE.error_check(result)
      self
    end

    def append_metric_memory_barrier
      result = ZE.zetCommandListAppendMetricMemoryBarrier(@handle)
      ZE.error_check(result)
      self
    end

    private
    def _create_event_list(wait_events)
      count = 0
      ph_wait_events = nil
      if wait_events
        events = [wait_events].flatten
        count = events.length
        if count > 0
          ph_wait_events = MemoryPointer::new(:ze_event_handle_t, count)
          ph_wait_events.write_array_of_ze_event_handle_t(events.collect(&:handle).collect(&:address))
        end
      end
      [count, ph_wait_events]
Brice Videau's avatar
Brice Videau committed
1275
    end
1276
1277
1278
1279
1280
1281
1282
1283
1284
1285
1286
1287
1288
1289
1290
1291
1292
1293
1294
1295
1296
1297
1298
1299

    def _create_image_region_from_ranges(ranges)
      ranges.collect! do r
        case r
        when Range
          [r.size, r.first]
        when Integer
          [r, 0]
        when Array
          [r.last, r.first]
        else
          ZE.error_check(:ZE_RESULT_ERROR_INVALID_ARGUMENT)
        end
      end
      ranges += [[1, 0]] * (3 - ranges.length) if ranges.length < 3
      region = ZEImageRegion::new
      region[:originX] = ranges[0][1]
      region[:originY] = ranges[1][1]
      region[:originZ] = ranges[2][1]
      region[:width] = ranges[0][0]
      region[:height] = ranges[1][0]
      region[:depth] = ranges[2][0]
      region
    end
Brice Videau's avatar
Brice Videau committed
1300
1301
  end

Brice Videau's avatar
WIP    
Brice Videau committed
1302
  class Image < ZEManagedObject
Brice Videau's avatar
Brice Videau committed
1303
1304
1305
    @destructor = :zeImageDestroy
  end

Brice Videau's avatar
WIP    
Brice Videau committed
1306
  class Module < ZEManagedObject
Brice Videau's avatar
Brice Videau committed
1307
1308
    @destructor = :zeModuleDestroy
    attr_reader :device
Brice Videau's avatar
WIP    
Brice Videau committed
1309
    class BuildLog < ZEManagedObject
Brice Videau's avatar
Brice Videau committed
1310
1311
1312
1313
1314
1315
1316
1317
1318
1319
1320
1321
1322
1323
      @destructor = :zeModuleBuildLogDestroy

      def string
        p_size = MemoryPointer::new(:size_t)
        result = ZE.zeModuleBuildLogGetString(@handle, p_size, nil)
        ZE.error_check(result)
        sz = p_size.read_size_t
        p_build_log = MemoryPointer::new(sz)
        result = ZE.zeModuleBuildLogGetString(@handle, p_size, p_build_log)
        ZE.error_check(result)
        return p_build_log.read_string(sz)
      end
      alias to_s string
    end
1324
    add_property :properties, memoize: false
Brice Videau's avatar
Brice Videau committed
1325
1326
1327
1328
1329
1330
1331
1332
1333
1334
1335
1336
1337
1338
1339
1340
1341
1342

    def initialize(handle, device)
      super(handle)
      @device = device
    end

    def native_binary
      p_size = MemoryPointer::new(:size_t)
      result = ZE.zeModuleGetNativeBinary(@handle, p_size, nil)
      ZE.error_check(result)
      sz = p_size.read_size_t
      p_binary = MemoryPointer::new(sz)
      result = ZE.zeModuleGetNativeBinary(@handle, p_size, p_binary)
      ZE.error_check(result)
      return p_binary.read_bytes(sz)
    end

    def global_pointer(name)
1343
1344
      pptr = MemoryPointer::(:pointer)
      p_size = MemoryPointer::(:size_t)
Brice Videau's avatar
Brice Videau committed
1345
      p_name = MemoryPointer::from_string(name)
1346
      result = ZE.zeModuleGetGlobalPointer(@handle, p_name, p_size, pptr)
Brice Videau's avatar
Brice Videau committed
1347
      ZE.error_check(result)
1348
      return pptr.read_pointer.slice(0, p_size.read_size_t)
Brice Videau's avatar
Brice Videau committed
1349
1350
1351
1352
1353
1354
1355
1356
1357
1358
1359
1360
1361
1362
1363
1364
1365
1366
1367
1368
1369
1370
1371
1372
1373
1374
1375
1376
1377
1378
1379
    end

    def kernel_names
      p_count = MemoryPointer::new(:uint32)
      result = ZE.zeModuleGetKernelNames(@handle, p_count, nil)
      ZE.error_check(result)
      count  = p_count.read_uint32
      p_names = MemoryPointer::new(:pointer, count)
      result = ZE.zeModuleGetKernelNames(@handle, p_count, p_names)
      ZE.error_check(result)
      p_names.read_array_of_pointer(count).collect { |ptr| ptr.read_string }
    end

    def kernel_create(name, flags: 0)
      desc = ZEKernelDesc::new
      desc[:flags] = flags
      desc[:pKernelName] = MemoryPointer::from_string(name)
      ph_kernel = MemoryPointer::new(:ze_kernel_handle_t)
      result = ZE.zeKernelCreate(@handle, desc, ph_kernel)
      ZE.error_check(result)
      Kernel::new(ph_kernel.read_ze_kernel_handle_t)
    end

    def function_pointer(name)
      pptr = MemoryPointer::new(:pointer)
      p_name = MemoryPointer::from_string(name)
      result = ZE.zeModuleGetFunctionPointer(@handle, p_name, pptr)
      ZE.error_check(result)
      return pptr.read_pointer
    end

Brice Videau's avatar
Brice Videau committed
1380
1381
1382
1383
1384
1385
1386
1387
1388
1389
    def debug_info(format: :ZET_MODULE_DEBUG_INFO_FORMAT_ELF_DWARF)
      p_size = MemoryPointer::new(:size_t)
      result = ZE.zetModuleGetDebugInfo(@handle, format, p_size, nil)
      ZE.error_check(result)
      p_debug_info = MemoryPointer::new(p_size.read_size_t)
      result = ZE.zetModuleGetDebugInfo(@handle, format, p_size, p_debug_info)
      ZE.error_check(result)
      p_debug_info.read_bytes(p_size.read_size_t)
    end

Brice Videau's avatar
Brice Videau committed
1390
1391
  end

Brice Videau's avatar
WIP    
Brice Videau committed
1392
  class Kernel < ZEManagedObject
Brice Videau's avatar
Brice Videau committed
1393
1394
1395
    @destructor = :zeKernelDestroy
    add_property :properties

1396
1397
    def set_group_size(group_size_x, group_size_y = 1, group_size_z = 1)
      result = ZE.zeKernelSetGroupSize(@handle, group_size_x, group_size_y, group_size_z)
Brice Videau's avatar
Brice Videau committed
1398
1399
1400
1401
      ZE.error_check(result)
      return self
    end

1402
    def suggest_group_size(global_size_x, global_size_y = 1, global_size_z = 1)
Brice Videau's avatar
Brice Videau committed
1403
      ptr = MemoryPointer::new(:uint32, 3)
1404
      result = ZE.zeKernelSuggestGroupSize(@handle, global_size_x, global_size_y, global_size_z, ptr, ptr + 4, ptr + 8)
Brice Videau's avatar
Brice Videau committed
1405
1406
1407
1408
1409
1410
1411
1412
1413
1414
1415
1416
1417
1418
1419
1420
1421
1422
1423
      ZE.error_check(result)
      ptr.read_array_of_uint32(3)
    end

    def suggest_max_cooperative_group_count
      ptr = MemoryPointer::new(:uint32)
      result = ZE.zeKernelSuggestMaxCooperativeGroupCount(@handle, ptr)
      ZE.error_check(result)
      ptr.read_uint32
    end

    def set_argument_value(index, ptr, size: nil)
      sz = size
      sz = ptr.to_ptr.size unless sz
      result = ZE.zeKernelSetArgumentValue(@handle, index, sz, ptr)
      ZE.error_check(result)
      self
    end

1424
1425
    def set_indirect_access(flags)
      result = ZE.zeKernelSetIndirectAccess(@handle, flags)
Brice Videau's avatar
Brice Videau committed
1426
1427
1428
      ZE.error_check(result)
      self
    end
1429
1430
1431
1432
    def indirect_access=(flags)
      set_indirect_access(flags)
      flags
    end
Brice Videau's avatar
Brice Videau committed
1433

1434
1435
1436
    def indirect_acess
      ptr = MemoryPointer::new(:ze_kernel_indirect_access_flags_t)
      result = ZE.zeKernelGetIndirectAccess(@handle, ptr)
Brice Videau's avatar
Brice Videau committed
1437
      ZE.error_check(result)
1438
1439
1440
1441
1442
1443
1444
1445
      ptr.read_ze_kernel_indirect_access_flags_t
    end

    def source_attributes
      ptr = MemoryPointer::new(:pointer)
      result = ZE.zeKernelGetSourceAttributes(@handle, 0, ptr)
      ZE.error_check(result)
      ptr.read_pointer.read_string.split(" ")
Brice Videau's avatar
Brice Videau committed
1446
1447
    end

1448
1449
    def set_cache_config(flags)
      result = ZE.zeKernelSetCacheConfig(@handle, flags)
Brice Videau's avatar
Brice Videau committed
1450
1451
1452
      ZE.error_check(result)
      self
    end
Brice Videau's avatar
Brice Videau committed
1453

1454
1455
1456
1457
1458
1459
1460
1461
1462
1463
1464
1465
1466
    def name
      @name ||= begin
        p_size = MemoryPointer::new(:size_t)
        result = zeKernelGetName(@handle, p_size, nil)
        ZE.error_check(result)
        p_name = MemoryPointer::new(:char, p_size.read_size_t)
        result = zeKernelGetName(@handle, p_size, p_name)
        ZE.error_check(result)
        p_name.read_string
      end
    end

    # This is not supported by intel and documentation makes no sense...
Brice Videau's avatar
Brice Videau committed
1467
    def profile_info
1468
      info = ZETProfileProperties::new
Brice Videau's avatar
Brice Videau committed
1469
1470
1471
1472
      result = ZE.zetKernelGetProfileInfo(@handle, info)
      ZE.error_check(result)
      info
    end
1473
1474
  end

Brice Videau's avatar
WIP    
Brice Videau committed
1475
  class EventPool < ZEManagedObject
Brice Videau's avatar
Brice Videau committed
1476
1477
1478
1479
1480
1481
1482
1483
1484
1485
    @destructor = :zeEventPoolDestroy

    def event_create(index, signal: 0, wait: 0)
      desc = ZEEventDesc::new
      desc[:index] = index
      desc[:signal] = signal
      desc[:wait] = wait
      ph_event = MemoryPointer::new(:ze_event_handle_t)
      result = ZE.zeEventCreate(@handle, desc, ph_event)
      ZE.error_check(result)
Brice Videau's avatar
Brice Videau committed
1486
      Event::new(ph_event.read_ze_event_handle_t)