Skip to content

Commit 44db66a

Browse files
wenju-hewenju-he
committed
check device code
1 parent b42fd22 commit 44db66a

File tree

2 files changed

+105
-0
lines changed

2 files changed

+105
-0
lines changed

‎llvm/test/SYCLLowerIR/group_local_memory_inline.ll‎

Lines changed: 65 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,65 @@
1+
; RUN: opt < %s -passes=sycllowerwglocalmemory -S | FileCheck %s
2+
3+
; Check group_local_memory_for_overwrite and group_local_memory functions are inlined.
4+
; Check __sycl_allocateLocalMemory calls are lowered to four separate allocations.
5+
6+
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-G1"
7+
target triple = "spir64-unknown-unknown"
8+
9+
%"class.sycl::_V1::multi_ptr" = type { ptr addrspace(3) }
10+
%"class.sycl::_V1::group" = type { %"class.sycl::_V1::range", %"class.sycl::_V1::range", %"class.sycl::_V1::range", %"class.sycl::_V1::id" }
11+
%"class.sycl::_V1::range" = type { %"class.sycl::_V1::detail::array" }
12+
%"class.sycl::_V1::detail::array" = type { [1 x i64] }
13+
%"class.sycl::_V1::id" = type { %"class.sycl::_V1::detail::array" }
14+
15+
; CHECK: @WGLocalMem{{.*}} = internal addrspace(3) global [4 x i8] undef, align 4
16+
; CHECK: @WGLocalMem{{.*}} = internal addrspace(3) global [4 x i8] undef, align 4
17+
; CHECK: @WGLocalMem{{.*}} = internal addrspace(3) global [4 x i8] undef, align 4
18+
; CHECK: @WGLocalMem{{.*}} = internal addrspace(3) global [4 x i8] undef, align 4
19+
20+
; Function Attrs: alwaysinline
21+
define internal spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_ENKUlNS0_7nd_itemILi1EEEE_clES5_() #0 {
22+
entry:
23+
; CHECK: define internal spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_ENKUlNS0_7nd_itemILi1EEEE_clES5_(
24+
; CHECK: store ptr addrspace(3) @WGLocalMem{{.*}}, ptr addrspace(4) %AllocatedMem{{.*}}, align 8
25+
; CHECK: store ptr addrspace(3) @WGLocalMem{{.*}}, ptr addrspace(4) %AllocatedMem{{.*}}, align 8
26+
; CHECK: store ptr addrspace(3) @WGLocalMem{{.*}}, ptr addrspace(4) %AllocatedMem{{.*}}, align 8
27+
; CHECK: store ptr addrspace(3) @WGLocalMem{{.*}}, ptr addrspace(4) %AllocatedMem{{.*}}, align 8
28+
29+
%Ptr = alloca %"class.sycl::_V1::multi_ptr", align 8
30+
%agg = alloca %"class.sycl::_V1::group", align 8
31+
%Ptr.ascast = addrspacecast ptr %Ptr to ptr addrspace(4)
32+
call spir_func void @_ZN4sycl3_V13ext6oneapi32group_local_memory_for_overwriteIiNS0_5groupILi1EEEEENSt9enable_ifIXaasr3stdE27is_trivially_destructible_vIT_Esr4sycl6detail8is_groupIT0_EE5valueENS0_9multi_ptrIS7_LNS0_6access13address_spaceE3ELNSA_9decoratedE2EEEE4typeES8_(ptr addrspace(4) sret(%"class.sycl::_V1::multi_ptr") align 8 %Ptr.ascast, ptr byval(%"class.sycl::_V1::group") align 8 %agg)
33+
call spir_func void @_ZN4sycl3_V13ext6oneapi32group_local_memory_for_overwriteIiNS0_5groupILi1EEEEENSt9enable_ifIXaasr3stdE27is_trivially_destructible_vIT_Esr4sycl6detail8is_groupIT0_EE5valueENS0_9multi_ptrIS7_LNS0_6access13address_spaceE3ELNSA_9decoratedE2EEEE4typeES8_(ptr addrspace(4) sret(%"class.sycl::_V1::multi_ptr") align 8 %Ptr.ascast, ptr byval(%"class.sycl::_V1::group") align 8 %agg)
34+
call spir_func void @_ZN4sycl3_V13ext6oneapi18group_local_memoryIiNS0_5groupILi1EEEJEEENSt9enable_ifIXaasr3stdE27is_trivially_destructible_vIT_Esr4sycl6detail8is_groupIT0_EE5valueENS0_9multi_ptrIS7_LNS0_6access13address_spaceE3ELNSA_9decoratedE2EEEE4typeES8_DpOT1_(ptr addrspace(4) sret(%"class.sycl::_V1::multi_ptr") align 8 %Ptr.ascast, ptr byval(%"class.sycl::_V1::group") align 8 %agg)
35+
call spir_func void @_ZN4sycl3_V13ext6oneapi18group_local_memoryIiNS0_5groupILi1EEEJEEENSt9enable_ifIXaasr3stdE27is_trivially_destructible_vIT_Esr4sycl6detail8is_groupIT0_EE5valueENS0_9multi_ptrIS7_LNS0_6access13address_spaceE3ELNSA_9decoratedE2EEEE4typeES8_DpOT1_(ptr addrspace(4) sret(%"class.sycl::_V1::multi_ptr") align 8 %Ptr.ascast, ptr byval(%"class.sycl::_V1::group") align 8 %agg)
36+
ret void
37+
}
38+
39+
; CHECK-NOT: define {{.*}} @_ZN4sycl3_V13ext6oneapi32group_local_memory_for_overwriteIiNS0_5groupILi1EEEEENSt9enable_ifIXaasr3stdE27is_trivially_destructible_vIT_Esr4sycl6detail8is_groupIT0_EE5valueENS0_9multi_ptrIS7_LNS0_6access13address_spaceE3ELNSA_9decoratedE2EEEE4typeES8_(
40+
41+
; Function Attrs: alwaysinline
42+
define spir_func void @_ZN4sycl3_V13ext6oneapi32group_local_memory_for_overwriteIiNS0_5groupILi1EEEEENSt9enable_ifIXaasr3stdE27is_trivially_destructible_vIT_Esr4sycl6detail8is_groupIT0_EE5valueENS0_9multi_ptrIS7_LNS0_6access13address_spaceE3ELNSA_9decoratedE2EEEE4typeES8_(ptr addrspace(4) sret(%"class.sycl::_V1::multi_ptr") align 8 %agg.result, ptr byval(%"class.sycl::_V1::group") align 8 %g) #0 {
43+
entry:
44+
%AllocatedMem = alloca ptr addrspace(3), align 8
45+
%AllocatedMem.ascast = addrspacecast ptr %AllocatedMem to ptr addrspace(4)
46+
%call = call spir_func ptr addrspace(3) @__sycl_allocateLocalMemory(i64 4, i64 4)
47+
store ptr addrspace(3) %call, ptr addrspace(4) %AllocatedMem.ascast, align 8
48+
ret void
49+
}
50+
51+
; CHECK-NOT: define {{.*}} @_ZN4sycl3_V13ext6oneapi18group_local_memoryIiNS0_5groupILi1EEEJEEENSt9enable_ifIXaasr3stdE27is_trivially_destructible_vIT_Esr4sycl6detail8is_groupIT0_EE5valueENS0_9multi_ptrIS7_LNS0_6access13address_spaceE3ELNSA_9decoratedE2EEEE4typeES8_DpOT1_(
52+
53+
; Function Attrs: alwaysinline
54+
define spir_func void @_ZN4sycl3_V13ext6oneapi18group_local_memoryIiNS0_5groupILi1EEEJEEENSt9enable_ifIXaasr3stdE27is_trivially_destructible_vIT_Esr4sycl6detail8is_groupIT0_EE5valueENS0_9multi_ptrIS7_LNS0_6access13address_spaceE3ELNSA_9decoratedE2EEEE4typeES8_DpOT1_(ptr addrspace(4) sret(%"class.sycl::_V1::multi_ptr") align 8 %agg.result, ptr byval(%"class.sycl::_V1::group") align 8 %g) #0 {
55+
entry:
56+
%AllocatedMem = alloca ptr addrspace(3), align 8
57+
%AllocatedMem.ascast = addrspacecast ptr %AllocatedMem to ptr addrspace(4)
58+
%call = call spir_func ptr addrspace(3) @__sycl_allocateLocalMemory(i64 4, i64 4)
59+
store ptr addrspace(3) %call, ptr addrspace(4) %AllocatedMem.ascast, align 8
60+
ret void
61+
}
62+
63+
declare spir_func ptr addrspace(3) @__sycl_allocateLocalMemory(i64 noundef, i64 noundef)
64+
65+
attributes #0 = { alwaysinline }

