opencl_model.rb 39.7 KB
Newer Older
Brice Videau's avatar
Brice Videau committed
1
2
3
require 'nokogiri'
require 'yaml'

Brice Videau's avatar
WIP.    
Brice Videau committed
4
5
6
7
8
9
if ENV["SRC_DIR"]
  SRC_DIR = ENV["SRC_DIR"]
else
  SRC_DIR = "."
end

10
11
12
13
START = "entry"
STOP = "exit"
SUFFIXES = { "start" => START, "stop" => STOP }

Brice Videau's avatar
Brice Videau committed
14
15
MEMBER_SEPARATOR = "__"

Brice Videau's avatar
Brice Videau committed
16
17
GENERATE_ENUMS_TRACEPOINTS = false

18
19
HOST_PROFILE = true

Brice Videau's avatar
Brice Videau committed
20
21
WINDOWS = /D3D|DX9/

22
VENDOR_EXT = /QCOM$|INTEL$|ARM$|APPLE$|IMG$|OCLICD$/
Brice Videau's avatar
Brice Videau committed
23
24
25

ABSENT_FUNCTIONS = /^clIcdGetPlatformIDsKHR$|^clCreateProgramWithILKHR$|^clTerminateContextKHR$|^clCreateCommandQueueWithPropertiesKHR$|^clEnqueueMigrateMemObjectEXT$/

26
27
EXTENSION_FUNCTIONS = /KHR$|EXT$|GL/

Brice Videau's avatar
WIP.    
Brice Videau committed
28
SUPPORTED_EXTENSION_FUNCTIONS = /#{YAML::load_file(File.join(SRC_DIR,"supported_extensions.yaml")).join("|")}/
29

Brice Videau's avatar
Brice Videau committed
30
31
INIT_FUNCTIONS = /clGetPlatformIDs|clGetPlatformInfo|clGetDeviceIDs|clCreateContext|clCreateContextFromType|clUnloadPlatformCompiler|clGetExtensionFunctionAddressForPlatform|clGetExtensionFunctionAddress|clGetGLContextInfoKHR/

32
LTTNG_AVAILABLE_PARAMS = 25
33
34
LTTNG_USABLE_PARAMS = LTTNG_AVAILABLE_PARAMS - 1

35
36
37
38
ENUMS = {}
ENUM_PARAM_NAME_MAP = {}
ENUM_TYPES = []

Brice Videau's avatar
Brice Videau committed
39
40
#map = Hash::new { |h, k| h[k] = [] }

Brice Videau's avatar
Brice Videau committed
41
doc = Nokogiri::XML(open("cl.xml.patched"))
Brice Videau's avatar
Brice Videau committed
42
43
44
45
46
funcs_e = doc.xpath("//commands/command").reject do |l|
  name = l.search("proto/name").text
  name.match(VENDOR_EXT) || name.match(ABSENT_FUNCTIONS) || name.match(WINDOWS)
end.collect

47
48
49
50
51
ext_funcs_e = doc.xpath("//commands/command").select do |l|
  name = l.search("proto/name").text
  name.match(SUPPORTED_EXTENSION_FUNCTIONS)
end.collect

52
53
54
55
typedef_e = doc.xpath("//types/type").select do |l|
  l["category"] == "define" && l.search("type").size > 0
end.collect

Brice Videau's avatar
Brice Videau committed
56
57
58
59
struct_e = doc.xpath("//types/type").select do |l|
  l["category"] == "struct"
end.collect

60
$constants = doc.xpath("//enums/enum").collect { |n|
61
62
63
64
65
66
67
  if n["value"]
    [n["name"], n["value"]]
  elsif n["bitpos"]
    [n["name"], "(1 << #{n["bitpos"]})"]
  end
}.to_h

Brice Videau's avatar
Brice Videau committed
68
69
CL_OBJECTS = ["cl_platform_id", "cl_device_id", "cl_context", "cl_command_queue", "cl_mem", "cl_program", "cl_kernel", "cl_event", "cl_sampler"]

Brice Videau's avatar
Brice Videau committed
70
CL_EXT_OBJECTS = ["cl_GLsync", "CLeglImageKHR", "CLeglDisplayKHR", "CLeglSyncKHR", "cl_accelerator_intel"]
Brice Videau's avatar
Brice Videau committed
71

72
CL_INT_SCALARS = ["unsigned int", "int","uintptr_t", "intptr_t", "size_t", "cl_int", "cl_uint", "cl_long", "cl_ulong", "cl_short", "cl_ushort", "cl_char", "cl_uchar"]
73
CL_FLOAT_SCALARS = ["cl_half", "cl_float", "cl_double"]
Brice Videau's avatar
Brice Videau committed
74
CL_FLOAT_SCALARS_MAP = {"cl_half" => "cl_ushort", "cl_float" => "cl_uint", "cl_double" => "cl_ulong"}
75
76
77
78
79
CL_BASE_TYPES = CL_INT_SCALARS + CL_FLOAT_SCALARS

CL_TYPE_MAP = typedef_e.collect { |l|
  [l.search("name").text, l.search("type").text]
}.to_h
Brice Videau's avatar
Brice Videau committed
80
81
82
(CL_BASE_TYPES + CL_EXT_OBJECTS + CL_OBJECTS).each { |t|
  CL_TYPE_MAP.delete(t)
}
83

Brice Videau's avatar
Brice Videau committed
84
err = false
85
CL_TYPE_MAP.transform_values! { |v|
Brice Videau's avatar
Brice Videau committed
86
87
88
  counter = 0
  until CL_BASE_TYPES.include?(v) || counter > 10
    counter += 1
89
90
    v = CL_TYPE_MAP[v]
  end
Brice Videau's avatar
Brice Videau committed
91
  err = true if counter > 10
92
93
  v
}
Brice Videau's avatar
Brice Videau committed
94
95
96
97
98
99
if err
  CL_TYPE_MAP.each { |k, v|
    $stderr.puts "#{k}" unless v
  }
  raise "Failed to achieve transitive closure!"
end
100

Brice Videau's avatar
Brice Videau committed
101
FFI_BASE_TYPES = ["ffi_type_int", "ffi_type_uint", "ffi_type_uint8", "ffi_type_sint8", "ffi_type_uint16", "ffi_type_sint16", "ffi_type_uint32", "ffi_type_sint32", "ffi_type_uint64", "ffi_type_sint64", "ffi_type_float", "ffi_type_double", "ffi_type_void", "ffi_type_pointer"]
102
FFI_TYPE_MAP =  {
Brice Videau's avatar
Brice Videau committed
103
104
 "int" => "ffi_type_int",
 "unsigned int" => "ffi_type_uint",
105
106
107
108
109
110
111
112
113
114
115
 "uint8_t" => "ffi_type_uint8",
 "int8_t" => "ffi_type_sint8",
 "uint16_t" => "ffi_type_uint16",
 "int16_t" => "ffi_type_sint16",
 "uint32_t" => "ffi_type_uint32",
 "int32_t" => "ffi_type_sint32",
 "uint64_t" => "ffi_type_uint64",
 "int64_t" => "ffi_type_sint64",
 "float" => "ffi_type_float",
 "double" => "ffi_type_double",
 "intptr_t" => "ffi_type_pointer",
116
 "uintptr_t" => "ffi_type_pointer",
117
118
119
120
121
122
123
124
125
126
127
128
129
130
 "size_t" => "ffi_type_pointer",
 "cl_double" => "double",
 "cl_float" => "float",
 "cl_char" => "int8_t",
 "cl_uchar" => "uint8_t",
 "cl_short" => "int16_t",
 "cl_ushort" => "uint16_t",
 "cl_int" => "int32_t",
 "cl_uint" => "uint32_t",
 "cl_long" => "int64_t",
 "cl_ulong" => "uint64_t",
 "cl_half" => "uint8_t"
}

