clang/test/CodeGenCUDA/lambda-reference-var.cu

   1 // RUN: %clang_cc1 -x hip -emit-llvm -std=c++11 %s -o - \
   2 // RUN:   -triple x86_64-linux-gnu \
   3 // RUN:   | FileCheck -check-prefix=HOST %s
   4 // RUN: %clang_cc1 -x hip -emit-llvm -std=c++11 %s -o - \
   5 // RUN:   -triple amdgcn-amd-amdhsa -fcuda-is-device \
   6 // RUN:   | FileCheck -check-prefix=DEV %s
   7
   8 #include "Inputs/cuda.h"
   9
  10 // HOST: %[[T1:.*]] = type <{ ptr, i32, [4 x i8] }>
  11 // HOST: %[[T2:.*]] = type { ptr, ptr }
  12 // HOST: %[[T3:.*]] = type <{ ptr, i32, [4 x i8] }>
  13 // DEV: %[[T1:.*]] = type { ptr }
  14 // DEV: %[[T2:.*]] = type { ptr }
  15 // DEV: %[[T3:.*]] = type <{ ptr, i32, [4 x i8] }>
  16 int global_host_var;
  17 __device__ int global_device_var;
  18
  19 template<class F>
  20 __global__ void kern(F f) { f(); }
  21
  22 // DEV-LABEL: @_ZZ27dev_capture_dev_ref_by_copyPiENKUlvE_clEv(
  23 // DEV: %[[VAL:.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr)
  24 // DEV: store i32 %[[VAL]]
  25 __device__ void dev_capture_dev_ref_by_copy(int *out) {
  26   int &ref = global_device_var;
  27   [=](){ *out = ref;}();
  28 }
  29
  30 // DEV-LABEL: @_ZZ28dev_capture_dev_rval_by_copyPiENKUlvE_clEv(
  31 // DEV: store i32 3
  32 __device__ void dev_capture_dev_rval_by_copy(int *out) {
  33   constexpr int a = 1;
  34   constexpr int b = 2;
  35   constexpr int c = a + b;
  36   [=](){ *out = c;}();
  37 }
  38
  39 // DEV-LABEL: @_ZZ26dev_capture_dev_ref_by_refPiENKUlvE_clEv(
  40 // DEV: %[[VAL:.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr)
  41 // DEV: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1
  42 // DEV: store i32 %[[VAL2]], ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr)
  43 // DEV: %[[VAL:.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr)
  44 // DEV: store i32 %[[VAL]]
  45 __device__ void dev_capture_dev_ref_by_ref(int *out) {
  46   int &ref = global_device_var;
  47   [&](){ ref++; *out = ref;}();
  48 }
  49
  50 // DEV-LABEL: define{{.*}} void @_Z7dev_refPi(
  51 // DEV: %[[VAL:.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr)
  52 // DEV: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1
  53 // DEV: store i32 %[[VAL2]], ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr)
  54 // DEV: %[[VAL:.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr)
  55 // DEV: store i32 %[[VAL]]
  56 __device__ void dev_ref(int *out) {
  57   int &ref = global_device_var;
  58   ref++;
  59   *out = ref;
  60 }
  61
  62 // DEV-LABEL: @_ZZ14dev_lambda_refPiENKUlvE_clEv(
  63 // DEV: %[[VAL:.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr)
  64 // DEV: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1
  65 // DEV: store i32 %[[VAL2]], ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr)
  66 // DEV: %[[VAL:.*]] = load i32, ptr addrspacecast (ptr addrspace(1) @global_device_var to ptr)
  67 // DEV: store i32 %[[VAL]]
  68 __device__ void dev_lambda_ref(int *out) {
  69   [=](){
  70     int &ref = global_device_var;
  71     ref++;
  72     *out = ref;
  73   }();
  74 }
  75
  76 // HOST-LABEL: @_ZZ29host_capture_host_ref_by_copyPiENKUlvE_clEv(
  77 // HOST: %[[VAL:.*]] = load i32, ptr @global_host_var
  78 // HOST: store i32 %[[VAL]]
  79 void host_capture_host_ref_by_copy(int *out) {
  80   int &ref = global_host_var;
  81   [=](){ *out = ref;}();
  82 }
  83
  84 // HOST-LABEL: @_ZZ28host_capture_host_ref_by_refPiENKUlvE_clEv(
  85 // HOST: %[[CAP:.*]] = getelementptr inbounds %[[T2]], ptr %this1, i32 0, i32 0
  86 // HOST: %[[REF:.*]] = load ptr, ptr %[[CAP]]
  87 // HOST: %[[VAL:.*]] = load i32, ptr %[[REF]]
  88 // HOST: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1
  89 // HOST: store i32 %[[VAL2]], ptr %[[REF]]
  90 // HOST: %[[VAL:.*]] = load i32, ptr @global_host_var
  91 // HOST: store i32 %[[VAL]]
  92 void host_capture_host_ref_by_ref(int *out) {
  93   int &ref = global_host_var;
  94   [&](){ ref++; *out = ref;}();
  95 }
  96
  97 // HOST-LABEL: define{{.*}} void @_Z8host_refPi(
  98 // HOST: %[[VAL:.*]] = load i32, ptr @global_host_var
  99 // HOST: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1
 100 // HOST: store i32 %[[VAL2]], ptr @global_host_var
 101 // HOST: %[[VAL:.*]] = load i32, ptr @global_host_var
 102 // HOST: store i32 %[[VAL]]
 103 void host_ref(int *out) {
 104   int &ref = global_host_var;
 105   ref++;
 106   *out = ref;
 107 }
 108
 109 // HOST-LABEL: @_ZZ15host_lambda_refPiENKUlvE_clEv(
 110 // HOST: %[[VAL:.*]] = load i32, ptr @global_host_var
 111 // HOST: %[[VAL2:.*]] = add nsw i32 %[[VAL]], 1
 112 // HOST: store i32 %[[VAL2]], ptr @global_host_var
 113 // HOST: %[[VAL:.*]] = load i32, ptr @global_host_var
 114 // HOST: store i32 %[[VAL]]
 115 void host_lambda_ref(int *out) {
 116   [=](){
 117     int &ref = global_host_var;
 118     ref++;
 119     *out = ref;
 120   }();
 121 }
 122
 123 // HOST-LABEL: define{{.*}} void @_Z28dev_capture_host_ref_by_copyPi(
 124 // HOST: %[[CAP:.*]] = getelementptr inbounds %[[T3]], ptr %{{.*}}, i32 0, i32 1
 125 // HOST: %[[VAL:.*]] = load i32, ptr @global_host_var
 126 // HOST: store i32 %[[VAL]], ptr %[[CAP]]
 127 // DEV-LABEL: define internal void @_ZZ28dev_capture_host_ref_by_copyPiENKUlvE_clEv(
 128 // DEV: %[[CAP:.*]] = getelementptr inbounds %[[T3]], ptr %this1, i32 0, i32 1
 129 // DEV: %[[VAL:.*]] = load i32, ptr %[[CAP]]
 130 // DEV: store i32 %[[VAL]]
 131 void dev_capture_host_ref_by_copy(int *out) {
 132   int &ref = global_host_var;
 133   kern<<<1, 1>>>([=]__device__() { *out = ref;});
 134 }
 135