‎sycl/test/check_device_code/extensions/group_local_memory.cpp‎

Lines changed: 40 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,40 @@
1+
// RUN: %clangxx -fsycl -fsycl-device-only -S -emit-llvm %s -o - | FileCheck %s
2+
// RUN: %clangxx -fsycl -fsycl-device-only -S -emit-llvm %s -fno-sycl-early-optimizations -o - | FileCheck %s
3+
// RUN: %clangxx -fsycl -fsycl-device-only -S -emit-llvm %s -O0 -o - | FileCheck %s
4+
5+
// The test checks that multiple calls to the same template instantiation of a
6+
// group local memory function result in separate allocations.
7+
8+
#include <sycl/detail/core.hpp>
9+
#include <sycl/ext/oneapi/group_local_memory.hpp>
10+
#include <sycl/usm.hpp>
11+
12+
using namespace sycl;
13+
14+
int main() {
15+
queue Q;
16+
17+
int **Out = malloc_shared<int *>(4, Q);
18+
19+
Q.submit([&](handler &Cgh) {
20+
Cgh.parallel_for(nd_range<1>({1}, {1}), [=](nd_item<1> Item) {
21+
auto Ptr0 = ext::oneapi::group_local_memory_for_overwrite<int>(
22+
Item.get_group());
23+
auto Ptr1 = ext::oneapi::group_local_memory_for_overwrite<int>(
24+
Item.get_group());
25+
auto Ptr2 = ext::oneapi::group_local_memory<int>(
26+
Item.get_group());
27+
auto Ptr3 = ext::oneapi::group_local_memory<int>(
28+
Item.get_group());
29+
Out[0] = Ptr0;
30+
Out[1] = Ptr1;
31+
Out[2] = Ptr2;
32+
Out[3] = Ptr3;
33+
});
34+
});
35+
}
36+
37+
// CHECK: @WGLocalMem{{.*}} = internal addrspace(3) global [4 x i8] undef, align 4
38+
// CHECK-NEXT: @WGLocalMem{{.*}} = internal addrspace(3) global [4 x i8] undef, align 4
39+
// CHECK-NEXT: @WGLocalMem{{.*}} = internal addrspace(3) global [4 x i8] undef, align 4
40+
// CHECK-NEXT: @WGLocalMem{{.*}} = internal addrspace(3) global [4 x i8] undef, align 4

0 commit comments

Comments
 (0)