Skip to content

Commit

Permalink
[SYCL][Bindless][E2E] Test normalized usm bindless images
Browse files Browse the repository at this point in the history 
Modify pre-existing read_norm_types test to add testing for 2D
USM bindless images.
  • Loading branch information
@Seanst98
Seanst98 committed Sep 11, 2024
1 parent 4f2ad4f commit 9d86971
Show file tree
Hide file tree
Showing 2 changed files with 57 additions and 19 deletions.
14 changes: 7 additions & 7 deletions sycl/cmake/modules/FetchUnifiedRuntime.cmake
View file
Original file line number Diff line number Diff line change
Expand Up @@ -116,14 +116,14 @@ if(SYCL_UR_USE_FETCH_CONTENT)
CACHE PATH "Path to external '${name}' adapter source dir" FORCE)
endfunction()

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

set(UMF_BUILD_EXAMPLES OFF CACHE INTERNAL "EXAMPLES")
# Due to the use of dependentloadflag and no installer for UMF and hwloc we need
Expand Down
62 changes: 50 additions & 12 deletions sycl/test-e2e/bindless_images/read_norm_types.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -6,6 +6,7 @@
#include <iostream>
#include <limits>
#include <sycl/detail/core.hpp>
#include <sycl/usm.hpp>

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

std::vector<InputType> dataIn(numElems, InputType((DType)dtypeMaxVal));
std::vector<OutputType> dataOut(numElems);
std::vector<OutputType> expected(numElems, OutputType(1.f));
std::vector<OutputType> expected(numElems, OutputType(2.f));

try {

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

auto imgIn = syclexp::create_image(imgMemIn, sampler, descIn, q);
auto imgIn1 = syclexp::create_image(imgMemIn, sampler, descIn, q);
auto imgOut = syclexp::create_image(imgMemOut, descOut, q);

void *allocUSM = nullptr;
syclexp::image_mem_handle allocMem;
syclexp::sampled_image_handle imgIn2;

if constexpr (NDims == 2) {
size_t pitch = 0;
allocUSM = syclexp::pitched_alloc_device(&pitch, descIn, q);

if (allocUSM == nullptr) {
std::cerr << "Error allocating 2D USM memory!" << std::endl;
return false;
}
imgIn2 = syclexp::create_image(allocUSM, pitch, sampler, descIn, q);
q.ext_oneapi_copy(dataIn.data(), allocUSM, descIn, pitch);

} else {
allocMem = syclexp::alloc_image_mem(descIn, q);
imgIn2 = syclexp::create_image(allocMem, sampler, descIn, q);
q.ext_oneapi_copy(dataIn.data(), allocMem, descIn);
}

q.ext_oneapi_copy(dataIn.data(), imgMemIn, descIn);
q.wait_and_throw();

Expand All @@ -60,27 +82,35 @@ bool run_test(sycl::range<NDims> globalSize, sycl::range<NDims> localSize) {
if constexpr (NDims == 1) {
size_t dim0 = it.get_global_id(0);
float fdim0 = dim0 / globalSize[0];
OutputType pixel =
syclexp::sample_image<OutputType>(imgIn, fdim0);
syclexp::write_image(imgOut, int(dim0), pixel);
OutputType pixel1 =
syclexp::sample_image<OutputType>(imgIn1, fdim0);
OutputType pixel2 =
syclexp::sample_image<OutputType>(imgIn2, fdim0);
syclexp::write_image(imgOut, int(dim0), pixel1 + pixel2);
} else if constexpr (NDims == 2) {
size_t dim0 = it.get_global_id(0);
size_t dim1 = it.get_global_id(1);
float fdim0 = dim0 / globalSize[0];
float fdim1 = dim1 / globalSize[1];
OutputType pixel = syclexp::sample_image<OutputType>(
imgIn, sycl::float2(fdim0, fdim1));
syclexp::write_image(imgOut, sycl::int2(dim0, dim1), pixel);
OutputType pixel1 = syclexp::sample_image<OutputType>(
imgIn1, sycl::float2(fdim0, fdim1));
OutputType pixel2 = syclexp::sample_image<OutputType>(
imgIn2, sycl::float2(fdim0, fdim1));
syclexp::write_image(imgOut, sycl::int2(dim0, dim1),
pixel1 + pixel2);
} else if constexpr (NDims == 3) {
size_t dim0 = it.get_global_id(0);
size_t dim1 = it.get_global_id(1);
size_t dim2 = it.get_global_id(2);
float fdim0 = dim0 / globalSize[0];
float fdim1 = dim1 / globalSize[1];
float fdim2 = dim2 / globalSize[2];
OutputType pixel = syclexp::sample_image<OutputType>(
imgIn, sycl::float3(fdim0, fdim1, fdim2));
syclexp::write_image(imgOut, sycl::int3(dim0, dim1, dim2), pixel);
OutputType pixel1 = syclexp::sample_image<OutputType>(
imgIn1, sycl::float3(fdim0, fdim1, fdim2));
OutputType pixel2 = syclexp::sample_image<OutputType>(
imgIn2, sycl::float3(fdim0, fdim1, fdim2));
syclexp::write_image(imgOut, sycl::int3(dim0, dim1, dim2),
pixel1 + pixel2);
}
});
});
Expand All @@ -89,12 +119,20 @@ bool run_test(sycl::range<NDims> globalSize, sycl::range<NDims> localSize) {
q.ext_oneapi_copy(imgMemOut, dataOut.data(), descOut);
q.wait_and_throw();

syclexp::destroy_image_handle(imgIn, q);
syclexp::destroy_image_handle(imgIn1, q);
syclexp::destroy_image_handle(imgIn2, q);
syclexp::destroy_image_handle(imgOut, q);

syclexp::free_image_mem(imgMemIn, syclexp::image_type::standard, dev, ctxt);
syclexp::free_image_mem(imgMemOut, syclexp::image_type::standard, dev,
ctxt);

if constexpr (NDims == 2) {
sycl::free(allocUSM, ctxt);
} else {
syclexp::free_image_mem(allocMem, syclexp::image_type::standard, dev,
ctxt);
}
} catch (sycl::exception e) {
std::cerr << "SYCL exception caught! : " << e.what() << "\n";
return false;
Expand Down

0 comments on commit 9d86971

Please sign in to comment.