Skip to content

Commit 87d4137

Browse files
author
Artem Gindinson
committed
Add LIT tests
Signed-off-by: Artem Gindinson <artem.gindinson@intel.com>
1 parent 4d1c6be commit 87d4137

File tree

3 files changed

+777
-0
lines changed

3 files changed

+777
-0
lines changed

‎llvm/test/SYCLLowerIR/printf_generic_as.ll‎

Lines changed: 101 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,101 @@
1+
;; This tests replacement of string literal address space for __spirv_ocl_printf
2+
;; at the regular O2 optimization level.
3+
4+
;; Compiled with the following command (custom build of SYCL Clang with
5+
;; SYCLMutatePrintfAddrspacePass turned off):
6+
;; clang++ -fsycl -fsycl-device-only experimental-printf.cpp -S -D__SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__
7+
;;
8+
;; // experimental-printf.cpp
9+
;; #include <CL/sycl.hpp>
10+
;; using namespace sycl;
11+
;; int main() {
12+
;; queue q;
13+
;; q.submit([&](handler &cgh) {
14+
;; cgh.single_task([=]() {
15+
;; ext::oneapi::experimental::printf("String No. %f\n", 1.0f);
16+
;; const char *IntFormatString = "String No. %i\n";
17+
;; ext::oneapi::experimental::printf(IntFormatString, 2);
18+
;; ext::oneapi::experimental::printf(IntFormatString, 3);
19+
;; });
20+
;; });
21+
;; return 0;
22+
;; }
23+
24+
; RUN: opt < %s --SYCLMutatePrintfAddrspace -S | FileCheck %s
25+
26+
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
27+
target triple = "spir64-unknown-unknown"
28+
29+
%"struct.cl::sycl::detail::AssertHappened" = type { i32, [257 x i8], [257 x i8], [129 x i8], i32, i64, i64, i64, i64, i64, i64 }
30+
%"class.cl::sycl::range" = type { %"class.cl::sycl::detail::array" }
31+
%"class.cl::sycl::detail::array" = type { [1 x i64] }
32+
%"class.cl::sycl::id" = type { %"class.cl::sycl::detail::array" }
33+
34+
$_ZTSN2cl4sycl6detail23__sycl_service_kernel__16AssertInfoCopierE = comdat any
35+
36+
$_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_EUlvE_ = comdat any
37+
38+
; CHECK-DAG: @.str._AS2 = internal addrspace(2) constant [15 x i8] c"String No. %f\0A\00", align 1
39+
@.str = private unnamed_addr addrspace(1) constant [15 x i8] c"String No. %f\0A\00", align 1
40+
; CHECK-DAG: @.str.1._AS2 = internal addrspace(2) constant [15 x i8] c"String No. %i\0A\00", align 1
41+
@.str.1 = private unnamed_addr addrspace(1) constant [15 x i8] c"String No. %i\0A\00", align 1
42+
43+
; Function Attrs: convergent norecurse
44+
define weak_odr dso_local spir_kernel void @_ZTSN2cl4sycl6detail23__sycl_service_kernel__16AssertInfoCopierE(%"struct.cl::sycl::detail::AssertHappened" addrspace(1)* %_arg_, %"class.cl::sycl::range"* byval(%"class.cl::sycl::range") align 8 %_arg_1, %"class.cl::sycl::range"* byval(%"class.cl::sycl::range") align 8 %_arg_2, %"class.cl::sycl::id"* byval(%"class.cl::sycl::id") align 8 %_arg_3) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !5 {
45+
entry:
46+
%0 = getelementptr inbounds %"class.cl::sycl::id", %"class.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0
47+
%1 = addrspacecast i64* %0 to i64 addrspace(4)*
48+
%2 = load i64, i64 addrspace(4)* %1, align 8
49+
%add.ptr.i = getelementptr inbounds %"struct.cl::sycl::detail::AssertHappened", %"struct.cl::sycl::detail::AssertHappened" addrspace(1)* %_arg_, i64 %2
50+
%3 = bitcast %"struct.cl::sycl::detail::AssertHappened" addrspace(1)* %add.ptr.i to i8 addrspace(1)*
51+
%4 = addrspacecast i8 addrspace(1)* %3 to i8 addrspace(4)*
52+
tail call spir_func void @__devicelib_assert_read(i8 addrspace(4)* %4) #3
53+
ret void
54+
}
55+
56+
; Function Attrs: convergent
57+
declare extern_weak dso_local spir_func void @__devicelib_assert_read(i8 addrspace(4)*) local_unnamed_addr #1
58+
59+
; Function Attrs: convergent mustprogress norecurse
60+
define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_EUlvE_() local_unnamed_addr #2 comdat !kernel_arg_buffer_location !6 {
61+
entry:
62+
; In particular, make sure that no argument promotion has been done for float
63+
; upon variadic redeclaration:
64+
; CHECK: tail call spir_func i32 (i8 addrspace(2)*, ...) @_Z18__spirv_ocl_printfPU3AS2Kcz(i8 addrspace(2)* getelementptr inbounds ([15 x i8], [15 x i8] addrspace(2)* @.str._AS2, i32 0, i32 0), float 1.000000e+00) #3
65+
%call.i.i = tail call spir_func i32 @_Z18__spirv_ocl_printfIJfEEiPKcDpT_(i8 addrspace(4)* getelementptr inbounds ([15 x i8], [15 x i8] addrspace(4)* addrspacecast ([15 x i8] addrspace(1)* @.str to [15 x i8] addrspace(4)*), i64 0, i64 0), float 1.000000e+00) #3
66+
; CHECK-NEXT: tail call spir_func i32 (i8 addrspace(2)*, ...) @_Z18__spirv_ocl_printfPU3AS2Kcz(i8 addrspace(2)* getelementptr inbounds ([15 x i8], [15 x i8] addrspace(2)* @.str.1._AS2, i32 0, i32 0), i32 2) #3
67+
%call.i1.i = tail call spir_func i32 @_Z18__spirv_ocl_printfIJiEEiPKcDpT_(i8 addrspace(4)* getelementptr inbounds ([15 x i8], [15 x i8] addrspace(4)* addrspacecast ([15 x i8] addrspace(1)* @.str.1 to [15 x i8] addrspace(4)*), i64 0, i64 0), i32 2) #3
68+
; CHECK-NEXT: tail call spir_func i32 (i8 addrspace(2)*, ...) @_Z18__spirv_ocl_printfPU3AS2Kcz(i8 addrspace(2)* getelementptr inbounds ([15 x i8], [15 x i8] addrspace(2)* @.str.1._AS2, i32 0, i32 0), i32 3) #3
69+
%call.i2.i = tail call spir_func i32 @_Z18__spirv_ocl_printfIJiEEiPKcDpT_(i8 addrspace(4)* getelementptr inbounds ([15 x i8], [15 x i8] addrspace(4)* addrspacecast ([15 x i8] addrspace(1)* @.str.1 to [15 x i8] addrspace(4)*), i64 0, i64 0), i32 3) #3
70+
ret void
71+
}
72+
73+
; Make sure the non-variadic declarations have been wiped out
74+
; in favor of the single variadic one:
75+
; CHECK-NOT: declare dso_local spir_func i32 @_Z18__spirv_ocl_printf{{.*}}(i8 addrspace(4)*, float)
76+
; CHECK-NOT: declare dso_local spir_func i32 @_Z18__spirv_ocl_printf{{.*}}(i8 addrspace(4)*, i32)
77+
; CHECK: declare dso_local spir_func i32 @_Z18__spirv_ocl_printfPU3AS2Kcz(i8 addrspace(2)*, ...) #1
78+
79+
; Function Attrs: convergent
80+
declare dso_local spir_func i32 @_Z18__spirv_ocl_printfIJfEEiPKcDpT_(i8 addrspace(4)*, float) local_unnamed_addr #1
81+
82+
; Function Attrs: convergent
83+
declare dso_local spir_func i32 @_Z18__spirv_ocl_printfIJiEEiPKcDpT_(i8 addrspace(4)*, i32) local_unnamed_addr #1
84+
85+
attributes #0 = { convergent norecurse "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="experimental-printf.cpp" "uniform-work-group-size"="true" }
86+
attributes #1 = { convergent "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
87+
attributes #2 = { convergent mustprogress norecurse "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="experimental-printf.cpp" "uniform-work-group-size"="true" }
88+
attributes #3 = { convergent }
89+
90+
!llvm.module.flags = !{!0, !1}
91+
!opencl.spir.version = !{!2}
92+
!spirv.Source = !{!3}
93+
!llvm.ident = !{!4}
94+
95+
!0 = !{i32 1, !"wchar_size", i32 4}
96+
!1 = !{i32 7, !"frame-pointer", i32 2}
97+
!2 = !{i32 1, i32 2}
98+
!3 = !{i32 4, i32 100000}
99+
!4 = !{!"clang version 14.0.0"}
100+
!5 = !{i32 -1, i32 -1, i32 -1, i32 -1}
101+
!6 = !{}

0 commit comments

Comments
 (0)