Skip to content

Commit a2998c5

Browse files
victor-edsvictor-eds
committed
Improve testing
1 parent 606be48 commit a2998c5

File tree

9 files changed

+62
-93
lines changed

9 files changed

+62
-93
lines changed

‎clang/include/clang/Basic/DiagnosticSemaKinds.td

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -179,7 +179,7 @@ def err_intel_sycl_ptr_annotation_mismatch
179179

180180
def err_intel_sycl_alloca_no_alias
181181
: Error<"__builtin_intel_sycl_alloca cannot be used in source code. "
182-
"Use the private_alloca alias instead.">;
182+
"Use the private_alloca alias instead">;
183183
def err_intel_sycl_alloca_wrong_arg_count
184184
: Error<"__builtin_intel_sycl_alloca expects to be passed a single "
185185
"argument. Got %0">;

‎clang/lib/CodeGen/CGBuiltin.cpp

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -23727,9 +23727,8 @@ CodeGenFunction::EmitIntelSYCLAllocaBuiltin(const CallExpr *E,
2372723727
}
2372823728

2372923729
// If no slot is provided, simply return allocation.
23730-
if (ReturnValue.isNull()) {
23730+
if (ReturnValue.isNull())
2373123731
return RValue::get(Allocation);
23732-
}
2373323732

2373423733
// If a slot is provided, store pointer there.
2373523734
Builder.CreateStore(Allocation, ReturnValue.getValue());

‎clang/lib/Sema/SemaChecking.cpp

