From fa64381fa6d11c9d4ec16647f9d2ae88a8aab508 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Pekka=20J=C3=A4=C3=A4skel=C3=A4inen?= Date: Thu, 9 Nov 2023 17:12:12 +0200 Subject: [PATCH 01/17] README.md: missing dynamic/static arg in the llvm installer example --- README.md | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/README.md b/README.md index 0cac38807..7adc8cc2e 100644 --- a/README.md +++ b/README.md @@ -39,9 +39,9 @@ It's recommended to use the chipStar forks of LLVM and SPIRV-LLVM-Translator. For this you can use a script included in the chipStar repository: ```bash -# chipStar/scripts/configure_llvm.sh -chipStar/scripts/configure_llvm.sh 17 /opt/install/llvm/17.0 -cd ./llvm-project/llvm/build_17 +# chipStar/scripts/configure_llvm.sh +chipStar/scripts/configure_llvm.sh 17 /opt/install/llvm/17.0 dynamic +cd llvm-project/llvm/build_17 make -j 16 make install ``` From b37403deffe68756b83be67d720ecc8934c04dfc Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Pekka=20J=C3=A4=C3=A4skel=C3=A4inen?= Date: Fri, 10 Nov 2023 07:49:50 +0200 Subject: [PATCH 02/17] llvm config: Actually link dynamically when using 'dynamic'. (#679) Lacking this caused static linking of the LLVM binaries anyhow leading to massive disk space (e.g. opt 1.5GB) usage on my poor laptop. --- scripts/check.py | 2 +- scripts/configure_llvm.sh | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/scripts/check.py b/scripts/check.py index 5e5aa6b41..c205b18a6 100755 --- a/scripts/check.py +++ b/scripts/check.py @@ -84,7 +84,7 @@ def run_cmd(cmd): cmd = f"{modules} {env_vars} ./samples/hipInfo/hipInfo" out, _ = run_cmd(cmd) -texture_support = 0 < int(out.split("maxTexture1DLinear:")[1].split("\n")[0].strip()) +texture_support = "maxTexture1DLinear:" in out and 0 < int(out.split("maxTexture1DLinear:")[1].split("\n")[0].strip()) if not texture_support: texture_cmd = "|[Tt]ex" else: diff --git a/scripts/configure_llvm.sh b/scripts/configure_llvm.sh index 8fbfb9635..753119a78 100755 --- a/scripts/configure_llvm.sh +++ b/scripts/configure_llvm.sh @@ -71,7 +71,7 @@ elif [ "$BUILD_TYPE" == "dynamic" ]; then -DCMAKE_INSTALL_RPATH=${INSTALL_DIR}/lib \ -DLLVM_ENABLE_PROJECTS="clang;openmp" \ -DLLVM_TARGETS_TO_BUILD=host \ - -DLLVM_LINK_LLVM_DYLIB=OFF \ + -DLLVM_LINK_LLVM_DYLIB=ON \ -DLLVM_BUILD_LLVM_DYLIB=ON \ -DLLVM_PARALLEL_LINK_JOBS=2 \ -DCMAKE_BUILD_TYPE=RelWithDebInfo \ From 15ed7bcf69021cc814dc3ad6ea0125745b41e7a1 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Pekka=20J=C3=A4=C3=A4skel=C3=A4inen?= Date: Fri, 10 Nov 2023 15:09:17 +0200 Subject: [PATCH 03/17] README.md: Updated the outdated check.py instructions --- README.md | 8 ++++++-- 1 file changed, 6 insertions(+), 2 deletions(-) diff --git a/README.md b/README.md index 7adc8cc2e..22b803c2d 100644 --- a/README.md +++ b/README.md @@ -113,12 +113,16 @@ NOTE: If you don't have libOpenCL.so (for example from the `ocl-icd-opencl-dev` There's a script `check.py` which can be used to run unit tests and which filters out known failing tests for different platforms. Its usage is as follows. ```bash -# BACKEND={opencl/level0/pocl} # Which backend/driver you wish to test, "opencl" = Intel OpenCL runtime, "level0" = Intel LevelZero runtime, "pocl" = PoCL OpenCL runtime +# BACKEND={opencl/level0-{reg,imm}/pocl} +# ^ Which backend/driver/platform you wish to test: +# "opencl" = Intel OpenCL runtime, "level0" = Intel LevelZero runtime with regular command lists (reg) or immediate command lists (imm), "pocl" = PoCL OpenCL runtime # DEVICE={cpu,igpu,dgpu} # What kind of device to test. +# ^ This selects the expected test pass lists. +# 'igpu' is a Intel Iris Xe iGPU, 'dgpu' a typical recent Intel dGPU such as Data Center GPU Max series or an Arc. # PARALLEL={N} # How many tests to run in parallel. # export CHIP_PLATFORM=N # If there are multiple OpenCL platforms present on the system, selects which one to use -python3 $SOURCE_DIR/scripts/check.py --num-threads $PARALLEL $BUILD_DIR $DEVICE $BACKEND +python3 $SOURCE_DIR/scripts/check.py -m off --num-threads $PARALLEL $BUILD_DIR $DEVICE $BACKEND ``` Please refer to the [user documentation](docs/Using.md) for instructions on how to use the installed chipStar to build CUDA/HIP programs. From 0387f643ab46a22d8fd2a38346eb61f85094fc51 Mon Sep 17 00:00:00 2001 From: Michal Babej Date: Thu, 16 Nov 2023 11:02:30 +0200 Subject: [PATCH 04/17] Minor fixes to release notes & README.md (#684) --- README.md | 4 ++-- docs/release_notes/chipStar_1.1.rst | 4 +++- 2 files changed, 5 insertions(+), 3 deletions(-) diff --git a/README.md b/README.md index 22b803c2d..09d8fb401 100644 --- a/README.md +++ b/README.md @@ -135,8 +135,8 @@ CHIP_PLATFORM= # If there are multiple platform CHIP_DEVICE= # If there are multiple devices present on the system, selects which one to use. Defaults to 0 CHIP_LOGLEVEL= # Sets the log level. If compiled in RELEASE, only err/crit are available CHIP_DUMP_SPIRV= # Dumps the generated SPIR-V code to a file -CHIP_JIT_FLAGS= # String to override the default JIT flags. Defaults to -x spir -cl-kernel-arg-info -cl-std=CL3.0 -CHIP_L0_COLLECT_EVENTS_TIMEOUT= # Timeout in milliseconds for collecting Level Zero events +CHIP_JIT_FLAGS= # String to override the default JIT flags. Defaults to -cl-kernel-arg-info -cl-std=CL3.0 +CHIP_L0_COLLECT_EVENTS_TIMEOUT= # Timeout in seconds for collecting Level Zero events CHIP_L0_IMM_CMD_LISTS= # Use immediate command lists in Level Zero ``` diff --git a/docs/release_notes/chipStar_1.1.rst b/docs/release_notes/chipStar_1.1.rst index be9f210f3..0ca45cd0b 100644 --- a/docs/release_notes/chipStar_1.1.rst +++ b/docs/release_notes/chipStar_1.1.rst @@ -58,7 +58,9 @@ Previous versions of chipStar used command queue barriers excessively for synchronization, which led to limited opportunities for asynchronous execution. In chipStar 1.1, command queue synchronization is done using event dependencies, which leads to more task parallelism opportunities presented to the drivers, -speeding up various workloads significantly. +speeding up various workloads significantly. Workloads that do not exploit +parallelism but enqueue a lot of very small kernels (in the 10's of microseconds +range) may also benefit as the barrier itself could dominate the execution time. ============== Minor Features From 68370c29af2cb924fd2688c9437e1ed4f002d23b Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Wed, 22 Nov 2023 00:09:32 -0600 Subject: [PATCH 05/17] fix Intel CPU runner --- .github/workflows/presubmit.yml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/.github/workflows/presubmit.yml b/.github/workflows/presubmit.yml index cc9c88f17..bd18ebc5c 100644 --- a/.github/workflows/presubmit.yml +++ b/.github/workflows/presubmit.yml @@ -185,7 +185,7 @@ jobs: name: Build and test chipStar on ${{ matrix.backend }} (llvm-${{ matrix.version }}) runs-on: ubuntu-latest env: - OPENCL_ENV: ${{ matrix.backend == 'intel' && 'LD_LIBRARY_PATH=$(dirname $(find /opt/intel | grep libsvml.so)) OCL_ICD_FILENAMES=$(cat /etc/OpenCL/vendors/intel64.icd)' || 'OPENCL_VENDOR_PATH=$HOME/opt/pocl/4.0/etc/OpenCL/vendors/' }} + OPENCL_ENV: ${{ matrix.backend == 'intel' && 'LD_LIBRARY_PATH=$(dirname $(find /opt/intel -name libsvml.so)):$(dirname $(find /opt/intel/oneapi/tbb -name libtbb.so* -print -quit)):$LD_LIBRARY_PATH OCL_ICD_FILENAMES=$(cat /etc/OpenCL/vendors/intel64.icd)' || 'OPENCL_VENDOR_PATH=$HOME/opt/pocl/4.0/etc/OpenCL/vendors/' }} EXCLUDE: ${{ matrix.backend == 'intel' && '"`cat ./test_lists/cpu_opencl_failed_tests.txt`"' || '"`cat ./test_lists/cpu_pocl_failed_tests.txt`"' }} strategy: matrix: From 3ff3979df0823e4c6a91ca82f016b27746d13d4e Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Thu, 23 Nov 2023 10:44:54 +0200 Subject: [PATCH 06/17] Hipcc rewrite (#696) * Add test for -MD -MT options * Update HIPCC submodule to re-written HIPCC which uses loop parsing instead of regex --- HIPCC | 2 +- tests/compiler/CMakeLists.txt | 1 + tests/compiler/TestHipcc692Regression.bash | 38 ++++++++++++++++++++++ 3 files changed, 40 insertions(+), 1 deletion(-) create mode 100644 tests/compiler/TestHipcc692Regression.bash diff --git a/HIPCC b/HIPCC index 68df7aac2..f43fdf9bb 160000 --- a/HIPCC +++ b/HIPCC @@ -1 +1 @@ -Subproject commit 68df7aac29bccb4fbd3e306d45f4f2b9692d55b8 +Subproject commit f43fdf9bbda1fdb21962e83382ea166efcf1361a diff --git a/tests/compiler/CMakeLists.txt b/tests/compiler/CMakeLists.txt index b9ee474d6..de2499412 100644 --- a/tests/compiler/CMakeLists.txt +++ b/tests/compiler/CMakeLists.txt @@ -104,6 +104,7 @@ add_shell_test(TestRDCWithMultipleHipccCmds.bash) add_shell_test(TestWholeProgramCompilation.bash) add_shell_test(TestHipccDashX.bash) add_shell_test(TestHipccFp16Include.bash) +add_shell_test(TestHipcc692Regression.bash) add_test(NAME "TestHipccMultiSource" COMMAND ${CMAKE_BINARY_DIR}/bin/hipcc ${CMAKE_CURRENT_SOURCE_DIR}/TestHipccCompileThenLinkMain.cpp ${CMAKE_CURRENT_SOURCE_DIR}/TestHipccCompileThenLinkKernel.cpp -o TestHipccMultiSource) diff --git a/tests/compiler/TestHipcc692Regression.bash b/tests/compiler/TestHipcc692Regression.bash new file mode 100644 index 000000000..b2bba848d --- /dev/null +++ b/tests/compiler/TestHipcc692Regression.bash @@ -0,0 +1,38 @@ +#!/bin/bash +# Check that -MD -MT options are respected +set -eu + +BIN_DIR=@CMAKE_BINARY_DIR@ +SRC_DIR=@CMAKE_SOURCE_DIR@ +OUT_DIR=@CMAKE_CURRENT_BINARY_DIR@/@TEST_NAME@.d +HIPCC=@CMAKE_BINARY_DIR@/bin/hipcc + +mkdir -p ${OUT_DIR} + +touch ${OUT_DIR}/test_hipcub_basic.cpp +mkdir -p ${OUT_DIR}/CMakeFiles/test_hipcub_basic.dir + +export HIPCC_VERBOSE=7 + +${HIPCC} \ +-DGTEST_LINKED_AS_SHARED_LIBRARY=1 \ +-O3 -DNDEBUG \ +-x hip -D__HIP_PLATFORM_SPIRV__= --offload=spirv64 -nohipwrapperinc --hip-path=${BIN_DIR} --target=x86_64-unknown-linux-gnu \ +-include ${SRC_DIR}/include/hip/spirv_fixups.h -std=c++14 \ +-c ${OUT_DIR}/test_hipcub_basic.cpp \ +-MD -MT ${OUT_DIR}/CMakeFiles/test_hipcub_basic.dir/test_hipcub_basic.cpp.o \ +-MF ${OUT_DIR}/CMakeFiles/test_hipcub_basic.dir/test_hipcub_basic.cpp.o.d \ +-o ${OUT_DIR}/CMakeFiles/test_hipcub_basic.dir/test_hipcub_basic.cpp.o \ +| tee TestHipcc692ReggressionOutput.log +# make sure that -MT is respected +grep -- "-MT ${OUT_DIR}/CMakeFiles/test_hipcub_basic.dir/test_hipcub_basic.cpp.o" ./TestHipcc692ReggressionOutput.log +if [ $? -ne 0 ]; then + echo "-MT not respected" + exit 1 +fi + +grep -- "-MF ${OUT_DIR}/CMakeFiles/test_hipcub_basic.dir/test_hipcub_basic.cpp.o.d" ./TestHipcc692ReggressionOutput.log +if [ $? -ne 0 ]; then + echo "-MF not respected" + exit 1 +fi From b44bf5827ab5016aa0ff330081a21bc8ddf8f519 Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Fri, 24 Nov 2023 01:39:49 -0600 Subject: [PATCH 07/17] switch unit tests to single thread --- scripts/unit_tests.sh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/scripts/unit_tests.sh b/scripts/unit_tests.sh index 69241a4c1..d36274d6c 100755 --- a/scripts/unit_tests.sh +++ b/scripts/unit_tests.sh @@ -36,7 +36,7 @@ shift # Set the number of tries based on the argument or default to 1 num_tries=1 -num_threads=8 +num_threads=1 timeout=200 for arg in "$@" do From 0ec18bc84b461dc1bf356ded11c89d58d955324c Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Mon, 27 Nov 2023 12:40:46 +0200 Subject: [PATCH 08/17] Merge pull request #707 from CHIP-SPV/single-thread-tests Identify tests which timeout when run in parallel --- cmake/UnitTests.cmake | 60 +++++++++++++++++++++++++++++++++---------- scripts/check.py | 37 +++++++++++++++++++++----- scripts/unit_tests.sh | 6 ++--- 3 files changed, 81 insertions(+), 22 deletions(-) diff --git a/cmake/UnitTests.cmake b/cmake/UnitTests.cmake index fc68ffdfa..d4a50a71f 100644 --- a/cmake/UnitTests.cmake +++ b/cmake/UnitTests.cmake @@ -11,6 +11,47 @@ list(APPEND DGPU_LEVEL0_RCL_FAILED_TESTS " ") list(APPEND DGPU_LEVEL0_ICL_FAILED_TESTS " ") list(APPEND CPU_POCL_FAILED_TESTS " ") list(APPEND GPU_POCL_FAILED_TESTS " ") # TODO +list(APPEND NON_PARALLEL_TESTS " ") + +list(APPEND NON_PARALLEL_TESTS "clock") +list(APPEND NON_PARALLEL_TESTS "Unit_hipMemcpyWithStream_TestwithTwoStream") +list(APPEND NON_PARALLEL_TESTS "Unit_hipMemcpyWithStream_TestDtoDonSameDevice") +list(APPEND NON_PARALLEL_TESTS "Unit_hipMemcpy_H2H-H2D-D2H-H2PinMem - int") +list(APPEND NON_PARALLEL_TESTS "Unit_hipMemcpy_H2H-H2D-D2H-H2PinMem - float") +list(APPEND NON_PARALLEL_TESTS "Unit_hipMemcpy_H2H-H2D-D2H-H2PinMem - double") +list(APPEND NON_PARALLEL_TESTS "broadcast2") +list(APPEND NON_PARALLEL_TESTS "Unit_hipMemsetFunctional_SmallSize_3D") +list(APPEND NON_PARALLEL_TESTS "Unit_hipMemcpy_KernelLaunch - int") +list(APPEND NON_PARALLEL_TESTS "Unit_hipMemcpy_KernelLaunch - float") +list(APPEND NON_PARALLEL_TESTS "Unit_hipMemcpy_KernelLaunch - double") +list(APPEND NON_PARALLEL_TESTS "fp16") +list(APPEND NON_PARALLEL_TESTS "SimpleConvolution") +list(APPEND NON_PARALLEL_TESTS "Unit_hipHostMalloc_Basic") +list(APPEND NON_PARALLEL_TESTS "Unit_hipMalloc_LoopRegressionAllocFreeCycles") +list(APPEND NON_PARALLEL_TESTS "Unit_hipMultiThreadStreams1_AsyncAsync") +list(APPEND NON_PARALLEL_TESTS "cuda-scan") +list(APPEND NON_PARALLEL_TESTS "Unit_hipMemset3DAsync_ConcurrencyMthread") +list(APPEND NON_PARALLEL_TESTS "cuda-bandwidthTest") +list(APPEND NON_PARALLEL_TESTS "Unit_hipMemsetAsync_QueueJobsMultithreaded") +list(APPEND NON_PARALLEL_TESTS "Unit_hipMemset2DAsync_MultiThread") +list(APPEND NON_PARALLEL_TESTS "DCT") +list(APPEND NON_PARALLEL_TESTS "Unit_hipMemsetFunctional_ZeroSize_3D") +list(APPEND NON_PARALLEL_TESTS "cuda-matrixMul") +list(APPEND NON_PARALLEL_TESTS "cuda-FDTD3d") +list(APPEND NON_PARALLEL_TESTS "Unit_hipMultiThreadStreams1_AsyncSync") +list(APPEND NON_PARALLEL_TESTS "FastWalshTransform") +list(APPEND NON_PARALLEL_TESTS "Unit_hipMultiStream_sameDevice") +list(APPEND NON_PARALLEL_TESTS "Unit_hipStreamCreate_MultistreamBasicFunctionalities") +list(APPEND NON_PARALLEL_TESTS "Unit_hipMemcpy_MultiThreadWithSerialization") +list(APPEND NON_PARALLEL_TESTS "dwtHaar1D") +list(APPEND NON_PARALLEL_TESTS "Unit_hipHostRegister_Memcpy - int") +list(APPEND NON_PARALLEL_TESTS "Unit_hipHostRegister_Memcpy - float") +list(APPEND NON_PARALLEL_TESTS "Unit_hipHostRegister_Memcpy - double") +list(APPEND NON_PARALLEL_TESTS "TestStlFunctionsDouble") +list(APPEND NON_PARALLEL_TESTS "Unit_hipMemset_SetMemoryWithOffset") +list(APPEND NON_PARALLEL_TESTS "Unit_hipMemsetAsync_SetMemoryWithOffset") +list(APPEND NON_PARALLEL_TESTS "BitonicSort") +list(APPEND NON_PARALLEL_TESTS "FloydWarshall") # This test gets enabled only if LLVM' FileCheck tool is found in PATH. # It fails with "error: cannot find ROCm device library; @@ -732,6 +773,9 @@ list(APPEND IGPU_OPENCL_FAILED_TESTS "cuda-simpleCallback") # SEGFAULT list(APPEND IGPU_OPENCL_FAILED_TESTS "Unit_hipGraphMemcpyNodeSetParams_Functional") # Subprocess aborted # dGPU OpenCL Unit Test Failures +list(APPEND DGPU_OPENCL_FAILED_TESTS "Unit_hipGraphEventRecordNodeSetEvent_SetEventProperty") # flaky +list(APPEND DGPU_OPENCL_FAILED_TESTS "Unit_hipGraphAddEventRecordNode_Functional_ElapsedTime") # flaky +list(APPEND DGPU_OPENCL_FAILED_TESTS "Unit_hipEventRecord") # flaky list(APPEND DGPU_OPENCL_FAILED_TESTS "Unit_hipMultiThreadStreams1_AsyncSame") # invalid free() list(APPEND DGPU_OPENCL_FAILED_TESTS "Unit_hipStreamPerThread_MultiThread") list(APPEND DGPU_OPENCL_FAILED_TESTS "Unit_hipStreamPerThread_DeviceReset_1") @@ -2061,8 +2105,7 @@ string(REGEX REPLACE ";" "\$|" IGPU_LEVEL0_RCL_FAILED_TESTS_STR "${IGPU_LEVEL0_R string(REGEX REPLACE ";" "\$|" IGPU_LEVEL0_ICL_FAILED_TESTS_STR "${IGPU_LEVEL0_ICL_FAILED_TESTS}") string(REGEX REPLACE ";" "\$|" CPU_POCL_FAILED_TESTS_STR "${CPU_POCL_FAILED_TESTS}") string(REGEX REPLACE ";" "\$|" ALL_FAILED_TESTS_STR "${ALL_FAILED_TESTS}") - -add_custom_target(check COMMAND ${CMAKE_CTEST_COMMAND} ${TEST_OPTIONS} -E ${ALL_FAILED_TESTS_STR} VERBATIM) +string(REGEX REPLACE ";" "\$|" NON_PARALLEL_TESTS_STR "${NON_PARALLEL_TESTS}") string(CONCAT DGPU_OPENCL_FAILED_TESTS_STR ${DGPU_OPENCL_FAILED_TESTS_STR} "\$|") string(CONCAT IGPU_OPENCL_FAILED_TESTS_STR ${IGPU_OPENCL_FAILED_TESTS_STR} "\$|") @@ -2073,6 +2116,7 @@ string(CONCAT IGPU_LEVEL0_RCL_FAILED_TESTS_STR ${IGPU_LEVEL0_RCL_FAILED_TESTS_ST string(CONCAT IGPU_LEVEL0_ICL_FAILED_TESTS_STR ${IGPU_LEVEL0_ICL_FAILED_TESTS_STR} "\$|") string(CONCAT CPU_POCL_FAILED_TESTS_STR ${CPU_POCL_FAILED_TESTS_STR} "\$|") string(CONCAT ALL_FAILED_TESTS_STR ${ALL_FAILED_TESTS_STR} "\$|") +string(CONCAT NON_PARALLEL_TESTS_STR ${NON_PARALLEL_TESTS_STR} "\$|") FILE(WRITE "${CMAKE_BINARY_DIR}/test_lists/dgpu_opencl_failed_tests.txt" "\"${DGPU_OPENCL_FAILED_TESTS_STR}\"") FILE(WRITE "${CMAKE_BINARY_DIR}/test_lists/igpu_opencl_failed_tests.txt" "\"${IGPU_OPENCL_FAILED_TESTS_STR}\"") @@ -2083,14 +2127,4 @@ FILE(WRITE "${CMAKE_BINARY_DIR}/test_lists/igpu_level0_failed_reg_tests.txt" "\" FILE(WRITE "${CMAKE_BINARY_DIR}/test_lists/igpu_level0_failed_imm_tests.txt" "\"${IGPU_LEVEL0_ICL_FAILED_TESTS_STR}\"") FILE(WRITE "${CMAKE_BINARY_DIR}/test_lists/cpu_pocl_failed_tests.txt" "\"${CPU_POCL_FAILED_TESTS_STR}\"") FILE(WRITE "${CMAKE_BINARY_DIR}/test_lists/all_failed_tests.txt" "\"${ALL_FAILED_TESTS_STR}\"") - -# TODO fix-254 how do I make these read from the environment? -# MULTI_TESTS_REPEAT=33 make multi_tests -# Preferably without an additional reconfigure. Every way I tried escaping ${MULTI_TESTS_REPEAT} results in something undesirable like \${MULTI_TESTS_REPEAT} -set(FLAKY_TESTS_REPEAT 100) -set(MULTI_TESTS_REPEAT 10) -set(PARALLEL_TESTS 1) - -set(TEST_OPTIONS -j ${PARALLEL_TESTS} --timeout 120 --output-on-failure) -add_custom_target(flaky_tests COMMAND ${CMAKE_CTEST_COMMAND} ${TEST_OPTIONS} -R ${FLAKY_TESTS} --repeat until-fail:${FLAKY_TESTS_REPEAT} USES_TERMINAL VERBATIM) -add_custom_target(multi_tests COMMAND ${CMAKE_CTEST_COMMAND} ${TEST_OPTIONS} -R "[Aa]sync|[Mm]ulti[Tt]hread|[Mm]ulti[Ss]tream|[Tt]hread|[Ss]tream" --repeat until-fail:${MULTI_TESTS_REPEAT} USES_TERMINAL VERBATIM) +FILE(WRITE "${CMAKE_BINARY_DIR}/test_lists/non_parallel_tests.txt" "\"${NON_PARALLEL_TESTS_STR}\"") diff --git a/scripts/check.py b/scripts/check.py index c205b18a6..fceb59e09 100755 --- a/scripts/check.py +++ b/scripts/check.py @@ -18,6 +18,8 @@ parser.add_argument('--timeout', type=int, nargs='?', default=200, help='Timeout in seconds (default: 200)') parser.add_argument('-m', '--modules', type=str, choices=['on', 'off'], default="on", help='load modulefiles automatically (default: on)') parser.add_argument('-v', '--verbose', action='store_true', help='verbose output') +parser.add_argument('-d', '--dry-run', '-N', action='store_true', help='dry run') +parser.add_argument('-c', '--categories', action='store_true', help='run tests by categories, including running a set of tests in a single thread') args = parser.parse_args() @@ -90,10 +92,33 @@ def run_cmd(cmd): else: texture_cmd = "" -cmd = f"{modules} {env_vars} ctest --output-on-failure --timeout {args.timeout} --repeat until-fail:{args.num_tries} -j {args.num_threads} -E \"`cat ./test_lists/{args.device_type}_{args.backend}_failed_{level0_cmd_list}tests.txt`{texture_cmd}\" -O checkpy_{args.device_type}_{args.backend}.txt" -res, ctest_return_code = run_cmd(cmd) -# check if "0 tests failed" is in the output, if so return 0 -if "0 tests failed" in res: - exit(0) +if args.categories: + cmd_deviceFunc = f"{modules} {env_vars} ctest --output-on-failure --timeout {args.timeout} --repeat until-fail:{args.num_tries} -j 100 -E \"`cat ./test_lists/{args.device_type}_{args.backend}_failed_{level0_cmd_list}tests.txt`{texture_cmd}\" -R deviceFunc -O checkpy_{args.device_type}_{args.backend}_device.txt" + cmd_graph = f"{modules} {env_vars} ctest --output-on-failure --timeout {args.timeout} --repeat until-fail:{args.num_tries} -j 100 -E \"`cat ./test_lists/{args.device_type}_{args.backend}_failed_{level0_cmd_list}tests.txt`{texture_cmd}\" -R \"[Gg]raph\" -O checkpy_{args.device_type}_{args.backend}_graph.txt" + cmd_single = f"{modules} {env_vars} ctest --output-on-failure --timeout {args.timeout} --repeat until-fail:{args.num_tries} -j 1 -E \"`cat ./test_lists/{args.device_type}_{args.backend}_failed_{level0_cmd_list}tests.txt`{texture_cmd}\" -R \"`cat ./test_lists/non_parallel_tests.txt`\" -O checkpy_{args.device_type}_{args.backend}_single.txt" + cmd_other = f"{modules} {env_vars} ctest --output-on-failure --timeout {args.timeout} --repeat until-fail:{args.num_tries} -j {args.num_threads} -E \"`cat ./test_lists/{args.device_type}_{args.backend}_failed_{level0_cmd_list}tests.txt`{texture_cmd}|deviceFunc|[Gg]raph|`cat ./test_lists/non_parallel_tests.txt`\" -O checkpy_{args.device_type}_{args.backend}_other.txt" + if(args.dry_run): + print(cmd_deviceFunc) + print(cmd_graph) + print(cmd_single) + print(cmd_other) + exit(0) + res_deviceFunc, err = run_cmd(cmd_deviceFunc) + res_graph, err = run_cmd(cmd_graph) + res_single, err = run_cmd(cmd_single) + res_other, err = run_cmd(cmd_other) + + if "0 tests failed" in res_deviceFunc and "0 tests failed" in res_graph and "0 tests failed" in res_single and "0 tests failed" in res_other: + exit(0) + else: + exit(1) else: - exit(ctest_return_code) + cmd = f"{modules} {env_vars} ctest --output-on-failure --timeout {args.timeout} --repeat until-fail:{args.num_tries} -j {args.num_threads} -E \"`cat ./test_lists/{args.device_type}_{args.backend}_failed_{level0_cmd_list}tests.txt`{texture_cmd}\" -O checkpy_{args.device_type}_{args.backend}.txt" + if(args.dry_run): + print(cmd) + exit(0) + res, err = run_cmd(cmd) + if "0 tests failed" in res: + exit(0) + else: + exit(1) diff --git a/scripts/unit_tests.sh b/scripts/unit_tests.sh index d36274d6c..4d0e48f48 100755 --- a/scripts/unit_tests.sh +++ b/scripts/unit_tests.sh @@ -36,7 +36,7 @@ shift # Set the number of tries based on the argument or default to 1 num_tries=1 -num_threads=1 +num_threads=24 timeout=200 for arg in "$@" do @@ -183,7 +183,7 @@ echo "end dgpu_level0_failed_imm_tests" echo "begin igpu_opencl_failed_tests" # module load opencl/igpu # module list -../scripts/check.py ./ igpu opencl --num-threads=${num_threads} --timeout=$timeout --num-tries=$num_tries | tee igpu_opencl_make_check_result.txt +../scripts/check.py ./ igpu opencl --num-threads=${num_threads} --timeout=$timeout --num-tries=$num_tries --categories | tee igpu_opencl_make_check_result.txt # ctest --timeout $timeout --repeat until-fail:${num_tries} $(ctest_j_option 4) --output-on-failure -E "`cat ./test_lists/igpu_opencl_failed_tests.txt`" | tee igpu_opencl_make_check_result.txt #pushd ${LIBCEED_DIR} #make FC= CC=clang CXX=clang++ BACKENDS="/gpu/hip/ref /gpu/hip/shared /gpu/hip/gen" prove --repeat until-fail:${num_tries} $(ctest_j_option 12) PROVE_OPS="-j" | tee igpu_opencl_make_check_result.txt @@ -196,7 +196,7 @@ echo "begin dgpu_opencl_failed_tests" # module load intel/opencl # sets ICD # module load opencl/dgpu # sets CHIP_BE # module list -../scripts/check.py ./ dgpu opencl --num-threads=${num_threads} --timeout=$timeout --num-tries=$num_tries | tee dgpu_opencl_make_check_result.txt +../scripts/check.py ./ dgpu opencl --num-threads=${num_threads} --timeout=$timeout --num-tries=$num_tries --categories | tee dgpu_opencl_make_check_result.txt # ctest --timeout $timeout --repeat until-fail:${num_tries} $(ctest_j_option 8) --output-on-failure -E "`cat ./test_lists/dgpu_opencl_failed_tests.txt`" | tee dgpu_opencl_make_check_result.txt # pushd ${LIBCEED_DIR} # HIP_DIR=${CHIPSTAR_INSTALL_DIR} make FC= CC=clang CXX=clang++ BACKENDS="/gpu/hip/ref /gpu/hip/shared /gpu/hip/gen" prove --repeat until-fail:${num_tries} $(ctest_j_option 12) PROVE_OPS="-j" | tee dgpu_opencl_make_check_result.txt From 99320fd1e0105526d036378f80f58c26e34ca540 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Henry=20Linjam=C3=A4ki?= Date: Thu, 23 Nov 2023 13:37:04 +0200 Subject: [PATCH 09/17] Resurrect __bitinsert_u32, implement __bitinsert_u64 Note that they are made available by passing -DCHIP_ENABLE_NON_COMPLIANT_DEVICELIB_CODE to HIP compilation. Resolves #701. --- bitcode/devicelib.cl | 18 ++++++--- .../hip/devicelib/integer/int_intrinsics.hh | 29 +++++++++++++++ tests/runtime/CMakeLists.txt | 5 +++ tests/runtime/TestBitInsert.hip | 37 +++++++++++++++++++ 4 files changed, 84 insertions(+), 5 deletions(-) create mode 100644 tests/runtime/TestBitInsert.hip diff --git a/bitcode/devicelib.cl b/bitcode/devicelib.cl index e9c03334b..b9589ab5a 100644 --- a/bitcode/devicelib.cl +++ b/bitcode/devicelib.cl @@ -181,11 +181,19 @@ EXPORT unsigned /* long */ long __bitextract_u64(unsigned /* long */ long src0, return width == 0 ? 0 : (src0 << (64 - offset - width)) >> (64 - width); } -EXPORT unsigned int __bitinsert_u32(unsigned int src0, unsigned int src1, - unsigned int src2, unsigned int src3) { - unsigned long offset = src2 & 31; - unsigned long width = src3 & 31; - unsigned long mask = (1 << width) - 1; +EXPORT unsigned int __chip_bitinsert_u32(uint src0, uint src1, uint raw_offset, + uint raw_width) { + uint offset = raw_offset & 31u; + uint width = raw_width & 31u; + uint mask = (1u << width) - 1u; + return ((src0 & ~(mask << offset)) | ((src1 & mask) << offset)); +} + +EXPORT ulong __chip_bitinsert_u64(ulong src0, ulong src1, ulong raw_offset, + ulong raw_width) { + ulong offset = raw_offset & 63ul; + ulong width = raw_width & 63ul; + ulong mask = (1ul << width) - 1ul; return ((src0 & ~(mask << offset)) | ((src1 & mask) << offset)); } diff --git a/include/hip/devicelib/integer/int_intrinsics.hh b/include/hip/devicelib/integer/int_intrinsics.hh index 9b689758d..85fc7ef4b 100644 --- a/include/hip/devicelib/integer/int_intrinsics.hh +++ b/include/hip/devicelib/integer/int_intrinsics.hh @@ -25,6 +25,35 @@ #include +#ifdef CHIP_ENABLE_NON_COMPLIANT_DEVICELIB_CODE +// __bitinsert_* intrinsics are not found in the HIP programming +// manual but they are provided by hipamd. +// +// They replace a 'width' bits sized block in 'src0' staring at 'offset' +// with least significant bits extracted from 'src1'. + +extern "C" __device__ unsigned int __chip_bitinsert_u32(unsigned int src0, + unsigned int src1, + unsigned int offset, + unsigned int width); +extern "C++" inline __device__ unsigned int +__bitinsert_u32(unsigned int src0, unsigned int src1, unsigned int offset, + unsigned int width) { + return __chip_bitinsert_u32(src0, src1, offset, width); +} + +extern "C" __device__ uint64_t __chip_bitinsert_u64(uint64_t src0, + uint64_t src1, + uint64_t offset, + uint64_t width); +extern "C++" inline __device__ uint64_t __bitinsert_u64(uint64_t src0, + uint64_t src1, + uint64_t offset, + uint64_t width) { + return __chip_bitinsert_u64(src0, src1, offset, width); +} +#endif // CHIP_ENABLE_NON_COMPLIANT_DEVICELIB_CODE + extern "C" __device__ unsigned int __chip_brev(unsigned int x); // Custom extern "C++" inline __device__ unsigned int __brev(unsigned int x) { return __chip_brev(x); diff --git a/tests/runtime/CMakeLists.txt b/tests/runtime/CMakeLists.txt index 52d6ed13a..b5870174d 100644 --- a/tests/runtime/CMakeLists.txt +++ b/tests/runtime/CMakeLists.txt @@ -88,3 +88,8 @@ add_shell_test(TestRuntimeWarnings.bash) add_hip_runtime_test(TestAPIs.hip) add_hip_runtime_test(TestMemFunctions.hip) add_hip_runtime_test(TestAlignAttrRuntime.hip) + +add_hip_runtime_test(TestBitInsert.hip) +# Imports __bitinsert_*. +target_compile_definitions(TestBitInsert + PRIVATE CHIP_ENABLE_NON_COMPLIANT_DEVICELIB_CODE) diff --git a/tests/runtime/TestBitInsert.hip b/tests/runtime/TestBitInsert.hip new file mode 100644 index 000000000..f1dc05feb --- /dev/null +++ b/tests/runtime/TestBitInsert.hip @@ -0,0 +1,37 @@ +#include + +__global__ void bitInsert(unsigned *Dst, unsigned Src0, unsigned Src1, + unsigned Offset, unsigned Width) { + *Dst = __bitinsert_u32(Src0, Src1, Offset, Width); +} + +__global__ void bitInsert(uint64_t *Dst, uint64_t Src0, uint64_t Src1, + uint64_t Offset, uint64_t Width) { + *Dst = __bitinsert_u64(Src0, Src1, Offset, Width); +} + +int main() { + unsigned *Dst1D, Dst1H = 0; + uint64_t *Dst2D, Dst2H = 0; + + (void)hipMalloc(&Dst1D, sizeof(unsigned)); + (void)hipMalloc(&Dst2D, sizeof(uint64_t)); + + (void)bitInsert<<<1, 1>>>(Dst1D, 0xffff4321u, 0xffff8765ul, + // Arguments + a garbage bit that should be ignored. + 16u + 32u, 16u + 32u); + hipMemcpy(&Dst1H, Dst1D, sizeof(unsigned), hipMemcpyDeviceToHost); + if (Dst1H != 0x87654321) + return 1; + + (void)bitInsert<<<1, 1>>>(Dst2D, 0xffffffff44332211ul, 0xffffffff88776655ul, + // Arguments + a garbage bit that should be ignored. + 32u + 64u, 32u + 64u); + (void)hipMemcpy(&Dst2H, Dst2D, sizeof(uint64_t), hipMemcpyDeviceToHost); + if (Dst2H != 0x8877665544332211ul) + return 2; + + (void)hipFree(Dst1D); + (void)hipFree(Dst2D); + return 0; +} From 2de6f4de6e2029d065dcd5bc1809b054ab8895e2 Mon Sep 17 00:00:00 2001 From: Michal Babej Date: Fri, 24 Nov 2023 12:34:47 +0200 Subject: [PATCH 10/17] more fixes to documentation --- docs/Development.md | 5 +++++ docs/Device_API_support_matrix.md | 6 +++--- 2 files changed, 8 insertions(+), 3 deletions(-) diff --git a/docs/Development.md b/docs/Development.md index 6efb2a4b6..f41a29fdb 100644 --- a/docs/Development.md +++ b/docs/Development.md @@ -74,7 +74,12 @@ There are several transformations (LLVM passes) done on the LLVM IR of the devic * HipAbort.cpp - special handling for abort() calls from the device side (to cause a host abort currently). * HipDefrost.cpp - removes freeze from instructions (workaround for the llvm-spirv translator). * HipDynMem.cpp - replaces dynamically sized shared-memory variables (`extern __shared__ type variable[];`) with a kernel argument. This is because in OpenCL, dynamically-sized local memory can only be passed as kernel argument. +* HipEmitLoweredNames.cpp - required processing for hiprtcGetLoweredName() * HipGlobalVariable.cpp - creates special kernels that handle access and modification of global scope variables. +* HipKernelArgSpiller.cpp - Reduces the size of large kernel parameter lists by spilling them into a device buffer +* HipLowerSwitch.cpp - Lowers switch instructions with a "non-standard" integer bitwidth (e.g. i4) to bitwidth supported by SPIRV-LLVM-Translator +* HipLowerZeroLengthArrays.cpp - Lowers occurrences of zero length array types (unsupported by SPIRV-LLVM-Translator) +* HipSanityChecks.cpp - sanity checks on the LLVM IR just before HIP-to-SPIR-V lowering * HipPasses.cpp - defines a pass plugin that runs a collection of LLVM passes (= rest of the files in this directory). * HipPrintf.cpp - pass to convert calls to the CUDA/HIP printf() to OpenCL/SPIR-V compatible printf() calls. * HipStripUsedIntrinsics.cpp - pass to remove llvm.used and llvm.compiler.used intrinsic variables. diff --git a/docs/Device_API_support_matrix.md b/docs/Device_API_support_matrix.md index d74b9cd3b..4b3fac93b 100644 --- a/docs/Device_API_support_matrix.md +++ b/docs/Device_API_support_matrix.md @@ -7,9 +7,9 @@ | Feature | HIP API # of funcs | # of impl in chipStar | chipStar notes | |-------------------------------|---------------------|-----------------------|---------------------------| | Coordinate Built-Ins | 12 | 12 | | -| Warp Size variable | supported | supported | chipStar support probably low effort, but requires guarantee from driver side to respect warpSize (cl_intel_required_subgroup_size) | -| Timer functions | 2 | 0 | missing: clock, clock64; seems already available in intel GPU hardware & driver (TODO: unclear about HW clock bit width), possibly needs software (SPIR-V) support | -| Atomic functions | ~30 | ~30 | all supported, but a few (on float/double types) are emulated, proper impl requires OpenCL/driver/HW support | +| Warp Size variable | supported | supported | implemented, but requires support from driver side to respect warpSize (= cl_intel_required_subgroup_size extension) | +| Timer functions | 2 | 2 | currently only fallback implementations of clock, clock64 are available | +| Atomic functions | ~30 | ~30 | all supported; the implementation is efficient only if cl_ext_float_atomics is present & supported by backend & HW| | Vector Types | 48 | 48 | | | Memory-Fence Instructions | 3 | 2 | \_\_threadfence_system is unsupported | | Synchronization Functions | 4 | 4 | | From 314daccfe6c247a6fc1b32d0a26dda0bfd41325f Mon Sep 17 00:00:00 2001 From: Michal Babej Date: Mon, 27 Nov 2023 19:36:49 +0200 Subject: [PATCH 11/17] Fix issue #708 tests fail with "undefined reference to pthread_create" --- CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index c8fd67be1..40f54e400 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -340,7 +340,7 @@ include(UnitTests) target_compile_definitions( CHIP PRIVATE ${HIP_ENABLE_SPIRV} ${CHIP_SPV_DEFINITIONS}) -target_link_libraries(CHIP INTERFACE ${CHIP_INTERFACE_LIBS}) +target_link_libraries(CHIP PUBLIC ${CHIP_INTERFACE_LIBS}) if(HAS_EXPERIMENTAL_FILESYSTEM) target_link_libraries(CHIP PUBLIC stdc++fs) From 7fc2b7bc323bb6b0339319bfb6652a4e29fdc963 Mon Sep 17 00:00:00 2001 From: Michal Babej Date: Tue, 28 Nov 2023 16:53:59 +0200 Subject: [PATCH 12/17] Fix issue #699 disable OCML_BASIC_ROUNDED_OPERATIONS --- CMakeLists.txt | 2 -- chipStarConfig.hh.in | 3 ++- 2 files changed, 2 insertions(+), 3 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 40f54e400..51666663c 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -217,8 +217,6 @@ option(CHIP_EXT_FLOAT_ATOMICS "Use cl_ext_float_atomics." ON) # https://github.com/CHIP-SPV/chipStar/issues/458 https://github.com/intel/compute-runtime/issues/631 option(CHIP_L0_FIRST_TOUCH "First-touch workaround for Level Zero." OFF) -option(OCML_BASIC_ROUNDED_OPERATIONS "Use OCML implementations for devicelib functions with explicit rounding mode such as __dadd_rd. Otherwise, rounding mode will be ignored" OFF) - # Warpsize would optimally be a device-specific, queried and made # effective at runtime. However, we need to fix the warpsize since SPIR-Vs need # to be portable across multiple devices. It should be more portable to diff --git a/chipStarConfig.hh.in b/chipStarConfig.hh.in index cdc9d8b4e..76d2dd277 100644 --- a/chipStarConfig.hh.in +++ b/chipStarConfig.hh.in @@ -23,7 +23,8 @@ #ifndef CHIP_CONFIG_H #define CHIP_CONFIG_H -#cmakedefine OCML_BASIC_ROUNDED_OPERATIONS "@OCML_BASIC_ROUNDED_OPERATIONS@" +// not implemented yet +#undef OCML_BASIC_ROUNDED_OPERATIONS #cmakedefine CHIP_SOURCE_DIR "@CHIP_SOURCE_DIR@" From f2fbbd98c802474a2489f937af4d850304255174 Mon Sep 17 00:00:00 2001 From: Michal Babej Date: Tue, 28 Nov 2023 16:58:13 +0200 Subject: [PATCH 13/17] fix issue #705 - update ROCM submodule fails to build when /usr/bin/clang doesn't exist, or if it points to a different Clang version --- bitcode/ROCm-Device-Libs | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/bitcode/ROCm-Device-Libs b/bitcode/ROCm-Device-Libs index 7eca6d212..bbc22e330 160000 --- a/bitcode/ROCm-Device-Libs +++ b/bitcode/ROCm-Device-Libs @@ -1 +1 @@ -Subproject commit 7eca6d2125b7e8a1738313326a2f874ce945bb61 +Subproject commit bbc22e330dc12d52dfe6924f8c5d34e81809b2cf From 1da1ae764a490263841002d2dde67d0a7e2fcbbd Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Pekka=20J=C3=A4=C3=A4skel=C3=A4inen?= Date: Fri, 1 Dec 2023 13:07:11 +0200 Subject: [PATCH 14/17] Add format-diff.sh Formats the uncommitted changes (copied from PoCL). --- scripts/clang-format-diff.py | 135 +++++++++++++++++++++++++++++++++++ scripts/format-diff.sh | 23 ++++++ 2 files changed, 158 insertions(+) create mode 100755 scripts/clang-format-diff.py create mode 100755 scripts/format-diff.sh diff --git a/scripts/clang-format-diff.py b/scripts/clang-format-diff.py new file mode 100755 index 000000000..ea483f59e --- /dev/null +++ b/scripts/clang-format-diff.py @@ -0,0 +1,135 @@ +#!/usr/bin/env python3 +# +#===- clang-format-diff.py - ClangFormat Diff Reformatter ----*- python -*--===# +# +# 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 +# +#===------------------------------------------------------------------------===# + +""" +This script reads input from a unified diff and reformats all the changed +lines. This is useful to reformat all the lines touched by a specific patch. +Example usage for git/svn users: + + git diff -U0 --no-color --relative HEAD^ | clang-format-diff.py -p1 -i + svn diff --diff-cmd=diff -x-U0 | clang-format-diff.py -i + +It should be noted that the filename contained in the diff is used unmodified +to determine the source file to update. Users calling this script directly +should be careful to ensure that the path in the diff is correct relative to the +current working directory. +""" +from __future__ import absolute_import, division, print_function + +import argparse +import difflib +import re +import subprocess +import sys + +if sys.version_info.major >= 3: + from io import StringIO +else: + from io import BytesIO as StringIO + + +def main(): + parser = argparse.ArgumentParser(description=__doc__, + formatter_class= + argparse.RawDescriptionHelpFormatter) + parser.add_argument('-i', action='store_true', default=False, + help='apply edits to files instead of displaying a diff') + parser.add_argument('-p', metavar='NUM', default=0, + help='strip the smallest prefix containing P slashes') + parser.add_argument('-regex', metavar='PATTERN', default=None, + help='custom pattern selecting file paths to reformat ' + '(case sensitive, overrides -iregex)') + parser.add_argument('-iregex', metavar='PATTERN', default= + r'.*\.(cpp|cc|c\+\+|cxx|c|cl|h|hh|hpp|hxx|m|mm|inc|js|ts' + r'|proto|protodevel|java|cs|json)', + help='custom pattern selecting file paths to reformat ' + '(case insensitive, overridden by -regex)') + parser.add_argument('-sort-includes', action='store_true', default=False, + help='let clang-format sort include blocks') + parser.add_argument('-v', '--verbose', action='store_true', + help='be more verbose, ineffective without -i') + parser.add_argument('-style', + help='formatting style to apply (LLVM, GNU, Google, Chromium, ' + 'Microsoft, Mozilla, WebKit)') + parser.add_argument('-binary', default='clang-format', + help='location of binary to use for clang-format') + args = parser.parse_args() + + # Extract changed lines for each file. + filename = None + lines_by_file = {} + for line in sys.stdin: + match = re.search(r'^\+\+\+\ (.*?/){%s}(\S*)' % args.p, line) + if match: + filename = match.group(2) + if filename is None: + continue + + if args.regex is not None: + if not re.match('^%s$' % args.regex, filename): + continue + else: + if not re.match('^%s$' % args.iregex, filename, re.IGNORECASE): + continue + + match = re.search(r'^@@.*\+(\d+)(,(\d+))?', line) + if match: + start_line = int(match.group(1)) + line_count = 1 + if match.group(3): + line_count = int(match.group(3)) + if line_count == 0: + continue + end_line = start_line + line_count - 1 + lines_by_file.setdefault(filename, []).extend( + ['-lines', str(start_line) + ':' + str(end_line)]) + + # Reformat files containing changes in place. + for filename, lines in lines_by_file.items(): + if args.i and args.verbose: + print('Formatting {}'.format(filename)) + command = [args.binary, filename] + if args.i: + command.append('-i') + if args.sort_includes: + command.append('-sort-includes') + command.extend(lines) + if args.style: + command.extend(['-style', args.style]) + + try: + p = subprocess.Popen(command, + stdout=subprocess.PIPE, + stderr=None, + stdin=subprocess.PIPE, + universal_newlines=True) + except OSError as e: + # Give the user more context when clang-format isn't + # found/isn't executable, etc. + raise RuntimeError( + 'Failed to run "%s" - %s"' % (" ".join(command), e.strerror)) + + stdout, stderr = p.communicate() + if p.returncode != 0: + sys.exit(p.returncode) + + if not args.i: + with open(filename) as f: + code = f.readlines() + formatted_code = StringIO(stdout).readlines() + diff = difflib.unified_diff(code, formatted_code, + filename, filename, + '(before formatting)', '(after formatting)') + diff_string = ''.join(diff) + if len(diff_string) > 0: + sys.stdout.write(diff_string) + +if __name__ == '__main__': + main() diff --git a/scripts/format-diff.sh b/scripts/format-diff.sh new file mode 100755 index 000000000..7e0672d27 --- /dev/null +++ b/scripts/format-diff.sh @@ -0,0 +1,23 @@ +#!/usr/bin/env bash + +GITROOT=$(git rev-parse --show-toplevel 2>/dev/null) +if [ $? -ne 0 ]; then + echo "must be run in git repo" + exit 1 +fi + +SCRIPTPATH=$( realpath "$0" ) +RELPATH=$(dirname "$SCRIPTPATH") + +# cd to root directory of the git repo +pushd ${GITROOT} > /dev/null + +PATCHY=$(mktemp /tmp/pocl.XXXXXXXX.patch) +trap "rm -f $PATCHY" EXIT + +git diff -U0 --no-color >$PATCHY + +$RELPATH/clang-format-diff.py -regex '(.*(\.hpp$|\.cl$|\.hh$|\.cc$|\.cpp$))|(lib/llvmopencl/.*)|(lib/CL/devices/tce/.*)' -i -p1 -style LLVM <$PATCHY + +# cd back wherever we were previously +popd > /dev/null From 8cb46ca1ac7828a06d4a049ec6b022978b55ab4a Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Pekka=20J=C3=A4=C3=A4skel=C3=A4inen?= Date: Fri, 1 Dec 2023 13:06:25 +0200 Subject: [PATCH 15/17] Require CG SVM / USM from the OpenCL device Without this, it lets such a device pass until the first SVM allocation, which the crashes uglily. --- bitcode/ROCm-Device-Libs | 2 +- src/backend/OpenCL/CHIPBackendOpenCL.cc | 41 +++++++++++++++++-------- 2 files changed, 30 insertions(+), 13 deletions(-) diff --git a/bitcode/ROCm-Device-Libs b/bitcode/ROCm-Device-Libs index bbc22e330..7eca6d212 160000 --- a/bitcode/ROCm-Device-Libs +++ b/bitcode/ROCm-Device-Libs @@ -1 +1 @@ -Subproject commit bbc22e330dc12d52dfe6924f8c5d34e81809b2cf +Subproject commit 7eca6d2125b7e8a1738313326a2f874ce945bb61 diff --git a/src/backend/OpenCL/CHIPBackendOpenCL.cc b/src/backend/OpenCL/CHIPBackendOpenCL.cc index 35ce4db06..324f5534c 100644 --- a/src/backend/OpenCL/CHIPBackendOpenCL.cc +++ b/src/backend/OpenCL/CHIPBackendOpenCL.cc @@ -1579,36 +1579,53 @@ void CHIPBackendOpenCL::initializeImpl(std::string CHIPPlatformStr, StrStream.str(""); - StrStream << "OpenCL Devices of type " << CHIPDeviceTypeStr - << " with SPIR-V_1 support:\n"; - std::vector SpirvDevices; + StrStream << "OpenCL Devices of type " << CHIPDeviceTypeStr << ":\n"; + + std::vector SupportedDevices; std::vector Dev; Err = SelectedPlatform.getDevices(SelectedDevType, &Dev); CHIPERR_CHECK_LOG_AND_THROW(Err, CL_SUCCESS, hipErrorInitializationError); for (auto D : Dev) { - std::string Ver = D.getInfo(&Err); - if ((Err == CL_SUCCESS) && (Ver.rfind("SPIR-V_1.", 0) == 0)) { - std::string DeviceName = D.getInfo(); - StrStream << DeviceName << "\n"; - SpirvDevices.push_back(D); + + std::string DeviceName = D.getInfo(); + + StrStream << DeviceName << " "; + std::string SPIRVVer = D.getInfo(&Err); + if ((Err != CL_SUCCESS) || + (SPIRVVer.rfind("SPIR-V_1.", 0) == std::string::npos)) { + StrStream << " no SPIR-V support.\n"; + continue; } + + // We require at least CG SVM or Device USM. + std::string DevExts = D.getInfo(); + cl_device_svm_capabilities SVMCapabilities = + D.getInfo(); + + if ((SVMCapabilities & CL_DEVICE_SVM_COARSE_GRAIN_BUFFER) == 0 && + DevExts.find("cl_intel_unified_shared_memory") == std::string::npos) { + StrStream << " no SVM/USM support.\n"; + continue; + } + + StrStream << " is supported.\n"; + SupportedDevices.push_back(D); } logTrace("{}", StrStream.str()); int SelectedDeviceIdx = atoi(CHIPDeviceStr.c_str()); - if (SelectedDeviceIdx >= SpirvDevices.size()) { + if (SelectedDeviceIdx >= SupportedDevices.size()) { logCritical("Selected OpenCL device {} is out of range", SelectedDeviceIdx); std::exit(1); } - auto Device = SpirvDevices[SelectedDeviceIdx]; + auto Device = SupportedDevices[SelectedDeviceIdx]; logDebug("CHIP_DEVICE={} Selected OpenCL device {}", SelectedDeviceIdx, Device.getInfo()); // Create context which has devices // Create queues that have devices each of which has an associated context - // TODO Change this to spirv_enabled_devices - cl::Context Ctx(SpirvDevices); + cl::Context Ctx(SupportedDevices); CHIPContextOpenCL *ChipContext = new CHIPContextOpenCL(Ctx, Device, SelectedPlatform); ::Backend->addContext(ChipContext); From f52f306aaeea392eace1a64e9c8d9e862faae566 Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Wed, 6 Dec 2023 07:19:18 -0600 Subject: [PATCH 16/17] modify configure_llvm.sh enable configuring llvm without extensions --- scripts/configure_llvm.sh | 69 ++++++++++++++++++++++++--------------- 1 file changed, 43 insertions(+), 26 deletions(-) diff --git a/scripts/configure_llvm.sh b/scripts/configure_llvm.sh index c87d59001..8b37d5e9a 100755 --- a/scripts/configure_llvm.sh +++ b/scripts/configure_llvm.sh @@ -1,10 +1,14 @@ #!/bin/bash +# if an error is enountered, exit +set -e + # check arguments -if [ $# -ne 3 ]; then - echo "Usage: $0 " +if [ $# -ne 4 ]; then + echo "Usage: $0 " echo "version: LLVM version 15, 16, 17" - echo "build_type: static or dynamic" + echo "link_type: static or dynamic" + echo "only_necessary_spirv_exts: on or off" exit 1 fi @@ -14,48 +18,61 @@ if [ "$1" != "15" ] && [ "$1" != "16" ] && [ "$1" != "17" ]; then exit 1 fi -# check build_type argument +# check link_type argument if [ "$3" != "static" ] && [ "$3" != "dynamic" ]; then - echo "Invalid build_type. Must be 'static' or 'dynamic'." + echo "Invalid link_type. Must be 'static' or 'dynamic'." exit 1 fi -# if an error is enountered, exit -set -e +# check only-necessary-spirv-exts argument +if [ "$4" != "on" ] && [ "$4" != "off" ]; then + echo "Invalid only_necessary_spirv_exts. Must be 'on' or 'off'." + exit 1 +fi VERSION=$1 INSTALL_DIR=$2 -BUILD_TYPE=$3 +LINK_TYPE=$3 + +# set the brach name for checkuot based on only-necessary-spirv-exts +if [ "$4" == "on" ]; then + LLVM_BRANCH="spirv-ext-fixes-${VERSION}" + TRANSLATOR_BRANCH="chipStar-llvm-${VERSION}" +else + LLVM_BRANCH="chipStar-llvm-${VERSION}" + TRANSLATOR_BRANCH="chipStar-llvm-${VERSION}" +fi + export LLVM_DIR=`pwd`/llvm-project/llvm # check if llvm-project exists, if not clone it if [ ! -d llvm-project ]; then - git clone https://github.com/CHIP-SPV/llvm-project.git -b chipStar-llvm-${VERSION} --depth 1 + git clone https://github.com/CHIP-SPV/llvm-project.git -b ${LLVM_BRANCH} --depth 1 cd ${LLVM_DIR}/projects - git clone https://github.com/CHIP-SPV/SPIRV-LLVM-Translator.git -b chipStar-llvm-${VERSION} --depth 1 + git clone https://github.com/CHIP-SPV/SPIRV-LLVM-Translator.git -b ${TRANSLATOR_BRANCH} --depth 1 cd ${LLVM_DIR} else # Warn the user, error out echo "llvm-project directory already exists. Assuming it's cloned from chipStar." cd ${LLVM_DIR} # check if already on the desired branch - if [ `git branch --show-current` == "chipStar-llvm-${VERSION}" ]; then - echo "Already on branch chipStar-llvm-${VERSION}" + if [ `git branch --show-current` == "${LLVM_BRANCH}" ]; then + echo "Already on branch ${LLVM_BRANCH}" else - echo "Switching to branch chipStar-llvm-${VERSION}" - git br -D chipStar-llvm-${VERSION} &> /dev/null - git fetch origin chipStar-llvm-${VERSION}:chipStar-llvm-${VERSION} - git checkout chipStar-llvm-${VERSION} + echo "Switching to branch ${LLVM_BRANCH}" + git br -D ${LLVM_BRANCH} &> /dev/null + git fetch origin ${LLVM_BRANCH}:${LLVM_BRANCH} + git checkout ${LLVM_BRANCH} fi cd ${LLVM_DIR}/projects/SPIRV-LLVM-Translator # check if already on the desired branch - if [ `git branch --show-current` == "chipStar-llvm-${VERSION}" ]; then - echo "Already on branch chipStar-llvm-${VERSION}" + if [ `git branch --show-current` == "${TRANSLATOR_BRANCH}" ]; then + echo "Already on branch ${TRANSLATOR_BRANCH}" else - echo "Switching to branch chipStar-llvm-${VERSION}" - git br -D chipStar-llvm-${VERSION} &> /dev/null - git fetch origin chipStar-llvm-${VERSION}:chipStar-llvm-${VERSION} - git checkout chipStar-llvm-${VERSION} + echo "Switching to branch ${TRANSLATOR_BRANCH}" + git br -D ${TRANSLATOR_BRANCH} &> /dev/null + git fetch origin ${TRANSLATOR_BRANCH}:${TRANSLATOR_BRANCH} + git checkout ${TRANSLATOR_BRANCH} fi cd ${LLVM_DIR} fi @@ -71,13 +88,13 @@ else fi # Add build type condition -if [ "$BUILD_TYPE" == "static" ]; then +if [ "$LINK_TYPE" == "static" ]; then cmake ../ \ -DCMAKE_INSTALL_PREFIX=${INSTALL_DIR} \ -DCMAKE_BUILD_TYPE=Release \ -DLLVM_ENABLE_PROJECTS="clang;openmp" \ -DLLVM_TARGETS_TO_BUILD=host -elif [ "$BUILD_TYPE" == "dynamic" ]; then +elif [ "$LINK_TYPE" == "dynamic" ]; then cmake ../ \ -DCMAKE_INSTALL_PREFIX=${INSTALL_DIR} \ -DCMAKE_INSTALL_RPATH=${INSTALL_DIR}/lib \ @@ -86,9 +103,9 @@ elif [ "$BUILD_TYPE" == "dynamic" ]; then -DLLVM_LINK_LLVM_DYLIB=ON \ -DLLVM_BUILD_LLVM_DYLIB=ON \ -DLLVM_PARALLEL_LINK_JOBS=2 \ - -DCMAKE_BUILD_TYPE=RelWithDebInfo \ + -DCMAKE_BUILD_TYPE=Release \ -DLLVM_ENABLE_ASSERTIONS=On else - echo "Invalid build_type. Must be 'static' or 'dynamic'." + echo "Invalid link_type. Must be 'static' or 'dynamic'." exit 1 fi From 7bb9b3f61967937dd6a41c546a181551ef793f44 Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Wed, 6 Dec 2023 12:53:21 -0600 Subject: [PATCH 17/17] additional non-parallel tests --- cmake/UnitTests.cmake | 14 ++++++++++++++ 1 file changed, 14 insertions(+) diff --git a/cmake/UnitTests.cmake b/cmake/UnitTests.cmake index b18ac6a68..0106f719a 100644 --- a/cmake/UnitTests.cmake +++ b/cmake/UnitTests.cmake @@ -13,6 +13,20 @@ list(APPEND CPU_POCL_FAILED_TESTS " ") list(APPEND GPU_POCL_FAILED_TESTS " ") # TODO list(APPEND NON_PARALLEL_TESTS " ") +list(APPEND NON_PARALLEL_TESTS "TestLargeGlobalVar") +list(APPEND NON_PARALLEL_TESTS "cuda-asyncAPI") +list(APPEND NON_PARALLEL_TESTS "Unit_hipMemcpy_Negative") +list(APPEND NON_PARALLEL_TESTS "firstTouch") +list(APPEND NON_PARALLEL_TESTS "Unit_hipMemcpy_HalfMemCopy") +list(APPEND NON_PARALLEL_TESTS "Unit_hipStreamBeginCapture_ColligatedStrmCapture_defaultflag") +list(APPEND NON_PARALLEL_TESTS "Unit_hipMemcpyWithStream_TestkindDtoH") +list(APPEND NON_PARALLEL_TESTS "Unit_hipMemcpyWithStream_TestkindDefault") +list(APPEND NON_PARALLEL_TESTS "Unit_hipMemsetFunctional_ZeroValue_2D") +list(APPEND NON_PARALLEL_TESTS "Unit_hipHostMalloc_NonCoherent") +list(APPEND NON_PARALLEL_TESTS "Unit_hipStreamAddCallback_WithCreatedStream") +list(APPEND NON_PARALLEL_TESTS "cuda-sortnet") +list(APPEND NON_PARALLEL_TESTS "Unit_hipMemset3DAsync_SeekSetArrayPortion") +list(APPEND NON_PARALLEL_TESTS "Unit_hipMemcpyToFromSymbol_SyncAndAsync") list(APPEND NON_PARALLEL_TESTS "MatrixMultiply") list(APPEND NON_PARALLEL_TESTS "Unit_hipMemcpy2DFromArray_PinnedMemSameGPU") list(APPEND NON_PARALLEL_TESTS "Unit_hipMemset3D_SeekSetArrayPortion")