llvm/test/CodeGen/AMDGPU/attr-amdgpu-max-num-workgroups.ll

   1 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 < %s | FileCheck %s
   2
   3 ; Attribute not specified.
   4 ; CHECK-LABEL: {{^}}empty_no_attribute:
   5 define amdgpu_kernel void @empty_no_attribute() {
   6 entry:
   7   ret void
   8 }
   9
  10 ; Ignore if number of work groups for x dimension is 0.
  11 ; CHECK-LABEL: {{^}}empty_max_num_workgroups_x0:
  12 define amdgpu_kernel void @empty_max_num_workgroups_x0() #0 {
  13 entry:
  14   ret void
  15 }
  16 attributes #0 = {"amdgpu-max-num-workgroups"="0,2,3"}
  17
  18 ; Ignore if number of work groups for y dimension is 0.
  19 ; CHECK-LABEL: {{^}}empty_max_num_workgroups_y0:
  20 define amdgpu_kernel void @empty_max_num_workgroups_y0() #1 {
  21 entry:
  22   ret void
  23 }
  24 attributes #1 = {"amdgpu-max-num-workgroups"="1,0,3"}
  25
  26 ; Ignore if number of work groups for z dimension is 0.
  27 ; CHECK-LABEL: {{^}}empty_max_num_workgroups_z0:
  28 define amdgpu_kernel void @empty_max_num_workgroups_z0() #2 {
  29 entry:
  30   ret void
  31 }
  32 attributes #2 = {"amdgpu-max-num-workgroups"="1,2,0"}
  33
  34 ; CHECK-LABEL: {{^}}empty_max_num_workgroups_1_2_3:
  35 define amdgpu_kernel void @empty_max_num_workgroups_1_2_3() #3 {
  36 entry:
  37   ret void
  38 }
  39 attributes #3 = {"amdgpu-max-num-workgroups"="1,2,3"}
  40
  41 ; CHECK-LABEL: {{^}}empty_max_num_workgroups_1024_1024_1024:
  42 define amdgpu_kernel void @empty_max_num_workgroups_1024_1024_1024() #4 {
  43 entry:
  44   ret void
  45 }
  46 attributes #4 = {"amdgpu-max-num-workgroups"="1024,1024,1024"}
  47
  48
  49
  50 ; Ignore if number of work groups for x dimension is 0.
  51 ; CHECK-LABEL: {{^}}empty_max_num_workgroups_x_max:
  52 define amdgpu_kernel void @empty_max_num_workgroups_x_max() #5 {
  53 entry:
  54   ret void
  55 }
  56 attributes #5 = {"amdgpu-max-num-workgroups"="4294967295,2,3"}
  57
  58 ; Ignore if number of work groups for y dimension is 0.
  59 ; CHECK-LABEL: {{^}}empty_max_num_workgroups_y_max:
  60 define amdgpu_kernel void @empty_max_num_workgroups_y_max() #6 {
  61 entry:
  62   ret void
  63 }
  64 attributes #6 = {"amdgpu-max-num-workgroups"="1,4294967295,3"}
  65
  66 ; Ignore if number of work groups for z dimension is 0.
  67 ; CHECK-LABEL: {{^}}empty_max_num_workgroups_z_max:
  68 define amdgpu_kernel void @empty_max_num_workgroups_z_max() #7 {
  69 entry:
  70   ret void
  71 }
  72 attributes #7 = {"amdgpu-max-num-workgroups"="1,2,4294967295"}
  73
  74
  75 ; CHECK: .amdgpu_metadata
  76 ; CHECK: - .args:
  77 ; CHECK:        .max_flat_workgroup_size: 1024
  78 ; CHECK-NEXT:   .name:           empty_no_attribute
  79 ; CHECK-NEXT:   .private_segment_fixed_size: 0
  80
  81 ; CHECK: - .args:
  82 ; CHECK:        .max_flat_workgroup_size: 1024
  83 ; CHECK-NEXT:   .max_num_workgroups_y: 2
  84 ; CHECK-NEXT:   .max_num_workgroups_z: 3
  85 ; CHECK-NEXT:   .name:           empty_max_num_workgroups_x0
  86 ; CHECK-NEXT:   .private_segment_fixed_size: 0
  87
  88 ; CHECK: - .args:
  89 ; CHECK:        .max_flat_workgroup_size: 1024
  90 ; CHECK-NEXT:   .max_num_workgroups_x: 1
  91 ; CHECK-NEXT:   .max_num_workgroups_z: 3
  92 ; CHECK-NEXT:   .name:           empty_max_num_workgroups_y0
  93 ; CHECK-NEXT:   .private_segment_fixed_size: 0
  94
  95 ; CHECK: - .args:
  96 ; CHECK:        .max_flat_workgroup_size: 1024
  97 ; CHECK-NEXT:   .max_num_workgroups_x: 1
  98 ; CHECK-NEXT:   .max_num_workgroups_y: 2
  99 ; CHECK-NEXT:   .name:           empty_max_num_workgroups_z0
 100 ; CHECK-NEXT:   .private_segment_fixed_size: 0
 101
 102 ; CHECK: - .args:
 103 ; CHECK:        .max_flat_workgroup_size: 1024
 104 ; CHECK-NEXT:   .max_num_workgroups_x: 1
 105 ; CHECK-NEXT:   .max_num_workgroups_y: 2
 106 ; CHECK-NEXT:   .max_num_workgroups_z: 3
 107 ; CHECK-NEXT:   .name:           empty_max_num_workgroups_1_2_3
 108 ; CHECK-NEXT:   .private_segment_fixed_size: 0
 109
 110 ; CHECK: - .args:
 111 ; CHECK:        .max_flat_workgroup_size: 1024
 112 ; CHECK-NEXT:   .max_num_workgroups_x: 1024
 113 ; CHECK-NEXT:   .max_num_workgroups_y: 1024
 114 ; CHECK-NEXT:   .max_num_workgroups_z: 1024
 115 ; CHECK-NEXT:   .name:           empty_max_num_workgroups_1024_1024_1024
 116 ; CHECK-NEXT:   .private_segment_fixed_size: 0
 117
 118
 119 ; CHECK: - .args:
 120 ; CHECK:        .max_flat_workgroup_size: 1024
 121 ; CHECK-NEXT:   .max_num_workgroups_y: 2
 122 ; CHECK-NEXT:   .max_num_workgroups_z: 3
 123 ; CHECK-NEXT:   .name:           empty_max_num_workgroups_x_max
 124 ; CHECK-NEXT:   .private_segment_fixed_size: 0
 125
 126 ; CHECK: - .args:
 127 ; CHECK:        .max_flat_workgroup_size: 1024
 128 ; CHECK-NEXT:   .max_num_workgroups_x: 1
 129 ; CHECK-NEXT:   .max_num_workgroups_z: 3
 130 ; CHECK-NEXT:   .name:           empty_max_num_workgroups_y_max
 131 ; CHECK-NEXT:   .private_segment_fixed_size: 0
 132
 133 ; CHECK: - .args:
 134 ; CHECK:        .max_flat_workgroup_size: 1024
 135 ; CHECK-NEXT:   .max_num_workgroups_x: 1
 136 ; CHECK-NEXT:   .max_num_workgroups_y: 2
 137 ; CHECK-NEXT:   .name:           empty_max_num_workgroups_z_max
 138 ; CHECK-NEXT:   .private_segment_fixed_size: 0