diff --git a/projects/hip-tests/catch/include/utils.hh b/projects/hip-tests/catch/include/utils.hh index f025768c14..3855308a42 100644 --- a/projects/hip-tests/catch/include/utils.hh +++ b/projects/hip-tests/catch/include/utils.hh @@ -170,7 +170,7 @@ inline bool DeviceAttributesSupport(const int device, Attributes... attributes) return (... && DeviceAttributeSupport(device, attributes)); } -inline int GetDeviceAttribute(int device, const hipDeviceAttribute_t attr) { +inline int GetDeviceAttribute(const hipDeviceAttribute_t attr, int device) { int value = 0; HIP_CHECK(hipDeviceGetAttribute(&value, attr, device)); return value; diff --git a/projects/hip-tests/catch/unit/CMakeLists.txt b/projects/hip-tests/catch/unit/CMakeLists.txt index 6b63292c91..6b38e58ff6 100644 --- a/projects/hip-tests/catch/unit/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/CMakeLists.txt @@ -41,7 +41,6 @@ add_subdirectory(device_memory) add_subdirectory(warp) add_subdirectory(dynamicLoading) add_subdirectory(g++) -add_subdirectory(module) add_subdirectory(channelDescriptor) add_subdirectory(executionControl) add_subdirectory(math) diff --git a/projects/hip-tests/catch/unit/executionControl/CMakeLists.txt b/projects/hip-tests/catch/unit/executionControl/CMakeLists.txt index a27f9dc4f1..877addd79b 100644 --- a/projects/hip-tests/catch/unit/executionControl/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/executionControl/CMakeLists.txt @@ -4,6 +4,7 @@ set(TEST_SRC hipFuncSetSharedMemConfig.cc hipFuncSetAttribute.cc hipFuncGetAttributes.cc + hipLaunchKernel.cc hipLaunchCooperativeKernel.cc hipLaunchCooperativeKernelMultiDevice.cc ) @@ -12,6 +13,7 @@ if(HIP_PLATFORM MATCHES "amd") set(TEST_SRC ${TEST_SRC} hipExtLaunchKernel.cc hipExtLaunchMultiKernelMultiDevice.cc + launch_api.cc ) endif() diff --git a/projects/hip-tests/catch/unit/executionControl/hipExtLaunchKernel.cc b/projects/hip-tests/catch/unit/executionControl/hipExtLaunchKernel.cc index 8b85507de5..1b336b4d74 100644 --- a/projects/hip-tests/catch/unit/executionControl/hipExtLaunchKernel.cc +++ b/projects/hip-tests/catch/unit/executionControl/hipExtLaunchKernel.cc @@ -49,19 +49,19 @@ TEST_CASE("Unit_hipExtLaunchKernel_Positive_Basic") { TEST_CASE("Unit_hipExtLaunchKernel_Positive_Parameters") { SECTION("blockDim.x == maxBlockDimX") { - const unsigned int x = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimX); + const unsigned int x = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimX, 0); HIP_CHECK(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{x, 1, 1}, nullptr, 0, nullptr, nullptr, nullptr, 0u)); } SECTION("blockDim.y == maxBlockDimY") { - const unsigned int y = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimY); + const unsigned int y = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimY, 0); HIP_CHECK(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{y, 1, 1}, nullptr, 0, nullptr, nullptr, nullptr, 0u)); } SECTION("blockDim.z == maxBlockDimZ") { - const unsigned int z = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimZ); + const unsigned int z = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimZ, 0); HIP_CHECK(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{z, 1, 1}, nullptr, 0, nullptr, nullptr, nullptr, 0u)); } @@ -111,28 +111,28 @@ TEST_CASE("Unit_hipExtLaunchKernel_Negative_Parameters") { } SECTION("blockDim.x > maxBlockDimX") { - const unsigned int x = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimX) + 1u; + const unsigned int x = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimX, 0) + 1u; HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{x, 1, 1}, nullptr, 0, nullptr, nullptr, nullptr, 0u), hipErrorInvalidConfiguration); } SECTION("blockDim.y > maxBlockDimY") { - const unsigned int y = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimY) + 1u; + const unsigned int y = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimY, 0) + 1u; HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{1, y, 1}, nullptr, 0, nullptr, nullptr, nullptr, 0u), hipErrorInvalidConfiguration); } SECTION("blockDim.z > maxBlockDimZ") { - const unsigned int z = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimZ) + 1u; + const unsigned int z = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimZ, 0) + 1u; HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{1, 1, z}, nullptr, 0, nullptr, nullptr, nullptr, 0u), hipErrorInvalidConfiguration); } SECTION("blockDim.x * blockDim.y * blockDim.z > maxThreadsPerBlock") { - const unsigned int max = GetDeviceAttribute(0, hipDeviceAttributeMaxThreadsPerBlock); + const unsigned int max = GetDeviceAttribute(hipDeviceAttributeMaxThreadsPerBlock, 0); const unsigned int dim = std::ceil(std::cbrt(max)); HIP_CHECK_ERROR( hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{dim, dim, dim}, @@ -141,7 +141,7 @@ TEST_CASE("Unit_hipExtLaunchKernel_Negative_Parameters") { } SECTION("sharedMemBytes > maxSharedMemoryPerBlock") { - const unsigned int max = GetDeviceAttribute(0, hipDeviceAttributeMaxSharedMemoryPerBlock) + 1u; + const unsigned int max = GetDeviceAttribute(hipDeviceAttributeMaxSharedMemoryPerBlock, 0) + 1u; HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{1, 1, 1}, nullptr, max, nullptr, nullptr, nullptr, 0u), hipErrorOutOfMemory); diff --git a/projects/hip-tests/catch/unit/executionControl/hipFuncGetAttributes.cc b/projects/hip-tests/catch/unit/executionControl/hipFuncGetAttributes.cc index e97f44300e..c3ce1c835e 100644 --- a/projects/hip-tests/catch/unit/executionControl/hipFuncGetAttributes.cc +++ b/projects/hip-tests/catch/unit/executionControl/hipFuncGetAttributes.cc @@ -35,8 +35,8 @@ TEST_CASE("Unit_hipFuncGetAttributes_Positive_Basic") { SECTION("binaryVersion") { #if HT_NVIDIA - const auto major = GetDeviceAttribute(0, hipDeviceAttributeComputeCapabilityMajor); - const auto minor = GetDeviceAttribute(0, hipDeviceAttributeComputeCapabilityMinor); + const auto major = GetDeviceAttribute(hipDeviceAttributeComputeCapabilityMajor, 0); + const auto minor = GetDeviceAttribute(hipDeviceAttributeComputeCapabilityMinor, 0); REQUIRE(attr.binaryVersion == major * 10 + minor); #elif HT_AMD REQUIRE(attr.binaryVersion > 0); @@ -48,7 +48,7 @@ TEST_CASE("Unit_hipFuncGetAttributes_Positive_Basic") { SECTION("constSizeBytes") { REQUIRE(attr.constSizeBytes == kConstSizeBytes); } SECTION("maxThreadsPerBlock") { - REQUIRE(attr.maxThreadsPerBlock == GetDeviceAttribute(0, hipDeviceAttributeMaxThreadsPerBlock)); + REQUIRE(attr.maxThreadsPerBlock == GetDeviceAttribute(hipDeviceAttributeMaxThreadsPerBlock, 0)); } SECTION("numRegs") { REQUIRE(attr.numRegs >= 0); } @@ -57,7 +57,7 @@ TEST_CASE("Unit_hipFuncGetAttributes_Positive_Basic") { SECTION("sharedSizeBytes") { REQUIRE(attr.sharedSizeBytes <= - GetDeviceAttribute(0, hipDeviceAttributeMaxSharedMemoryPerBlock)); + GetDeviceAttribute(hipDeviceAttributeMaxSharedMemoryPerBlock, 0)); } } diff --git a/projects/hip-tests/catch/unit/executionControl/hipLaunchCooperativeKernel.cc b/projects/hip-tests/catch/unit/executionControl/hipLaunchCooperativeKernel.cc index eb7eb2293f..5beeed4621 100644 --- a/projects/hip-tests/catch/unit/executionControl/hipLaunchCooperativeKernel.cc +++ b/projects/hip-tests/catch/unit/executionControl/hipLaunchCooperativeKernel.cc @@ -61,19 +61,19 @@ TEST_CASE("Unit_hipLaunchCooperativeKernel_Positive_Parameters") { } SECTION("blockDim.x == maxBlockDimX") { - const unsigned int x = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimX); + const unsigned int x = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimX, 0); HIP_CHECK(hipLaunchCooperativeKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{x, 1, 1}, nullptr, 0, nullptr)); } SECTION("blockDim.y == maxBlockDimY") { - const unsigned int y = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimY); + const unsigned int y = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimY, 0); HIP_CHECK(hipLaunchCooperativeKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{y, 1, 1}, nullptr, 0, nullptr)); } SECTION("blockDim.z == maxBlockDimZ") { - const unsigned int z = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimZ); + const unsigned int z = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimZ, 0); HIP_CHECK(hipLaunchCooperativeKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{z, 1, 1}, nullptr, 0, nullptr)); } @@ -128,28 +128,28 @@ TEST_CASE("Unit_hipLaunchCooperativeKernel_Negative_Parameters") { } SECTION("blockDim.x > maxBlockDimX") { - const unsigned int x = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimX) + 1u; + const unsigned int x = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimX, 0) + 1u; HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{x, 1, 1}, nullptr, 0, nullptr), hipErrorInvalidConfiguration); } SECTION("blockDim.y > maxBlockDimY") { - const unsigned int y = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimY) + 1u; + const unsigned int y = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimY, 0) + 1u; HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{1, y, 1}, nullptr, 0, nullptr), hipErrorInvalidConfiguration); } SECTION("blockDim.z > maxBlockDimZ") { - const unsigned int z = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimZ) + 1u; + const unsigned int z = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimZ, 0) + 1u; HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{1, 1, z}, nullptr, 0, nullptr), hipErrorInvalidConfiguration); } SECTION("blockDim.x * blockDim.y * blockDim.z > maxThreadsPerBlock") { - const unsigned int max = GetDeviceAttribute(0, hipDeviceAttributeMaxThreadsPerBlock); + const unsigned int max = GetDeviceAttribute(hipDeviceAttributeMaxThreadsPerBlock, 0); const unsigned int dim = std::ceil(std::cbrt(max)); HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{dim, dim, dim}, nullptr, 0, nullptr), @@ -163,7 +163,7 @@ TEST_CASE("Unit_hipLaunchCooperativeKernel_Negative_Parameters") { HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&max_blocks, reinterpret_cast(kernel), 1, 0)); const unsigned int multiproc_count = - GetDeviceAttribute(0, hipDeviceAttributeMultiprocessorCount); + GetDeviceAttribute(hipDeviceAttributeMultiprocessorCount, 0); const unsigned int dim = std::ceil(std::cbrt(max_blocks * multiproc_count)); HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast(kernel), dim3{dim, dim, dim}, dim3{1, 1, 1}, nullptr, 0, nullptr), @@ -171,7 +171,7 @@ TEST_CASE("Unit_hipLaunchCooperativeKernel_Negative_Parameters") { } SECTION("sharedMemBytes > maxSharedMemoryPerBlock") { - const unsigned int max = GetDeviceAttribute(0, hipDeviceAttributeMaxSharedMemoryPerBlock) + 1u; + const unsigned int max = GetDeviceAttribute(hipDeviceAttributeMaxSharedMemoryPerBlock, 0) + 1u; HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{1, 1, 1}, nullptr, max, nullptr), hipErrorCooperativeLaunchTooLarge); diff --git a/projects/hip-tests/catch/unit/executionControl/hipLaunchKernel.cc b/projects/hip-tests/catch/unit/executionControl/hipLaunchKernel.cc new file mode 100644 index 0000000000..d9272107eb --- /dev/null +++ b/projects/hip-tests/catch/unit/executionControl/hipLaunchKernel.cc @@ -0,0 +1,156 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "execution_control_common.hh" + +#include +#include +#include +#include + +TEST_CASE("Unit_hipLaunchKernel_Positive_Basic") { + SECTION("Kernel with no arguments") { + HIP_CHECK(hipLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{1, 1, 1}, + nullptr, 0, nullptr)); + HIP_CHECK(hipDeviceSynchronize()); + } + + SECTION("Kernel with arguments using kernelParams") { + LinearAllocGuard result_dev(LinearAllocs::hipMalloc, sizeof(int)); + HIP_CHECK(hipMemset(result_dev.ptr(), 0, sizeof(*result_dev.ptr()))); + int* result_ptr = result_dev.ptr(); + void* kernel_args[1] = {&result_ptr}; + HIP_CHECK(hipLaunchKernel(reinterpret_cast(kernel_42), dim3{1, 1, 1}, dim3{1, 1, 1}, + kernel_args, 0, nullptr)); + int result = 0; + HIP_CHECK(hipMemcpy(&result, result_dev.ptr(), sizeof(result), hipMemcpyDefault)); + REQUIRE(result == 42); + } +} + +TEST_CASE("Unit_hipLaunchKernel_Positive_Parameters") { + SECTION("blockDim.x == maxBlockDimX") { + const unsigned int x = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimX, 0); + HIP_CHECK(hipLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{x, 1, 1}, + nullptr, 0, nullptr)); + } + + SECTION("blockDim.y == maxBlockDimY") { + const unsigned int y = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimY, 0); + HIP_CHECK(hipLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{y, 1, 1}, + nullptr, 0, nullptr)); + } + + SECTION("blockDim.z == maxBlockDimZ") { + const unsigned int z = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimZ, 0); + HIP_CHECK(hipLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{z, 1, 1}, + nullptr, 0, nullptr)); + } +} + +TEST_CASE("Unit_hipLaunchKernel_Negative_Parameters") { + SECTION("f == nullptr") { + HIP_CHECK_ERROR(hipLaunchKernel(nullptr, dim3{1, 1, 1}, dim3{1, 1, 1}, nullptr, 0, nullptr), + hipErrorInvalidDeviceFunction); + } + + SECTION("gridDim.x == 0") { + HIP_CHECK_ERROR(hipLaunchKernel(reinterpret_cast(kernel), dim3{0, 1, 1}, dim3{1, 1, 1}, + nullptr, 0, nullptr), + hipErrorInvalidValue); + } + + SECTION("gridDim.y == 0") { + HIP_CHECK_ERROR(hipLaunchKernel(reinterpret_cast(kernel), dim3{1, 0, 1}, dim3{1, 1, 1}, + nullptr, 0, nullptr), + hipErrorInvalidValue); + } + + SECTION("gridDim.z == 0") { + HIP_CHECK_ERROR(hipLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 0}, dim3{1, 1, 1}, + nullptr, 0, nullptr), + hipErrorInvalidValue); + } + + SECTION("blockDim.x == 0") { + HIP_CHECK_ERROR(hipLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{0, 1, 1}, + nullptr, 0, nullptr), + hipErrorInvalidValue); + } + + SECTION("blockDim.y == 0") { + HIP_CHECK_ERROR(hipLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{1, 0, 1}, + nullptr, 0, nullptr), + hipErrorInvalidValue); + } + + SECTION("blockDim.z == 0") { + HIP_CHECK_ERROR(hipLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{1, 1, 0}, + nullptr, 0, nullptr), + hipErrorInvalidValue); + } + + SECTION("blockDim.x > maxBlockDimX") { + const unsigned int x = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimX, 0) + 1u; + HIP_CHECK_ERROR(hipLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{x, 1, 1}, + nullptr, 0, nullptr), + hipErrorInvalidConfiguration); + } + + SECTION("blockDim.y > maxBlockDimY") { + const unsigned int y = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimY, 0) + 1u; + HIP_CHECK_ERROR(hipLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{1, y, 1}, + nullptr, 0, nullptr), + hipErrorInvalidConfiguration); + } + + SECTION("blockDim.z > maxBlockDimZ") { + const unsigned int z = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimZ, 0) + 1u; + HIP_CHECK_ERROR(hipLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{1, 1, z}, + nullptr, 0, nullptr), + hipErrorInvalidConfiguration); + } + + SECTION("blockDim.x * blockDim.y * blockDim.z > maxThreadsPerBlock") { + const unsigned int max = GetDeviceAttribute(hipDeviceAttributeMaxThreadsPerBlock, 0); + const unsigned int dim = std::ceil(std::cbrt(max)); + HIP_CHECK_ERROR(hipLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, + dim3{dim, dim, dim}, nullptr, 0, nullptr), + hipErrorInvalidConfiguration); + } + + SECTION("sharedMemBytes > maxSharedMemoryPerBlock") { + const unsigned int max = GetDeviceAttribute(hipDeviceAttributeMaxSharedMemoryPerBlock, 0) + 1u; + HIP_CHECK_ERROR(hipLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{1, 1, 1}, + nullptr, max, nullptr), + hipErrorOutOfMemory); + } + + SECTION("Invalid stream") { + hipStream_t stream = nullptr; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK_ERROR(hipLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{1, 1, 1}, + nullptr, 0, stream), + hipErrorInvalidValue); + } +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/executionControl/launch_api.cc b/projects/hip-tests/catch/unit/executionControl/launch_api.cc new file mode 100644 index 0000000000..64cdcf8266 --- /dev/null +++ b/projects/hip-tests/catch/unit/executionControl/launch_api.cc @@ -0,0 +1,69 @@ +/* +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "execution_control_common.hh" + +#include +#include + +TEST_CASE("Unit_hipLaunchByPtr_Positive_Basic") { + LinearAllocGuard alloc(LinearAllocs::hipMallocManaged, 4); + + SECTION("hipConfigureCall") { HIP_CHECK(hipConfigureCall(dim3{1}, dim3{1}, 0, nullptr)); } + + SECTION("__hipPushCallConfiguration") { + HIP_CHECK(__hipPushCallConfiguration(dim3{1}, dim3{1}, 0, nullptr)); + } + + int* arg = alloc.ptr(); + HIP_CHECK(hipSetupArgument(&arg, sizeof(int*), 0)); + + HIP_CHECK(hipLaunchByPtr(reinterpret_cast(kernel_42))); + HIP_CHECK(hipDeviceSynchronize()); + + REQUIRE(alloc.ptr()[0] == 42); +} + +TEST_CASE("Unit_hipLaunchByPtr_Negative_Parameters") { + HIP_CHECK(hipConfigureCall(dim3{1}, dim3{1}, 0, nullptr)); + HIP_CHECK_ERROR(hipLaunchByPtr(nullptr), hipErrorInvalidDeviceFunction); +} + +TEST_CASE("Unit___hipPushCallConfiguration_Positive_Basic") { + StreamGuard stream_guard(Streams::created); + HIP_CHECK(__hipPushCallConfiguration(dim3{1, 2, 3}, dim3{3, 2, 1}, 1024, stream_guard.stream())); + + dim3 grid; + dim3 block; + size_t shmem; + hipStream_t stream; + HIP_CHECK(__hipPopCallConfiguration(&grid, &block, &shmem, &stream)); + + REQUIRE(grid.x == 1); + REQUIRE(grid.y == 2); + REQUIRE(grid.z == 3); + REQUIRE(block.x == 3); + REQUIRE(block.y == 2); + REQUIRE(block.z == 1); + REQUIRE(shmem == 1024); + REQUIRE(stream == stream_guard.stream()); +} \ No newline at end of file