AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge 'develop' into 'amd-staging'

Browse Source 
Change-Id: Idd8cc63ca258dc2a72126f6894cb5a135432f3ed
This commit is contained in:
Jenkins
2023-03-14 23:10:37 +00:00
parent 8ba8ba51e0 ac4267e21e
commit a28856619c
13 changed files with 46 additions and 39 deletions
Download Patch File Download Diff File Expand all files Collapse all files
catch/include/utils.hh
+1 -1
View File
@@ -40,7 +40,7 @@ void ArrayMismatch(T* const expected, T* const actual, const size_t num_elements
template <typename It, typename T> void ArrayFindIfNot(It begin, It end, const T expected_value) {
const auto it = std::find_if_not(
begin, end, [expected_value](const int elem) { return expected_value == elem; });
begin, end, [expected_value](const T elem) { return expected_value == elem; });
if (it != end) {
const auto idx = std::distance(begin, it);
catch/unit/callback/hipGetStreamDeviceId.cc
+1 -1
View File
@@ -70,7 +70,7 @@ TEST_CASE("Unit_hipGetStreamDeviceId_Positive_Multithreaded_Basic") {
const int device_count = HipTest::getDeviceCount();
auto thread_function = [&]() {
for(unsigned int id = 0; id < device_count; ++id) {
for(int id = 0; id < device_count; ++id) {
HIP_CHECK_THREAD(hipSetDevice(id));
StreamGuard stream_guard{Streams::perThread};
catch/unit/deviceLib/SimpleAtomicsTest.cc
+7 -7
View File
@@ -108,9 +108,9 @@ template<typename T, typename enable_if<is_integral<T>{}>::type* = nullptr>
bool VerifyIntegral(T* gpuData, int len) {
// atomic Max
T val = 0;
for (T i = 0; i < len; ++i) {
for (int i = 0; i < len; ++i) {
// fourth element should be len-1
val = max(val, i);
val = max(val, static_cast<T>(i));
}
REQUIRE(val == gpuData[3]);
@@ -118,14 +118,14 @@ bool VerifyIntegral(T* gpuData, int len) {
// atomic Min
val = 1 << 8;
for (T i = 0; i < len; ++i) {
val = min(val, i);
for (int i = 0; i < len; ++i) {
val = min(val, static_cast<T>(i));
}
REQUIRE(val == gpuData[4]);
// atomic Inc
int limit = 17;
T limit = 17;
val = 0;
for (int i = 0; i < len; ++i) {
@@ -145,9 +145,9 @@ bool VerifyIntegral(T* gpuData, int len) {
REQUIRE(val == gpuData[6]);
// atomic CAS
for (T i = 0; i < len; ++i) {
for (int i = 0; i < len; ++i) {
// eighth element should be a member of [0, len)
if (i == gpuData[7]) {
if (static_cast<T>(i) == gpuData[7]) {
return true;
break;
}
catch/unit/deviceLib/hipStdComplex.cc
+5 -3
View File
@@ -24,7 +24,6 @@ THE SOFTWARE.
// Tolerance for error
const double tolerance = 1e-6;
const bool verbose = false;
#define LEN 64
@@ -117,14 +116,15 @@ void test() {
hipLaunchKernelGGL(kernel<FloatT>, dim3(1), dim3(LEN), 0, 0,
Ad, Bd, Cd, CK);
HIP_CHECK(hipMemcpy(C, Cd, sizeof(ComplexT)*LEN, hipMemcpyDeviceToHost));
bool pass = true;
for (int i = 0; i < LEN; i++) {
ComplexT Expected = calc(A[i], B[i], CK);
FloatT error = abs(C[i] - Expected);
if (abs(Expected) > tolerance)
error /= abs(Expected);
bool pass = error < tolerance;
pass &= error < tolerance;
}
return true;
return pass;
};
#define OP(x) assert(test_fun(CK_##x));
@@ -140,6 +140,7 @@ void test() {
delete[] D;
}
#if HT_AMD
TEST_CASE("Unit_StdComplex") {
SECTION("Test run with float") {
test<float>();
@@ -148,3 +149,4 @@ TEST_CASE("Unit_StdComplex") {
test<double>();
}
}
#endif
catch/unit/deviceLib/hipTestAtomicAdd.cc
+1 -3
View File
@@ -56,13 +56,11 @@ bool p_atomicNoRet = false;
template <typename T>
__global__ void atomicnoret_manywaves(T* C_d) {
size_t tid = (blockIdx.x * blockDim.x + threadIdx.x);
atomicAddNoRet(C_d, INCREMENT_VALUE);
}
template <typename T>
__global__ void atomic_manywaves(T* C_d) {
size_t tid = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
atomicAdd(C_d, INCREMENT_VALUE);
}
@@ -95,7 +93,7 @@ bool atomictest_manywaves(const T& initial_val) {
// Copy result from device to host
HIP_CHECK(hipMemcpy(hOData, dOData, memSize, hipMemcpyDeviceToHost));
REQUIRE(hOData[0] == initial_val+
(INCREMENT_VALUE*(ThreadsperBlock*numBlocks)));
static_cast<T>(INCREMENT_VALUE*(ThreadsperBlock*numBlocks)));
// Cleanup memory
free(hOData);
catch/unit/graph/hipLaunchHostFunc.cc
+2 -2
View File
@@ -117,7 +117,7 @@ TEST_CASE("Unit_hipLaunchHostFunc_Positive_Functional") {
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
// Replay the recorded sequence multiple times
for (int i = 0; i < kLaunchIters; i++) {
for (size_t i = 0; i < kLaunchIters; i++) {
std::fill_n(A_h.host_ptr(), 1, static_cast<float>(i));
HIP_CHECK(hipGraphLaunch(graphExec, stream));
HIP_CHECK(hipStreamSynchronize(stream));
@@ -171,7 +171,7 @@ TEST_CASE("Unit_hipLaunchHostFunc_Positive_Thread") {
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
// Replay the recorded sequence multiple times
for (int i = 0; i < kLaunchIters; i++) {
for (size_t i = 0; i < kLaunchIters; i++) {
std::fill_n(A_h.host_ptr(), 1, static_cast<float>(i));
HIP_CHECK(hipGraphLaunch(graphExec, stream));
HIP_CHECK(hipStreamSynchronize(stream));
catch/unit/graph/hipStreamBeginCapture.cc
+10 -10
View File
@@ -84,7 +84,7 @@ void captureStreamAndLaunchGraph(F graphFunc, hipStreamCaptureMode mode, hipStre
REQUIRE(graphExec != nullptr);
// Replay the recorded sequence multiple times
for (int i = 0; i < kLaunchIters; i++) {
for (size_t i = 0; i < kLaunchIters; i++) {
std::fill_n(A_h.host_ptr(), N, static_cast<float>(i));
HIP_CHECK(hipGraphLaunch(graphExec, stream));
HIP_CHECK(hipStreamSynchronize(stream));
@@ -243,7 +243,7 @@ static void interStrmEventSyncCapture(const hipStream_t& stream1, const hipStrea
REQUIRE(graphExec2 != nullptr);
// Replay the recorded sequence multiple times
for (int i = 0; i < kLaunchIters; i++) {
for (size_t i = 0; i < kLaunchIters; i++) {
// Execute the Graphs
HIP_CHECK(hipGraphLaunch(graphExec1, stream1));
HIP_CHECK(hipGraphLaunch(graphExec2, stream2));
@@ -286,7 +286,7 @@ static void colligatedStrmCapture(const hipStream_t& stream1, const hipStream_t&
REQUIRE(graphExec2 != nullptr);
// Replay the recorded sequence multiple times
for (int i = 0; i < kLaunchIters; i++) {
for (size_t i = 0; i < kLaunchIters; i++) {
// Execute the Graphs
HIP_CHECK(hipGraphLaunch(graphExec1, stream1));
HIP_CHECK(hipGraphLaunch(graphExec2, stream2));
@@ -340,15 +340,15 @@ static void colligatedStrmCaptureFunc(const hipStream_t& stream1, const hipStrea
REQUIRE(graphExec2 != nullptr);
// Execute the Graphs
for (int iter = 0; iter < kLaunchIters; iter++) {
for (size_t iter = 0; iter < kLaunchIters; iter++) {
std::fill_n(A_h.host_ptr(), N, iter);
std::fill_n(C_h.host_ptr(), N, iter);
HIP_CHECK(hipGraphLaunch(graphExec1, stream1));
HIP_CHECK(hipGraphLaunch(graphExec2, stream2));
HIP_CHECK(hipStreamSynchronize(stream1));
HIP_CHECK(hipStreamSynchronize(stream2));
ArrayFindIfNot(B_h.host_ptr(), iter * iter, N);
ArrayFindIfNot(D_h.host_ptr(), iter * iter, N);
ArrayFindIfNot(B_h.host_ptr(), static_cast<int>(iter * iter), N);
ArrayFindIfNot(D_h.host_ptr(), static_cast<int>(iter * iter), N);
}
// Free
@@ -407,15 +407,15 @@ static void multithreadedTest(hipStreamCaptureMode mode) {
REQUIRE(graphExec2 != nullptr);
// Execute the Graphs
for (int iter = 0; iter < kLaunchIters; iter++) {
for (size_t iter = 0; iter < kLaunchIters; iter++) {
std::fill_n(A_h.host_ptr(), N, iter);
std::fill_n(C_h.host_ptr(), N, iter);
HIP_CHECK(hipGraphLaunch(graphExec1, stream1));
HIP_CHECK(hipGraphLaunch(graphExec2, stream2));
HIP_CHECK(hipStreamSynchronize(stream1));
HIP_CHECK(hipStreamSynchronize(stream2));
ArrayFindIfNot(B_h.host_ptr(), iter * iter, N);
ArrayFindIfNot(D_h.host_ptr(), iter * iter, N);
ArrayFindIfNot(B_h.host_ptr(), static_cast<int>(iter * iter), N);
ArrayFindIfNot(D_h.host_ptr(), static_cast<int>(iter * iter), N);
}
// Free
@@ -1212,7 +1212,7 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_captureComplexGraph") {
hipGraphExec_t graphExec{nullptr};
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
// Verify graph
for (int iter = 0; iter < kLaunchIters; iter++) {
for (size_t iter = 0; iter < kLaunchIters; iter++) {
std::fill_n(Ah.host_ptr(), N, iter);
std::fill_n(Bh.host_ptr(), N, iter);
HIP_CHECK(hipGraphLaunch(graphExec, streams[0]));
catch/unit/graph/hipStreamEndCapture.cc
+1 -1
View File
@@ -194,7 +194,7 @@ TEST_CASE("Unit_hipStreamEndCapture_Positive_Thread") {
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
// Replay the recorded sequence multiple times
for (int i = 0; i < kLaunchIters; i++) {
for (size_t i = 0; i < kLaunchIters; i++) {
std::fill_n(A_h.host_ptr(), N, static_cast<float>(i));
HIP_CHECK(hipGraphLaunch(graphExec, stream));
HIP_CHECK(hipStreamSynchronize(stream));
catch/unit/graph/hipStreamGetCaptureInfo.cc
+4 -2
View File
@@ -71,7 +71,7 @@ void checkStreamCaptureInfo(hipStreamCaptureMode mode, hipStream_t stream) {
REQUIRE(graphExec != nullptr);
// Replay the recorded sequence multiple times
for (int i = 0; i < kLaunchIters; i++) {
for (size_t i = 0; i < kLaunchIters; i++) {
std::fill_n(A_h.host_ptr(), N, static_cast<float>(i));
HIP_CHECK(hipGraphLaunch(graphExec, stream));
HIP_CHECK(hipStreamSynchronize(stream));
@@ -166,7 +166,9 @@ TEST_CASE("Unit_hipStreamGetCaptureInfo_Positive_UniqueID") {
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_hipStreamGetCaptureInfo_Negative_Parameters") {
#if HT_NVIDIA
hipStreamCaptureStatus cStatus;
#endif
unsigned long long capSequenceID; // NOLINT
const auto stream_type = GENERATE(Streams::perThread, Streams::created);
StreamGuard stream_guard(stream_type);
@@ -194,4 +196,4 @@ TEST_CASE("Unit_hipStreamGetCaptureInfo_Negative_Parameters") {
hipErrorContextIsDestroyed);
}
#endif
}
}
catch/unit/graph/hipStreamGetCaptureInfo_v2.cc
+5 -3
View File
@@ -40,7 +40,7 @@ void checkStreamCaptureInfo_v2(hipStreamCaptureMode mode, hipStream_t stream) {
hipGraph_t graph{nullptr}, capInfoGraph{nullptr};
hipGraphExec_t graphExec{nullptr};
const hipGraphNode_t* nodelist{};
int numDepsCreated = 0;
size_t numDepsCreated = 0;
hipStreamCaptureStatus captureStatus{hipStreamCaptureStatusNone};
hipGraphNodeType type(hipGraphNodeTypeEmpty);
unsigned long long capSequenceID = 0; // NOLINT
@@ -119,7 +119,7 @@ void checkStreamCaptureInfo_v2(hipStreamCaptureMode mode, hipStream_t stream) {
REQUIRE(graphExec != nullptr);
// Replay the recorded sequence multiple times
for (int i = 0; i < kLaunchIters; i++) {
for (size_t i = 0; i < kLaunchIters; i++) {
std::fill_n(A_h.host_ptr(), N, static_cast<float>(i));
HIP_CHECK(hipGraphLaunch(graphExec, stream));
HIP_CHECK(hipStreamSynchronize(stream));
@@ -221,7 +221,9 @@ TEST_CASE("Unit_hipStreamGetCaptureInfo_v2_Positive_UniqueID") {
*/
TEST_CASE("Unit_hipStreamGetCaptureInfo_v2_Negative_Parameters") {
hipGraph_t capInfoGraph{};
#if HT_NVIDIA
hipStreamCaptureStatus captureStatus;
#endif
unsigned long long capSequenceID; // NOLINT
size_t numDependencies;
const hipGraphNode_t* nodelist{};
@@ -256,4 +258,4 @@ TEST_CASE("Unit_hipStreamGetCaptureInfo_v2_Negative_Parameters") {
hipErrorContextIsDestroyed);
}
#endif
}
}
catch/unit/graph/hipStreamIsCapturing.cc
+2 -2
View File
@@ -133,7 +133,7 @@ void checkStreamCaptureStatus(hipStreamCaptureMode mode, hipStream_t stream) {
REQUIRE(graphExec != nullptr);
// Replay the recorded sequence multiple times
for (int i = 0; i < kLaunchIters; i++) {
for (size_t i = 0; i < kLaunchIters; i++) {
std::fill_n(A_h.host_ptr(), N, static_cast<float>(i));
HIP_CHECK(hipGraphLaunch(graphExec, stream));
HIP_CHECK(hipStreamSynchronize(stream));
@@ -209,4 +209,4 @@ TEST_CASE("Unit_hipStreamIsCapturing_Positive_Thread") {
HIP_CHECK(hipStreamEndCapture(stream, &graph));
HIP_CHECK(hipGraphDestroy(graph));
}
}
catch/unit/graph/hipStreamUpdateCaptureDependencies.cc
+4 -4
View File
@@ -34,7 +34,7 @@ THE SOFTWARE.
* update the set of dependencies in a capturing stream
*/
static __global__ void vectorSet(const float* A_d, float* B_d, int64_t NELEM) {
static __global__ void vectorSet(const float* A_d, float* B_d, size_t NELEM) {
size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
size_t stride = blockDim.x * gridDim.x;
@@ -159,7 +159,7 @@ static void UpdateStreamCaptureDependenciesSet(hipStream_t stream,
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
// Replay the recorded sequence multiple times
for (int i = 0; i < kLaunchIters; i++) {
for (size_t i = 0; i < kLaunchIters; i++) {
std::fill_n(A_h.host_ptr(), N, static_cast<float>(i));
std::fill_n(C_h.host_ptr(), N, static_cast<float>(i));
HIP_CHECK(hipGraphLaunch(graphExec, stream));
@@ -275,7 +275,7 @@ static void UpdateStreamCaptureDependenciesAdd(hipStream_t stream,
HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
// Replay the recorded sequence multiple times
for (int i = 0; i < kLaunchIters; i++) {
for (size_t i = 0; i < kLaunchIters; i++) {
std::fill_n(A_h.host_ptr(), N, static_cast<float>(i));
std::fill_n(C_h.host_ptr(), N, static_cast<float>(i));
HIP_CHECK(hipGraphLaunch(graphExec, stream));
@@ -367,7 +367,7 @@ TEST_CASE("Unit_hipStreamUpdateCaptureDependencies_Positive_Parameters") {
const hipStreamUpdateCaptureDependenciesFlags flag =
GENERATE(hipStreamAddCaptureDependencies, hipStreamSetCaptureDependencies);
HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal));
HIP_CHECK(hipStreamBeginCapture(stream, captureMode)); //hipStreamCaptureModeGlobal));
HIP_CHECK(hipStreamUpdateCaptureDependencies(stream, nullptr, 0, flag));
catch/unit/graph/stream_capture_common.hh
+3
View File
@@ -40,6 +40,8 @@ void captureSequenceSimple(T* hostMem1, T* devMem1, T* hostMem2, size_t N,
template <typename T>
void captureSequenceLinear(T* hostMem1, T* devMem1, T* hostMem2, T* devMem2, size_t N,
hipStream_t captureStream) {
{(void)(hostMem2);} //unused hostMem2
size_t Nbytes = N * sizeof(T);
HIP_CHECK(hipMemcpyAsync(devMem1, hostMem1, Nbytes, hipMemcpyHostToDevice, captureStream));
@@ -51,6 +53,7 @@ template <typename T>
void captureSequenceBranched(T* hostMem1, T* devMem1, T* hostMem2, T* devMem2, size_t N,
hipStream_t captureStream, std::vector<hipStream_t>& streams,
std::vector<hipEvent_t>& events) {
{(void)(hostMem2);} //unused hostMem2
size_t Nbytes = N * sizeof(T);
HIP_CHECK(hipEventRecord(events[0], captureStream));