Lines changed: 2 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -7578,11 +7578,10 @@ bool Sema::CheckIntelSYCLAllocaBuiltinFunctionCall(unsigned, CallExpr *Call) {
75787578
if (CheckSize(getASTContext(), CST)) {
75797579
TemplateArgument TA = CST->get(1);
75807580
QualType Ty = TA.getNonTypeTemplateArgumentType();
7581-
if (Ty.isNull()) {
7581+
if (Ty.isNull())
75827582
Diag(Loc, diag::err_intel_sycl_alloca_no_size) << TA;
7583-
} else {
7583+
else
75847584
Diag(Loc, diag::err_intel_sycl_alloca_wrong_size) << TA << Ty;
7585-
}
75867585
return true;
75877586
}
75887587

‎clang/test/CodeGenSYCL/builtin-alloca.cpp

Lines changed: 26 additions & 64 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
11
// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s \
2-
// RUN: | FileCheck %s --check-prefix=CHECK-DEVICE
3-
// RUN: %clang_cc1 -fsycl-is-host -triple x86_64-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s -fcxx-exceptions\
4-
// RUN: | FileCheck %s --check-prefix=CHECK-HOST
2+
// RUN: | FileCheck %s
3+
4+
// Test codegen for __builtin_intel_sycl_alloca.
55

66
#include <stddef.h>
77

@@ -17,70 +17,32 @@ struct myStruct {
1717

1818
constexpr sycl::specialization_id<size_t> size(1);
1919
constexpr sycl::specialization_id<int> intSize(-1);
20-
constexpr sycl::specialization_id<unsigned short> shortSize(1);
21-
22-
// COM: On the device, for each call, we should generate a chain of: 'call @sycl.alloca.<ty>' + ('addrspacecast') + 'store'.
23-
// COM: The 'addrspacecast' will only appear when the pointer is not decorated, i.e., `DecorateAddress == sycl::access::decorated::no`.
24-
25-
// CHECK-DEVICE-LABEL: define dso_local spir_func void @_Z4testRN4sycl3_V114kernel_handlerE(
26-
// CHECK-DEVICE-SAME: ptr addrspace(4) noundef align 1 dereferenceable(1) [[KH:%.*]])
27-
// CHECK-DEVICE-NEXT: entry:
28-
// CHECK-DEVICE-NEXT: [[KH_ADDR:%.*]] = alloca ptr addrspace(4), align 8
29-
// CHECK-DEVICE-NEXT: [[PTR0:%.*]] = alloca %"class.sycl::_V1::multi_ptr", align 8
30-
// CHECK-DEVICE-NEXT: [[TMP0:%.*]] = call ptr @llvm.sycl.alloca.p0.p4.p4.p4.f64(ptr addrspace(4) addrspacecast (ptr {{.*}} to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) {{.*}} to ptr addrspace(4)), ptr addrspace(4) null, double 0.000000e+00, i64 8)
31-
// CHECK-DEVICE-NEXT: [[PTR1:%.*]] = alloca %"class.sycl::_V1::multi_ptr.0", align 8
32-
// CHECK-DEVICE-NEXT: [[TMP2:%.*]] = call ptr @llvm.sycl.alloca.p0.p4.p4.p4.i32(ptr addrspace(4) addrspacecast (ptr {{.*}} to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) {{.*}} to ptr addrspace(4)), ptr addrspace(4) null, i32 0, i64 4)
33-
// CHECK-DEVICE-NEXT: [[PTR2:%.*]] = alloca %"class.sycl::_V1::multi_ptr.2", align 8
34-
// CHECK-DEVICE-NEXT: [[TMP4:%.*]] = call ptr @llvm.sycl.alloca.p0.p4.p4.p4.s_struct.myStructs(ptr addrspace(4) addrspacecast (ptr {{.*}} to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) {{.*}} to ptr addrspace(4)), ptr addrspace(4) null, %struct.myStruct zeroinitializer, i64 1)
35-
// CHECK-DEVICE-NEXT: [[KH_ADDR_ASCAST:%.*]] = addrspacecast ptr [[KH_ADDR]] to ptr addrspace(4)
36-
// CHECK-DEVICE-NEXT: [[PTR0_ASCAST:%.*]] = addrspacecast ptr [[PTR0]] to ptr addrspace(4)
37-
// CHECK-DEVICE-NEXT: [[PTR1_ASCAST:%.*]] = addrspacecast ptr [[PTR1]] to ptr addrspace(4)
38-
// CHECK-DEVICE-NEXT: [[PTR2_ASCAST:%.*]] = addrspacecast ptr [[PTR2]] to ptr addrspace(4)
39-
// CHECK-DEVICE-NEXT: [[TMP5:%.*]] = addrspacecast ptr [[TMP4]] to ptr addrspace(4)
40-
// CHECK-DEVICE-NEXT: store ptr addrspace(4) [[KH]], ptr addrspace(4) [[KH_ADDR_ASCAST]], align 8
41-
// CHECK-DEVICE-NEXT: store ptr [[TMP0]], ptr addrspace(4) [[PTR0_ASCAST]], align 8
42-
// CHECK-DEVICE-NEXT: store ptr [[TMP2]], ptr addrspace(4) [[PTR1_ASCAST]], align 8
43-
// CHECK-DEVICE-NEXT: store ptr addrspace(4) [[TMP5]], ptr addrspace(4) [[PTR2_ASCAST]], align 8
44-
// CHECK-DEVICE-NEXT: ret void
4520

46-
// COM: On the host, each call should be materialized...
47-
48-
// CHECK-HOST-LABEL: define dso_local void @_Z4testRN4sycl3_V114kernel_handlerE(
49-
// CHECK-HOST-SAME: ptr noundef nonnull align 1 dereferenceable(1) [[KH:%.*]])
50-
// CHECK-HOST-NEXT: entry:
51-
// CHECK-HOST-NEXT: [[KH_ADDR:%.*]] = alloca ptr, align 8
52-
// CHECK-HOST-NEXT: [[PTR0:%.*]] = alloca %"class.sycl::_V1::multi_ptr", align 8
53-
// CHECK-HOST-NEXT: [[PTR1:%.*]] = alloca %"class.sycl::_V1::multi_ptr.0", align 8
54-
// CHECK-HOST-NEXT: [[PTR2:%.*]] = alloca %"class.sycl::_V1::multi_ptr.1", align 8
55-
// CHECK-HOST-NEXT: store ptr [[KH]], ptr [[KH_ADDR]], align 8
56-
// CHECK-HOST-NEXT: [[TMP0:%.*]] = load ptr, ptr [[KH_ADDR]], align 8
57-
// CHECK-HOST-NEXT: [[CALL:%.*]] = call ptr @_ZN4sycl3_V13ext6oneapi12experimental14private_allocaIdTnRDaL_ZL4sizeELNS0_6access9decoratedE1EEENS0_9multi_ptrIT_LNS6_13address_spaceE0EXT1_EEERNS0_14kernel_handlerE(ptr noundef nonnull align 1 dereferenceable(1) [[TMP0]])
58-
// CHECK-HOST-NEXT: [[COERCE_DIVE:%.*]] = getelementptr inbounds %"class.sycl::_V1::multi_ptr", ptr [[PTR0]], i32 0, i32 0
59-
// CHECK-HOST-NEXT: store ptr [[CALL]], ptr [[COERCE_DIVE]], align 8
60-
// CHECK-HOST-NEXT: [[TMP1:%.*]] = load ptr, ptr [[KH_ADDR]], align 8
61-
// CHECK-HOST-NEXT: [[CALL1:%.*]] = call ptr @_ZN4sycl3_V13ext6oneapi12experimental14private_allocaIiTnRDaL_ZL7intSizeELNS0_6access9decoratedE2EEENS0_9multi_ptrIT_LNS6_13address_spaceE0EXT1_EEERNS0_14kernel_handlerE(ptr noundef nonnull align 1 dereferenceable(1) [[TMP1]])
62-
// CHECK-HOST-NEXT: [[COERCE_DIVE2:%.*]] = getelementptr inbounds %"class.sycl::_V1::multi_ptr.0", ptr [[PTR1]], i32 0, i32 0
63-
// CHECK-HOST-NEXT: store ptr [[CALL1]], ptr [[COERCE_DIVE2]], align 8
64-
// CHECK-HOST-NEXT: [[TMP2:%.*]] = load ptr, ptr [[KH_ADDR]], align 8
65-
// CHECK-HOST-NEXT: [[CALL3:%.*]] = call ptr @_ZN4sycl3_V13ext6oneapi12experimental14private_allocaI8myStructTnRDaL_ZL7intSizeELNS0_6access9decoratedE0EEENS0_9multi_ptrIT_LNS7_13address_spaceE0EXT1_EEERNS0_14kernel_handlerE(ptr noundef nonnull align 1 dereferenceable(1) [[TMP2]])
66-
// CHECK-HOST-NEXT: [[COERCE_DIVE4:%.*]] = getelementptr inbounds %"class.sycl::_V1::multi_ptr.1", ptr [[PTR2]], i32 0, i32 0
67-
// CHECK-HOST-NEXT: store ptr [[CALL3]], ptr [[COERCE_DIVE4]], align 8
68-
// CHECK-HOST-NEXT: ret void
69-
//
21+
// For each call, we should generate a chain of: 'call @llvm.sycl.alloca.<ty>' + ('addrspacecast') + 'store'.
22+
// The 'addrspacecast' will only appear when the pointer is not decorated, i.e., `DecorateAddress == sycl::access::decorated::no`.
23+
24+
// CHECK-LABEL: define dso_local spir_func void @_Z4testRN4sycl3_V114kernel_handlerE(
25+
// CHECK-SAME: ptr addrspace(4) noundef align 1 dereferenceable(1) [[KH:%.*]])
26+
// CHECK-NEXT: entry:
27+
// CHECK-NEXT: [[KH_ADDR:%.*]] = alloca ptr addrspace(4), align 8
28+
// CHECK-NEXT: [[PTR0:%.*]] = alloca %"class.sycl::_V1::multi_ptr", align 8
29+
// CHECK-NEXT: [[TMP0:%.*]] = call ptr @llvm.sycl.alloca.p0.p4.p4.p4.f64(ptr addrspace(4) addrspacecast (ptr {{.*}} to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) {{.*}} to ptr addrspace(4)), ptr addrspace(4) null, double 0.000000e+00, i64 8)
30+
// CHECK-NEXT: [[PTR1:%.*]] = alloca %"class.sycl::_V1::multi_ptr.0", align 8
31+
// CHECK-NEXT: [[TMP2:%.*]] = call ptr @llvm.sycl.alloca.p0.p4.p4.p4.i32(ptr addrspace(4) addrspacecast (ptr {{.*}} to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) {{.*}} to ptr addrspace(4)), ptr addrspace(4) null, i32 0, i64 4)
32+
// CHECK-NEXT: [[PTR2:%.*]] = alloca %"class.sycl::_V1::multi_ptr.2", align 8
33+
// CHECK-NEXT: [[TMP4:%.*]] = call ptr @llvm.sycl.alloca.p0.p4.p4.p4.s_struct.myStructs(ptr addrspace(4) addrspacecast (ptr {{.*}} to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) {{.*}} to ptr addrspace(4)), ptr addrspace(4) null, %struct.myStruct zeroinitializer, i64 1)
34+
// CHECK-NEXT: [[KH_ADDR_ASCAST:%.*]] = addrspacecast ptr [[KH_ADDR]] to ptr addrspace(4)
35+
// CHECK-NEXT: [[PTR0_ASCAST:%.*]] = addrspacecast ptr [[PTR0]] to ptr addrspace(4)
36+
// CHECK-NEXT: [[PTR1_ASCAST:%.*]] = addrspacecast ptr [[PTR1]] to ptr addrspace(4)
37+
// CHECK-NEXT: [[PTR2_ASCAST:%.*]] = addrspacecast ptr [[PTR2]] to ptr addrspace(4)
38+
// CHECK-NEXT: [[TMP5:%.*]] = addrspacecast ptr [[TMP4]] to ptr addrspace(4)
39+
// CHECK-NEXT: store ptr addrspace(4) [[KH]], ptr addrspace(4) [[KH_ADDR_ASCAST]], align 8
40+
// CHECK-NEXT: store ptr [[TMP0]], ptr addrspace(4) [[PTR0_ASCAST]], align 8
41+
// CHECK-NEXT: store ptr [[TMP2]], ptr addrspace(4) [[PTR1_ASCAST]], align 8
42+
// CHECK-NEXT: store ptr addrspace(4) [[TMP5]], ptr addrspace(4) [[PTR2_ASCAST]], align 8
43+
// CHECK-NEXT: ret void
7044
SYCL_EXTERNAL void test(sycl::kernel_handler &kh) {
7145
auto ptr0 = sycl::ext::oneapi::experimental::private_alloca<double, size, sycl::access::decorated::yes>(kh);
7246
auto ptr1 = sycl::ext::oneapi::experimental::private_alloca<int, intSize, sycl::access::decorated::legacy>(kh);
7347
auto ptr2 = sycl::ext::oneapi::experimental::private_alloca<myStruct, intSize, sycl::access::decorated::no>(kh);
7448
}
75-
76-
// COM: And the body function should be simply a throw
77-
78-
// CHECK-HOST-LABEL: define internal ptr @_ZN4sycl3_V13ext6oneapi12experimental14private_allocaIdTnRDaL_ZL4sizeELNS0_6access9decoratedE1EEENS0_9multi_ptrIT_LNS6_13address_spaceE0EXT1_EEERNS0_14kernel_handlerE
79-
// CHECK-HOST-SAME: ptr noundef nonnull align 1 dereferenceable(1) [[H:%.*]])
80-
// CHECK-HOST-NEXT: entry:
81-
// CHECK-HOST-NEXT: [[H_ADDR:%.*]] = alloca ptr, align 8
82-
// CHECK-HOST-NEXT: store ptr [[H]], ptr [[H_ADDR]], align 8
83-
// CHECK-HOST-NEXT: [[EXCEPTION:%.*]] = call ptr @__cxa_allocate_exception(i64 8)
84-
// CHECK-HOST-NEXT: store ptr @.str, ptr [[EXCEPTION]], align 16
85-
// CHECK-HOST-NEXT: call void @__cxa_throw(ptr [[EXCEPTION]], ptr @_ZTIPKc, ptr null)
86-
// CHECK-HOST-NEXT: unreachable

‎clang/test/SemaSYCL/Inputs/private_alloca.hpp

Lines changed: 20 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,20 @@
1+
#pragma once
2+
3+
#include "./sycl.hpp"
4+
5+
namespace sycl {
6+
inline namespace _V1 {
7+
namespace ext {
8+
namespace oneapi {
9+
namespace experimental {
10+
11+
template <typename ElementType, auto &Size, access::decorated DecorateAddress>
12+
__SYCL_BUILTIN_ALIAS(__builtin_intel_sycl_alloca)
13+
multi_ptr<ElementType, access::address_space::private_space,
14+
DecorateAddress> private_alloca(kernel_handler &h);
15+
16+
} // namespace experimental
17+
} // namesapce oneapi
18+
} // namespace ext
19+
} // namespace _V1
20+
} // namespace sycl

‎clang/test/SemaSYCL/Inputs/sycl.hpp

Lines changed: 0 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -480,25 +480,6 @@ class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_ptr) annotated_p
480480
#endif
481481
};
482482

