Skip to content

Commit 9d86971

Browse files
Seanst98Seanst98
committed
[SYCL][Bindless][E2E] Test normalized usm bindless images
Modify pre-existing read_norm_types test to add testing for 2D USM bindless images.
1 parent 4f2ad4f commit 9d86971

File tree

2 files changed

+57
-19
lines changed

2 files changed

+57
-19
lines changed

‎sycl/cmake/modules/FetchUnifiedRuntime.cmake‎

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -116,14 +116,14 @@ if(SYCL_UR_USE_FETCH_CONTENT)
116116
CACHE PATH "Path to external '${name}' adapter source dir" FORCE)
117117
endfunction()
118118

119-
set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
120-
# commit eb63d1a21729f6928bb6cccc5f92856b0690aca6
121-
# Merge: e26bba51 45a781f4
119+
set(UNIFIED_RUNTIME_REPO "https://github.com/Seanst98/unified-runtime.git")
120+
# commit d0a50523006fa6f283da6a36811081add3bb22fc
121+
# Merge: 804851e4 04deb8b3
122122
# Author: Omar Ahmed <[email protected]>
123-
# Date: Tue Sep 10 12:08:57 2024 +0100
124-
# Merge pull request #1796 from GeorgeWeb/georgi/ur_kernel_max_active_wgs
125-
# [CUDA] Implement urKernelSuggestMaxCooperativeGroupCountExp for Cuda
126-
set(UNIFIED_RUNTIME_TAG eb63d1a21729f6928bb6cccc5f92856b0690aca6)
123+
# Date: Tue Aug 20 16:28:30 2024 +0100
124+
# Merge pull request #1940 from RossBrunton/ross/urcall
125+
# [XPTI] Use `ur.call` rather than `ur` in XPTI
126+
set(UNIFIED_RUNTIME_TAG 94cb7b07e5dc5712432d9793b2879916dc9b8653)
127127

128128
set(UMF_BUILD_EXAMPLES OFF CACHE INTERNAL "EXAMPLES")
129129
# Due to the use of dependentloadflag and no installer for UMF and hwloc we need

‎sycl/test-e2e/bindless_images/read_norm_types.cpp‎

Lines changed: 50 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -6,6 +6,7 @@
66
#include <iostream>
77
#include <limits>
88
#include <sycl/detail/core.hpp>
9+
#include <sycl/usm.hpp>
910

