Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Integrate chipstar #459

Open
wants to merge 19 commits into
base: develop
Choose a base branch
from
Open

Integrate chipstar #459

wants to merge 19 commits into from

Conversation

pvelesko
Copy link
Contributor

@pvelesko pvelesko commented Mar 11, 2024

This PR makes the necessary changes to enable these tests to work with chipStar HIP runtime.

  • Import old CatchAddTests.cmake and add CATCH2_DISCOVER_TESTS_COMPILE_TIME CMake option to fix the compile-time test discovery.
  • Remove unused code related to compile time test discovery
  • Fix some compilation issues related to template resolution
  • Disable/skip tests that don't work with chipStar

Comment on lines +327 to +338
// template <typename... Typenames, typename K, typename Dim, typename... Args>
// void launchKernel(K kernel, Dim numBlocks, Dim numThreads, std::uint32_t memPerBlock,
// hipStream_t stream, Args&&... packedArgs) {
// #ifndef RTC_TESTING
// validateArguments(kernel, packedArgs...);
// kernel<<<numBlocks, numThreads, memPerBlock, stream>>>(std::forward<Args>(packedArgs)...);
// #else
// launchRTCKernel<Typenames...>(kernel, numBlocks, numThreads, memPerBlock, stream,
// std::forward<Args>(packedArgs)...);
// #endif
// HIP_CHECK(hipGetLastError());
// }
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This wasn't working - fix implemented below. With the following changes I was able to compile using LLVM17.

This solution seems not to respect const. Does the original implementation (can't test since it doesn't compile) ?

Comment on lines +1 to +6
#ifndef KERNELS_PATH_H
#define KERNELS_PATH_H

#define KERNELS_PATH "@CMAKE_CURRENT_SOURCE_DIR@/kernels/"

#endif
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This part also didn't compile. Tried escaping in the CMakeLists.txt but didn't work.

@pvelesko
Copy link
Contributor Author

pvelesko commented Mar 11, 2024

Compilation error in original code

