EXSWHTEC-115 - Implement tests for hipLaunchCooperativeKernel APIs #59
Change-Id: I5bda5ee3787a73aeeab5c25c05278e7aa2f8bfa2
[ROCm/hip-tests commit: 380ace735f]
This commit is contained in:
Mirza Halilčević
committed by
Rakesh Roy
Rakesh Roy
parent
b66fc2b672
commit
f3a4512a78
@@ -48,11 +48,6 @@
|
||||
"Unit_hipFuncSetAttribute_Positive_PreferredSharedMemoryCarveout",
|
||||
"Unit_hipFuncSetAttribute_Positive_Parameters",
|
||||
"Unit_hipFuncSetAttribute_Negative_Parameters",
|
||||
"NOTE: The following 4 tests are disabled due to defect - EXSWHTEC-240",
|
||||
"Unit_hipFuncSetCacheConfig_Negative_Not_Supported",
|
||||
"Unit_hipFuncSetSharedMemConfig_Negative_Not_Supported",
|
||||
"Unit_hipFuncSetAttribute_Positive_MaxDynamicSharedMemorySize_Not_Supported",
|
||||
"Unit_hipFuncSetAttribute_Positive_PreferredSharedMemoryCarveout_Not_Supported",
|
||||
"NOTE: The following test is disabled due to defect - EXSWHTEC-243",
|
||||
"Unit_hipExtLaunchKernel_Negative_Parameters",
|
||||
"NOTE: The following test is disabled due to defect - EXSWHTEC-244",
|
||||
|
||||
@@ -115,11 +115,6 @@
|
||||
"Unit_hipEventCreateWithFlags_DefaultFlg_NonCohHstMem",
|
||||
"Unit_hipEventCreateWithFlags_DisableSystemFence_CohHstMem",
|
||||
"Unit_hipEventCreateWithFlags_DefaultFlg_CohHstMem",
|
||||
"NOTE: The following 4 tests are disabled due to defect - EXSWHTEC-240",
|
||||
"Unit_hipFuncSetCacheConfig_Negative_Not_Supported",
|
||||
"Unit_hipFuncSetSharedMemConfig_Negative_Not_Supported",
|
||||
"Unit_hipFuncSetAttribute_Positive_MaxDynamicSharedMemorySize_Not_Supported",
|
||||
"Unit_hipFuncSetAttribute_Positive_PreferredSharedMemoryCarveout_Not_Supported",
|
||||
"Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_Negative_Parameters",
|
||||
"Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_Negative_Parameters",
|
||||
"Unit_hipModuleOccupancyMaxPotentialBlockSizeWithFlags_Negative_Parameters",
|
||||
|
||||
@@ -3,6 +3,9 @@ set(TEST_SRC
|
||||
hipFuncSetCacheConfig.cc
|
||||
hipFuncSetSharedMemConfig.cc
|
||||
hipFuncSetAttribute.cc
|
||||
hipFuncGetAttributes.cc
|
||||
hipLaunchCooperativeKernel.cc
|
||||
hipLaunchCooperativeKernelMultiDevice.cc
|
||||
)
|
||||
|
||||
if(HIP_PLATFORM MATCHES "amd")
|
||||
@@ -15,4 +18,4 @@ endif()
|
||||
hip_add_exe_to_target(NAME ExecutionControlTest
|
||||
TEST_SRC ${TEST_SRC}
|
||||
TEST_TARGET_NAME build_tests
|
||||
COMPILE_OPTIONS -std=c++17)
|
||||
COMPILE_OPTIONS -std=c++17)
|
||||
|
||||
@@ -23,5 +23,15 @@ THE SOFTWARE.
|
||||
#include "execution_control_common.hh"
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip/hip_cooperative_groups.h>
|
||||
|
||||
__global__ void kernel() {}
|
||||
__global__ void kernel() {}
|
||||
|
||||
__global__ void kernel2() {}
|
||||
|
||||
__global__ void kernel_42(int* val) { *val = 42; }
|
||||
|
||||
__global__ void coop_kernel() {
|
||||
cooperative_groups::grid_group grid = cooperative_groups::this_grid();
|
||||
grid.sync();
|
||||
}
|
||||
@@ -22,4 +22,10 @@ THE SOFTWARE.
|
||||
|
||||
#pragma once
|
||||
|
||||
__global__ void kernel();
|
||||
__global__ void kernel();
|
||||
|
||||
__global__ void kernel2();
|
||||
|
||||
__global__ void kernel_42(int* val);
|
||||
|
||||
__global__ void coop_kernel();
|
||||
@@ -0,0 +1,188 @@
|
||||
/*
|
||||
Copyright (c) 2022 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 <hip_test_common.hh>
|
||||
#include <hip/hip_runtime_api.h>
|
||||
#include <resource_guards.hh>
|
||||
#include <utils.hh>
|
||||
|
||||
TEST_CASE("Unit_hipLaunchCooperativeKernel_Positive_Basic") {
|
||||
if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) {
|
||||
HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported");
|
||||
return;
|
||||
}
|
||||
|
||||
SECTION("Cooperative kernel with no arguments") {
|
||||
HIP_CHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(coop_kernel), dim3{2, 2, 1},
|
||||
dim3{1, 1, 1}, nullptr, 0, nullptr));
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
}
|
||||
|
||||
SECTION("Kernel with arguments using kernelParams") {
|
||||
LinearAllocGuard<int> 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(hipLaunchCooperativeKernel(reinterpret_cast<void*>(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_hipLaunchCooperativeKernel_Positive_Parameters") {
|
||||
if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) {
|
||||
HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported");
|
||||
return;
|
||||
}
|
||||
|
||||
SECTION("blockDim.x == maxBlockDimX") {
|
||||
const unsigned int x = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimX);
|
||||
HIP_CHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
|
||||
dim3{x, 1, 1}, nullptr, 0, nullptr));
|
||||
}
|
||||
|
||||
SECTION("blockDim.y == maxBlockDimY") {
|
||||
const unsigned int y = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimY);
|
||||
HIP_CHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
|
||||
dim3{y, 1, 1}, nullptr, 0, nullptr));
|
||||
}
|
||||
|
||||
SECTION("blockDim.z == maxBlockDimZ") {
|
||||
const unsigned int z = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimZ);
|
||||
HIP_CHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
|
||||
dim3{z, 1, 1}, nullptr, 0, nullptr));
|
||||
}
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipLaunchCooperativeKernel_Negative_Parameters") {
|
||||
if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) {
|
||||
HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported");
|
||||
return;
|
||||
}
|
||||
|
||||
SECTION("f == nullptr") {
|
||||
HIP_CHECK_ERROR(hipLaunchCooperativeKernel(static_cast<void*>(nullptr), dim3{1, 1, 1},
|
||||
dim3{1, 1, 1}, nullptr, 0, nullptr),
|
||||
hipErrorInvalidDeviceFunction);
|
||||
}
|
||||
|
||||
SECTION("gridDim.x == 0") {
|
||||
HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast<void*>(kernel), dim3{0, 1, 1},
|
||||
dim3{1, 1, 1}, nullptr, 0, nullptr),
|
||||
hipErrorInvalidConfiguration);
|
||||
}
|
||||
|
||||
SECTION("gridDim.y == 0") {
|
||||
HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast<void*>(kernel), dim3{1, 0, 1},
|
||||
dim3{1, 1, 1}, nullptr, 0, nullptr),
|
||||
hipErrorInvalidConfiguration);
|
||||
}
|
||||
|
||||
SECTION("gridDim.z == 0") {
|
||||
HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 0},
|
||||
dim3{1, 1, 1}, nullptr, 0, nullptr),
|
||||
hipErrorInvalidConfiguration);
|
||||
}
|
||||
|
||||
SECTION("blockDim.x == 0") {
|
||||
HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
|
||||
dim3{0, 1, 1}, nullptr, 0, nullptr),
|
||||
hipErrorInvalidConfiguration);
|
||||
}
|
||||
|
||||
SECTION("blockDim.y == 0") {
|
||||
HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
|
||||
dim3{1, 0, 1}, nullptr, 0, nullptr),
|
||||
hipErrorInvalidConfiguration);
|
||||
}
|
||||
|
||||
SECTION("blockDim.z == 0") {
|
||||
HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
|
||||
dim3{1, 1, 0}, nullptr, 0, nullptr),
|
||||
hipErrorInvalidConfiguration);
|
||||
}
|
||||
|
||||
SECTION("blockDim.x > maxBlockDimX") {
|
||||
const unsigned int x = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimX) + 1u;
|
||||
HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast<void*>(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;
|
||||
HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast<void*>(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;
|
||||
HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast<void*>(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 dim = std::ceil(std::cbrt(max));
|
||||
HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
|
||||
dim3{dim, dim, dim}, nullptr, 0, nullptr),
|
||||
hipErrorInvalidConfiguration);
|
||||
}
|
||||
|
||||
SECTION(
|
||||
"gridDim.x * gridDim.y * gridDim.z > maxActiveBlocksPerMultiprocessor * "
|
||||
"multiProcessorCount") {
|
||||
int max_blocks;
|
||||
HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&max_blocks,
|
||||
reinterpret_cast<void*>(kernel), 1, 0));
|
||||
const unsigned int multiproc_count =
|
||||
GetDeviceAttribute(0, hipDeviceAttributeMultiprocessorCount);
|
||||
const unsigned int dim = std::ceil(std::cbrt(max_blocks * multiproc_count));
|
||||
HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast<void*>(kernel), dim3{dim, dim, dim},
|
||||
dim3{1, 1, 1}, nullptr, 0, nullptr),
|
||||
hipErrorCooperativeLaunchTooLarge);
|
||||
}
|
||||
|
||||
SECTION("sharedMemBytes > maxSharedMemoryPerBlock") {
|
||||
const unsigned int max = GetDeviceAttribute(0, hipDeviceAttributeMaxSharedMemoryPerBlock) + 1u;
|
||||
HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
|
||||
dim3{1, 1, 1}, nullptr, max, nullptr),
|
||||
hipErrorCooperativeLaunchTooLarge);
|
||||
}
|
||||
|
||||
SECTION("Invalid stream") {
|
||||
hipStream_t stream = nullptr;
|
||||
HIP_CHECK(hipStreamCreate(&stream));
|
||||
HIP_CHECK(hipStreamDestroy(stream));
|
||||
HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast<void*>(kernel), dim3{1, 1, 1},
|
||||
dim3{1, 1, 1}, nullptr, 0, stream),
|
||||
hipErrorContextIsDestroyed);
|
||||
}
|
||||
}
|
||||
+159
@@ -0,0 +1,159 @@
|
||||
/*
|
||||
Copyright (c) 2022 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 <hip_test_common.hh>
|
||||
#include <hip/hip_runtime_api.h>
|
||||
#include <resource_guards.hh>
|
||||
#include <utils.hh>
|
||||
|
||||
TEST_CASE("Unit_hipLaunchCooperativeKernelMultiDevice_Positive_Basic") {
|
||||
if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) {
|
||||
HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported");
|
||||
return;
|
||||
}
|
||||
|
||||
const auto device_count = HipTest::getDeviceCount();
|
||||
|
||||
std::vector<hipLaunchParams> params_list(device_count);
|
||||
|
||||
int device = 0;
|
||||
for (auto& params : params_list) {
|
||||
params.func = reinterpret_cast<void*>(coop_kernel);
|
||||
params.gridDim = dim3{1, 1, 1};
|
||||
params.blockDim = dim3{1, 1, 1};
|
||||
params.args = nullptr;
|
||||
params.sharedMem = 0;
|
||||
HIP_CHECK(hipSetDevice(device++));
|
||||
HIP_CHECK(hipStreamCreate(¶ms.stream));
|
||||
}
|
||||
|
||||
HIP_CHECK(hipLaunchCooperativeKernelMultiDevice(params_list.data(), device_count, 0u));
|
||||
|
||||
for (const auto params : params_list) {
|
||||
HIP_CHECK(hipStreamSynchronize(params.stream));
|
||||
}
|
||||
|
||||
for (const auto params : params_list) {
|
||||
HIP_CHECK(hipStreamDestroy(params.stream));
|
||||
}
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipLaunchCooperativeKernelMultiDevice_Negative_Parameters") {
|
||||
if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) {
|
||||
HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported");
|
||||
return;
|
||||
}
|
||||
|
||||
const auto device_count = HipTest::getDeviceCount();
|
||||
|
||||
std::vector<hipLaunchParams> params_list(device_count);
|
||||
|
||||
int device = 0;
|
||||
for (auto& params : params_list) {
|
||||
params.func = reinterpret_cast<void*>(coop_kernel);
|
||||
params.gridDim = dim3{1, 1, 1};
|
||||
params.blockDim = dim3{1, 1, 1};
|
||||
params.args = nullptr;
|
||||
params.sharedMem = 0;
|
||||
HIP_CHECK(hipSetDevice(device++));
|
||||
HIP_CHECK(hipStreamCreate(¶ms.stream));
|
||||
}
|
||||
|
||||
SECTION("launchParamsList == nullptr") {
|
||||
HIP_CHECK_ERROR(hipLaunchCooperativeKernelMultiDevice(nullptr, device_count, 0u),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("numDevices == 0") {
|
||||
HIP_CHECK_ERROR(hipLaunchCooperativeKernelMultiDevice(params_list.data(), 0, 0u),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("numDevices > device count") {
|
||||
HIP_CHECK_ERROR(hipLaunchCooperativeKernelMultiDevice(params_list.data(), device_count + 1, 0u),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("invalid flags") {
|
||||
HIP_CHECK_ERROR(hipLaunchCooperativeKernelMultiDevice(params_list.data(), device_count, 999),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
if (device_count > 1) {
|
||||
SECTION("launchParamsList.func doesn't match across all devices") {
|
||||
params_list[1].func = reinterpret_cast<void*>(kernel);
|
||||
HIP_CHECK_ERROR(hipLaunchCooperativeKernelMultiDevice(params_list.data(), device_count, 0u),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("launchParamsList.gridDim doesn't match across all kernels") {
|
||||
params_list[1].gridDim = dim3{2, 2, 2};
|
||||
HIP_CHECK_ERROR(hipLaunchCooperativeKernelMultiDevice(params_list.data(), device_count, 0u),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("launchParamsList.blockDim doesn't match across all kernels") {
|
||||
params_list[1].blockDim = dim3{2, 2, 2};
|
||||
HIP_CHECK_ERROR(hipLaunchCooperativeKernelMultiDevice(params_list.data(), device_count, 0u),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("launchParamsList.sharedMem doesn't match across all kernels") {
|
||||
params_list[1].sharedMem = 1024;
|
||||
HIP_CHECK_ERROR(hipLaunchCooperativeKernelMultiDevice(params_list.data(), device_count, 0u),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
}
|
||||
|
||||
for (const auto params : params_list) {
|
||||
HIP_CHECK(hipStreamDestroy(params.stream));
|
||||
}
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipLaunchCooperativeKernelMultiDevice_Negative_MultiKernelSameDevice") {
|
||||
if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) {
|
||||
HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported");
|
||||
return;
|
||||
}
|
||||
|
||||
HIP_CHECK(hipSetDevice(0));
|
||||
|
||||
std::vector<hipLaunchParams> params_list(2);
|
||||
|
||||
for (auto& params : params_list) {
|
||||
params.func = reinterpret_cast<void*>(coop_kernel);
|
||||
params.gridDim = dim3{1, 1, 1};
|
||||
params.blockDim = dim3{1, 1, 1};
|
||||
params.args = nullptr;
|
||||
params.sharedMem = 0;
|
||||
HIP_CHECK(hipStreamCreate(¶ms.stream));
|
||||
}
|
||||
|
||||
HIP_CHECK_ERROR(hipLaunchCooperativeKernelMultiDevice(params_list.data(), 2, 0u),
|
||||
hipErrorInvalidValue);
|
||||
|
||||
for (const auto params : params_list) {
|
||||
HIP_CHECK(hipStreamDestroy(params.stream));
|
||||
}
|
||||
}
|
||||
Reference in New Issue
Block a user