SWDEV-1 - Fix warnings while building tests (#201)
Change-Id: I7faf3446f8c9dafbbfc6bab4c6c1b756f0e4219d
This commit is contained in:
@@ -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;
|
||||
}
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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);
|
||||
|
||||
Reference in New Issue
Block a user