Brice Videau's avatar
Brice Videau committed
131
FFI_TYPE_MAP.merge! CL_TYPE_MAP
132
133
134

FFI_TYPE_MAP.transform_values! { |v|
  until FFI_BASE_TYPES.include? v
Brice Videau's avatar
Brice Videau committed
135
    ov = v
136
    v = FFI_TYPE_MAP[v]
Brice Videau's avatar
Brice Videau committed
137
138
    $stderr.puts ov unless v
    $stderr.puts FFI_BASE_TYPES unless v
139
140
141
142
143
    exit unless v
  end
  v
}

Brice Videau's avatar
Brice Videau committed
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
class CLXML

  attr_reader :__node

  def initialize(node)
    @__node = node
  end

  def inspect
    str = "#<#{self.class}:#{(object_id << 1).to_s(16)} "
    str << instance_variables.reject { |v| v == :@__node }.collect { |v| "#{v.to_s}=#{instance_variable_get(v).inspect}" }.join(", ")
    str << ">"
    str
  end

end

161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
class Require < CLXML
  attr_reader :comment

  def initialize(node)
    super
    @comment = node["comment"]
  end

  def bitfield?
    @comment.match("bitfield")
  end

  def enums
    @__node.search("enum").collect { |e| e["name"] }
  end

end

179
180
$requires = (doc.xpath("//feature/require").to_a + doc.xpath("//extensions/extension/require").to_a).collect { |r| Require::new(r) }

Brice Videau's avatar
Brice Videau committed
181
if GENERATE_ENUMS_TRACEPOINTS
Brice Videau's avatar
WIP.    
Brice Videau committed
182
  enums = YAML::load_file(File.join(SRC_DIR,"supported_enums.yaml"))
