Skip to content

Commit f7dc00d

Browse files
victor-edsvictor-eds
authored
[SYCL][SCLA] Drop unneeded addrspacecast instruction (#13739)
As pointed out in #13514 (comment), `sycl::multi_ptr` pointer member is always decorated in DPC++. Drop unneeeded `addrspacecast` to the generic address space when `[aligned_]private_alloca` returns a `raw_private_ptr`. Signed-off-by: Victor Perez <[email protected]>
1 parent 190b15e commit f7dc00d

File tree

5 files changed

+12
-61
lines changed

5 files changed

+12
-61
lines changed

‎clang/lib/CodeGen/CGBuiltin.cpp

Lines changed: 0 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -23990,22 +23990,6 @@ RValue CodeGenFunction::EmitIntelSYCLAllocaBuiltin(
2399023990
return CI;
2399123991
}();
2399223992

23993-
// Perform AS cast if needed.
23994-
23995-
constexpr int NoDecorated = 0;
23996-
llvm::APInt Decorated = TAL->get(DecorateAddressIndex).getAsIntegral();
23997-
// Both 'sycl::access::decorated::{yes and legacy}' lead to decorated (private
23998-
// AS) pointer type. Perform cast if 'sycl::access::decorated::no'.
23999-
if (Decorated == NoDecorated) {
24000-
IRBuilderBase::InsertPointGuard IPG(Builder);
24001-
Builder.SetInsertPoint(getPostAllocaInsertPoint());
24002-
unsigned DestAddrSpace =
24003-
getContext().getTargetAddressSpace(LangAS::Default);
24004-
llvm::PointerType *DestTy =
24005-
llvm::PointerType::get(Builder.getContext(), DestAddrSpace);
24006-
Allocation = Builder.CreateAddrSpaceCast(Allocation, DestTy);
24007-
}
24008-
2400923993
// If no slot is provided, simply return allocation.
2401023994
if (ReturnValue.isNull())
2401123995
return RValue::get(Allocation);

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

Lines changed: 1 addition & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -178,29 +178,16 @@ struct DecoratedType<ElementType, access::address_space::constant_space> {
178178
#endif
179179
};
180180

181-
// Equivalent to std::conditional
182-
template <bool B, class T, class F>
183-
struct conditional { using type = T; };
184-
185-
template <class T, class F>
186-
struct conditional<false, T, F> { using type = F; };
187-
188-
template <bool B, class T, class F>
189-
using conditional_t = typename conditional<B, T, F>::type;
190-
191181
template <typename T, access::address_space AS,
192182
access::decorated DecorateAddress = access::decorated::legacy>
193183
class __SYCL_TYPE(multi_ptr) multi_ptr {
194-
static constexpr bool is_decorated =
195-
DecorateAddress == access::decorated::yes;
196-
197184
using decorated_type = typename DecoratedType<T, AS>::type;
198185

199186
static_assert(DecorateAddress != access::decorated::legacy);
200187
static_assert(AS != access::address_space::constant_space);
201188

202189
public:
203-
using pointer = conditional_t<is_decorated, decorated_type *, T *>;
190+
using pointer = decorated_type *;
204191

205192
multi_ptr(typename multi_ptr<T, AS, access::decorated::yes>::pointer Ptr)
206193
: m_Pointer((pointer)(Ptr)) {} // #MultiPtrConstructor

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

Lines changed: 3 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -18,8 +18,7 @@ struct myStruct {
1818
constexpr sycl::specialization_id<size_t> size(1);
1919
constexpr sycl::specialization_id<int> intSize(1);
2020

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`.
21+
// For each call, we should generate a chain of: 'call @llvm.sycl.alloca.<ty>' + 'store'.
2322

2423
// CHECK-LABEL: define dso_local spir_func void @_Z4testRN4sycl3_V114kernel_handlerE(
2524
// CHECK-SAME: ptr addrspace(4) noundef align 1 dereferenceable(1) [[KH:%.*]])
@@ -35,11 +34,10 @@ constexpr sycl::specialization_id<int> intSize(1);
3534
// CHECK-NEXT: [[PTR0_ASCAST:%.*]] = addrspacecast ptr [[PTR0]] to ptr addrspace(4)
3635
// CHECK-NEXT: [[PTR1_ASCAST:%.*]] = addrspacecast ptr [[PTR1]] to ptr addrspace(4)
3736
// CHECK-NEXT: [[PTR2_ASCAST:%.*]] = addrspacecast ptr [[PTR2]] to ptr addrspace(4)
38-
// CHECK-NEXT: [[TMP5:%.*]] = addrspacecast ptr [[TMP4]] to ptr addrspace(4)
3937
// CHECK-NEXT: store ptr addrspace(4) [[KH]], ptr addrspace(4) [[KH_ADDR_ASCAST]], align 8
4038
// CHECK-NEXT: store ptr [[TMP0]], ptr addrspace(4) [[PTR0_ASCAST]], align 8
4139
// 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
40+
// CHECK-NEXT: store ptr [[TMP4]], ptr addrspace(4) [[PTR2_ASCAST]], align 8
4341
// CHECK-NEXT: ret void
4442
SYCL_EXTERNAL void test(sycl::kernel_handler &kh) {
4543
auto ptr0 = sycl::ext::oneapi::experimental::private_alloca<double, size, sycl::access::decorated::yes>(kh);
@@ -67,11 +65,10 @@ SYCL_EXTERNAL void test(sycl::kernel_handler &kh) {
6765
// CHECK-NEXT: [[PTR0_ASCAST:%.*]] = addrspacecast ptr [[PTR0]] to ptr addrspace(4)
6866
// CHECK-NEXT: [[PTR1_ASCAST:%.*]] = addrspacecast ptr [[PTR1]] to ptr addrspace(4)
6967
// CHECK-NEXT: [[PTR2_ASCAST:%.*]] = addrspacecast ptr [[PTR2]] to ptr addrspace(4)
70-
// CHECK-NEXT: [[TMP5:%.*]] = addrspacecast ptr [[TMP4]] to ptr addrspace(4)
7168
// CHECK-NEXT: store ptr addrspace(4) [[KH]], ptr addrspace(4) [[KH_ADDR_ASCAST]], align 8
7269
// CHECK-NEXT: store ptr [[TMP0]], ptr addrspace(4) [[PTR0_ASCAST]], align 8
7370
// CHECK-NEXT: store ptr [[TMP2]], ptr addrspace(4) [[PTR1_ASCAST]], align 8
74-
// CHECK-NEXT: store ptr addrspace(4) [[TMP5]], ptr addrspace(4) [[PTR2_ASCAST]], align 8
71+
// CHECK-NEXT: store ptr [[TMP4]], ptr addrspace(4) [[PTR2_ASCAST]], align 8
7572
// CHECK-NEXT: ret void
7673
SYCL_EXTERNAL void test_aligned(sycl::kernel_handler &kh) {
7774
auto ptr0 = sycl::ext::oneapi::experimental::aligned_private_alloca<double, alignof(double) * 2, size, sycl::access::decorated::yes>(kh);

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

Lines changed: 1 addition & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -409,29 +409,16 @@ struct DecoratedType<ElementType, access::address_space::global_space> {
409409
using type = __attribute__((opencl_global)) ElementType;
410410
};
411411

412-
// Equivalent to std::conditional
413-
template <bool B, class T, class F>
414-
struct conditional { using type = T; };
415-
416-
template <class T, class F>
417-
struct conditional<false, T, F> { using type = F; };
418-
419-
template <bool B, class T, class F>
420-
using conditional_t = typename conditional<B, T, F>::type;
421-
422412
template <typename T, access::address_space AS,
423413
access::decorated DecorateAddress = access::decorated::legacy>
424414
class __SYCL_TYPE(multi_ptr) multi_ptr {
425-
static constexpr bool is_decorated =
426-
DecorateAddress == access::decorated::yes;
427-
428415
using decorated_type = typename DecoratedType<T, AS>::type;
429416

430417
static_assert(DecorateAddress != access::decorated::legacy);
431418
static_assert(AS != access::address_space::constant_space);
432419

433420
public:
434-
using pointer = conditional_t<is_decorated, decorated_type *, T *>;
421+
using pointer = decorated_type *;
435422

436423
multi_ptr(typename multi_ptr<T, AS, access::decorated::yes>::pointer Ptr)
437424
: m_Pointer((pointer)(Ptr)) {}

‎sycl/test/extensions/private_alloca.cpp

Lines changed: 7 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -7,13 +7,11 @@
77
// Check SPIR-V code generation for 'sycl_ext_oneapi_private_alloca'. Each call
88
// to the extension API is annotated as follows for future reference:
99
//
10-
// <NAME>: storage_class=<sc>, element_type=<et>, alignment=<align>
10+
// <NAME>: element_type=<et>, alignment=<align>
1111
//
1212
// - <NAME>: Variable name in the test below. These will be the result of
1313
// bitcasting a variable to a different pointer type. We use this instead of the
1414
// variable due to FileCheck limitations.
15-
// - <sc>: 'generic' if <NAME> is casted to generic before being stored in the
16-
// multi_ptr or 'function' otherwise.
1715
// - <et>: element type. 'Bitcast X <NAME> Y' will originate value <NAME>, being
1816
// X a pointer to <et> and storage class function.
1917
// - <align>: alignment. <NAME> will appear in a 'Decorage <NAME> Aligment
@@ -44,17 +42,17 @@ SYCL_EXTERNAL void test(sycl::kernel_handler &kh) {
4442
keep(/*B0: storage_class=function, element_type=f32, alignment=4*/
4543
sycl::ext::oneapi::experimental::private_alloca<
4644
float, int8_id, sycl::access::decorated::yes>(kh),
47-
/*B1: storage_class=generic, element_type=f64, alignment=8*/
45+
/*B1: element_type=f64, alignment=8*/
4846
sycl::ext::oneapi::experimental::private_alloca<
4947
double, uint32_id, sycl::access::decorated::no>(kh),
50-
/*B2: storage_class=function, element_type=i32, alignment=4*/
48+
/*B2: element_type=i32, alignment=4*/
5149
sycl::ext::oneapi::experimental::private_alloca<
5250
int, int16_id, sycl::access::decorated::legacy>(kh),
53-
/*B3: storage_class=generic, element_type=i64, alignment=16*/
51+
/*B3: element_type=i64, alignment=16*/
5452
sycl::ext::oneapi::experimental::aligned_private_alloca<
5553
int64_t, alignof(int64_t) * 2, uint64_id,
5654
sycl::access::decorated::no>(kh),
57-
/*B4: storage_class=function, element_type=composite, alignment=32*/
55+
/*B4: element_type=composite, alignment=32*/
5856
sycl::ext::oneapi::experimental::aligned_private_alloca<
5957
composite, alignof(composite) * 8, int32_id,
6058
sycl::access::decorated::yes>(kh));
@@ -120,15 +118,13 @@ SYCL_EXTERNAL void test(sycl::kernel_handler &kh) {
120118
// CHECK-SPV-DAG: Store {{.*}} [[#B0]]
121119
// CHECK-SPV-DAG: Variable [[#ARRF64PTRTY]] [[#V1:]] [[#FUNCTIONSTORAGE]]
122120
// CHECK-SPV-DAG: Bitcast [[#F64PTRTY]] [[#B1:]] [[#V1]]
123-
// CHECK-SPV-DAG: PtrCastToGeneric {{.*}} [[#G1:]] [[#B1]]
124-
// CHECK-SPV-DAG: Store {{.*}} [[#G1]]
121+
// CHECK-SPV-DAG: Store {{.*}} [[#B1]]
125122
// CHECK-SPV-DAG: Variable [[#ARRI32PTRTY]] [[#V2:]] [[#FUNCTIONSTORAGE]]
126123
// CHECK-SPV-DAG: Bitcast [[#I32PTRTY]] [[#B2:]] [[#V2]]
127124
// CHECK-SPV-DAG: Store {{.*}} [[#B2]]
128125
// CHECK-SPV-DAG: Variable [[#ARRI64PTRTY]] [[#V3:]] [[#FUNCTIONSTORAGE]]
129126
// CHECK-SPV-DAG: Bitcast [[#I64PTRTY]] [[#B3:]] [[#V3]]
130-
// CHECK-SPV-DAG: PtrCastToGeneric {{.*}} [[#G3:]] [[#B3]]
131-
// CHECK-SPV-DAG: Store {{.*}} [[#G3]]
127+
// CHECK-SPV-DAG: Store {{.*}} [[#B3]]
132128
// CHECK-SPV-DAG: Variable [[#ARRCOMPPTRTY]] [[#V4:]] [[#FUNCTIONSTORAGE]]
133129
// CHECK-SPV-DAG: Bitcast [[#COMPPTRTY]] [[#B4:]] [[#V4]]
134130
// CHECK-SPV-DAG: Store {{.*}} [[#B4]]

0 commit comments

Comments
 (0)