hip-tests-update/hip-tests/catch/unit/occupancy/hipOccupancyMaxPotentialBlockSize.cc
[build] /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/unit/occupancy/hipOccupancyMaxPotentialBlockSize.cc:84:18: error: no matching function for call to 'hipOccupancyMaxPotentialBlockSize'
[build]    84 |           return hipOccupancyMaxPotentialBlockSize<void (*)(int*)>(gridSize, blockSize, f2<int>, 0, 0);
[build]       |                  ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
[build] /space/pvelesko/chipStar/hip-tests-update/HIP/include/hip/hip_runtime_api.h:8489:35: note: candidate function template not viable: no overload of 'f2' matching 'void (*)(int *)' for 3rd argument
[build]  8490 | static hipError_t __host__ inline hipOccupancyMaxPotentialBlockSize(int* gridSize, int* blockSize,
[build]       |                                   ^
[build]  8491 |     T f, size_t dynSharedMemPerBlk = 0, int blockSizeLimit = 0) {
[build]       |     ~~~
[build] /space/pvelesko/chipStar/hip-tests-update/HIP/include/hip/hip_runtime_api.h:8775:19: note: candidate function template not viable: no overload of 'f2' matching 'void (*)(int *)' for 3rd argument
[build]  8776 | inline hipError_t hipOccupancyMaxPotentialBlockSize(int* gridSize, int* blockSize,
[build]       |                   ^
[build]  8777 |                                                     F kernel, size_t dynSharedMemPerBlk, uint32_t blockSizeLimit) {
[build]       |                                                     ~~~~~~~~
[build] /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/unit/occupancy/hipOccupancyMaxPotentialBlockSize.cc:92:18: error: no matching function for call to 'hipOccupancyMaxPotentialBlockSize'
[build]    92 |           return hipOccupancyMaxPotentialBlockSize<void (*)(int*)>(
[build]       |                  ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
[build] /space/pvelesko/chipStar/hip-tests-update/HIP/include/hip/hip_runtime_api.h:8489:35: note: candidate function template not viable: no overload of 'f2' matching 'void (*)(int *)' for 3rd argument
[build]  8490 | static hipError_t __host__ inline hipOccupancyMaxPotentialBlockSize(int* gridSize, int* blockSize,
[build]       |                                   ^
[build]  8491 |     T f, size_t dynSharedMemPerBlk = 0, int blockSizeLimit = 0) {
[build]       |     ~~~
[build] /space/pvelesko/chipStar/hip-tests-update/HIP/include/hip/hip_runtime_api.h:8775:19: note: candidate function template not viable: no overload of 'f2' matching 'void (*)(int *)' for 3rd argument
[build]  8776 | inline hipError_t hipOccupancyMaxPotentialBlockSize(int* gridSize, int* blockSize,
[build]       |                   ^
[build]  8777 |                                                     F kernel, size_t dynSharedMemPerBlk, uint32_t blockSizeLimit) {
[build]       |                                                     ~~~~~~~~
[build] 2 errors generated when compiling for .
[build] 

@pvelesko
Copy link
Contributor Author

[build] FAILED: catch/catch_tests/unit/memory/CMakeFiles/MemoryTest1.dir/hipGetSymbolSizeAddress.cc.o 
[build] /space/pvelesko/chipStar/hip-tests-update/build/bin/hipcc  -I/space/pvelesko/chipStar/hip-tests-update/CHIP -I/space/pvelesko/chipStar/hip-tests-update/PUBLIC -I/usr/include/level_zero -I/space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/external/Catch2 -I/space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/./include -I/space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/./kernels -I/space/pvelesko/chipStar/hip-tests-update/build/include -I/space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/external/picojson -I/space/pvelesko/chipStar/hip-tests-update/build/catch -Wno-duplicate-decl-specifier -Wno-tautological-constant-compare  -Wno-c++20-extensions -Wno-unused-result -Wno-delete-abstract-non-virtual-dtor -Wno-deprecated-declarations -Wunused-command-line-argument --std=c++17 -g -fPIE -Wno-format-extra-args -mf16c -Wall -O1 -std=c++17 -MD -MT catch/catch_tests/unit/memory/CMakeFiles/MemoryTest1.dir/hipGetSymbolSizeAddress.cc.o -MF catch/catch_tests/unit/memory/CMakeFiles/MemoryTest1.dir/hipGetSymbolSizeAddress.cc.o.d -o catch/catch_tests/unit/memory/CMakeFiles/MemoryTest1.dir/hipGetSymbolSizeAddress.cc.o -c /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/unit/memory/hipGetSymbolSizeAddress.cc
[build] /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/unit/memory/hipGetSymbolSizeAddress.cc:91:20: error: no matching function for call to 'HipGetSymbolSizeAddressTest'
[build]    91 |   SECTION("int") { HIP_GET_SYMBOL_SIZE_ADDRESS_TEST(int); }
[build]       |                    ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
[build] /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/unit/memory/hipGetSymbolSizeAddress.cc:86:3: note: expanded from macro 'HIP_GET_SYMBOL_SIZE_ADDRESS_TEST'
[build]    86 |   HipGetSymbolSizeAddressTest<type, 1, type##_var_address_validation_kernel>(SYMBOL(type##_var));  \
[build]       |   ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
[build] /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/unit/memory/hipGetSymbolSizeAddress.cc:49:13: note: candidate template ignored: invalid explicitly-specified argument for template parameter 'validation_kernel'
[build]    49 | static void HipGetSymbolSizeAddressTest(const void* symbol) {
[build]       |             ^
[build] /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/unit/memory/hipGetSymbolSizeAddress.cc:91:20: error: no matching function for call to 'HipGetSymbolSizeAddressTest'
[build]    91 |   SECTION("int") { HIP_GET_SYMBOL_SIZE_ADDRESS_TEST(int); }
[build]       |                    ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
[build] /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/unit/memory/hipGetSymbolSizeAddress.cc:87:3: note: expanded from macro 'HIP_GET_SYMBOL_SIZE_ADDRESS_TEST'
[build]    87 |   HipGetSymbolSizeAddressTest<type, kArraySize, type##_arr_address_validation_kernel>(             \
[build]       |   ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
[build] /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/unit/memory/hipGetSymbolSizeAddress.cc:49:13: note: candidate template ignored: invalid explicitly-specified argument for template parameter 'validation_kernel'
[build]    49 | static void HipGetSymbolSizeAddressTest(const void* symbol) {
[build]       |             ^
[build] /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/unit/memory/hipGetSymbolSizeAddress.cc:92:22: error: no matching function for call to 'HipGetSymbolSizeAddressTest'
[build]    92 |   SECTION("float") { HIP_GET_SYMBOL_SIZE_ADDRESS_TEST(float); }
[build]       |                      ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
[build] /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/unit/memory/hipGetSymbolSizeAddress.cc:86:3: note: expanded from macro 'HIP_GET_SYMBOL_SIZE_ADDRESS_TEST'
[build]    86 |   HipGetSymbolSizeAddressTest<type, 1, type##_var_address_validation_kernel>(SYMBOL(type##_var));  \
[build]       |   ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
[build] /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/unit/memory/hipGetSymbolSizeAddress.cc:49:13: note: candidate template ignored: invalid explicitly-specified argument for template parameter 'validation_kernel'
[build]    49 | static void HipGetSymbolSizeAddressTest(const void* symbol) {
[build]       |             ^
[build] /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/unit/memory/hipGetSymbolSizeAddress.cc:92:22: error: no matching function for call to 'HipGetSymbolSizeAddressTest'
[build]    92 |   SECTION("float") { HIP_GET_SYMBOL_SIZE_ADDRESS_TEST(float); }
[build]       |                      ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
[build] /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/unit/memory/hipGetSymbolSizeAddress.cc:87:3: note: expanded from macro 'HIP_GET_SYMBOL_SIZE_ADDRESS_TEST'
[build]    87 |   HipGetSymbolSizeAddressTest<type, kArraySize, type##_arr_address_validation_kernel>(             \
[build]       |   ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
[build] /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/unit/memory/hipGetSymbolSizeAddress.cc:49:13: note: candidate template ignored: invalid explicitly-specified argument for template parameter 'validation_kernel'
[build]    49 | static void HipGetSymbolSizeAddressTest(const void* symbol) {
[build]       |             ^
[build] /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/unit/memory/hipGetSymbolSizeAddress.cc:93:21: error: no matching function for call to 'HipGetSymbolSizeAddressTest'
[build]    93 |   SECTION("char") { HIP_GET_SYMBOL_SIZE_ADDRESS_TEST(char); }
[build]       |                     ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
[build] /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/unit/memory/hipGetSymbolSizeAddress.cc:86:3: note: expanded from macro 'HIP_GET_SYMBOL_SIZE_ADDRESS_TEST'
[build]    86 |   HipGetSymbolSizeAddressTest<type, 1, type##_var_address_validation_kernel>(SYMBOL(type##_var));  \
[build]       |   ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
[build] /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/unit/memory/hipGetSymbolSizeAddress.cc:49:13: note: candidate template ignored: invalid explicitly-specified argument for template parameter 'validation_kernel'
[build]    49 | static void HipGetSymbolSizeAddressTest(const void* symbol) {
[build]       |             ^
[build] /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/unit/memory/hipGetSymbolSizeAddress.cc:93:21: error: no matching function for call to 'HipGetSymbolSizeAddressTest'
[build]    93 |   SECTION("char") { HIP_GET_SYMBOL_SIZE_ADDRESS_TEST(char); }
[build]       |                     ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
[build] /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/unit/memory/hipGetSymbolSizeAddress.cc:87:3: note: expanded from macro 'HIP_GET_SYMBOL_SIZE_ADDRESS_TEST'
[build]    87 |   HipGetSymbolSizeAddressTest<type, kArraySize, type##_arr_address_validation_kernel>(             \
[build]       |   ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
[build] /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/unit/memory/hipGetSymbolSizeAddress.cc:49:13: note: candidate template ignored: invalid explicitly-specified argument for template parameter 'validation_kernel'
[build]    49 | static void HipGetSymbolSizeAddressTest(const void* symbol) {
[build]       |             ^
[build] /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/unit/memory/hipGetSymbolSizeAddress.cc:94:23: error: no matching function for call to 'HipGetSymbolSizeAddressTest'
[build]    94 |   SECTION("double") { HIP_GET_SYMBOL_SIZE_ADDRESS_TEST(double); }
[build]       |                       ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
[build] /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/unit/memory/hipGetSymbolSizeAddress.cc:86:3: note: expanded from macro 'HIP_GET_SYMBOL_SIZE_ADDRESS_TEST'
[build]    86 |   HipGetSymbolSizeAddressTest<type, 1, type##_var_address_validation_kernel>(SYMBOL(type##_var));  \
[build]       |   ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
[build] /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/unit/memory/hipGetSymbolSizeAddress.cc:49:13: note: candidate template ignored: invalid explicitly-specified argument for template parameter 'validation_kernel'
[build]    49 | static void HipGetSymbolSizeAddressTest(const void* symbol) {
[build]       |             ^
[build] /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/unit/memory/hipGetSymbolSizeAddress.cc:94:23: error: no matching function for call to 'HipGetSymbolSizeAddressTest'
[build]    94 |   SECTION("double") { HIP_GET_SYMBOL_SIZE_ADDRESS_TEST(double); }
[build]       |                       ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
[build] /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/unit/memory/hipGetSymbolSizeAddress.cc:87:3: note: expanded from macro 'HIP_GET_SYMBOL_SIZE_ADDRESS_TEST'
[build]    87 |   HipGetSymbolSizeAddressTest<type, kArraySize, type##_arr_address_validation_kernel>(             \
[build]       |   ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
[build] /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/unit/memory/hipGetSymbolSizeAddress.cc:49:13: note: candidate template ignored: invalid explicitly-specified argument for template parameter 'validation_kernel'
[build]    49 | static void HipGetSymbolSizeAddressTest(const void* symbol) {
[build]       |             ^
[build] 8 errors generated when compiling for .

@pvelesko
Copy link
Contributor Author

Compilation error with -DRTC_TESTING=ON

In file included from /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/unit/threadfence/__threadfence_block.cc:24:
In file included from /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/./include/hip_test_common.hh:39:
/space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/./include/hip_test_rtc.hh:294:5: error: no matching function for call to 'launchRTCKernel'
  294 |     launchRTCKernel(kernelName, dim3(numBlocks), dim3(numThreads), memPerBlock, stream,
      |     ^~~~~~~~~~~~~~~
/space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/./include/hip_test_common.hh:357:3: note: in instantiation of function template specialization 'HipTest::launchRTCKernel<void (*)(int *, int *), int *, int *>' requested here
  357 |   launchRTCKernel<Typenames...>(kernel, numBlocks, numThreads, memPerBlock, stream,
      |   ^
/space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/unit/threadfence/__threadfence_block.cc:57:14: note: in instantiation of function template specialization 'HipTest::launchKernel<void (*)(int *, int *), int, int *, int *>' requested here
   57 |     HipTest::launchKernel(ThreadfenceTestKernel<ThreadfenceScope::kBlock, true>, 1, 2,
      |              ^
/space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/./include/hip_test_rtc.hh:289:6: note: candidate function template not viable: no known conversion from 'dim3' to 'int' for 2nd argument
  289 | void launchRTCKernel(KernelFunc kernelFunc, int numBlocks, int numThreads,
      |      ^                                      ~~~~~~~~~~~~~
/space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/./include/hip_test_rtc.hh:217:6: note: candidate function template not viable: no known conversion from 'std::string' (aka 'basic_string<char>') to 'std::string (*)()' (aka 'basic_string<char> (*)()') for 1st argument
  217 | void launchRTCKernel(std::string (*getKernelName)(), dim3 numBlocks, dim3 numThreads,
      |      ^               ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
1 warning generated when compiling for .
1 warning and 1 error generated when compiling for .

@pvelesko
Copy link
Contributor Author

@gargrahul

* device-side mallloc and free, wall_clock placeholders were implemented
@rakesroy
Copy link
Contributor

!verify

@pvelesko
Copy link
Contributor Author

pvelesko commented May 6, 2024

@rakesroy Any information regarding intention to merge?

@rakesroy
Copy link
Contributor

rakesroy commented May 6, 2024

Hi @pvelesko,
We have cloned this PR to internal repo for review & to check PSDB.
I'll update you on its progress.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

3 participants