Lines Matching refs:gpu

1 // RUN: mlir-opt -split-input-file -convert-gpu-to-spirv="use-64bit-index=false" %s -o - | FileChec…
2 // RUN: mlir-opt -split-input-file -convert-gpu-to-spirv="use-64bit-index=true" %s -o - | FileCheck…
5   gpu.container_module,
10     gpu.launch_func @kernels::@builtin_workgroup_id_x
19   gpu.module @kernels {
20     gpu.func @builtin_workgroup_id_x() kernel
26       %0 = gpu.block_id x
27       gpu.return
35   gpu.container_module,
41     gpu.launch_func @kernels::@builtin_workgroup_id_y
49   gpu.module @kernels {
50     gpu.func @builtin_workgroup_id_y() kernel
55       %0 = gpu.block_id y
56       gpu.return
64   gpu.container_module,
69     gpu.launch_func @kernels::@builtin_workgroup_id_z
76   gpu.module @kernels {
77     gpu.func @builtin_workgroup_id_z() kernel
82       %0 = gpu.block_id z
83       gpu.return
91   gpu.container_module,
96     gpu.launch_func @kernels::@builtin_workgroup_size_x
102   gpu.module @kernels {
103     gpu.func @builtin_workgroup_size_x() kernel
106       // Note that this ignores the workgroup size specification in gpu.launch.
107       // We may want to define gpu.workgroup_size and convert it to the entry
110       %0 = gpu.block_dim x
111       gpu.return
119   gpu.container_module,
124     gpu.launch_func @kernels::@builtin_workgroup_size_y
130   gpu.module @kernels {
131     gpu.func @builtin_workgroup_size_y() kernel
135       %0 = gpu.block_dim y
136       gpu.return
144   gpu.container_module,
149     gpu.launch_func @kernels::@builtin_workgroup_size_z
155   gpu.module @kernels {
156     gpu.func @builtin_workgroup_size_z() kernel
160       %0 = gpu.block_dim z
161       gpu.return
169   gpu.container_module,
174     gpu.launch_func @kernels::@builtin_local_id_x
181   gpu.module @kernels {
182     gpu.func @builtin_local_id_x() kernel
187       %0 = gpu.thread_id x
188       gpu.return
196   gpu.container_module,
201     gpu.launch_func @kernels::@builtin_num_workgroups_x
208   gpu.module @kernels {
209     gpu.func @builtin_num_workgroups_x() kernel
214       %0 = gpu.grid_dim x
215       gpu.return
223   gpu.container_module,
228   gpu.module @kernels {
229     gpu.func @builtin_subgroup_id() kernel
233       %0 = gpu.subgroup_id : index
234       gpu.return
242   gpu.container_module,
247   gpu.module @kernels {
248     gpu.func @builtin_num_subgroups() kernel
252       %0 = gpu.num_subgroups : index
253       gpu.return
261   gpu.container_module,
266     gpu.launch_func @kernels::@builtin_workgroup_size_x
273   gpu.module @kernels {
274     gpu.func @builtin_workgroup_size_x() kernel
279       %0 = gpu.block_dim x
280       gpu.return
288   gpu.container_module,
293     gpu.launch_func @kernels::@builtin_workgroup_size_y
300   gpu.module @kernels {
301     gpu.func @builtin_workgroup_size_y() kernel
306       %0 = gpu.block_dim y
307       gpu.return
315   gpu.container_module,
320     gpu.launch_func @kernels::@builtin_workgroup_size_z
327   gpu.module @kernels {
328     gpu.func @builtin_workgroup_size_z() kernel
333       %0 = gpu.block_dim z
334       gpu.return
342   gpu.container_module,
347     gpu.launch_func @kernels::@builtin_global_id_x
354   gpu.module @kernels {
355     gpu.func @builtin_global_id_x() kernel
360       %0 = gpu.global_id x
361       gpu.return
369   gpu.container_module,
374     gpu.launch_func @kernels::@builtin_global_id_y
381   gpu.module @kernels {
382     gpu.func @builtin_global_id_y() kernel
387       %0 = gpu.global_id y
388       gpu.return
396   gpu.container_module,
401     gpu.launch_func @kernels::@builtin_global_id_z
408   gpu.module @kernels {
409     gpu.func @builtin_global_id_z() kernel
414       %0 = gpu.global_id z
415       gpu.return
424   gpu.container_module,
431   gpu.module @kernels {
432     gpu.func @builtin_subgroup_size() kernel
437       %0 = gpu.subgroup_size : index
438       gpu.return