-
Notifications
You must be signed in to change notification settings - Fork 769
[SYCL][SCLA] Add sycl::aspect::ext_oneapi_private_alloca
#13181
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
Changes from all commits
00563d6
d528078
2586560
b1e9fff
02fff2b
cb3022f
File filter
Filter by extension
Conversations
Jump to
Diff view
Diff view
There are no files selected for viewing
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,24 @@ | ||
; RUN: sycl-post-link -spec-const=emulation %s 2>&1 | FileCheck %s | ||
|
||
; This test checks the `-spec-const` pass on SPIR-V targets and emulation mode, | ||
; i.e., on AOT SPIR-V targets. In this scenario, 'llvm.sycl.alloca' intrinsics | ||
; must be left unmodified. | ||
|
||
; Note that coming from clang this case should never be reached. | ||
|
||
; CHECK: sycl-post-link NOTE: no modifications to the input LLVM IR have been made | ||
|
||
target triple = "spir64_x86_64" | ||
|
||
%"class.sycl::_V1::specialization_id" = type { i64 } | ||
|
||
@size_i64 = addrspace(1) constant %"class.sycl::_V1::specialization_id" { i64 10 }, align 8 | ||
|
||
@size_i64_stable_name = private unnamed_addr constant [36 x i8] c"_ZTS14name_generatorIL_Z8size_i64EE\00", align 1 | ||
|
||
define dso_local void @private_alloca() { | ||
call ptr @llvm.sycl.alloca.p0.p4.p4.p4.f64(ptr addrspace(4) addrspacecast (ptr @size_i64_stable_name to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @size_i64 to ptr addrspace(4)), ptr addrspace(4) null, double 0.000000e+00, i64 8) | ||
ret void | ||
} | ||
|
||
declare ptr @llvm.sycl.alloca.p0.p4.p4.p4.f32(ptr addrspace(4), ptr addrspace(4), ptr addrspace(4), float, i64) |
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -1,7 +1,10 @@ | ||
; RUN: sycl-post-link -spec-const=native < %s -S -o %t.table | ||
; RUN: FileCheck %s -check-prefixes=CHECK-RT < %t_0.ll | ||
; RUN: FileCheck %s -check-prefixes=CHECK,CHECK-RT < %t_0.ll | ||
; RUN: FileCheck %s --check-prefixes=CHECK-PROPS < %t_0.prop | ||
|
||
; RUN: sycl-post-link -spec-const=emulation < %s -S -o %t.table | ||
; RUN: FileCheck %s -check-prefixes=CHECK,CHECK-EMULATION < %t_0.ll | ||
|
||
; This test checks that the post link tool is able to correctly transform | ||
; SYCL alloca intrinsics in SPIR-V devices. | ||
|
||
|
@@ -10,9 +13,9 @@ | |
%"class.sycl::_V1::specialization_id.1" = type { i16 } | ||
%my_range = type { ptr addrspace(4), ptr addrspace(4) } | ||
|
||
@size_i64 = internal addrspace(1) constant %"class.sycl::_V1::specialization_id" { i64 10 }, align 8 | ||
@size_i32 = internal addrspace(1) constant %"class.sycl::_V1::specialization_id.0" { i32 120 }, align 4 | ||
@size_i16 = internal addrspace(1) constant %"class.sycl::_V1::specialization_id.1" { i16 1 }, align 2 | ||
@size_i64 = addrspace(1) constant %"class.sycl::_V1::specialization_id" { i64 10 }, align 8 | ||
@size_i32 = addrspace(1) constant %"class.sycl::_V1::specialization_id.0" { i32 120 }, align 4 | ||
@size_i16 = addrspace(1) constant %"class.sycl::_V1::specialization_id.1" { i16 1 }, align 2 | ||
Comment on lines
-13
to
+18
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. Needed to not drop |
||
|
||
; Check that the following globals are preserved: even though they are not used | ||
; in the module anymore, they could still be referenced by debug info metadata | ||
|
@@ -30,10 +33,19 @@ | |
define dso_local void @private_alloca() { | ||
; CHECK-RT: [[LENGTH:%.*]] = call i32 @_Z20__spirv_SpecConstantii(i32 1, i32 120) | ||
; CHECK-RT: {{.*}} = alloca double, i32 [[LENGTH]], align 8 | ||
|
||
; CHECK-EMULATION: alloca double, i32 120, align 8 | ||
call ptr @llvm.sycl.alloca.p0.p4.p4.p4.f64(ptr addrspace(4) addrspacecast (ptr @size_i32_stable_name to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @size_i32 to ptr addrspace(4)), ptr addrspace(4) null, double 0.000000e+00, i64 8) | ||
; CHECK-RT: [[LENGTH:%.*]] = call i64 @_Z20__spirv_SpecConstantix(i32 0, i64 10) | ||
; CHECK-RT: {{.*}} = alloca float, i64 [[LENGTH]], align 8 | ||
|
||
; CHECK-EMULATION: alloca float, i64 10, align 8 | ||
call ptr @llvm.sycl.alloca.p0.p4.p4.p4.f32(ptr addrspace(4) addrspacecast (ptr @size_i64_stable_name to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @size_i64 to ptr addrspace(4)), ptr addrspace(4) null, float 0.000000e+00, i64 8) | ||
|
||
; CHECK-RT: %[[LENGTH:.*]] = call i16 @_Z20__spirv_SpecConstantis(i32 2, i16 1) | ||
; CHECK-RT: {{.*}} = alloca %my_range, i16 %[[LENGTH]], align 64 | ||
Comment on lines
+45
to
+46
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. Previously missing |
||
|
||
; CHECK-EMULATION: alloca %my_range, i16 1, align 64 | ||
call ptr @llvm.sycl.alloca.p0.p4.p4.p4.s_my_range(ptr addrspace(4) addrspacecast (ptr @size_i16_stable_name to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @size_i16 to ptr addrspace(4)), ptr addrspace(4) null, %my_range zeroinitializer, i64 64) | ||
ret void | ||
} | ||
|
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,32 @@ | ||
; RUN: sycl-post-link -split=auto -spec-const=native -S -o %t.table %s -generate-device-image-default-spec-consts | ||
; RUN: FileCheck %s -input-file %t_1.ll --implicit-check-not="SpecConst" | ||
|
||
; This test checks that the post link tool is able to correctly transform | ||
; SYCL alloca intrinsics in SPIR-V devices when using default values. | ||
|
||
%"class.sycl::_V1::specialization_id" = type { i64 } | ||
%"class.sycl::_V1::specialization_id.0" = type { i32 } | ||
%"class.sycl::_V1::specialization_id.1" = type { i16 } | ||
%my_range = type { ptr addrspace(4), ptr addrspace(4) } | ||
|
||
@size_i64 = internal addrspace(1) constant %"class.sycl::_V1::specialization_id" { i64 10 }, align 8 | ||
@size_i32 = internal addrspace(1) constant %"class.sycl::_V1::specialization_id.0" { i32 120 }, align 4 | ||
@size_i16 = internal addrspace(1) constant %"class.sycl::_V1::specialization_id.1" { i16 1 }, align 2 | ||
|
||
@size_i64_stable_name = private unnamed_addr constant [36 x i8] c"_ZTS14name_generatorIL_Z8size_i64EE\00", align 1 | ||
@size_i32_stable_name = private unnamed_addr constant [36 x i8] c"_ZTS14name_generatorIL_Z8size_i32EE\00", align 1 | ||
@size_i16_stable_name = private unnamed_addr constant [36 x i8] c"_ZTS14name_generatorIL_Z8size_i16EE\00", align 1 | ||
|
||
define dso_local void @private_alloca() { | ||
; CHECK: alloca double, i32 120, align 8 | ||
call ptr @llvm.sycl.alloca.p0.p4.p4.p4.f64(ptr addrspace(4) addrspacecast (ptr @size_i32_stable_name to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @size_i32 to ptr addrspace(4)), ptr addrspace(4) null, double 0.000000e+00, i64 8) | ||
; CHECK: alloca float, i64 10, align 8 | ||
call ptr @llvm.sycl.alloca.p0.p4.p4.p4.f32(ptr addrspace(4) addrspacecast (ptr @size_i64_stable_name to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @size_i64 to ptr addrspace(4)), ptr addrspace(4) null, float 0.000000e+00, i64 8) | ||
; CHECK: alloca %my_range, i16 1, align 64 | ||
call ptr @llvm.sycl.alloca.p0.p4.p4.p4.s_my_range(ptr addrspace(4) addrspacecast (ptr @size_i16_stable_name to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @size_i16 to ptr addrspace(4)), ptr addrspace(4) null, %my_range zeroinitializer, i64 64) | ||
ret void | ||
} | ||
|
||
declare ptr @llvm.sycl.alloca.p0.p4.p4.p4.f32(ptr addrspace(4), ptr addrspace(4), ptr addrspace(4), float, i64) | ||
declare ptr @llvm.sycl.alloca.p0.p4.p4.p4.f64(ptr addrspace(4), ptr addrspace(4), ptr addrspace(4), double, i64) | ||
declare ptr @llvm.sycl.alloca.p0.p4.p4.p4.s_my_range(ptr addrspace(4), ptr addrspace(4), ptr addrspace(4), %my_range, i64) |
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,40 @@ | ||
// RUN: %{build} -o %t.out | ||
// RUN: %{run} %t.out | ||
|
||
// Check that an exception with an exception with the `errc::invalid` error code | ||
// thrown when trying to use `sycl_ext_oneapi_private_alloca` and no device | ||
// supports the aspect. | ||
|
||
#include <sycl/detail/core.hpp> | ||
|
||
#include <sycl/ext/oneapi/experimental/alloca.hpp> | ||
#include <sycl/specialization_id.hpp> | ||
#include <sycl/usm.hpp> | ||
|
||
class Kernel; | ||
|
||
constexpr sycl::specialization_id<int> Size(10); | ||
|
||
static std::error_code test() { | ||
sycl::queue Queue; | ||
sycl::buffer<int> B(10); | ||
|
||
try { | ||
Queue.submit([&](sycl::handler &Cgh) { | ||
sycl::accessor Acc(B, Cgh, sycl::write_only, sycl::no_init); | ||
Cgh.parallel_for<Kernel>(10, [=](sycl::id<1>, sycl::kernel_handler Kh) { | ||
sycl::ext::oneapi::experimental::private_alloca< | ||
int, Size, sycl::access::decorated::no>(Kh); | ||
}); | ||
}); | ||
} catch (sycl::exception &Exception) { | ||
return Exception.code(); | ||
} | ||
assert(false && "Exception not thrown"); | ||
} | ||
|
||
int main() { | ||
assert(test() == sycl::errc::invalid && "Unexpected error code"); | ||
|
||
return 0; | ||
} |
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1 @@ | ||
config.required_features += ['!aspect-ext_oneapi_private_alloca'] |
Original file line number | Diff line number | Diff line change |
---|---|---|
|
@@ -2,9 +2,10 @@ | |
|
||
// Template for private alloca tests. | ||
|
||
#include <sycl/sycl.hpp> | ||
#include <sycl/detail/core.hpp> | ||
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. Also did this in all the tests |
||
|
||
#include <sycl/ext/oneapi/experimental/alloca.hpp> | ||
#include <sycl/specialization_id.hpp> | ||
|
||
template <typename ElementType, typename SizeType, | ||
sycl::access::decorated DecorateAddress> | ||
|
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1 @@ | ||
config.required_features += ['aspect-ext_oneapi_private_alloca'] |
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -1,6 +1,5 @@ | ||
// RUN: %{build} -w -o %t.out | ||
// RUN: echo 1 | %{run} %t.out | ||
// UNSUPPORTED: cuda || hip | ||
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. Not needed, using |
||
|
||
// Test checking size of 'bool' type. This is not expected to be ever used, but, | ||
// as 'bool' is an integral type, it is a possible scenario. | ||
|
Uh oh!
There was an error while loading. Please reload this page.