483-
#ifdef __SYCL_DEVICE_ONLY__
484-
485-
template <typename ElementType, auto &Size,
486-
access::decorated DecorateAddress>
487-
__SYCL_BUILTIN_ALIAS(__builtin_intel_sycl_alloca)
488-
multi_ptr<ElementType, access::address_space::private_space, DecorateAddress>
489-
private_alloca(kernel_handler &h);
490-
491-
#else
492-
493-
template <typename ElementType, auto &Size,
494-
access::decorated DecorateAddress>
495-
multi_ptr<ElementType, access::address_space::private_space, DecorateAddress>
496-
private_alloca(kernel_handler &h) {
497-
throw "sycl::ext::oneapi::experimental::private_alloca is not supported in the host";
498-
}
499-
500-
#endif
501-
502483
} // namespace experimental
503484
} // namespace oneapi
504485
} // namespace ext

‎clang/test/SemaSYCL/builtin-alloca-errors-device.cpp

Lines changed: 5 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,8 +1,11 @@
1-
// RUN: %clang_cc1 -fsyntax-only -fsycl-is-device -triple spir64-unknown-unknown -verify -Wpedantic %s
1+
// RUN: %clang_cc1 -fsyntax-only -fsycl-is-device -triple spir64-unknown-unknown -verify %s
2+
3+
// Test errors of __builtin_intel_sycl_alloca when used in SYCL device code.
24

