diff --git a/tests/catch/unit/stream/hipStreamGetCUMask.cc b/tests/catch/unit/stream/hipStreamGetCUMask.cc index 49f21bdbda..f63acc212a 100644 --- a/tests/catch/unit/stream/hipStreamGetCUMask.cc +++ b/tests/catch/unit/stream/hipStreamGetCUMask.cc @@ -117,9 +117,21 @@ TEST_CASE("Unit_hipExtStreamGetCUMask_verifyDefaultAndCustomMask") { } SECTION("Verify with custom mask set") { + hipDeviceProp_t props; std::vector customMask(defaultCUMask); hipStream_t stream; - customMask[0] = 0xe; + int deviceId; + + HIP_CHECK(hipGetDevice(&deviceId)); + HIP_CHECK(hipGetDeviceProperties(&props, deviceId)); + + if (props.major >= 10) { + // For gfx >= 10, one work group processor encompasses 2 CUs & + // hence the CUs need to be enabled in pair + customMask[0] = 0xc; + } else { + customMask[0] = 0xe; + } HIP_CHECK(hipExtStreamCreateWithCUMask(&stream, customMask.size(), customMask.data())); diff --git a/tests/catch/unit/stream/hipStreamWithCUMask.cc b/tests/catch/unit/stream/hipStreamWithCUMask.cc index ea6b033d9c..adf1db3e80 100644 --- a/tests/catch/unit/stream/hipStreamWithCUMask.cc +++ b/tests/catch/unit/stream/hipStreamWithCUMask.cc @@ -204,21 +204,27 @@ TEST_CASE("Unit_hipExtStreamCreateWithCUMask_Functionality") { hA[i] = CONSTANT + i; } + int cuCountPerGroup = 1; + if (props.major >= 10) { + cuCountPerGroup = 2; // For gfx >= 10, one work group processor encompasses 2 CUs + } + + unsigned long mask = pow(2, cuCountPerGroup) - 1; for (int np = 0; np < KNumPartition; np++) { HIP_CHECK(hipMalloc(&dA[np], Nbytes)); HIP_CHECK(hipMalloc(&dC[np], Nbytes)); // make unique CU masks in the multiple of dwords for each stream uint32_t temp = 0; - uint32_t bit_index = np; - for (int i = np; i < props.multiProcessorCount; i = i + 4) { - temp |= 1UL << bit_index; + uint32_t bit_index = cuCountPerGroup * np; + for (int i = np; i < props.multiProcessorCount; i = i + cuCountPerGroup * 4) { + temp |= mask << bit_index; if (bit_index >= 32) { cuMasks[np].push_back(temp); temp = 0; - bit_index = np; - temp |= 1UL << bit_index; + bit_index = cuCountPerGroup * np; + temp |= mask << bit_index; } - bit_index += 4; + bit_index += cuCountPerGroup * 4; } if (bit_index != 0) { cuMasks[np].push_back(temp); @@ -229,9 +235,9 @@ TEST_CASE("Unit_hipExtStreamCreateWithCUMask_Functionality") { HIP_CHECK(hipMemcpy(dA[np], hA, Nbytes, hipMemcpyHostToDevice)); - ss[np] << std::hex; + ss[np] << std::hex << std::setfill('0'); for (int i = cuMasks[np].size() - 1; i >= 0; i--) { - ss[np] << cuMasks[np][i]; + ss[np] << std::setw(8) << cuMasks[np][i]; } } diff --git a/tests/src/runtimeApi/stream/hipStreamGetCUMask.cpp b/tests/src/runtimeApi/stream/hipStreamGetCUMask.cpp index 3468a1fa53..ec3722f91d 100644 --- a/tests/src/runtimeApi/stream/hipStreamGetCUMask.cpp +++ b/tests/src/runtimeApi/stream/hipStreamGetCUMask.cpp @@ -124,7 +124,13 @@ int main(int argc, char* argv[]) { cout << "info: CU mask for the default stream is: 0x" << ss.str().c_str() << endl; vector cuMask1(defaultCUMask); - cuMask1[0] = 0xe; + if (props.major >= 10) { + // For gfx >= 10, one work group processor encompasses 2 CUs & + // hence the CUs need to be enabled in pair + cuMask1[0] = 0xc; + } else { + cuMask1[0] = 0xe; + } HIPCHECK(hipExtStreamCreateWithCUMask(&stream, cuMask1.size(), cuMask1.data())); ss.str(""); diff --git a/tests/src/runtimeApi/stream/hipStreamWithCUMask.cpp b/tests/src/runtimeApi/stream/hipStreamWithCUMask.cpp index efe98a8bae..07ac4c527b 100644 --- a/tests/src/runtimeApi/stream/hipStreamWithCUMask.cpp +++ b/tests/src/runtimeApi/stream/hipStreamWithCUMask.cpp @@ -72,6 +72,13 @@ int main(int argc, char* argv[]) { hA[i] = 1.618f + i; } + int cuCountPerGroup = 1; + if (props.major >= 10) { + cuCountPerGroup = 2; // For gfx >= 10, one work group processor encompasses 2 CUs + } + + unsigned long mask = pow(2, cuCountPerGroup) - 1; + for (int np = 0; np < numPartition; np++) { HIPCHECK(hipMalloc(&dA[np], Nbytes)); @@ -79,16 +86,16 @@ int main(int argc, char* argv[]) { // make unique CU masks in the multiple of dwords for each stream uint32_t temp = 0; - uint32_t bit_index = np; - for (int i = np; i < props.multiProcessorCount; i = i + 4) { - temp |= 1UL << bit_index; + uint32_t bit_index = cuCountPerGroup * np; + for (int i = np; i < props.multiProcessorCount; i = i + cuCountPerGroup * 4) { + temp |= mask << bit_index; if (bit_index >= 32) { cuMasks[np].push_back(temp); temp = 0; - bit_index = np; - temp |= 1UL << bit_index; + bit_index = cuCountPerGroup * np; + temp |= mask << bit_index; } - bit_index += 4; + bit_index += cuCountPerGroup * 4; } if (bit_index != 0) { cuMasks[np].push_back(temp); @@ -98,9 +105,9 @@ int main(int argc, char* argv[]) { HIPCHECK(hipMemcpy(dA[np], hA, Nbytes, hipMemcpyHostToDevice)); - ss[np] << std::hex; + ss[np] << std::hex << std::setfill('0'); for (int i = cuMasks[np].size() - 1; i >= 0; i--) { - ss[np] << cuMasks[np][i]; + ss[np] << std::setw(8) << cuMasks[np][i]; } }