// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx700 -show-encoding %s | FileCheck --check-prefix=CHECK %s
// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx800 -show-encoding %s | FileCheck --check-prefix=CHECK %s
// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx900 -show-encoding %s | FileCheck --check-prefix=CHECK %s

; CHECK:      	.amdgpu_metadata
; CHECK:      amdhsa.kernels:  
; CHECK-NEXT:   - .args:           
; CHECK-NEXT:       - .offset:         1
; CHECK-NEXT:         .size:           1
; CHECK-NEXT:         .type_name:      char
; CHECK-NEXT:         .value_kind:     by_value
; CHECK-NEXT:         .value_type:     i8
; CHECK-NEXT:       - .offset:         8
; CHECK-NEXT:         .size:           8
; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
; CHECK-NEXT:         .value_type:     i64
; CHECK-NEXT:       - .offset:         8
; CHECK-NEXT:         .size:           8
; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
; CHECK-NEXT:         .value_type:     i64
; CHECK-NEXT:       - .offset:         8
; CHECK-NEXT:         .size:           8
; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
; CHECK-NEXT:         .value_type:     i64
; CHECK-NEXT:       - .address_space:  global
; CHECK-NEXT:         .offset:         8
; CHECK-NEXT:         .size:           8
; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
; CHECK-NEXT:         .value_type:     i8
; CHECK-NEXT:     .group_segment_fixed_size: 16
; CHECK-NEXT:     .kernarg_segment_align: 64
; CHECK-NEXT:     .kernarg_segment_size: 8
; CHECK-NEXT:     .language:       OpenCL C
; CHECK-NEXT:     .language_version: 
; CHECK-NEXT:       - 2
; CHECK-NEXT:       - 0
; CHECK-NEXT:     .max_flat_workgroup_size: 256
; CHECK-NEXT:     .name:           test_kernel
; CHECK-NEXT:     .private_segment_fixed_size: 32
; CHECK-NEXT:     .sgpr_count:     14
; CHECK-NEXT:     .symbol:         'test_kernel@kd'
; CHECK-NEXT:     .uses_dynamic_stack: true
; CHECK-NEXT:     .vgpr_count:     40
; CHECK-NEXT:     .wavefront_size: 128
; CHECK-NEXT: amdhsa.printf:   
; CHECK-NEXT:   - '1:1:4:%d\n'
; CHECK-NEXT:   - '2:1:8:%g\n'
; CHECK-NEXT: amdhsa.version:  
; CHECK-NEXT:   - 1
; CHECK-NEXT:   - 0
; CHECK:      	.end_amdgpu_metadata
.amdgpu_metadata
  amdhsa.version:
    - 1
    - 0
  amdhsa.printf:
    - '1:1:4:%d\n'
    - '2:1:8:%g\n'
  amdhsa.kernels:
    - .name:            test_kernel
      .symbol:      test_kernel@kd
      .language:        OpenCL C
      .language_version:
        - 2
        - 0
      .kernarg_segment_size: 8
      .group_segment_fixed_size: 16
      .private_segment_fixed_size: 32
      .uses_dynamic_stack: true
      .kernarg_segment_align: 64
      .wavefront_size: 128
      .sgpr_count: 14
      .vgpr_count: 40
      .max_flat_workgroup_size: 256
      .args:
        - .type_name:      char
          .size:          1
          .offset:         1
          .value_kind:     by_value
          .value_type:     i8
        - .size:          8
          .offset:         8
          .value_kind:     hidden_global_offset_x
          .value_type:     i64
        - .size:          8
          .offset:         8
          .value_kind:     hidden_global_offset_y
          .value_type:     i64
        - .size:          8
          .offset:         8
          .value_kind:     hidden_global_offset_z
          .value_type:     i64
        - .size:          8
          .offset:         8
          .value_kind:     hidden_printf_buffer
          .value_type:     i8
          .address_space: global
.end_amdgpu_metadata