Brice Videau's avatar
Brice Videau committed
183
184
185


  enums.each { |e|
186
    vals = $requires.select { |r|
Brice Videau's avatar
Brice Videau committed
187
188
189
190
      r.comment && r.comment.match(/#{e["name"]}(\z| )/)
    }.collect { |r|
      r.enums
    }.reduce(:+).collect { |v|
191
     [v, $constants[v]]
Brice Videau's avatar
Brice Videau committed
192
193
194
195
196
197
198
    }.to_h
    ENUMS[e["name"]] = { "values" => vals, "trace_name" => e["trace_name"], "type_name" => e["type_name"] }
    ENUM_PARAM_NAME_MAP[e["trace_name"]] = e["type_name"]
    ENUM_TYPES.push(e["type_name"] ? e["type_name"] : e["name"])
  }
  ENUM_TYPES.push "cl_bool"
end
199

Brice Videau's avatar
Brice Videau committed
200
class Declaration < CLXML
Brice Videau's avatar
Brice Videau committed
201
202
203
204
205
206
207
  attr_reader :type
  attr_reader :name

  def initialize(param)
    super
    @name = param.search("name").text
    @type = param.search("type").text
208
    @type += "*" if decl.match?(/\*\*/)
Brice Videau's avatar
Brice Videau committed
209
    @type += "*" if decl.match?(/\[\]/)
Brice Videau's avatar
Brice Videau committed
210
211
212
213
214
215
216
217
    @__callback = nil
  end

  def decl
    @__node.children.collect(&:text).join(" ").squeeze(" ")
  end

  def decl_pointer
Brice Videau's avatar
Brice Videau committed
218
    @__node.children.collect { |n| "#{n.name == "name" ? "" : n.text}" }.join(" ").squeeze(" ").strip
Brice Videau's avatar
Brice Videau committed
219
220
221
222
  end

  def pointer?
    @__pointer if !@__pointer.nil?
Brice Videau's avatar
Brice Videau committed
223
    @__pointer = false
Brice Videau's avatar
Brice Videau committed
224
225
    @__node.children.collect { |n|
      break if n.name == "name"
Brice Videau's avatar
Brice Videau committed
226
227
      if n.text.match("\\*")
        @__pointer = true
Brice Videau's avatar
Brice Videau committed
228
229
        break
      end
Brice Videau's avatar
Brice Videau committed
230
231
232
233
    }
    @__pointer
  end

Brice Videau's avatar
Brice Videau committed
234
235
236
end

class Member < Declaration
237
  def initialize(command, member, prefix, dir = "start")
Brice Videau's avatar
Brice Videau committed
238
    super(member)
Brice Videau's avatar
Brice Videau committed
239
    name = "#{prefix}#{MEMBER_SEPARATOR}#{@name}"
Brice Videau's avatar
Brice Videau committed
240
    expr = "#{prefix} != NULL ? #{prefix}->#{@name} : 0"
Brice Videau's avatar
Brice Videau committed
241
    @dir = dir
242
    @lttng_type = ["ctf_integer_hex", "uintptr_t", name, "(uintptr_t)(#{expr})"] if pointer?
Brice Videau's avatar
Brice Videau committed
243
244
245
246
    t = @type
    t = CL_TYPE_MAP[@type] if CL_TYPE_MAP[@type]
    case t
    when *CL_OBJECTS, *CL_EXT_OBJECTS
247
      @lttng_type = ["ctf_integer_hex", "uintptr_t", name, "(uintptr_t)(#{expr})"]
Brice Videau's avatar
Brice Videau committed
248
    when *CL_INT_SCALARS
249
      @lttng_type = ["ctf_integer", t, name, expr]
Brice Videau's avatar
Brice Videau committed
250
    when *CL_FLOAT_SCALARS
251
      @lttng_type = ["ctf_float", t, name, expr]
Brice Videau's avatar
Brice Videau committed
252
253
254
255
    end
   end

   def lttng_in_type
256
     @dir == "start" ? @lttng_type : nil
Brice Videau's avatar
Brice Videau committed
257
258
259
   end

   def lttng_out_type
260
     @dir == "start" ? nil : @lttng_type
Brice Videau's avatar
Brice Videau committed
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
   end

end

CL_STRUCT_MAP = struct_e.collect { |s|
  members = s.search("member")
  [s["name"], members]
}.to_h

CL_STRUCTS = CL_STRUCT_MAP.keys

class Parameter < Declaration

  def initialize(param)
    super
    @__callback = nil
  end

  def callback?
    @__callback if !@__callback.nil?
    @__callback = false
    @__node.children.collect { |n| @__callback = true if n.text.match("CL_CALLBACK") }
    @__callback
  end

  def pointer?
    return true if callback?
    super
  end

Brice Videau's avatar
Brice Videau committed
291
292
  def lttng_in_type
    if pointer?
293
      return ["ctf_integer_hex", "uintptr_t", @name, "(uintptr_t)#{@name}"]
Brice Videau's avatar
Brice Videau committed
294
    end
295
296
    t = @type
    t = CL_TYPE_MAP[@type] if CL_TYPE_MAP[@type]
297
    if ENUM_TYPES.include? @type
298
      return ["ctf_enum", "lttng_ust_opencl", @type, t, @name, @name]
299
300
301
    else
      case t
      when *CL_OBJECTS, *CL_EXT_OBJECTS
302
        return ["ctf_integer_hex", "uintptr_t", @name, "(uintptr_t)#{@name}"]
303
      when *CL_INT_SCALARS
304
        return ["ctf_integer", t, @name, @name]
305
      when *CL_FLOAT_SCALARS
306
        return ["ctf_float", t, @name, @name]
307
      end
Brice Videau's avatar
Brice Videau committed
308
309
310
311
    end
    nil
  end

312
313
314
315
  def void?
    decl.strip == "void"
  end

Brice Videau's avatar
Brice Videau committed
316
  def lttng_out_type
Brice Videau's avatar
Brice Videau committed
317
    nil
Brice Videau's avatar
Brice Videau committed
318
319
  end

320
321
322
323
324
325
  def ffi_type
    return "ffi_type_pointer" if pointer? || CL_OBJECTS.include?(type) || CL_EXT_OBJECTS.include?(type)
    return "ffi_type_void" if void?
    return FFI_TYPE_MAP[type]
  end

Brice Videau's avatar
Brice Videau committed
326
327
328
329
330
331
332
333
334
335
336
end

class Prototype < CLXML

  attr_reader :return_type
  attr_reader :name

  def has_return_type?
    return_type != "void"
  end

337
338
339
340
341
342
  def ffi_return_type
    return "ffi_type_void" unless has_return_type?
    return "ffi_type_pointer" if return_type.match(/\*/) || CL_OBJECTS.include?(return_type) || CL_EXT_OBJECTS.include?(return_type)
    FFI_TYPE_MAP[return_type]
  end

Brice Videau's avatar
Brice Videau committed
343
344
345
346
347
348
349
350
351
352
  def initialize(proto)
    super
    @name = proto.search("name").text
    @return_type = @__node.children.reject { |c| c.name == "name" }.collect(&:text).join(" ").squeeze(" ").strip
  end

  def decl
    @__node.children.collect { |n| "#{n.name == "name" ? "CL_API_CALL " : ""}#{n.text}" }.join(" ").squeeze(" ")
  end

Brice Videau's avatar
Brice Videau committed
353
354
  def decl_pointer(type: false)
    @__node.children.collect { |n| "#{n.name == "name" ? "(CL_API_CALL *#{type ? pointer_type_name : pointer_name})" : n.text}" }.join(" ").squeeze(" ")
Brice Videau's avatar
Brice Videau committed
355
356
357
358
359
360
  end

  def pointer_name
    @name + "_ptr"
  end

361
362
363
364
  def ffi_function_name
    @name + "_ffi"
  end

Brice Videau's avatar
Brice Videau committed
365
366
367
368
  def pointer_type_name
    @name + "_t"
  end

Brice Videau's avatar
Brice Videau committed
369
370
  def lttng_return_type
    if @return_type.match("\\*")
371
      return ["ctf_integer_hex", "uintptr_t", "_retval", "(uintptr_t)_retval"]
Brice Videau's avatar
Brice Videau committed
372
373
374
    end
    case @return_type
    when "cl_int"
Brice Videau's avatar
Brice Videau committed
375
      if GENERATE_ENUMS_TRACEPOINTS
376
        return ["ctf_enum", "lttng_ust_opencl", "cl_errcode", "cl_int", "errcode_ret_val", "_retval"]
Brice Videau's avatar
Brice Videau committed
377
      else
378
        return ["ctf_integer", "cl_int", "errcode_ret_val", "_retval"]
Brice Videau's avatar
Brice Videau committed
379
      end
Brice Videau's avatar
Brice Videau committed
380
    when *CL_OBJECTS
381
      return ["ctf_integer_hex", "uintptr_t", @return_type.gsub(/^cl_/,""), "(uintptr_t)_retval"]
Brice Videau's avatar
Brice Videau committed
382
    when *CL_EXT_OBJECTS
383
      return ["ctf_integer_hex", "uintptr_t", @return_type.gsub(/^CL/,"").gsub(/KHR$/,""), "(uintptr_t)_retval"]
Brice Videau's avatar
Brice Videau committed
384
    when "void*"
385
      return ["ctf_integer_hex", "uintptr_t", "ret_ptr", "(uintptr_t)_retval"]
Brice Videau's avatar
Brice Videau committed
386
387
388
389
    end
    nil
  end

Brice Videau's avatar
Brice Videau committed
390
391
end

Brice Videau's avatar
Brice Videau committed
392
class MetaParameter
Brice Videau's avatar
Brice Videau committed
393
394
395
396
397
  def initialize(command, name)
    @command = command
    @name = name
  end

398
399
400
401
402
403
404
405
406
407
  def lttng_array_type_broker(type, name, size, stype = nil)
    type = CL_TYPE_MAP[type] if CL_TYPE_MAP[type]
    if stype
      stype = CL_TYPE_MAP[stype] if CL_TYPE_MAP[stype]
      lttng_arr_type = "sequence"
      lttng_args = [ stype, "#{name} == NULL ? 0 : #{size}" ]
    else
      lttng_arr_type = "array"
      lttng_args = [ size ]
    end
Brice Videau's avatar
Brice Videau committed
408
    expr = name
409
410
    case type
    when *CL_OBJECTS, *CL_EXT_OBJECTS
411
      lttng_type = ["ctf_#{lttng_arr_type}_hex", "uintptr_t"]
412
413
414
415
416
    when *CL_INT_SCALARS
      lttng_type = ["ctf_#{lttng_arr_type}", type]
    when *CL_FLOAT_SCALARS
      lttng_type = ["ctf_#{lttng_arr_type}_hex", CL_FLOAT_SCALARS_MAP[type]]
    when *CL_STRUCTS
417
      lttng_type = ["ctf_#{lttng_arr_type}_text", "uint8_t"]
Brice Videau's avatar
Brice Videau committed
418
    when "void"
419
      lttng_type = ["ctf_#{lttng_arr_type}_text", "uint8_t"]
420
    when /\*/
421
      lttng_type = ["ctf_#{lttng_arr_type}_hex", "uintptr_t"]
422
    else
Brice Videau's avatar
Brice Videau committed
423
      raise "Unknown Type: #{type.inspect} for #{name} in #{@command.prototype.name}!"
424
    end
Brice Videau's avatar
Brice Videau committed
425
    lttng_type += [ name+"_vals", expr ]
426
427
428
    lttng_type += lttng_args
  end

429
  def lttng_in_type
Brice Videau's avatar
Brice Videau committed
430
431
432
    nil
  end

433
  def lttng_out_type
Brice Videau's avatar
Brice Videau committed
434
435
    nil
  end
436
end
Brice Videau's avatar
Brice Videau committed
437

Brice Videau's avatar
Brice Videau committed
438
439
440
441
442
443
444
445
446
447
448
449
450
class OutMetaParameter < MetaParameter
  def lttng_out_type
    @lttng_out_type
  end
end

class InMetaParameter < MetaParameter
  def lttng_in_type
    @lttng_in_type
  end
end

class OutScalar < OutMetaParameter
451
  def initialize(command, name)
Brice Videau's avatar
Brice Videau committed
452
    super
Brice Videau's avatar
Brice Videau committed
453
    raise "Couldn't find variable #{name} for #{command.prototype.name}!" unless command[name]
454
    type = command[name].type.gsub("*", "")
Brice Videau's avatar
Brice Videau committed
455
    type = CL_TYPE_MAP[type] if CL_TYPE_MAP[type]
456
    if ENUM_PARAM_NAME_MAP[name]
457
      @lttng_out_type = ["ctf_enum", "lttng_ust_opencl", ENUM_PARAM_NAME_MAP[name], type, name+"_val", "#{name} == NULL ? 0 : *#{name}"]
Brice Videau's avatar
Brice Videau committed
458
    else
459
460
      case type
      when *CL_OBJECTS, *CL_EXT_OBJECTS
461
        @lttng_out_type = ["ctf_integer_hex", "uintptr_t", name+"_val", "(uintptr_t)(#{name} == NULL ? 0 : *#{name})"]
462
      when *CL_INT_SCALARS
463
        @lttng_out_type = ["ctf_integer", type, name+"_val", "#{name} == NULL ? 0 : *#{name}"]
464
      when *CL_FLOAT_SCALARS
465
        @lttng_out_type = ["ctf_float", type, name+"_val", "#{name} == NULL ? 0 : *#{name}"]
Brice Videau's avatar
Brice Videau committed
466
      when "void"
467
        @lttng_out_type = ["ctf_integer_hex", "uintptr_t", name+"_val", "(uintptr_t)(#{name} == NULL ? 0 : *#{name})"]
468
469
470
      else
        raise "Unknown Type: #{type.inspect}!"
      end
Brice Videau's avatar
Brice Videau committed
471
472
473
474
    end
  end
end

Brice Videau's avatar
Brice Videau committed
475
476
477
class InFixedArray  < InMetaParameter
  def initialize(command, name, count)
    super(command, name)
Brice Videau's avatar
Brice Videau committed
478
    raise "Couldn't find variable #{name} for #{command.prototype.name}!" unless command[name]
Brice Videau's avatar
Brice Videau committed
479
    type = command[name].type
480
    @lttng_in_type = lttng_array_type_broker(type, name, count)
Brice Videau's avatar
Brice Videau committed
481
482
483
  end
end

Brice Videau's avatar
Brice Videau committed
484
class OutArray < OutMetaParameter
485
  def initialize(command, name, sname = "num_entries")
Brice Videau's avatar
Brice Videau committed
486
    super(command, name)
Brice Videau's avatar
Brice Videau committed
487
    @sname = sname
Brice Videau's avatar
Brice Videau committed
488
    raise "Couldn't find variable #{name} for #{command.prototype.name}!" unless command[name]
489
    type = command[name].type
Brice Videau's avatar
Brice Videau committed
490
    raise "Couldn't find variable #{sname} for #{command.prototype.name}!" unless command[sname]
491
    stype = command[sname].type
492
    @lttng_out_type = lttng_array_type_broker(type, name, sname, stype)
Brice Videau's avatar
Brice Videau committed
493
494
495
496
  end
end

class InArray < InMetaParameter
497
  def initialize(command, name, sname = "num_entries")
Brice Videau's avatar
Brice Videau committed
498
    super(command, name)
Brice Videau's avatar
Brice Videau committed
499
    @sname = sname
Brice Videau's avatar
Brice Videau committed
500
    raise "Couldn't find variable #{name} for #{command.prototype.name}!" unless command[name]
501
    type = command[name].type
Brice Videau's avatar
Brice Videau committed
502
    raise "Couldn't find variable #{sname} for #{command.prototype.name}!" unless command[sname]
503
    stype = command[sname].type
504
    @lttng_in_type = lttng_array_type_broker(type, name, sname, stype)
Brice Videau's avatar
Brice Videau committed
505
506
507
  end
end

508
class DeviceFissionPropertyList < InArray
509
510
  def initialize(command, name)
    sname = "_#{name}_size"
Brice Videau's avatar
Brice Videau committed
511
    type = command[name].type
512
513
514
    command.tracepoint_parameters.push TracepointParameter::new(sname, "size_t", <<EOF)
  #{sname} = 0;
  if(#{name} != NULL) {
515
516
517
518
519
520
521
522
523
524
    while(#{name}[#{sname}++] != CL_PROPERTIES_LIST_END_EXT) {
      switch(#{name}[#{sname}]) {
      case CL_DEVICE_PARTITION_EQUALLY_EXT:
      case CL_DEVICE_PARTITION_BY_AFFINITY_DOMAIN_EXT:
        #{sname}++; //value
        break;
      case CL_DEVICE_PARTITION_BY_COUNTS_EXT:
        while(#{name}[#{sname}++] != CL_PARTITION_BY_COUNTS_LIST_END_EXT);
        break;
      case CL_DEVICE_PARTITION_BY_NAMES_EXT:
Brice Videau's avatar
Brice Videau committed
525
        while(#{name}[#{sname}] != (#{type})CL_PARTITION_BY_NAMES_LIST_END_EXT);
526
527
        break;
      }
528
    }
529
530
531
532
533
534
  }
EOF
    super(command, name, sname)
  end
end

535
# NULL terminated Key Value pairs
536
537
538
539
540
541
class InNullArray < InArray
  def initialize(command, name)
    sname = "_#{name}_size"
    command.tracepoint_parameters.push TracepointParameter::new(sname, "size_t", <<EOF)
  #{sname} = 0;
  if(#{name} != NULL) {
542
543
544
545
    while(#{name}[#{sname}] != 0) {
      #{sname} += 2;
    }
    #{sname} ++;
546
547
548
549
550
551
  }
EOF
    super(command, name, sname)
  end
end

Brice Videau's avatar
Brice Videau committed
552
553
554
class InString < InMetaParameter
  def initialize(command, name)
    super
555
    @lttng_in_type = ["ctf_string", name+"_val", name]
Brice Videau's avatar
Brice Videau committed
556
557
558
  end
end

Brice Videau's avatar
Brice Videau committed
559
class AutoMetaParameter
560
  def self.create_if_match(command)
Brice Videau's avatar
Brice Videau committed
561
562
563
564
    nil
  end
end

565
class EventWaitList < AutoMetaParameter
Brice Videau's avatar
Brice Videau committed
566
  def self.create_if_match(command)
Brice Videau's avatar
Brice Videau committed
567
    el = command.parameters.find { |p| p.name == "event_wait_list" }
Brice Videau's avatar
Brice Videau committed
568
    if el
569
      return InArray::new(command, "event_wait_list", "num_events_in_wait_list")
Brice Videau's avatar
Brice Videau committed
570
571
572
573
574
    end
    nil
  end
end

575
576
577
578
579
580
581
582
583
584
585
class AutoOutScalar
  def self.create(name)
    str = <<EOF
    Class::new(AutoMetaParameter) do
      def self.create_if_match(command)
        par = command.parameters.find { |p| p.name == "#{name}" && p.pointer? }
        if par
          return OutScalar::new(command, "#{name}")
        end
        nil
      end
Brice Videau's avatar
Brice Videau committed
586
    end
587
588
EOF
    eval str
Brice Videau's avatar
Brice Videau committed
589
590
591
  end
end

Brice Videau's avatar
Brice Videau committed
592
class ParamValue < AutoMetaParameter
Brice Videau's avatar
Brice Videau committed
593
  def self.create_if_match(command)
594
    return nil if command.prototype.name == "clSetKernelExecInfo"
Brice Videau's avatar
Brice Videau committed
595
    pv = command.parameters.find { |p| p.name == "param_value" }
596
597
    if pv
      return OutArray::new(command, "param_value", "param_value_size")
Brice Videau's avatar
Brice Videau committed
598
599
600
    end
    nil
  end
Brice Videau's avatar
Brice Videau committed
601
end
Brice Videau's avatar
Brice Videau committed
602

603
604
605
606
607
608
609
610
611
612
613
614
615
616
617
618
619
620
class TracepointParameter
  attr_reader :name
  attr_reader :type
  attr_reader :init

  def initialize(name, type, init)
    @name = name
    @type = type
    @init = init
  end
end

ErrCodeRet = AutoOutScalar::create("errcode_ret")

ParamValueSizeRet = AutoOutScalar::create("param_value_size_ret")

Event = AutoOutScalar::create("event")

Brice Videau's avatar
Brice Videau committed
621
def register_meta_parameter(method, type, *args)
Brice Videau's avatar
Brice Videau committed
622
  raise "Unknown method: #{method}!" unless OPENCL_COMMAND_NAMES.include?(method) || OPENCL_EXTENSION_COMMAND_NAMES.include?(method)
Brice Videau's avatar
Brice Videau committed
623
624
625
626
  META_PARAMETERS[method].push [type, args]
end

def register_meta_struct(method, name, type)
Brice Videau's avatar
Brice Videau committed
627
  raise "Unknown method: #{method}!" unless OPENCL_COMMAND_NAMES.include?(method) || OPENCL_EXTENSION_COMMAND_NAMES.include?(method)
Brice Videau's avatar
Brice Videau committed
628
629
630
631
  raise "Unknown struct: #{type}!" unless CL_STRUCTS.include?(type)
  CL_STRUCT_MAP[type].each { |m|
    META_PARAMETERS[method].push [Member, [m, name]]
  }
Brice Videau's avatar
Brice Videau committed
632
633
end

Brice Videau's avatar
Brice Videau committed
634

Brice Videau's avatar
Brice Videau committed
635
def register_prologue(method, code)
Brice Videau's avatar
Brice Videau committed
636
  raise "Unknown method: #{method}!" unless OPENCL_COMMAND_NAMES.include?(method) || OPENCL_EXTENSION_COMMAND_NAMES.include?(method)
Brice Videau's avatar
Brice Videau committed
637
  PROLOGUES[method].push(code)
Brice Videau's avatar
Brice Videau committed
638
639
640
end

def register_epilogue(method, code)
Brice Videau's avatar
Brice Videau committed
641
  raise "Unknown method: #{method}!" unless OPENCL_COMMAND_NAMES.include?(method) || OPENCL_EXTENSION_COMMAND_NAMES.include?(method)
Brice Videau's avatar
Brice Videau committed
642
643
644
  EPILOGUES[method].push(code)
end

Brice Videau's avatar
Brice Videau committed
645
AUTO_META_PARAMETERS = [EventWaitList, ErrCodeRet, ParamValueSizeRet, ParamValue, Event]
Brice Videau's avatar
Brice Videau committed
646
META_PARAMETERS = Hash::new { |h, k| h[k] = [] }
Brice Videau's avatar
Brice Videau committed
647
PROLOGUES = Hash::new { |h, k| h[k] = [] }
Brice Videau's avatar
Brice Videau committed
648
EPILOGUES = Hash::new { |h, k| h[k] = [] }
Brice Videau's avatar
Brice Videau committed
649

Brice Videau's avatar
Brice Videau committed
650
651
652
653
class Command < CLXML

  attr_reader :prototype
  attr_reader :parameters
654
  attr_reader :tracepoint_parameters
Brice Videau's avatar
Brice Videau committed
655
  attr_reader :meta_parameters
Brice Videau's avatar
Brice Videau committed
656
  attr_reader :prologues
Brice Videau's avatar
Brice Videau committed
657
  attr_reader :epilogues
Brice Videau's avatar
Brice Videau committed
658
659
660
661
662

  def initialize( command )
    super
    @prototype = Prototype::new( command.search("proto" ) )
    @parameters = command.search("param").collect { |p| Parameter::new(p) }
663
    @tracepoint_parameters = []
664
    @meta_parameters = AUTO_META_PARAMETERS.collect { |klass| klass.create_if_match(self) }.compact
665
666
667
    @meta_parameters += META_PARAMETERS[@prototype.name].collect { |type, args|
      type::new(self, *args)
    }
668
    @extension = @prototype.name.match(EXTENSION_FUNCTIONS)
Brice Videau's avatar
Brice Videau committed
669
    @init      = @prototype.name.match(INIT_FUNCTIONS)
Brice Videau's avatar
Brice Videau committed
670
    @prologues = PROLOGUES[@prototype.name]
Brice Videau's avatar
Brice Videau committed
671
    @epilogues = EPILOGUES[@prototype.name]
672
673
674
675
676
677
  end

  def [](name)
    res = @parameters.find { |p| p.name == name }
    return res if res
    @tracepoint_parameters.find { |p| p.name == name }
Brice Videau's avatar
Brice Videau committed
678
679
680
681
682
683
  end

  def decl
    "CL_API_ENTRY " + @prototype.decl + "(" + @parameters.collect(&:decl).join(", ") + ")"
  end

Brice Videau's avatar
Brice Videau committed
684
685
  def decl_pointer(type: false)
    "CL_API_ENTRY " + @prototype.decl_pointer(type: type) + "(" + @parameters.collect(&:decl_pointer).join(", ") + ")"
Brice Videau's avatar
Brice Videau committed
686
687
  end

688
  def decl_ffi_wrapper
689
    "void #{@prototype.ffi_function_name}(ffi_cif *cif, #{@prototype.return_type} *ffi_ret, void** args, #{@prototype.pointer_type_name} #{@prototype.pointer_name})"
690
691
  end

Brice Videau's avatar
Brice Videau committed
692
693
694
695
696
697
698
699
  def event?
    returns_event? || @parameters.find { |p| p.name == "event" && p.pointer? }
  end

  def returns_event?
    prototype.return_type == "cl_event"
  end

700
701
702
703
  def extension?
    return !!@extension
  end

Brice Videau's avatar
Brice Videau committed
704
705
706
707
  def init?
    return !!@init
  end

708
709
710
711
  def void_parameters?
    @parameters.size == 1 && @parameters.first.void?
  end

Brice Videau's avatar
Brice Videau committed
712
713
end

Brice Videau's avatar
Brice Videau committed
714
715
OPENCL_COMMAND_NAMES = funcs_e.collect { |c| Prototype::new( c.search("proto" ) ) }.collect { |p| p.name }
OPENCL_EXTENSION_COMMAND_NAMES = ext_funcs_e.collect { |c| Prototype::new( c.search("proto" ) ) }.collect { |p| p.name }
Brice Videau's avatar
Brice Videau committed
716

Brice Videau's avatar
WIP.    
Brice Videau committed
717
$meta_parameters = YAML::load_file(File.join(SRC_DIR,"opencl_meta_parameters.yaml"))
Brice Videau's avatar
Brice Videau committed
718
$meta_parameters["meta_parameters"].each  { |func, list|
Brice Videau's avatar
Brice Videau committed
719
720
721
722
  list.each { |type, *args|
    register_meta_parameter func, Kernel.const_get(type), *args
  }
}
Brice Videau's avatar
Brice Videau committed
723

Brice Videau's avatar
Brice Videau committed
724
$meta_parameters["meta_structs"].each { |func, list|
725
726
727
728
  list.each { |args|
    register_meta_struct func, *args
  }
}
Brice Videau's avatar
Brice Videau committed
729

Brice Videau's avatar
Brice Videau committed
730
731
732
733
734
735
736
737
$opencl_commands = funcs_e.collect { |func|
  Command::new(func)
}

$opencl_extension_commands = ext_funcs_e.collect { |func|
  Command::new(func)
}

738
739
740
741
742
743
744
745
$opencl_commands.each { |c|
  eval "$#{c.prototype.name} = c"
}

$opencl_extension_commands.each { |c|
  eval "$#{c.prototype.name} = c"
}

746
747
748
749
750
751
752
753
754
755
def upper_snake_case(str)
  str.gsub(/([A-Z][A-Z0-9]*)/, '_\1').upcase
end

OPENCL_POINTER_NAMES = ($opencl_commands.collect { |c|
  [c, upper_snake_case(c.prototype.pointer_name)]
} + $opencl_extension_commands.collect { |c|
  [c, c.prototype.pointer_name]
}).to_h

756
757
758
759
760
761
762
763
764
765
766
767
768
($opencl_commands+$opencl_extension_commands).select { |c|
   c.parameters.find { |p| p.name == "param_value_size_ret" && p.pointer? }
}.each { |c|
   c.prologues.push <<EOF
  size_t _new_param_value_size;
  if (!param_value_size_ret)
    param_value_size_ret = &_new_param_value_size;
EOF
  c.epilogues.push <<EOF
  param_value_size = (param_value_size <= *param_value_size_ret ? param_value_size : *param_value_size_ret );
EOF
}

769
buffer_create_info = InMetaParameter::new($clCreateSubBuffer, "buffer_create_info")
770
buffer_create_info.instance_variable_set(:@lttng_in_type, ["ctf_sequence_text", "uint8_t", "buffer_create_info_vals", "buffer_create_info", "size_t", "buffer_create_info == NULL ? 0 : (buffer_create_type == CL_BUFFER_CREATE_TYPE_REGION ? sizeof(cl_buffer_region) : 0)"])
Brice Videau's avatar
Brice Videau committed
771

772
$clCreateSubBuffer.meta_parameters.push buffer_create_info
773

Brice Videau's avatar
Brice Videau committed
774

775
($opencl_commands+$opencl_extension_commands).each { |c|
776
  if c.prototype.name.match "clEnqueue"
Brice Videau's avatar
Brice Videau committed
777
    c.prologues.push <<EOF
Brice Videau's avatar
Brice Videau committed
778
  int64_t _enqueue_counter = 0;
779
780
781
782
783
784
785
786
787
788
789
  if (do_dump) {
    pthread_mutex_lock(&enqueue_counter_mutex);
    _enqueue_counter = enqueue_counter;
    enqueue_counter++;
    pthread_mutex_unlock(&enqueue_counter_mutex);
    tracepoint(lttng_ust_opencl_dump, enqueue_counter, _enqueue_counter);
  }
EOF
  end
}

Brice Videau's avatar
Brice Videau committed
790
791
792
793
794
795
796
797
798
799
800
801
802
class ParamName < MetaParameter
  def initialize(c)
    super(c, "param_name")
    raise "Couldn't find variable param_name for #{c.prototype.name}!" unless c["param_name"]
    @type = c["param_name"].type.gsub("*", "")
  end

  def lttng_out_type
    ["ctf_integer_hex", @type, "_param_name", "param_name"]
  end
end

($opencl_commands+$opencl_extension_commands).each { |c|
803
  if c.prototype.name.match(/clGet(\w*?)Info/) && c["param_name"]
Brice Videau's avatar
Brice Videau committed
804
805
806
807
    c.meta_parameters.push(ParamName::new(c))
  end
}

808
809
810
811
812
813
814
815
register_epilogue "clCreateKernel", <<EOF
  if (do_dump && _retval != NULL) {
    add_kernel(_retval);
  }
EOF

register_epilogue "clSetKernelArg", <<EOF
  if (do_dump && _retval == CL_SUCCESS) {
816
817
818
819
820
821
822
823
824
825
826
827
828
829
830
831
832
833
834
    add_kernel_arg(kernel, arg_index, arg_size, arg_value, 0);
  }
EOF

register_epilogue "clSetKernelArgSVMPointer", <<EOF
  if (do_dump && _retval == CL_SUCCESS) {
    add_kernel_arg(kernel, arg_index, sizeof(arg_value), arg_value, 1);
  }
EOF

register_epilogue "clSVMAlloc", <<EOF
  if (do_dump && _retval != NULL) {
    add_svmptr(_retval, size);
  }
EOF

register_prologue "clSVMFree", <<EOF
  if (do_dump && svm_pointer != NULL) {
    remove_svmptr(svm_pointer);
835
836
837
  }
EOF

838
str = <<EOF
Brice Videau's avatar
Brice Videau committed
839
840
841
  int _dump_release_events = 0;
  int _dump_release_event = 0;
  cl_event extra_event;
842
  if (do_dump && command_queue != NULL && kernel != NULL && _enqueue_counter >= dump_start && _enqueue_counter <= dump_end) {
Brice Videau's avatar
Brice Videau committed
843
    cl_command_queue_properties properties;
844
    #{OPENCL_POINTER_NAMES[$clGetCommandQueueInfo]}(command_queue, CL_QUEUE_PROPERTIES, sizeof(cl_command_queue_properties), &properties, NULL);
Brice Videau's avatar
Brice Videau committed
845
846
847
848
849
    _dump_release_events = dump_kernel_args(command_queue, kernel, _enqueue_counter, properties, &num_events_in_wait_list, (cl_event **)&event_wait_list);
    if (properties | CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE && event == NULL) {
      event = &extra_event;
      _dump_release_event = 1;
    }
850
851
  }
EOF
852
853
register_prologue "clEnqueueNDRangeKernel", str
register_prologue "clEnqueueNDRangeKernelINTEL", str
854

855
str = <<EOF
Brice Videau's avatar
Brice Videau committed
856
  if (do_dump && _dump_release_events) {
Brice Videau's avatar
Brice Videau committed
857
    for (cl_uint event_index = 0; event_index < num_events_in_wait_list; event_index++) {
858
      #{OPENCL_POINTER_NAMES[$clReleaseEvent]}(event_wait_list[event_index]);
859
    }
Brice Videau's avatar
Brice Videau committed
860
    free((void *)event_wait_list);
861
862
  }
EOF
863
864
register_epilogue "clEnqueueNDRangeKernel", str
register_epilogue "clEnqueueNDRangeKernelINTEL", str
865
866
867
868
869
870
871
872
873
874
875
876
877
878

register_prologue "clCreateBuffer", <<EOF
  if (do_dump) {
    flags &= ~CL_MEM_HOST_WRITE_ONLY;
    flags &= ~CL_MEM_HOST_NO_ACCESS;
  }
EOF

register_epilogue "clCreateBuffer", <<EOF
  if (do_dump && _retval != NULL) {
    add_buffer(_retval, size);
  }
EOF

Brice Videau's avatar
Brice Videau committed
879
register_prologue "clCreateCommandQueue", <<EOF
880
  if (tracepoint_enabled(lttng_ust_opencl_profiling, event_profiling)) {
Brice Videau's avatar
Brice Videau committed
881
882
883
884
    properties |= CL_QUEUE_PROFILING_ENABLE;
  }
EOF

Brice Videau's avatar
Brice Videau committed
885
register_prologue "clCreateCommandQueueWithProperties", <<EOF
Brice Videau's avatar
Brice Videau committed
886
  cl_queue_properties *_profiling_properties = NULL;
887
  if (tracepoint_enabled(lttng_ust_opencl_profiling, event_profiling)) {
Brice Videau's avatar
Brice Videau committed
888
889
890
891
892
893
894
895
896
897
898
899
900
901
902
903
904
905
906
907
908
909
910
911
912
913
914
915
916
917
918
919
920
921
922
923
924
925
926
927
928
929
930
931
932
933
934
    int _found_queue_properties = 0;
    int _queue_properties_index = 0;
    int _properties_count = 0;
    if (properties) {
      while(properties[_properties_count]) {
        if (properties[_properties_count] == CL_QUEUE_PROPERTIES){
          _found_queue_properties = 1;
          _queue_properties_index = _properties_count;
        }
        _properties_count += 2;
      }
      _properties_count++;
      if (!_found_queue_properties)
        _properties_count +=2;
    } else
      _properties_count = 3;
    _profiling_properties = (cl_queue_properties *)malloc(_properties_count*sizeof(cl_queue_properties));
    if (_profiling_properties) {
      if (properties) {
        int _i = 0;
        while(properties[_i]) {
          _profiling_properties[_i] = properties[_i];
          _profiling_properties[_i+1] = properties[_i+1];
          _i += 2;
        }
        if (_found_queue_properties) {
          _profiling_properties[_queue_properties_index+1] |= CL_QUEUE_PROFILING_ENABLE;
          _profiling_properties[_i] = 0;
        } else {
          _profiling_properties[_i++] = CL_QUEUE_PROPERTIES;
          _profiling_properties[_i++] = CL_QUEUE_PROFILING_ENABLE;
          _profiling_properties[_i] = 0;
        }
      } else {
        _profiling_properties[0] = CL_QUEUE_PROPERTIES;
        _profiling_properties[1] = CL_QUEUE_PROFILING_ENABLE;
        _profiling_properties[2] = 0;
      }
      properties = _profiling_properties;
    }
  }
EOF

register_epilogue "clCreateCommandQueueWithProperties", <<EOF
  if (_profiling_properties) free(_profiling_properties);
EOF

Brice Videau's avatar
Brice Videau committed
935
register_prologue "clCreateProgramWithSource", <<EOF
936
  if (tracepoint_enabled(lttng_ust_opencl_source, program_string) && strings != NULL) {
Brice Videau's avatar
Brice Videau committed
937
    cl_uint index;
Brice Videau's avatar
Brice Videau committed
938
939
    for (index = 0; index < count; index++) {
      size_t length = 0;
940
941
      char path[sizeof(SOURCE_TEMPLATE)];
      strncpy(path, SOURCE_TEMPLATE, sizeof(path));
Brice Videau's avatar
Brice Videau committed
942
943
944
945
946
947
      if ( strings[index] != NULL ) {
        if (lengths == NULL || lengths[index] == 0)
          length = strlen(strings[index]);
        else
          length = lengths[index];
      }
948
      create_file_and_write(path, length, strings[index]);
949
      do_tracepoint(lttng_ust_opencl_source, program_string, index, length, path);
Brice Videau's avatar
Brice Videau committed
950
951
952
953
    }
  }
EOF

Brice Videau's avatar
Brice Videau committed
954
register_prologue "clCreateProgramWithBinary", <<EOF
955
  if (tracepoint_enabled(lttng_ust_opencl_source, program_binary) && binaries != NULL && lengths != NULL) {
Brice Videau's avatar
Brice Videau committed
956
    cl_uint index;
Brice Videau's avatar
Brice Videau committed
957
    for (index = 0; index < num_devices; index++) {
958
959
960
      char path[sizeof(BIN_SOURCE_TEMPLATE)];
      strncpy(path, BIN_SOURCE_TEMPLATE, sizeof(path));
      create_file_and_write(path, lengths[index], binaries[index]);
961
      do_tracepoint(lttng_ust_opencl_source, program_binary, index, lengths[index], path);
Brice Videau's avatar
Brice Videau committed
962
963
964
965
    }
  }
EOF

Brice Videau's avatar
Brice Videau committed
966
register_prologue "clCreateProgramWithIL", <<EOF
967
  if (tracepoint_enabled(lttng_ust_opencl_source, program_il) && il != NULL) {
968
969
970
    char path[sizeof(IL_SOURCE_TEMPLATE)];
    strncpy(path, IL_SOURCE_TEMPLATE, sizeof(path));
    create_file_and_write(path, length, il);
971
    do_tracepoint(lttng_ust_opencl_source, program_il, length, path);
Brice Videau's avatar
Brice Videau committed
972
973
974
  }
EOF

975
register_prologue "clCreateProgramWithILKHR", <<EOF
976
  if (tracepoint_enabled(lttng_ust_opencl_source, program_il) && il != NULL) {
977
978
979
    char path[sizeof(IL_SOURCE_TEMPLATE)];
    strncpy(path, IL_SOURCE_TEMPLATE, sizeof(path));
    create_file_and_write(path, length, il);
980
    do_tracepoint(lttng_ust_opencl_source, program_il, length, path);
981
982
983
  }
EOF

Brice Videau's avatar
Brice Videau committed
984
985
str = <<EOF
  int _free_options = 0;
986
  if (tracepoint_enabled(lttng_ust_opencl_arguments, argument_info)) {
987
988
989
990
991
992
993
994
995
    struct opencl_version version = {1, 0};
    get_program_platform_version(program, &version);
    if (compare_opencl_version(&version, &opencl_version_1_2) >= 0) {
      if (options) {
        if (!strstr(options, "-cl-kernel-arg-info")) {
          size_t sz = strlen(options) + strlen("-cl-kernel-arg-info") + 2;
          char * new_options = (char *)malloc(sz);
          if (new_options) {
            snprintf(new_options, sz, "%s %s", options, "-cl-kernel-arg-info");
Brice Videau's avatar
Brice Videau committed
996
            _free_options = 1;
997
998
999
1000
1001
1002
1003
1004
1005
            options = new_options;
          }
        }
      } else {
        options = "-cl-kernel-arg-info";
      }
    }
  }
EOF
Brice Videau's avatar
Brice Videau committed
1006
1007
1008
1009
1010
register_prologue "clBuildProgram", str
register_prologue "clCompileProgram", str
register_prologue "clLinkProgram", <<EOF
  int _free_options = 0;
  if (tracepoint_enabled(lttng_ust_opencl_arguments, argument_info) && input_programs && num_input_programs > 0) {
1011
    struct opencl_version version = {1, 0};
Brice Videau's avatar
Brice Videau committed
1012
    get_program_platform_version(*input_programs, &version);
1013
1014
1015
1016
1017
1018
1019
    if (compare_opencl_version(&version, &opencl_version_1_2) >= 0) {
      if (options) {
        if (!strstr(options, "-cl-kernel-arg-info")) {
          size_t sz = strlen(options) + strlen("-cl-kernel-arg-info") + 2;
          char * new_options = (char *)malloc(sz);
          if (new_options) {
            snprintf(new_options, sz, "%s %s", options, "-cl-kernel-arg-info");
Brice Videau's avatar
Brice Videau committed
1020
            _free_options = 1;
1021
1022
1023
1024
1025
1026
1027
1028
1029
1030
            options = new_options;
          }
        }
      } else {
        options = "-cl-kernel-arg-info";
      }
    }
  }
EOF

Brice Videau's avatar
Brice Videau committed
1031
1032
str = <<EOF
  if (_free_options)
1033
1034
    free((char *)options);
EOF
Brice Videau's avatar
Brice Videau committed
1035
1036
register_epilogue "clBuildProgram", str
register_epilogue "clCompileProgram", str
Brice Videau's avatar
Brice Videau committed
1037
register_epilogue "clLinkProgram", str
Brice Videau's avatar
Brice Videau committed
1038
1039
1040
1041

l = lambda { |func, name: "pfn_notify", extra_conditions: nil|
  register_prologue func, <<EOF
  struct #{func}_callback_payload *_payload = NULL;
1042
  if ((tracepoint_enabled(lttng_ust_opencl, #{func}_callback_#{START})#{extra_conditions ? " || #{extra_conditions.join(" || ")}" : ""}) && #{name}) {
Brice Videau's avatar
Brice Videau committed
1043
1044
1045
1046
1047
    _payload = (struct #{func}_callback_payload *)malloc(sizeof(struct #{func}_callback_payload));
    _payload->#{name} = #{name};
    _payload->user_data = user_data;
    user_data = (void *)_payload;
    #{name} = &#{func}_callback;
1048
1049
  }
EOF
Brice Videau's avatar
Brice Videau committed
1050
1051
1052
}
program_conditions = ["tracepoint_enabled(lttng_ust_opencl_build, binaries)", "tracepoint_enabled(lttng_ust_opencl_build, infos)"]
l.call("clBuildProgram", extra_conditions: program_conditions)
1053
1054
l.call("clCompileProgram", extra_conditions: ["tracepoint_enabled(lttng_ust_opencl_build, objects)", "tracepoint_enabled(lttng_ust_opencl_build, infos)"])
l.call("clLinkProgram", extra_conditions: program_conditions)
Brice Videau's avatar
Brice Videau committed
1055
1056
1057
1058
1059
1060
l.call("clCreateContext")
l.call("clCreateContextFromType")
l.call("clSetMemObjectDestructorCallback")
l.call("clSetProgramReleaseCallback")
l.call("clSetEventCallback")
l.call("clEnqueueSVMFree", name: "pfn_free_func")
1061

Brice Videau's avatar
Brice Videau committed
1062
1063
1064
1065
1066
1067
1068
1069
1070
1071
1072
1073
1074
str = <<EOF
  if (_payload && _retval != CL_SUCCESS)
    free(_payload);
EOF
register_epilogue "clBuildProgram", str
register_epilogue "clCompileProgram", str
register_epilogue "clSetMemObjectDestructorCallback", str
register_epilogue "clSetProgramReleaseCallback", str
register_epilogue "clSetEventCallback", str
register_epilogue "clEnqueueSVMFree", str
str = <<EOF
  if (_payload && !_retval)
    free(_payload);
1075
EOF
Brice Videau's avatar
Brice Videau committed
1076
1077
1078
register_epilogue "clLinkProgram", str
register_epilogue "clCreateContext", str
register_epilogue "clCreateContextFromType", str
1079

1080
register_epilogue "clBuildProgram", <<EOF
Brice Videau's avatar
Brice Videau committed
1081
1082
  if (tracepoint_enabled(lttng_ust_opencl_build, binaries) && !pfn_notify) {
    dump_program_binaries(program);
1083
1084
  }
EOF
1085
1086
1087
1088
1089
register_epilogue "clCompileProgram", <<EOF
  if (tracepoint_enabled(lttng_ust_opencl_build, objects) && !pfn_notify) {
    dump_program_objects(program);
  }
EOF
1090
register_epilogue "clLinkProgram", <<EOF
1091
  if (tracepoint_enabled(lttng_ust_opencl_build, binaries) && _retval && !pfn_notify) {
Brice Videau's avatar
Brice Videau committed
1092
    dump_program_binaries(_retval);
1093
1094
1095
  }
EOF

Brice Videau's avatar
Brice Videau committed
1096
1097
1098
1099
1100
1101
1102
str = <<EOF
  if (tracepoint_enabled(lttng_ust_opencl_build, infos) && !pfn_notify) {
    dump_program_build_infos(program);
  }
EOF
register_epilogue "clBuildProgram", str
register_epilogue "clCompileProgram", str
1103
register_epilogue "clLinkProgram", <<EOF
1104
  if (tracepoint_enabled(lttng_ust_opencl_build, infos) && _retval && !pfn_notify) {
1105
    dump_program_build_infos(_retval);
Brice Videau's avatar
Brice Videau committed
1106
1107
1108
  }
EOF

1109
register_epilogue "clCreateKernel", <<EOF
1110
  if (tracepoint_enabled(lttng_ust_opencl_arguments, kernel_info)) {
1111
1112
1113
1114
    dump_kernel_info(_retval);
  }
EOF

1115
register_epilogue "clCreateKernel", <<EOF
1116
  if (tracepoint_enabled(lttng_ust_opencl_arguments, argument_info) && _retval != NULL) {
1117
1118
1119
1120
1121
1122
    dump_kernel_arguments(program, _retval);
  }
EOF

register_prologue "clCreateKernelsInProgram", <<EOF
  cl_uint n_k = 0;
1123
  if (tracepoint_enabled(lttng_ust_opencl_arguments, kernel_info) && !num_kernels_ret && kernels) {
1124
1125
1126
1127
1128
    num_kernels_ret = &n_k;
  }
EOF

register_epilogue "clCreateKernelsInProgram", <<EOF