From 2e11a85d9f2dcbdfd7f4d52ae2db6624edfd67ce Mon Sep 17 00:00:00 2001 From: "Reyes Castro, Ruyman" Date: Sat, 10 May 2025 21:58:35 +0100 Subject: [PATCH 1/9] [SYCL][COMPAT] Fix error on headers, add helloworld test --- sycl/doc/syclcompat/README.md | 10 +- sycl/include/syclcompat/traits.hpp | 2 +- sycl/test-e2e/Compat/helloworld.cpp | 137 ++++++++++++++++++++++++++++ 3 files changed, 145 insertions(+), 4 deletions(-) create mode 100644 sycl/test-e2e/Compat/helloworld.cpp diff --git a/sycl/doc/syclcompat/README.md b/sycl/doc/syclcompat/README.md index 4f395ea26760a..5f83cf968a6af 100644 --- a/sycl/doc/syclcompat/README.md +++ b/sycl/doc/syclcompat/README.md @@ -3481,8 +3481,8 @@ using this library: #include #include -#include #include +#include /** * Slope intercept form of a straight line equation: Y = m * X + b @@ -3571,7 +3571,11 @@ int main(int argc, char **argv) { // Check output for (size_t i = 0; i < n_points; i++) { - assert(h_Y[i] - h_expected[i] < 1e-6); + if (std::abs(h_Y[i] - h_expected[i]) >= 1e-6) { + std::cerr << "Mismatch at index " << i << ": expected " << h_expected[i] + << ", but got " << h_Y[i] << std::endl; + exit(EXIT_FAILURE); + } } // Clean up memory @@ -3581,7 +3585,7 @@ int main(int argc, char **argv) { syclcompat::free(d_X); syclcompat::free(d_Y); - return 0; + return EXIT_SUCCESS; } ``` diff --git a/sycl/include/syclcompat/traits.hpp b/sycl/include/syclcompat/traits.hpp index 7ed4d765251bc..30719b376ae43 100644 --- a/sycl/include/syclcompat/traits.hpp +++ b/sycl/include/syclcompat/traits.hpp @@ -87,7 +87,7 @@ template struct range_to_item_map> { using ItemT = sycl::nd_item; }; template struct range_to_item_map> { - using ItemT = sycl::item; + using ItemT = sycl::item; }; template diff --git a/sycl/test-e2e/Compat/helloworld.cpp b/sycl/test-e2e/Compat/helloworld.cpp new file mode 100644 index 0000000000000..c00c4357c2d51 --- /dev/null +++ b/sycl/test-e2e/Compat/helloworld.cpp @@ -0,0 +1,137 @@ +/*************************************************************************** + * + * Copyright (C) Codeplay Software Ltd. + * + * Part of the LLVM Project, under the Apache License v2.0 with LLVM + * Exceptions. See https://llvm.org/LICENSE.txt for license information. + * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + * + * SYCLcompat API + * + * helloworld.cpp + * + * Description: + * Checks that the SYCLcompat example program compiles and runs + **************************************************************************/ + +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +#include +#include + +#include +#include + +#include +#include + +/** + * Slope intercept form of a straight line equation: Y = m * X + b + */ +template +void slope_intercept(float *Y, float *X, float m, float b, size_t n) { + + // Block index + size_t bx = syclcompat::work_group_id::x(); + // Thread index + size_t tx = syclcompat::local_id::x(); + + size_t i = bx * BLOCK_SIZE + tx; + // or i = syclcompat::global_id::x(); + if (i < n) + Y[i] = m * X[i] + b; +} + +void check_memory(void *ptr, std::string msg) { + if (ptr == nullptr) { + std::cerr << "Failed to allocate memory: " << msg << std::endl; + exit(EXIT_FAILURE); + } +} + +/** + * Program main + */ +int main(int argc, char **argv) { + std::cout << "Simple Kernel example" << std::endl; + + constexpr size_t n_points = 32; + constexpr float m = 1.5f; + constexpr float b = 0.5f; + + int block_size = 32; + if (block_size > syclcompat::get_current_device() + .get_info()) + block_size = 16; + + std::cout << "block_size = " << block_size << ", n_points = " << n_points + << std::endl; + + // Allocate host memory for vectors X and Y + size_t mem_size = n_points * sizeof(float); + float *h_X = (float *)syclcompat::malloc_host(mem_size); + float *h_Y = (float *)syclcompat::malloc_host(mem_size); + check_memory(h_X, "h_X allocation failed."); + check_memory(h_Y, "h_Y allocation failed."); + + // Alternative templated allocation for the expected output + float *h_expected = syclcompat::malloc_host(n_points); + check_memory(h_expected, "Not enough for h_expected."); + + // Initialize host memory & expected output + for (size_t i = 0; i < n_points; i++) { + h_X[i] = i + 1; + h_expected[i] = m * h_X[i] + b; + } + + // Allocate device memory + float *d_X = (float *)syclcompat::malloc(mem_size); + float *d_Y = (float *)syclcompat::malloc(mem_size); + check_memory(d_X, "d_X allocation failed."); + check_memory(d_Y, "d_Y allocation failed."); + + // copy host memory to device + syclcompat::memcpy(d_X, h_X, mem_size); + + size_t threads = block_size; + size_t grid = n_points / block_size; + + std::cout << "Computing result using SYCL Kernel... "; + if (block_size == 16) { + syclcompat::launch>(grid, threads, d_Y, d_X, m, b, + n_points); + } else { + syclcompat::launch>(grid, threads, d_Y, d_X, m, b, + n_points); + } + syclcompat::wait(); + std::cout << "DONE" << std::endl; + + // Async copy result from device to host + syclcompat::memcpy_async(h_Y, d_Y, mem_size).wait(); + + // Check output + for (size_t i = 0; i < n_points; i++) { + if (std::abs(h_Y[i] - h_expected[i]) >= 1e-6) { + std::cerr << "Mismatch at index " << i << ": expected " << h_expected[i] + << ", but got " << h_Y[i] << std::endl; + exit(EXIT_FAILURE); + } + } + + // Clean up memory + syclcompat::free(h_X); + syclcompat::free(h_Y); + syclcompat::free(h_expected); + syclcompat::free(d_X); + syclcompat::free(d_Y); + + return EXIT_SUCCESS; +} \ No newline at end of file From ec52be3787ea2d718c46dcab4eb5bcd5205e8868 Mon Sep 17 00:00:00 2001 From: Ruyman Date: Sat, 10 May 2025 22:41:30 +0100 Subject: [PATCH 2/9] Remove duplicated header Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> --- sycl/test-e2e/Compat/helloworld.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/sycl/test-e2e/Compat/helloworld.cpp b/sycl/test-e2e/Compat/helloworld.cpp index c00c4357c2d51..5970f02f56ed9 100644 --- a/sycl/test-e2e/Compat/helloworld.cpp +++ b/sycl/test-e2e/Compat/helloworld.cpp @@ -29,7 +29,6 @@ #include #include -#include #include /** From 9673473479c8220d45d4b66a427134102224a14f Mon Sep 17 00:00:00 2001 From: "Reyes Castro, Ruyman" Date: Mon, 12 May 2025 13:27:15 +0100 Subject: [PATCH 3/9] Addressing feedback from reviewers --- sycl/test-e2e/{Compat => syclcompat}/helloworld.cpp | 9 ++++----- 1 file changed, 4 insertions(+), 5 deletions(-) rename sycl/test-e2e/{Compat => syclcompat}/helloworld.cpp (98%) diff --git a/sycl/test-e2e/Compat/helloworld.cpp b/sycl/test-e2e/syclcompat/helloworld.cpp similarity index 98% rename from sycl/test-e2e/Compat/helloworld.cpp rename to sycl/test-e2e/syclcompat/helloworld.cpp index 5970f02f56ed9..da0365e27f579 100644 --- a/sycl/test-e2e/Compat/helloworld.cpp +++ b/sycl/test-e2e/syclcompat/helloworld.cpp @@ -26,11 +26,9 @@ #include #include -#include +#include #include -#include - /** * Slope intercept form of a straight line equation: Y = m * X + b */ @@ -67,8 +65,9 @@ int main(int argc, char **argv) { int block_size = 32; if (block_size > syclcompat::get_current_device() - .get_info()) + .get_info()) { block_size = 16; + } std::cout << "block_size = " << block_size << ", n_points = " << n_points << std::endl; @@ -133,4 +132,4 @@ int main(int argc, char **argv) { syclcompat::free(d_Y); return EXIT_SUCCESS; -} \ No newline at end of file +} From 535faa666486b5364eaa6358548c6b9f3867e0a2 Mon Sep 17 00:00:00 2001 From: "Reyes Castro, Ruyman" Date: Mon, 12 May 2025 21:53:42 +0100 Subject: [PATCH 4/9] README now links to the source code instead of copy/pasting the code itself --- sycl/doc/syclcompat/README.md | 116 +--------------------------------- 1 file changed, 2 insertions(+), 114 deletions(-) diff --git a/sycl/doc/syclcompat/README.md b/sycl/doc/syclcompat/README.md index 5f83cf968a6af..e3b4e31b53167 100644 --- a/sycl/doc/syclcompat/README.md +++ b/sycl/doc/syclcompat/README.md @@ -3474,120 +3474,8 @@ public: ## Sample Code -Below is a simple linear algebra sample, which computes `y = mx + b` implemented -using this library: - -``` c++ -#include -#include - -#include -#include - -/** - * Slope intercept form of a straight line equation: Y = m * X + b - */ -template -void slope_intercept(float *Y, float *X, float m, float b, size_t n) { - - // Block index - size_t bx = syclcompat::work_group_id::x(); - // Thread index - size_t tx = syclcompat::local_id::x(); - - size_t i = bx * BLOCK_SIZE + tx; - // or i = syclcompat::global_id::x(); - if (i < n) - Y[i] = m * X[i] + b; -} - -void check_memory(void *ptr, std::string msg) { - if (ptr == nullptr) { - std::cerr << "Failed to allocate memory: " << msg << std::endl; - exit(EXIT_FAILURE); - } -} - -/** - * Program main - */ -int main(int argc, char **argv) { - std::cout << "Simple Kernel example" << std::endl; - - constexpr size_t n_points = 32; - constexpr float m = 1.5f; - constexpr float b = 0.5f; - - int block_size = 32; - if (block_size > syclcompat::get_current_device() - .get_info()) - block_size = 16; - - std::cout << "block_size = " << block_size << ", n_points = " << n_points - << std::endl; - - // Allocate host memory for vectors X and Y - size_t mem_size = n_points * sizeof(float); - float *h_X = (float *)syclcompat::malloc_host(mem_size); - float *h_Y = (float *)syclcompat::malloc_host(mem_size); - check_memory(h_X, "h_X allocation failed."); - check_memory(h_Y, "h_Y allocation failed."); - - // Alternative templated allocation for the expected output - float *h_expected = syclcompat::malloc_host(n_points); - check_memory(h_expected, "Not enough for h_expected."); - - // Initialize host memory & expected output - for (size_t i = 0; i < n_points; i++) { - h_X[i] = i + 1; - h_expected[i] = m * h_X[i] + b; - } - - // Allocate device memory - float *d_X = (float *)syclcompat::malloc(mem_size); - float *d_Y = (float *)syclcompat::malloc(mem_size); - check_memory(d_X, "d_X allocation failed."); - check_memory(d_Y, "d_Y allocation failed."); - - // copy host memory to device - syclcompat::memcpy(d_X, h_X, mem_size); - - size_t threads = block_size; - size_t grid = n_points / block_size; - - std::cout << "Computing result using SYCL Kernel... "; - if (block_size == 16) { - syclcompat::launch>(grid, threads, d_Y, d_X, m, b, - n_points); - } else { - syclcompat::launch>(grid, threads, d_Y, d_X, m, b, - n_points); - } - syclcompat::wait(); - std::cout << "DONE" << std::endl; - - // Async copy result from device to host - syclcompat::memcpy_async(h_Y, d_Y, mem_size).wait(); - - // Check output - for (size_t i = 0; i < n_points; i++) { - if (std::abs(h_Y[i] - h_expected[i]) >= 1e-6) { - std::cerr << "Mismatch at index " << i << ": expected " << h_expected[i] - << ", but got " << h_Y[i] << std::endl; - exit(EXIT_FAILURE); - } - } - - // Clean up memory - syclcompat::free(h_X); - syclcompat::free(h_Y); - syclcompat::free(h_expected); - syclcompat::free(d_X); - syclcompat::free(d_Y); - - return EXIT_SUCCESS; -} -``` +The file [helloworld.cpp](sycl/test-e2e/syclcompat/helloworld.cpp) contains +a simple example which computes `y = mx + b` implemented using this library. ## Maintainers From 4a38d81cb0815efb762cf6f5efc1609dcdfc1433 Mon Sep 17 00:00:00 2001 From: "Reyes Castro, Ruyman" Date: Mon, 12 May 2025 21:54:52 +0100 Subject: [PATCH 5/9] Simplify check code and use \n instead of endl --- sycl/test-e2e/syclcompat/helloworld.cpp | 31 ++++++++++++------------- 1 file changed, 15 insertions(+), 16 deletions(-) diff --git a/sycl/test-e2e/syclcompat/helloworld.cpp b/sycl/test-e2e/syclcompat/helloworld.cpp index da0365e27f579..15101b8b83af5 100644 --- a/sycl/test-e2e/syclcompat/helloworld.cpp +++ b/sycl/test-e2e/syclcompat/helloworld.cpp @@ -29,6 +29,12 @@ #include #include +#define CHECK_MEMORY(ptr) \ + if ((ptr) == nullptr) { \ + std::cerr << "Failed to allocate memory: " << ( #ptr ) << "\n"; \ + exit(EXIT_FAILURE); \ + } + /** * Slope intercept form of a straight line equation: Y = m * X + b */ @@ -46,18 +52,11 @@ void slope_intercept(float *Y, float *X, float m, float b, size_t n) { Y[i] = m * X[i] + b; } -void check_memory(void *ptr, std::string msg) { - if (ptr == nullptr) { - std::cerr << "Failed to allocate memory: " << msg << std::endl; - exit(EXIT_FAILURE); - } -} - /** * Program main */ int main(int argc, char **argv) { - std::cout << "Simple Kernel example" << std::endl; + std::cout << "Simple Kernel example" << "\n"; constexpr size_t n_points = 32; constexpr float m = 1.5f; @@ -70,18 +69,18 @@ int main(int argc, char **argv) { } std::cout << "block_size = " << block_size << ", n_points = " << n_points - << std::endl; + << "\n"; // Allocate host memory for vectors X and Y size_t mem_size = n_points * sizeof(float); float *h_X = (float *)syclcompat::malloc_host(mem_size); float *h_Y = (float *)syclcompat::malloc_host(mem_size); - check_memory(h_X, "h_X allocation failed."); - check_memory(h_Y, "h_Y allocation failed."); + CHECK_MEMORY(h_X); + CHECK_MEMORY(h_Y); // Alternative templated allocation for the expected output float *h_expected = syclcompat::malloc_host(n_points); - check_memory(h_expected, "Not enough for h_expected."); + CHECK_MEMORY(h_expected); // Initialize host memory & expected output for (size_t i = 0; i < n_points; i++) { @@ -92,8 +91,8 @@ int main(int argc, char **argv) { // Allocate device memory float *d_X = (float *)syclcompat::malloc(mem_size); float *d_Y = (float *)syclcompat::malloc(mem_size); - check_memory(d_X, "d_X allocation failed."); - check_memory(d_Y, "d_Y allocation failed."); + CHECK_MEMORY(d_X); + CHECK_MEMORY(d_Y); // copy host memory to device syclcompat::memcpy(d_X, h_X, mem_size); @@ -110,7 +109,7 @@ int main(int argc, char **argv) { n_points); } syclcompat::wait(); - std::cout << "DONE" << std::endl; + std::cout << "DONE" << "\n"; // Async copy result from device to host syclcompat::memcpy_async(h_Y, d_Y, mem_size).wait(); @@ -119,7 +118,7 @@ int main(int argc, char **argv) { for (size_t i = 0; i < n_points; i++) { if (std::abs(h_Y[i] - h_expected[i]) >= 1e-6) { std::cerr << "Mismatch at index " << i << ": expected " << h_expected[i] - << ", but got " << h_Y[i] << std::endl; + << ", but got " << h_Y[i] << "\n"; exit(EXIT_FAILURE); } } From 9429d3d827a479a428554bde0ae4fb9199a946c0 Mon Sep 17 00:00:00 2001 From: "Reyes Castro, Ruyman" Date: Tue, 13 May 2025 13:05:31 +0100 Subject: [PATCH 6/9] apply clang-format --- sycl/test-e2e/syclcompat/helloworld.cpp | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/sycl/test-e2e/syclcompat/helloworld.cpp b/sycl/test-e2e/syclcompat/helloworld.cpp index 15101b8b83af5..0467e9286511e 100644 --- a/sycl/test-e2e/syclcompat/helloworld.cpp +++ b/sycl/test-e2e/syclcompat/helloworld.cpp @@ -29,10 +29,10 @@ #include #include -#define CHECK_MEMORY(ptr) \ - if ((ptr) == nullptr) { \ - std::cerr << "Failed to allocate memory: " << ( #ptr ) << "\n"; \ - exit(EXIT_FAILURE); \ +#define CHECK_MEMORY(ptr) \ + if ((ptr) == nullptr) { \ + std::cerr << "Failed to allocate memory: " << (#ptr) << "\n"; \ + exit(EXIT_FAILURE); \ } /** @@ -118,7 +118,7 @@ int main(int argc, char **argv) { for (size_t i = 0; i < n_points; i++) { if (std::abs(h_Y[i] - h_expected[i]) >= 1e-6) { std::cerr << "Mismatch at index " << i << ": expected " << h_expected[i] - << ", but got " << h_Y[i] << "\n"; + << ", but got " << h_Y[i] << "\n"; exit(EXIT_FAILURE); } } From fd55b72880861c4364cc28fcb0ecab41562095a2 Mon Sep 17 00:00:00 2001 From: "Reyes Castro, Ruyman" Date: Tue, 13 May 2025 13:08:29 +0100 Subject: [PATCH 7/9] Fixed: Path to test file from README.md --- sycl/doc/syclcompat/README.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/doc/syclcompat/README.md b/sycl/doc/syclcompat/README.md index e3b4e31b53167..00d99e6c919f7 100644 --- a/sycl/doc/syclcompat/README.md +++ b/sycl/doc/syclcompat/README.md @@ -3474,7 +3474,7 @@ public: ## Sample Code -The file [helloworld.cpp](sycl/test-e2e/syclcompat/helloworld.cpp) contains +The file [helloworld.cpp](../../test-e2e/syclcompat/helloworld.cpp) contains a simple example which computes `y = mx + b` implemented using this library. ## Maintainers From f506ac68a5be26734bd2f10e6f5894948f781f80 Mon Sep 17 00:00:00 2001 From: "Reyes Castro, Ruyman" Date: Tue, 13 May 2025 13:14:29 +0100 Subject: [PATCH 8/9] Use specific includes instead of the global compat and sycl ones --- sycl/test-e2e/syclcompat/helloworld.cpp | 12 ++++++++++-- 1 file changed, 10 insertions(+), 2 deletions(-) diff --git a/sycl/test-e2e/syclcompat/helloworld.cpp b/sycl/test-e2e/syclcompat/helloworld.cpp index 0467e9286511e..7b895b9a86284 100644 --- a/sycl/test-e2e/syclcompat/helloworld.cpp +++ b/sycl/test-e2e/syclcompat/helloworld.cpp @@ -23,8 +23,16 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out -#include -#include +#include + +// The example uses specific headers but the user can +// simple include to get all the +// functionality with a single header + +#include +#include +#include +#include #include #include From 375a86d6c4fbfd6fcb50f6e264fd75166408354d Mon Sep 17 00:00:00 2001 From: "Reyes Castro, Ruyman" Date: Tue, 13 May 2025 16:28:21 +0100 Subject: [PATCH 9/9] Fix clang format --- sycl/test-e2e/syclcompat/helloworld.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/test-e2e/syclcompat/helloworld.cpp b/sycl/test-e2e/syclcompat/helloworld.cpp index 7b895b9a86284..3e32f8a965eb6 100644 --- a/sycl/test-e2e/syclcompat/helloworld.cpp +++ b/sycl/test-e2e/syclcompat/helloworld.cpp @@ -111,10 +111,10 @@ int main(int argc, char **argv) { std::cout << "Computing result using SYCL Kernel... "; if (block_size == 16) { syclcompat::launch>(grid, threads, d_Y, d_X, m, b, - n_points); + n_points); } else { syclcompat::launch>(grid, threads, d_Y, d_X, m, b, - n_points); + n_points); } syclcompat::wait(); std::cout << "DONE" << "\n";