llvm/test/CodeGen/NVPTX/demote-vars.ll

   1 ; RUN: llc -o - %s -march=nvptx -mcpu=sm_20 | FileCheck %s
   2
   3 ; Check that we do global variable demotion when the symbols don't need to be
   4 ; shared across modules or functions.
   5
   6 ; Variable visible globally can't be demoted.
   7 ; CHECK: .visible .shared .align 4 .u64 external_global
   8 @external_global = addrspace(3) global i64 undef, align 4
   9
  10 ; Not externally visible global used in only one function => demote.
  11 ; It should live in @define_internal_global.
  12 @internal_global = internal addrspace(3) global i64 undef, align 4
  13
  14 ; Not externally visible global used in several functions => keep.
  15 ; CHECK: .shared .align 4 .u64 internal_global_used_in_different_fcts
  16 @internal_global_used_in_different_fcts = internal addrspace(3) global i64 undef, align 4
  17
  18 ; Not externally visible global used in only one function => demote.
  19 ; It should live in @define_private_global.
  20 @private_global = private addrspace(3) global i64 undef, align 4
  21
  22 ; Not externally visible global used in only one function
  23 ; (but with several uses) => demote.
  24 ; It should live in @define_private_global_more_than_one_use.
  25 @private_global_used_more_than_once_in_same_fct = private addrspace(3) global i64 undef, align 4
  26
  27 ; CHECK-LABEL: define_external_global(
  28 define void @define_external_global(i64 %val) {
  29   store i64 %val, ptr addrspace(3) @external_global
  30   ret void
  31 }
  32
  33 ; CHECK-LABEL: define_internal_global(
  34 ; Demoted `internal_global` should live here.
  35 ; CHECK: .shared .align 4 .u64 internal_global
  36 define void @define_internal_global(i64 %val) {
  37   store i64 %val, ptr addrspace(3) @internal_global
  38   ret void
  39 }
  40
  41 ; CHECK-LABEL: define_internal_global_with_different_fct1(
  42 ; CHECK-NOT: .shared .align 4 .u64 internal_global_used_in_different_fcts
  43 define void @define_internal_global_with_different_fct1(i64 %val) {
  44   store i64 %val, ptr addrspace(3) @internal_global_used_in_different_fcts
  45   ret void
  46 }
  47
  48 ; CHECK-LABEL: define_internal_global_with_different_fct2(
  49 ; CHECK-NOT: .shared .align 4 .u64 internal_global_used_in_different_fcts
  50 define void @define_internal_global_with_different_fct2(i64 %val) {
  51   store i64 %val, ptr addrspace(3) @internal_global_used_in_different_fcts
  52   ret void
  53 }
  54
  55 ; CHECK-LABEL: define_private_global(
  56 ; Demoted `private_global` should live here.
  57 ; CHECK: .shared .align 4 .u64 private_global
  58 define void @define_private_global(i64 %val) {
  59   store i64 %val, ptr addrspace(3) @private_global
  60   ret void
  61 }
  62
  63 ; CHECK-LABEL: define_private_global_more_than_one_use(
  64 ; Demoted `private_global_used_more_than_once_in_same_fct` should live here.
  65 ; CHECK: .shared .align 4 .u64 private_global_used_more_than_once_in_same_fct
  66 ;
  67 ; Also check that the if-then is still here, otherwise we may not be testing
  68 ; the "more-than-one-use" part.
  69 ; CHECK: st.shared.u64   [private_global_used_more_than_once_in_same_fct],
  70 ; CHECK: mov.u64 %[[VAR:.*]], 25
  71 ; CHECK: st.shared.u64   [private_global_used_more_than_once_in_same_fct], %[[VAR]]
  72 define void @define_private_global_more_than_one_use(i64 %val, i1 %cond) {
  73   store i64 %val, ptr addrspace(3) @private_global_used_more_than_once_in_same_fct
  74   br i1 %cond, label %then, label %end
  75
  76 then:
  77   store i64 25, ptr addrspace(3) @private_global_used_more_than_once_in_same_fct
  78   br label %end
  79
  80 end:
  81   ret void
  82 }