llvm/test/CodeGen/AMDGPU/attr-amdgpu-flat-work-group-size-v3.ll

   1 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx803 -verify-machineinstrs < %s | FileCheck --check-prefix=CHECK %s
   2 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx803 -verify-machineinstrs -amdgpu-verify-hsa-metadata -filetype=obj -o /dev/null < %s 2>&1 | FileCheck --check-prefix=PARSER %s
   3
   4 ; CHECK-LABEL: {{^}}min_64_max_64:
   5 ; CHECK: SGPRBlocks: 0
   6 ; CHECK: VGPRBlocks: 0
   7 ; CHECK: NumSGPRsForWavesPerEU: 1
   8 ; CHECK: NumVGPRsForWavesPerEU: 1
   9 define amdgpu_kernel void @min_64_max_64() #0 {
  10 entry:
  11   ret void
  12 }
  13 attributes #0 = {"amdgpu-flat-work-group-size"="64,64"}
  14
  15 ; CHECK-LABEL: {{^}}min_64_max_128:
  16 ; CHECK: SGPRBlocks: 0
  17 ; CHECK: VGPRBlocks: 0
  18 ; CHECK: NumSGPRsForWavesPerEU: 1
  19 ; CHECK: NumVGPRsForWavesPerEU: 1
  20 define amdgpu_kernel void @min_64_max_128() #1 {
  21 entry:
  22   ret void
  23 }
  24 attributes #1 = {"amdgpu-flat-work-group-size"="64,128"}
  25
  26 ; CHECK-LABEL: {{^}}min_128_max_128:
  27 ; CHECK: SGPRBlocks: 0
  28 ; CHECK: VGPRBlocks: 0
  29 ; CHECK: NumSGPRsForWavesPerEU: 1
  30 ; CHECK: NumVGPRsForWavesPerEU: 1
  31 define amdgpu_kernel void @min_128_max_128() #2 {
  32 entry:
  33   ret void
  34 }
  35 attributes #2 = {"amdgpu-flat-work-group-size"="128,128"}
  36
  37 ; CHECK-LABEL: {{^}}min_1024_max_1024
  38 ; CHECK: SGPRBlocks: 0
  39 ; CHECK: VGPRBlocks: 10
  40 ; CHECK: NumSGPRsForWavesPerEU: 2{{$}}
  41 ; CHECK: NumVGPRsForWavesPerEU: 43
  42 @var = addrspace(1) global float 0.0
  43 define amdgpu_kernel void @min_1024_max_1024() #3 {
  44   %val0 = load volatile float, ptr addrspace(1) @var
  45   %val1 = load volatile float, ptr addrspace(1) @var
  46   %val2 = load volatile float, ptr addrspace(1) @var
  47   %val3 = load volatile float, ptr addrspace(1) @var
  48   %val4 = load volatile float, ptr addrspace(1) @var
  49   %val5 = load volatile float, ptr addrspace(1) @var
  50   %val6 = load volatile float, ptr addrspace(1) @var
  51   %val7 = load volatile float, ptr addrspace(1) @var
  52   %val8 = load volatile float, ptr addrspace(1) @var
  53   %val9 = load volatile float, ptr addrspace(1) @var
  54   %val10 = load volatile float, ptr addrspace(1) @var
  55   %val11 = load volatile float, ptr addrspace(1) @var
  56   %val12 = load volatile float, ptr addrspace(1) @var
  57   %val13 = load volatile float, ptr addrspace(1) @var
  58   %val14 = load volatile float, ptr addrspace(1) @var
  59   %val15 = load volatile float, ptr addrspace(1) @var
  60   %val16 = load volatile float, ptr addrspace(1) @var
  61   %val17 = load volatile float, ptr addrspace(1) @var
  62   %val18 = load volatile float, ptr addrspace(1) @var
  63   %val19 = load volatile float, ptr addrspace(1) @var
  64   %val20 = load volatile float, ptr addrspace(1) @var
  65   %val21 = load volatile float, ptr addrspace(1) @var
  66   %val22 = load volatile float, ptr addrspace(1) @var
  67   %val23 = load volatile float, ptr addrspace(1) @var
  68   %val24 = load volatile float, ptr addrspace(1) @var
  69   %val25 = load volatile float, ptr addrspace(1) @var
  70   %val26 = load volatile float, ptr addrspace(1) @var
  71   %val27 = load volatile float, ptr addrspace(1) @var
  72   %val28 = load volatile float, ptr addrspace(1) @var
  73   %val29 = load volatile float, ptr addrspace(1) @var
  74   %val30 = load volatile float, ptr addrspace(1) @var
  75   %val31 = load volatile float, ptr addrspace(1) @var
  76   %val32 = load volatile float, ptr addrspace(1) @var
  77   %val33 = load volatile float, ptr addrspace(1) @var
  78   %val34 = load volatile float, ptr addrspace(1) @var
  79   %val35 = load volatile float, ptr addrspace(1) @var
  80   %val36 = load volatile float, ptr addrspace(1) @var
  81   %val37 = load volatile float, ptr addrspace(1) @var
  82   %val38 = load volatile float, ptr addrspace(1) @var
  83   %val39 = load volatile float, ptr addrspace(1) @var
  84   %val40 = load volatile float, ptr addrspace(1) @var
  85
  86   store volatile float %val0, ptr addrspace(1) @var
  87   store volatile float %val1, ptr addrspace(1) @var
  88   store volatile float %val2, ptr addrspace(1) @var
  89   store volatile float %val3, ptr addrspace(1) @var
  90   store volatile float %val4, ptr addrspace(1) @var
  91   store volatile float %val5, ptr addrspace(1) @var
  92   store volatile float %val6, ptr addrspace(1) @var
  93   store volatile float %val7, ptr addrspace(1) @var
  94   store volatile float %val8, ptr addrspace(1) @var
  95   store volatile float %val9, ptr addrspace(1) @var
  96   store volatile float %val10, ptr addrspace(1) @var
  97   store volatile float %val11, ptr addrspace(1) @var
  98   store volatile float %val12, ptr addrspace(1) @var
  99   store volatile float %val13, ptr addrspace(1) @var
 100   store volatile float %val14, ptr addrspace(1) @var
 101   store volatile float %val15, ptr addrspace(1) @var
 102   store volatile float %val16, ptr addrspace(1) @var
 103   store volatile float %val17, ptr addrspace(1) @var
 104   store volatile float %val18, ptr addrspace(1) @var
 105   store volatile float %val19, ptr addrspace(1) @var
 106   store volatile float %val20, ptr addrspace(1) @var
 107   store volatile float %val21, ptr addrspace(1) @var
 108   store volatile float %val22, ptr addrspace(1) @var
 109   store volatile float %val23, ptr addrspace(1) @var
 110   store volatile float %val24, ptr addrspace(1) @var
 111   store volatile float %val25, ptr addrspace(1) @var
 112   store volatile float %val26, ptr addrspace(1) @var
 113   store volatile float %val27, ptr addrspace(1) @var
 114   store volatile float %val28, ptr addrspace(1) @var
 115   store volatile float %val29, ptr addrspace(1) @var
 116   store volatile float %val30, ptr addrspace(1) @var
 117   store volatile float %val31, ptr addrspace(1) @var
 118   store volatile float %val32, ptr addrspace(1) @var
 119   store volatile float %val33, ptr addrspace(1) @var
 120   store volatile float %val34, ptr addrspace(1) @var
 121   store volatile float %val35, ptr addrspace(1) @var
 122   store volatile float %val36, ptr addrspace(1) @var
 123   store volatile float %val37, ptr addrspace(1) @var
 124   store volatile float %val38, ptr addrspace(1) @var
 125   store volatile float %val39, ptr addrspace(1) @var
 126   store volatile float %val40, ptr addrspace(1) @var
 127
 128   ret void
 129 }
 130 attributes #3 = {"amdgpu-flat-work-group-size"="1024,1024"}
 131
 132 !llvm.module.flags = !{!0}
 133 !0 = !{i32 1, !"amdgpu_code_object_version", i32 300}
 134
 135 ; CHECK: amdhsa.kernels:
 136 ; CHECK:   .max_flat_workgroup_size: 64
 137 ; CHECK:   .name:                 min_64_max_64
 138 ; CHECK:   .max_flat_workgroup_size: 128
 139 ; CHECK:   .name:                 min_64_max_128
 140 ; CHECK:   .max_flat_workgroup_size: 128
 141 ; CHECK:   .name:                 min_128_max_128
 142 ; CHECK:   .max_flat_workgroup_size: 1024
 143 ; CHECK:   .name:                 min_1024_max_1024
 144 ; CHECK: amdhsa.version:
 145 ; CHECK:   - 1
 146 ; CHECK:   - 0
 147
 148 ; PARSER: AMDGPU HSA Metadata Parser Test: PASS