test/MC/AMDGPU/hsa-metadata-kernel-attrs-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:      - .max_flat_workgroup_size: 256
   8 // CHECK:        .wavefront_size: 128
   9 // CHECK:        .symbol:      'test_kernel@kd'
  10 // CHECK:        .reqd_workgroup_size:
  11 // CHECK-NEXT:     - 1
  12 // CHECK-NEXT:     - 2
  13 // CHECK-NEXT:     - 4
  14 // CHECK:        .kernarg_segment_size: 8
  15 // CHECK:        .private_segment_fixed_size: 32
  16 // CHECK:        .workgroup_size_hint:
  17 // CHECK-NEXT:     - 8
  18 // CHECK-NEXT:     - 16
  19 // CHECK-NEXT:     - 32
  20 // CHECK:        .name:            test_kernel
  21 // CHECK:        .language:        OpenCL C
  22 // CHECK:        .sgpr_count:      14
  23 // CHECK:        .kernarg_segment_align: 64
  24 // CHECK:        .vgpr_count:      40
  25 // CHECK:        .language_version:
  26 // CHECK-NEXT:     - 2
  27 // CHECK-NEXT:     - 0
  28 // CHECK:        .vec_type_hint:       int
  29 // CHECK:    amdhsa.version:
  30 // CHECK-NEXT: - 1
  31 // CHECK-NEXT: - 0
  32 // CHECK:    amdhsa.printf:
  33 // CHECK:      - '1:1:4:%d\n'
  34 // CHECK:      - '2:1:8:%g\n'
  35 // CHECK: .end_amdgpu_metadata
  36 .amdgpu_metadata
  37   amdhsa.version:
  38     - 1
  39     - 0
  40   amdhsa.printf:
  41     - '1:1:4:%d\n'
  42     - '2:1:8:%g\n'
  43   amdhsa.kernels:
  44     - .name:            test_kernel
  45       .symbol:      test_kernel@kd
  46       .language:        OpenCL C
  47       .language_version:
  48         - 2
  49         - 0
  50       .kernarg_segment_size: 8
  51       .group_segment_fixed_size: 16
  52       .private_segment_fixed_size: 32
  53       .kernarg_segment_align: 64
  54       .wavefront_size: 128
  55       .sgpr_count: 14
  56       .vgpr_count: 40
  57       .max_flat_workgroup_size: 256
  58       .reqd_workgroup_size:
  59         - 1
  60         - 2
  61         - 4
  62       .workgroup_size_hint:
  63         - 8
  64         - 16
  65         - 32
  66       .vec_type_hint:       int
  67 .end_amdgpu_metadata