1011
#include "helpers/common.hpp"
1112
#include <sycl/ext/oneapi/bindless_images.hpp>
@@ -29,7 +30,7 @@ bool run_test(sycl::range<NDims> globalSize, sycl::range<NDims> localSize) {
2930

3031
std::vector<InputType> dataIn(numElems, InputType((DType)dtypeMaxVal));
3132
std::vector<OutputType> dataOut(numElems);
32-
std::vector<OutputType> expected(numElems, OutputType(1.f));
33+
std::vector<OutputType> expected(numElems, OutputType(2.f));
3334

3435
try {
3536

@@ -47,9 +48,30 @@ bool run_test(sycl::range<NDims> globalSize, sycl::range<NDims> localSize) {
4748
sycl::coordinate_normalization_mode::normalized,
4849
sycl::filtering_mode::nearest};
4950

50-
auto imgIn = syclexp::create_image(imgMemIn, sampler, descIn, q);
51+
auto imgIn1 = syclexp::create_image(imgMemIn, sampler, descIn, q);
5152
auto imgOut = syclexp::create_image(imgMemOut, descOut, q);
5253

54+
void *allocUSM = nullptr;
55+
syclexp::image_mem_handle allocMem;
56+
syclexp::sampled_image_handle imgIn2;
57+
58+
if constexpr (NDims == 2) {
59+
size_t pitch = 0;
60+
allocUSM = syclexp::pitched_alloc_device(&pitch, descIn, q);
61+
62+
if (allocUSM == nullptr) {
63+
std::cerr << "Error allocating 2D USM memory!" << std::endl;
64+
return false;
65+
}
66+
imgIn2 = syclexp::create_image(allocUSM, pitch, sampler, descIn, q);
67+
q.ext_oneapi_copy(dataIn.data(), allocUSM, descIn, pitch);
68+
69+
} else {
70+
allocMem = syclexp::alloc_image_mem(descIn, q);
71+
imgIn2 = syclexp::create_image(allocMem, sampler, descIn, q);
72+
q.ext_oneapi_copy(dataIn.data(), allocMem, descIn);
73+
}
74+
5375
q.ext_oneapi_copy(dataIn.data(), imgMemIn, descIn);
5476
q.wait_and_throw();
5577

@@ -60,27 +82,35 @@ bool run_test(sycl::range<NDims> globalSize, sycl::range<NDims> localSize) {
6082
if constexpr (NDims == 1) {
6183
size_t dim0 = it.get_global_id(0);
6284
float fdim0 = dim0 / globalSize[0];
63-
OutputType pixel =
64-
syclexp::sample_image<OutputType>(imgIn, fdim0);
65-
syclexp::write_image(imgOut, int(dim0), pixel);
85+
OutputType pixel1 =
86+
syclexp::sample_image<OutputType>(imgIn1, fdim0);
87+
OutputType pixel2 =
88+
syclexp::sample_image<OutputType>(imgIn2, fdim0);
89+
syclexp::write_image(imgOut, int(dim0), pixel1 + pixel2);
6690
} else if constexpr (NDims == 2) {
6791
size_t dim0 = it.get_global_id(0);
6892
size_t dim1 = it.get_global_id(1);
6993
float fdim0 = dim0 / globalSize[0];
7094
float fdim1 = dim1 / globalSize[1];
71-
OutputType pixel = syclexp::sample_image<OutputType>(
72-
imgIn, sycl::float2(fdim0, fdim1));
73-
syclexp::write_image(imgOut, sycl::int2(dim0, dim1), pixel);
95+
OutputType pixel1 = syclexp::sample_image<OutputType>(
96+
imgIn1, sycl::float2(fdim0, fdim1));
97+
OutputType pixel2 = syclexp::sample_image<OutputType>(
98+
imgIn2, sycl::float2(fdim0, fdim1));
99+
syclexp::write_image(imgOut, sycl::int2(dim0, dim1),
100+
pixel1 + pixel2);
74101
} else if constexpr (NDims == 3) {
75102
size_t dim0 = it.get_global_id(0);
76103
size_t dim1 = it.get_global_id(1);
77104
size_t dim2 = it.get_global_id(2);
78105
float fdim0 = dim0 / globalSize[0];
79106
float fdim1 = dim1 / globalSize[1];
80107
float fdim2 = dim2 / globalSize[2];
81-
OutputType pixel = syclexp::sample_image<OutputType>(
82-
imgIn, sycl::float3(fdim0, fdim1, fdim2));
83-
syclexp::write_image(imgOut, sycl::int3(dim0, dim1, dim2), pixel);
108+
OutputType pixel1 = syclexp::sample_image<OutputType>(
109+
imgIn1, sycl::float3(fdim0, fdim1, fdim2));
110+
OutputType pixel2 = syclexp::sample_image<OutputType>(
111+
imgIn2, sycl::float3(fdim0, fdim1, fdim2));
112+
syclexp::write_image(imgOut, sycl::int3(dim0, dim1, dim2),
113+
pixel1 + pixel2);
84114
}
85115
});
86116
});
@@ -89,12 +119,20 @@ bool run_test(sycl::range<NDims> globalSize, sycl::range<NDims> localSize) {
89119
q.ext_oneapi_copy(imgMemOut, dataOut.data(), descOut);
90120
q.wait_and_throw();
91121

92-
syclexp::destroy_image_handle(imgIn, q);
122+
syclexp::destroy_image_handle(imgIn1, q);
123+
syclexp::destroy_image_handle(imgIn2, q);
93124
syclexp::destroy_image_handle(imgOut, q);
94125

95126
syclexp::free_image_mem(imgMemIn, syclexp::image_type::standard, dev, ctxt);
96127
syclexp::free_image_mem(imgMemOut, syclexp::image_type::standard, dev,
97128
ctxt);
129+
130+
if constexpr (NDims == 2) {
131+
sycl::free(allocUSM, ctxt);
132+
} else {
133+
syclexp::free_image_mem(allocMem, syclexp::image_type::standard, dev,
134+
ctxt);
135+
}
98136
} catch (sycl::exception e) {
99137
std::cerr << "SYCL exception caught! : " << e.what() << "\n";
100138
return false;

0 commit comments

Comments
 (0)