clang/test/OpenMP/nvptx_target_printf_codegen.c

   1 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _
   2 // Test target codegen - host bc file has to be created first.
   3 // RUN: %clang_cc1 -verify -fopenmp -x c -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
   4 // RUN: %clang_cc1 -verify -fopenmp -x c -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=CHECK-64
   5 // RUN: %clang_cc1 -verify -fopenmp -x c -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
   6 // RUN: %clang_cc1 -verify -fopenmp -x c -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix=CHECK-32
   7 // expected-no-diagnostics
   8 extern int printf(const char *, ...);
   9
  10
  11 // Check a simple call to printf end-to-end.
  12 int CheckSimple(void) {
  13 #pragma omp target
  14   {
  15     // printf in master-only basic block.
  16     const char* fmt = "%d %lld %f";
  17
  18     printf(fmt, 1, 2ll, 3.0);
  19   }
  20
  21   return 0;
  22 }
  23
  24 void CheckNoArgs(void) {
  25 #pragma omp target
  26   {
  27     // printf in master-only basic block.
  28     printf("hello, world!");
  29   }
  30 }
  31
  32 // Check that printf's alloca happens in the entry block, not inside the if
  33 // statement.
  34 int foo;
  35 void CheckAllocaIsInEntryBlock(void) {
  36 #pragma omp target
  37   {
  38     if (foo) {
  39       printf("%d", 42);
  40     }
  41   }
  42 }
  43 // CHECK-64-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_CheckSimple_l13
  44 // CHECK-64-SAME: (ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR0:[0-9]+]] {
  45 // CHECK-64-NEXT:  entry:
  46 // CHECK-64-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8
  47 // CHECK-64-NEXT:    [[FMT:%.*]] = alloca ptr, align 8
  48 // CHECK-64-NEXT:    [[TMP:%.*]] = alloca [[PRINTF_ARGS:%.*]], align 8
  49 // CHECK-64-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8
  50 // CHECK-64-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_CheckSimple_l13_kernel_environment, ptr [[DYN_PTR]])
  51 // CHECK-64-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
  52 // CHECK-64-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
  53 // CHECK-64:       user_code.entry:
  54 // CHECK-64-NEXT:    store ptr @.str, ptr [[FMT]], align 8
  55 // CHECK-64-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[FMT]], align 8
  56 // CHECK-64-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [[PRINTF_ARGS]], ptr [[TMP]], i32 0, i32 0
  57 // CHECK-64-NEXT:    store i32 1, ptr [[TMP2]], align 4
  58 // CHECK-64-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [[PRINTF_ARGS]], ptr [[TMP]], i32 0, i32 1
  59 // CHECK-64-NEXT:    store i64 2, ptr [[TMP3]], align 8
  60 // CHECK-64-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [[PRINTF_ARGS]], ptr [[TMP]], i32 0, i32 2
  61 // CHECK-64-NEXT:    store double 3.000000e+00, ptr [[TMP4]], align 8
  62 // CHECK-64-NEXT:    [[TMP5:%.*]] = call i32 @__llvm_omp_vprintf(ptr [[TMP1]], ptr [[TMP]], i32 24)
  63 // CHECK-64-NEXT:    call void @__kmpc_target_deinit()
  64 // CHECK-64-NEXT:    ret void
  65 // CHECK-64:       worker.exit:
  66 // CHECK-64-NEXT:    ret void
  67 //
  68 //
  69 // CHECK-64-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_CheckNoArgs_l25
  70 // CHECK-64-SAME: (ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR0]] {
  71 // CHECK-64-NEXT:  entry:
  72 // CHECK-64-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8
  73 // CHECK-64-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8
  74 // CHECK-64-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_CheckNoArgs_l25_kernel_environment, ptr [[DYN_PTR]])
  75 // CHECK-64-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
  76 // CHECK-64-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
  77 // CHECK-64:       user_code.entry:
  78 // CHECK-64-NEXT:    [[TMP1:%.*]] = call i32 @__llvm_omp_vprintf(ptr @.str1, ptr null, i32 0)
  79 // CHECK-64-NEXT:    call void @__kmpc_target_deinit()
  80 // CHECK-64-NEXT:    ret void
  81 // CHECK-64:       worker.exit:
  82 // CHECK-64-NEXT:    ret void
  83 //
  84 //
  85 // CHECK-64-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_CheckAllocaIsInEntryBlock_l36
  86 // CHECK-64-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i64 noundef [[FOO:%.*]]) #[[ATTR0]] {
  87 // CHECK-64-NEXT:  entry:
  88 // CHECK-64-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 8
  89 // CHECK-64-NEXT:    [[FOO_ADDR:%.*]] = alloca i64, align 8
  90 // CHECK-64-NEXT:    [[TMP:%.*]] = alloca [[PRINTF_ARGS_0:%.*]], align 8
  91 // CHECK-64-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 8
  92 // CHECK-64-NEXT:    store i64 [[FOO]], ptr [[FOO_ADDR]], align 8
  93 // CHECK-64-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_CheckAllocaIsInEntryBlock_l36_kernel_environment, ptr [[DYN_PTR]])
  94 // CHECK-64-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
  95 // CHECK-64-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
  96 // CHECK-64:       user_code.entry:
  97 // CHECK-64-NEXT:    [[TMP1:%.*]] = load i32, ptr [[FOO_ADDR]], align 4
  98 // CHECK-64-NEXT:    [[TOBOOL:%.*]] = icmp ne i32 [[TMP1]], 0
  99 // CHECK-64-NEXT:    br i1 [[TOBOOL]], label [[IF_THEN:%.*]], label [[IF_END:%.*]]
 100 // CHECK-64:       if.then:
 101 // CHECK-64-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [[PRINTF_ARGS_0]], ptr [[TMP]], i32 0, i32 0
 102 // CHECK-64-NEXT:    store i32 42, ptr [[TMP2]], align 4
 103 // CHECK-64-NEXT:    [[TMP3:%.*]] = call i32 @__llvm_omp_vprintf(ptr @.str2, ptr [[TMP]], i32 4)
 104 // CHECK-64-NEXT:    br label [[IF_END]]
 105 // CHECK-64:       worker.exit:
 106 // CHECK-64-NEXT:    ret void
 107 // CHECK-64:       if.end:
 108 // CHECK-64-NEXT:    call void @__kmpc_target_deinit()
 109 // CHECK-64-NEXT:    ret void
 110 //
 111 //
 112 // CHECK-32-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_CheckSimple_l13
 113 // CHECK-32-SAME: (ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR0:[0-9]+]] {
 114 // CHECK-32-NEXT:  entry:
 115 // CHECK-32-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4
 116 // CHECK-32-NEXT:    [[FMT:%.*]] = alloca ptr, align 4
 117 // CHECK-32-NEXT:    [[TMP:%.*]] = alloca [[PRINTF_ARGS:%.*]], align 8
 118 // CHECK-32-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4
 119 // CHECK-32-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_CheckSimple_l13_kernel_environment, ptr [[DYN_PTR]])
 120 // CHECK-32-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
 121 // CHECK-32-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
 122 // CHECK-32:       user_code.entry:
 123 // CHECK-32-NEXT:    store ptr @.str, ptr [[FMT]], align 4
 124 // CHECK-32-NEXT:    [[TMP1:%.*]] = load ptr, ptr [[FMT]], align 4
 125 // CHECK-32-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [[PRINTF_ARGS]], ptr [[TMP]], i32 0, i32 0
 126 // CHECK-32-NEXT:    store i32 1, ptr [[TMP2]], align 4
 127 // CHECK-32-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [[PRINTF_ARGS]], ptr [[TMP]], i32 0, i32 1
 128 // CHECK-32-NEXT:    store i64 2, ptr [[TMP3]], align 8
 129 // CHECK-32-NEXT:    [[TMP4:%.*]] = getelementptr inbounds [[PRINTF_ARGS]], ptr [[TMP]], i32 0, i32 2
 130 // CHECK-32-NEXT:    store double 3.000000e+00, ptr [[TMP4]], align 8
 131 // CHECK-32-NEXT:    [[TMP5:%.*]] = call i32 @__llvm_omp_vprintf(ptr [[TMP1]], ptr [[TMP]], i32 24)
 132 // CHECK-32-NEXT:    call void @__kmpc_target_deinit()
 133 // CHECK-32-NEXT:    ret void
 134 // CHECK-32:       worker.exit:
 135 // CHECK-32-NEXT:    ret void
 136 //
 137 //
 138 // CHECK-32-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_CheckNoArgs_l25
 139 // CHECK-32-SAME: (ptr noalias noundef [[DYN_PTR:%.*]]) #[[ATTR0]] {
 140 // CHECK-32-NEXT:  entry:
 141 // CHECK-32-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4
 142 // CHECK-32-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4
 143 // CHECK-32-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_CheckNoArgs_l25_kernel_environment, ptr [[DYN_PTR]])
 144 // CHECK-32-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
 145 // CHECK-32-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
 146 // CHECK-32:       user_code.entry:
 147 // CHECK-32-NEXT:    [[TMP1:%.*]] = call i32 @__llvm_omp_vprintf(ptr @.str1, ptr null, i32 0)
 148 // CHECK-32-NEXT:    call void @__kmpc_target_deinit()
 149 // CHECK-32-NEXT:    ret void
 150 // CHECK-32:       worker.exit:
 151 // CHECK-32-NEXT:    ret void
 152 //
 153 //
 154 // CHECK-32-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_CheckAllocaIsInEntryBlock_l36
 155 // CHECK-32-SAME: (ptr noalias noundef [[DYN_PTR:%.*]], i32 noundef [[FOO:%.*]]) #[[ATTR0]] {
 156 // CHECK-32-NEXT:  entry:
 157 // CHECK-32-NEXT:    [[DYN_PTR_ADDR:%.*]] = alloca ptr, align 4
 158 // CHECK-32-NEXT:    [[FOO_ADDR:%.*]] = alloca i32, align 4
 159 // CHECK-32-NEXT:    [[TMP:%.*]] = alloca [[PRINTF_ARGS_0:%.*]], align 8
 160 // CHECK-32-NEXT:    store ptr [[DYN_PTR]], ptr [[DYN_PTR_ADDR]], align 4
 161 // CHECK-32-NEXT:    store i32 [[FOO]], ptr [[FOO_ADDR]], align 4
 162 // CHECK-32-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_target_init(ptr @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}_CheckAllocaIsInEntryBlock_l36_kernel_environment, ptr [[DYN_PTR]])
 163 // CHECK-32-NEXT:    [[EXEC_USER_CODE:%.*]] = icmp eq i32 [[TMP0]], -1
 164 // CHECK-32-NEXT:    br i1 [[EXEC_USER_CODE]], label [[USER_CODE_ENTRY:%.*]], label [[WORKER_EXIT:%.*]]
 165 // CHECK-32:       user_code.entry:
 166 // CHECK-32-NEXT:    [[TMP1:%.*]] = load i32, ptr [[FOO_ADDR]], align 4
 167 // CHECK-32-NEXT:    [[TOBOOL:%.*]] = icmp ne i32 [[TMP1]], 0
 168 // CHECK-32-NEXT:    br i1 [[TOBOOL]], label [[IF_THEN:%.*]], label [[IF_END:%.*]]
 169 // CHECK-32:       if.then:
 170 // CHECK-32-NEXT:    [[TMP2:%.*]] = getelementptr inbounds [[PRINTF_ARGS_0]], ptr [[TMP]], i32 0, i32 0
 171 // CHECK-32-NEXT:    store i32 42, ptr [[TMP2]], align 4
 172 // CHECK-32-NEXT:    [[TMP3:%.*]] = call i32 @__llvm_omp_vprintf(ptr @.str2, ptr [[TMP]], i32 4)
 173 // CHECK-32-NEXT:    br label [[IF_END]]
 174 // CHECK-32:       worker.exit:
 175 // CHECK-32-NEXT:    ret void
 176 // CHECK-32:       if.end:
 177 // CHECK-32-NEXT:    call void @__kmpc_target_deinit()
 178 // CHECK-32-NEXT:    ret void
 179 //