35
#include <stddef.h>
46

57
#include "Inputs/sycl.hpp"
8+
#include "Inputs/private_alloca.hpp"
69

710
constexpr sycl::specialization_id<size_t> size(1);
811
constexpr sycl::specialization_id<float> badsize(1);
@@ -64,7 +67,7 @@ void test(sycl::kernel_handler &h) {
6467
// expected-error@+1 {{builtin functions must be directly called}}
6568
auto funcPtr = sycl::ext::oneapi::experimental::private_alloca<float, size, sycl::access::decorated::no>;
6669

67-
// expected-error@+1 {{__builtin_intel_sycl_alloca cannot be used in source code. Use the private_alloca alias instead.}}
70+
// expected-error@+1 {{__builtin_intel_sycl_alloca cannot be used in source code. Use the private_alloca alias instead}}
6871
__builtin_intel_sycl_alloca(h);
6972

7073
// expected-error@+1 {{too few arguments to function call, expected 1, have 0}}

‎clang/test/SemaSYCL/builtin-alloca-errors-host.cpp

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,8 +1,11 @@
1-
// RUN: %clang_cc1 -fsyntax-only -fsycl-is-host -triple x86_64-unknown-unknown -verify -Wpedantic -fcxx-exceptions %s
1+
// RUN: %clang_cc1 -fsyntax-only -fsycl-is-host -triple x86_64-unknown-unknown -verify %s
2+
3+
// Test errors of __builtin_intel_sycl_alloca when used in targets other than SYCL devices.
24

35
#include <stddef.h>
46

57
#include "Inputs/sycl.hpp"
8+
#include "Inputs/private_alloca.hpp"
69

710
constexpr sycl::specialization_id<size_t> size(1);
811

‎clang/test/SemaSYCL/builtin-alloca.cpp

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,9 +1,11 @@
11
// RUN: %clang_cc1 -fsyntax-only -fsycl-is-device -triple spir64-unknown-unknown -verify -Wpedantic %s
2-
// RUN: %clang_cc1 -fsyntax-only -fsycl-is-host -triple x86_64-unknown-unknown -verify -Wpedantic -fcxx-exceptions %s
2+
3+
// Test verification of __builtin_intel_sycl_alloca when used in different valid ways.
34

45
#include <stddef.h>
56

67
#include "Inputs/sycl.hpp"
8+
#include "Inputs/private_alloca.hpp"
79

810
// expected-no-diagnostics
911

0 commit comments

Comments
 (0)