Skip to content

Commit

Permalink
[ESIMD][NFC][E2E] Fix 570 compilation warnings in ESIMD E2E tests (#1…
Browse files Browse the repository at this point in the history
…2748)

Warnings fixed:
- deprecated scatter_rgba
- deprecated get_cl_code
- deprecated lsc_fence
- deprecated uchar type usage
- deprecated get_access on HOST
- deprecated get_pointer
- usage of isfinite with -ffast-math
- deprecated dpas_argument_type::s1
- deprecated gpu_selector()

Also, the memory alloc/free in historgram*.cpp tests were updated to
simplify the potential memory leak avoidance.

Signed-off-by: Klochkov, Vyacheslav N <vyacheslav.n.klochkov@intel.com>
  • Loading branch information
v-klochkov authored Feb 20, 2024
1 parent 5fae0aa commit 436e687
Show file tree
Hide file tree
Showing 30 changed files with 219 additions and 273 deletions.
10 changes: 6 additions & 4 deletions sycl/test-e2e/ESIMD/accessor_global.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -2,7 +2,7 @@
// RUN: %{build} -fsycl-esimd-force-stateless-mem -o %t.out
// RUN: %{run} %t.out

// This test verifies usage of accessor methods operator[] and get_pointer().
// This test verifies usage of accessor methods operator[] and get_multi_ptr().

#include "esimd_test_utils.hpp"

Expand Down Expand Up @@ -39,7 +39,8 @@ bool test(queue Q, uint32_t LocalRange, uint32_t GlobalRange) {
for (int I = 0; I < VL; I++)
TmpAcc[GID * VL + I] = GID * 100 + I;
} else {
T *Ptr = TmpAcc.get_pointer();
T *Ptr =
TmpAcc.template get_multi_ptr<access::decorated::yes>().get();
simd<int, VL> IntValues(GID * 100, 1);
simd<T, VL> Values = IntValues;
block_store(Ptr + GID * VL, Values);
Expand All @@ -53,12 +54,13 @@ bool test(queue Q, uint32_t LocalRange, uint32_t GlobalRange) {
for (int I = 0; I < VL; I++)
Out[(GID + LID) * VL + I] = TmpAcc[(GID + LID) * VL + I];
} else {
T *Ptr = TmpAcc.get_pointer();
T *Ptr = TmpAcc.template get_multi_ptr<access::decorated::yes>()
.get();
simd<T, VL> Values = block_load<T, VL>(Ptr + (GID + LID) * VL);
Values.template copy_to(Out + (GID + LID) * VL);
}
} // end for (int LID = 0; LID < LocalRange; LID++)
} // end if (LID == 0)
} // end if (LID == 0)
});
}).wait();
} catch (sycl::exception const &e) {
Expand Down
8 changes: 5 additions & 3 deletions sycl/test-e2e/ESIMD/accessor_local.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -43,8 +43,10 @@ bool test(queue Q, uint32_t LocalRange, uint32_t GlobalRange) {
CGH.parallel_for(NDRange, [=](nd_item<1> Item) SYCL_ESIMD_KERNEL {
uint32_t GID = Item.get_global_id(0);
uint32_t LID = Item.get_local_id(0);
uint32_t LocalAccOffset = static_cast<uint32_t>(
reinterpret_cast<std::uintptr_t>(LocalAcc.get_pointer().get()));
uint32_t LocalAccOffset =
static_cast<uint32_t>(reinterpret_cast<std::uintptr_t>(
LocalAcc.template get_multi_ptr<access::decorated::yes>()
.get()));
if constexpr (TestSubscript) {
for (int I = 0; I < VL; I++)
LocalAcc[LID * VL + I] = GID * 100 + I;
Expand All @@ -67,7 +69,7 @@ bool test(queue Q, uint32_t LocalRange, uint32_t GlobalRange) {
ValuesFromSLM.copy_to(Out + (GID + LID) * VL);
}
} // end for (int LID = 0; LID < LocalRange; LID++)
} // end if (LID == 0)
} // end if (LID == 0)
});
}).wait();
} catch (sycl::exception const &e) {
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/ESIMD/api/simd_view_copy_move_assign.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -24,7 +24,7 @@ using namespace sycl::ext::intel::esimd;
template <unsigned VL, class T, class F>
bool test(queue q, std::string str, F funcUnderTest) {
std::cout << "Testing " << str << ", VL = " << VL << " ...\n";
size_t Size = 4 * VL;
constexpr size_t Size = 4 * VL;
T A[Size];
T B[Size];
constexpr unsigned HalfVL = VL > 1 ? (VL / 2) : 1;
Expand Down
10 changes: 3 additions & 7 deletions sycl/test-e2e/ESIMD/dpas/dpas_common.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -51,9 +51,7 @@ std::string toString(dpas_argument_type T) {
return "bf16";
case dpas_argument_type::tf32:
return "tf32";
case dpas_argument_type::s1:
case dpas_argument_type::u1:
case dpas_argument_type::Invalid:
default:
return "UNSUPPORTED";
}
return "UNRECOGNIZED";
Expand Down Expand Up @@ -127,9 +125,7 @@ template <dpas_argument_type T> constexpr int getBitSize() {
case dpas_argument_type::tf32:
return 32;

case dpas_argument_type::Invalid:
case dpas_argument_type::s1:
case dpas_argument_type::u1:
default:
break;
}
return 0;
Expand Down Expand Up @@ -405,7 +401,7 @@ bool test(queue &Q, bool Print) {
<< ") != expected (" << GoldRes << ")" << std::endl;
}
} // end for JJ
} // end for II
} // end for II

