test/MC/AMDGPU/hsa-metadata-kernel-args-v3.s

   1 // RUN: llvm-mc -mattr=+code-object-v3 -triple=amdgcn-amd-amdhsa -mcpu=gfx700 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX700 %s
   2 // RUN: llvm-mc -mattr=+code-object-v3 -triple=amdgcn-amd-amdhsa -mcpu=gfx800 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX800 %s
   3 // RUN: llvm-mc -mattr=+code-object-v3 -triple=amdgcn-amd-amdhsa -mcpu=gfx900 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX900 %s
   4
   5 ; CHECK:        .amdgpu_metadata
   6 ; CHECK:      amdhsa.kernels:
   7 ; CHECK-NEXT:   - .args:
   8 ; CHECK-NEXT:       - .offset:         1
   9 ; CHECK-NEXT:         .size:           1
  10 ; CHECK-NEXT:         .type_name:      char
  11 ; CHECK-NEXT:         .value_kind:     by_value
  12 ; CHECK-NEXT:         .value_type:     i8
  13 ; CHECK-NEXT:       - .offset:         8
  14 ; CHECK-NEXT:         .size:           8
  15 ; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
  16 ; CHECK-NEXT:         .value_type:     i64
  17 ; CHECK-NEXT:       - .offset:         8
  18 ; CHECK-NEXT:         .size:           8
  19 ; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
  20 ; CHECK-NEXT:         .value_type:     i64
  21 ; CHECK-NEXT:       - .offset:         8
  22 ; CHECK-NEXT:         .size:           8
  23 ; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
  24 ; CHECK-NEXT:         .value_type:     i64
  25 ; CHECK-NEXT:       - .address_space:  global
  26 ; CHECK-NEXT:         .offset:         8
  27 ; CHECK-NEXT:         .size:           8
  28 ; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
  29 ; CHECK-NEXT:         .value_type:     i8
  30 ; CHECK-NEXT:     .group_segment_fixed_size: 16
  31 ; CHECK-NEXT:     .kernarg_segment_align: 64
  32 ; CHECK-NEXT:     .kernarg_segment_size: 8
  33 ; CHECK-NEXT:     .language:       OpenCL C
  34 ; CHECK-NEXT:     .language_version:
  35 ; CHECK-NEXT:       - 2
  36 ; CHECK-NEXT:       - 0
  37 ; CHECK-NEXT:     .max_flat_workgroup_size: 256
  38 ; CHECK-NEXT:     .name:           test_kernel
  39 ; CHECK-NEXT:     .private_segment_fixed_size: 32
  40 ; CHECK-NEXT:     .sgpr_count:     14
  41 ; CHECK-NEXT:     .symbol:         'test_kernel@kd'
  42 ; CHECK-NEXT:     .vgpr_count:     40
  43 ; CHECK-NEXT:     .wavefront_size: 128
  44 ; CHECK-NEXT: amdhsa.printf:
  45 ; CHECK-NEXT:   - '1:1:4:%d\n'
  46 ; CHECK-NEXT:   - '2:1:8:%g\n'
  47 ; CHECK-NEXT: amdhsa.version:
  48 ; CHECK-NEXT:   - 1
  49 ; CHECK-NEXT:   - 0
  50 ; CHECK:        .end_amdgpu_metadata
  51 .amdgpu_metadata
  52   amdhsa.version:
  53     - 1
  54     - 0
  55   amdhsa.printf:
  56     - '1:1:4:%d\n'
  57     - '2:1:8:%g\n'
  58   amdhsa.kernels:
  59     - .name:            test_kernel
  60       .symbol:      test_kernel@kd
  61       .language:        OpenCL C
  62       .language_version:
  63         - 2
  64         - 0
  65       .kernarg_segment_size: 8
  66       .group_segment_fixed_size: 16
  67       .private_segment_fixed_size: 32
  68       .kernarg_segment_align: 64
  69       .wavefront_size: 128
  70       .sgpr_count: 14
  71       .vgpr_count: 40
  72       .max_flat_workgroup_size: 256
  73       .args:
  74         - .type_name:      char
  75           .size:          1
  76           .offset:         1
  77           .value_kind:     by_value
  78           .value_type:     i8
  79         - .size:          8
  80           .offset:         8
  81           .value_kind:     hidden_global_offset_x
  82           .value_type:     i64
  83         - .size:          8
  84           .offset:         8
  85           .value_kind:     hidden_global_offset_y
  86           .value_type:     i64
  87         - .size:          8
  88           .offset:         8
  89           .value_kind:     hidden_global_offset_z
  90           .value_type:     i64
  91         - .size:          8
  92           .offset:         8
  93           .value_kind:     hidden_printf_buffer
  94           .value_type:     i8
  95           .address_space: global
  96 .end_amdgpu_metadata