free(Res, Q);
free(APacked, Q);
Expand Down
4 changes: 4 additions & 0 deletions sycl/test-e2e/ESIMD/esimd_test_utils.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -27,6 +27,10 @@ using namespace sycl;

namespace esimd_test {

template <typename T>
using shared_allocator = sycl::usm_allocator<T, sycl::usm::alloc::shared>;
template <typename T> using shared_vector = std::vector<T, shared_allocator<T>>;

// This is the function provided to SYCL runtime by the application to decide
// on which device to run, or whether to run at all.
// When selecting a device, SYCL runtime first takes (1) a selector provided by
Expand Down
5 changes: 4 additions & 1 deletion sycl/test-e2e/ESIMD/ext_math.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -388,7 +388,10 @@ bool test(queue &Q, const std::string &Name, InitF Init = InitNarrow<T>{},
if constexpr (sizeof(T) <= 2)
delta = delta + delta;

bool BothFinite = std::isfinite(Test) && std::isfinite(Gold);
bool BothFinite = true;
#ifndef TEST_FAST_MATH
BothFinite = std::isfinite(Test) && std::isfinite(Gold);
#endif
if (BothFinite && std::abs(Test - Gold) > delta) {
if (++ErrCnt < 10) {
std::cout << " failed at index " << I << ", " << Test
Expand Down
16 changes: 2 additions & 14 deletions sycl/test-e2e/ESIMD/grf.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -71,13 +71,10 @@ int main(void) {
A[i] = i;
}

queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
esimd_test::printTestLabel(q);
try {
buffer<float, 1> bufa(A.data(), range<1>(Size));
queue q(gpu_selector{}, esimd_test::createExceptionHandler());

auto dev = q.get_device();
std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";

auto e = q.submit([&](handler &cgh) {
auto PA = bufa.get_access<access::mode::read_write>(cgh);
cgh.parallel_for<class SyclKernel>(Size,
Expand All @@ -98,11 +95,6 @@ int main(void) {

try {
buffer<float, 1> bufa(A.data(), range<1>(Size));
queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());

auto dev = q.get_device();
std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";

auto e = q.submit([&](handler &cgh) {
auto PA = bufa.get_access<access::mode::read_write>(cgh);
cgh.parallel_for<class EsimdKernel>(Size, [=](id<1> i) SYCL_ESIMD_KERNEL {
Expand All @@ -128,7 +120,6 @@ int main(void) {

try {
buffer<float, 1> bufa(A.data(), range<1>(Size));
queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
#ifdef USE_AUTO
sycl::ext::oneapi::experimental::properties prop{grf_size_automatic};
#elif defined(USE_NEW_API)
Expand All @@ -137,9 +128,6 @@ int main(void) {
sycl::ext::oneapi::experimental::properties prop{
register_alloc_mode<register_alloc_mode_enum::large>};
#endif
auto dev = q.get_device();
std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";

auto e = q.submit([&](handler &cgh) {
auto PA = bufa.get_access<access::mode::read_write>(cgh);
cgh.parallel_for<class EsimdKernelSpecifiedGRF>(
Expand Down
31 changes: 11 additions & 20 deletions sycl/test-e2e/ESIMD/histogram.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -85,15 +85,15 @@ int main(int argc, char *argv[]) {

// Allocate Input Buffer
queue q = esimd_test::createQueue();
esimd_test::printTestLabel(q);

auto dev = q.get_device();
unsigned char *srcY = malloc_shared<unsigned char>(width * height, q);
if (srcY == NULL) {
std::cerr << "Out of memory\n";
exit(1);
}
unsigned int *bins = malloc_shared<unsigned int>(NUM_BINS, q);
std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
esimd_test::shared_vector<uint8_t> srcY_vec(
width * height, esimd_test::shared_allocator<uint8_t>{q});
esimd_test::shared_vector<unsigned int> bins_vec(
NUM_BINS, esimd_test::shared_allocator<unsigned int>{q});
uint8_t *srcY = srcY_vec.data();
;
unsigned int *bins = bins_vec.data();

uint range_width = width / BLOCK_WIDTH;
uint range_height = height / BLOCK_HEIGHT;
Expand All @@ -106,16 +106,12 @@ int main(int argc, char *argv[]) {
FILE *f = fopen(input_file, "rb");
if (f == NULL) {
std::cerr << "Error opening file " << input_file;
free(srcY, q);
free(bins, q);
std::exit(1);
}

unsigned int cnt = fread(srcY, sizeof(unsigned char), input_size, f);
if (cnt != input_size) {
std::cerr << "Error reading input from " << input_file;
free(srcY, q);
free(bins, q);
std::exit(1);
}
} else {
Expand Down Expand Up @@ -171,18 +167,17 @@ int main(int argc, char *argv[]) {
uint h_pos = (tid % range_width) * BLOCK_WIDTH;
uint v_pos = (tid / range_width) * BLOCK_HEIGHT;

// Declare a 8x32 uchar matrix to store the input block pixel
// Declare a 8x32 uint8_t matrix to store the input block pixel
// value
simd<unsigned char, 8 * 32> in;
simd<uint8_t, 8 * 32> in;

// Declare a vector to store the local histogram
simd<unsigned int, NUM_BINS> histogram(0);

// Each thread handles BLOCK_HEIGHTxBLOCK_WIDTH pixel block
for (int y = 0; y < BLOCK_HEIGHT / 8; y++) {
// Perform 2D media block read to load 8x32 pixel block
in = media_block_load<unsigned char, 8, 32>(readAcc, h_pos,
v_pos);
in = media_block_load<uint8_t, 8, 32>(readAcc, h_pos, v_pos);

// Accumulate local histogram for each pixel value
#pragma unroll
Expand Down Expand Up @@ -236,8 +231,6 @@ int main(int argc, char *argv[]) {
// make sure data is given back to the host at the end of this scope
} catch (sycl::exception const &e) {
std::cerr << "SYCL exception caught: " << e.what() << '\n';
free(srcY, q);
free(bins, q);
return 1;
}

Expand All @@ -251,8 +244,6 @@ int main(int argc, char *argv[]) {
writeHist(cpuHistogram);
// Checking Histogram
bool Success = checkHistogram(cpuHistogram, bins);
free(srcY, q);
free(bins, q);

if (!Success) {
std::cerr << "FAILED\n";
Expand Down
28 changes: 12 additions & 16 deletions sycl/test-e2e/ESIMD/histogram_256_slm.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -45,7 +45,7 @@ ESIMD_INLINE void histogram_atomic(const uint32_t *input_ptr, uint32_t *output,
auto start_addr = ((unsigned int *)input_ptr) + start_off;
simd<uint, 32> data;
data.copy_from(start_addr);
auto in = data.bit_cast_view<uchar>();
auto in = data.bit_cast_view<uint8_t>();

#pragma unroll
for (int j = 0; j < BLOCK_WIDTH * sizeof(int); j += 16) {
Expand Down Expand Up @@ -75,7 +75,7 @@ void HistogramCPU(unsigned int size, unsigned int *src,
unsigned int *cpu_histogram) {
for (int i = 0; i < size; i++) {
unsigned int x = src[i];
cpu_histogram[(x)&0xFFU] += 1;
cpu_histogram[(x) & 0xFFU] += 1;
cpu_histogram[(x >> 8) & 0xFFU] += 1;
cpu_histogram[(x >> 16) & 0xFFU] += 1;
cpu_histogram[(x >> 24) & 0xFFU] += 1;
Expand Down Expand Up @@ -104,14 +104,18 @@ int CheckHistogram(unsigned int *cpu_histogram, unsigned int *gpu_histogram) {

int main() {
queue q = esimd_test::createQueue();
esimd_test::printTestLabel(q);

const char *input_file = nullptr;
unsigned int width = 1024;
unsigned int height = 1024;

// Initializes input.
unsigned int input_size = width * height;
unsigned int *input_ptr = malloc_shared<unsigned int>(input_size, q);

esimd_test::shared_vector<unsigned int> input_vec(
input_size, esimd_test::shared_allocator<unsigned int>{q});
unsigned int *input_ptr = input_vec.data();
printf("Processing %dx%d inputs\n", width, height);

srand(2009);
Expand All @@ -124,13 +128,8 @@ int main() {

// Allocates system memory for output buffer.
int buffer_size = sizeof(unsigned int) * NUM_BINS;
unsigned int *hist = new unsigned int[buffer_size];
if (hist == nullptr) {
free(input_ptr, q);
std::cerr << "Out of memory\n";
exit(1);
}
memset(hist, 0, buffer_size);
std::vector<unsigned int> hist_vec(buffer_size, 0);
unsigned int *hist = hist_vec.data();

// Uses the CPU to calculate the histogram output data.
unsigned int cpu_histogram[NUM_BINS];
Expand All @@ -141,7 +140,9 @@ int main() {
std::cout << "finish cpu_histogram\n";

// Uses the GPU to calculate the histogram output data.
unsigned int *output_surface = malloc_shared<unsigned int>(NUM_BINS, q);
esimd_test::shared_vector<unsigned int> output_vec(
NUM_BINS, esimd_test::shared_allocator<unsigned int>{q});
unsigned int *output_surface = output_vec.data();

unsigned int num_threads;
num_threads = width * height / (NUM_BLOCKS * BLOCK_WIDTH);
Expand Down Expand Up @@ -194,9 +195,6 @@ int main() {

memcpy(hist, output_surface, 4 * NUM_BINS);

free(output_surface, q);
free(input_ptr, q);

// Compares the CPU histogram output data with the
// GPU histogram output data.
// If there is no difference, the result is correct.
Expand All @@ -207,7 +205,5 @@ int main() {
else
std::cout << "FAILED\n";

delete[] hist;

return res ? 0 : -1;
}
Loading

0 comments on commit 436e687

Please sign in to comment.