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

SWDEV-289409 - Add first batch of device tests, add new exe, helper functions and update test Guidelines

Browse Source 
Change-Id: I71929caf44725ba2cab7a6f0224bc37b9d04bcbb
This commit is contained in:
cjatin
2021-06-15 18:07:58 +05:30
committed by Jatin Chaudhary
parent 6890dc3ea1
commit 068e1f5043
26 changed files with 1693 additions and 39 deletions
Download Patch File Download Diff File Expand all files Collapse all files
catch/unit/kernels/CMakeLists.txt → catch/ABM/AddKernels/CMakeLists.txt
+2 -2
View File
@@ -4,7 +4,7 @@ set(TEST_SRC
)
# Create shared lib of all tests
add_library(Kernels SHARED EXCLUDE_FROM_ALL ${TEST_SRC})
add_library(ABMAddKernels SHARED EXCLUDE_FROM_ALL ${TEST_SRC})
# Add dependency on build_tests to build it on this custom target
add_dependencies(build_tests Kernels)
add_dependencies(build_tests ABMAddKernels)
catch/unit/kernels/add.cc → catch/ABM/AddKernels/add.cc
+1 -1
View File
@@ -6,7 +6,7 @@ template <typename T> __global__ void add(T* a, T* b, T* c, size_t size) {
if (i < size) c[i] = a[i] + b[i];
}
TEMPLATE_TEST_CASE("Add Kernel", "[kernel][add]", int, long, float, long long, double) {
TEMPLATE_TEST_CASE("ABM_AddKernel_MultiTypeMultiSize", "", int, long, float, long long, double) {
auto size = GENERATE(as<size_t>{}, 100, 500, 1000);
TestType *d_a, *d_b, *d_c;
auto res = hipMalloc(&d_a, sizeof(TestType) * size);
catch/ABM/CMakeLists.txt
+1
View File
@@ -0,0 +1 @@
add_subdirectory(AddKernels)
catch/CMakeLists.txt
+1
View File
@@ -66,6 +66,7 @@ add_custom_target(build_tests)
# Tests folder
add_subdirectory(unit)
add_subdirectory(ABM)
add_subdirectory(hipTestMain)
add_subdirectory(stress)
catch/README.md
+15
View File
@@ -71,3 +71,18 @@ Catch2 allows multiple ways in which you can debug the test case.
## External Libs being used
- [Catch2](https://github.com/catchorg/Catch2) - Testing framework
- [picojson](https://github.com/kazuho/picojson) - For config file parsing
# Testing Guidelines
Tests fall in 5 categories and its file name prefix are as follows:
- Unit tests (Prefix: Unit_\*API\*_\*Optional Scenario\*, example : Unit_hipMalloc_Negative or Unit_hipMalloc): Unit Tests are simplest test for an API, the target here is to test the API with different types of input and different ways of calling.
- Application Behavior Modelling tests (Prefix: ABM_\*Intent\*_\*Optional Scenario\*, example: ABM_ModuleLoadAndRun): ABM tests are used to model a specific use case of HIP APIs, either seen in a customer app or a general purpose app. It mimics the calling behavior seen in aforementioned app.
- Stress/Scale tests (Prefix: Stress_\*API\*_\*Intent\*_\*Optional Scenario\*, example: Stress_hipMemset_ExhaustVRAM): These tests are used to see the behavior of HIP APIs in edge scenarios, for example what happens when we have exhausted vram and do a hipMalloc or run many instances of same API in parallel.
- Multi Process tests (Prefix: MultiProc_\*API\*_\*Optional Scenario\*, example: MultiProc_hipIPCMemHandle_GetDataFromProc): These tests are multi process tests and will only run on linux. They are used to test HIP APIs in multi process environment
- Performance tests(Prefix: Perf_\*Intent\*_\*Optional Scenario\*, example: Perf_DispatchLatenc y): Performance tests are used to get results of HIP APIs.
General Guidelines:
- Do not use the catch2 tags. Tags wont be used for filtering
- Add as many INFO() as you can in tests which prints state of the t est, this will help the debugger when the test fails (INFO macro only prints when the test fails)
- Check return of each HIP API and fail whenever there is a misma tch with hipSuccess or hiprtcSuccess.
- Each Category of test will hav e its own exe and catch_discover_test macro will be called on it to discover its tests
- Optional Scenario in test names are optional. For example you can test all Scenarios of hipMalloc API in one file, you can name the file Unit_hipMalloc, if you are having a file just for negative scenarios you can name it as Unit_hipMalloc_Negative.
catch/hipTestMain/CMakeLists.txt
+14 -3
View File
@@ -9,9 +9,8 @@ else()
target_compile_options(UnitTests PUBLIC -std=c++17)
endif()
target_link_libraries(UnitTests PRIVATE DeviceLibs
target_link_libraries(UnitTests PRIVATE UnitDeviceTests
MemoryTest
Kernels
stdc++fs)
# Add AMD Only Tests
@@ -20,7 +19,19 @@ if(HIP_PLATFORM MATCHES "amd")
endif()
catch_discover_tests(UnitTests PROPERTIES SKIP_REGULAR_EXPRESSION "HIP_SKIP_THIS_TEST")
add_dependencies(build_tests UnitTests)
# ABM exe
add_executable(ABMTests EXCLUDE_FROM_ALL main.cc hip_test_context.cc)
if(HIP_PLATFORM MATCHES "amd")
set_property(TARGET ABMTests PROPERTY CXX_STANDARD 17)
else()
target_compile_options(ABMTests PUBLIC -std=c++17)
endif()
target_link_libraries(ABMTests PRIVATE ABMAddKernels)
catch_discover_tests(ABMTests PROPERTIES SKIP_REGULAR_EXPRESSION "HIP_SKIP_THIS_TEST")
add_dependencies(build_tests UnitTests ABMTests)
# Add Multiproc tests as seperate binary
if(UNIX AND HIP_PLATFORM MATCHES "amd")
catch/include/hip_test_checkers.hh
+12 -12
View File
@@ -75,13 +75,13 @@ bool initArraysForHost(T** A_h, T** B_h, T** C_h, size_t N, bool usePinnedHost =
if (usePinnedHost) {
if (A_h) {
HIPCHECK(hipHostMalloc((void**)A_h, Nbytes));
HIP_CHECK(hipHostMalloc((void**)A_h, Nbytes));
}
if (B_h) {
HIPCHECK(hipHostMalloc((void**)B_h, Nbytes));
HIP_CHECK(hipHostMalloc((void**)B_h, Nbytes));
}
if (C_h) {
HIPCHECK(hipHostMalloc((void**)C_h, Nbytes));
HIP_CHECK(hipHostMalloc((void**)C_h, Nbytes));
}
} else {
if (A_h) {
@@ -110,13 +110,13 @@ bool initArrays(T** A_d, T** B_d, T** C_d, T** A_h, T** B_h, T** C_h, size_t N,
size_t Nbytes = N * sizeof(T);
if (A_d) {
HIPCHECK(hipMalloc(A_d, Nbytes));
HIP_CHECK(hipMalloc(A_d, Nbytes));
}
if (B_d) {
HIPCHECK(hipMalloc(B_d, Nbytes));
HIP_CHECK(hipMalloc(B_d, Nbytes));
}
if (C_d) {
HIPCHECK(hipMalloc(C_d, Nbytes));
HIP_CHECK(hipMalloc(C_d, Nbytes));
}
return initArraysForHost(A_h, B_h, C_h, N, usePinnedHost);
@@ -125,13 +125,13 @@ bool initArrays(T** A_d, T** B_d, T** C_d, T** A_h, T** B_h, T** C_h, size_t N,
template <typename T> bool freeArraysForHost(T* A_h, T* B_h, T* C_h, bool usePinnedHost) {
if (usePinnedHost) {
if (A_h) {
HIPCHECK(hipHostFree(A_h));
HIP_CHECK(hipHostFree(A_h));
}
if (B_h) {
HIPCHECK(hipHostFree(B_h));
HIP_CHECK(hipHostFree(B_h));
}
if (C_h) {
HIPCHECK(hipHostFree(C_h));
HIP_CHECK(hipHostFree(C_h));
}
} else {
if (A_h) {
@@ -150,13 +150,13 @@ template <typename T> bool freeArraysForHost(T* A_h, T* B_h, T* C_h, bool usePin
template <typename T>
bool freeArrays(T* A_d, T* B_d, T* C_d, T* A_h, T* B_h, T* C_h, bool usePinnedHost) {
if (A_d) {
HIPCHECK(hipFree(A_d));
HIP_CHECK(hipFree(A_d));
}
if (B_d) {
HIPCHECK(hipFree(B_d));
HIP_CHECK(hipFree(B_d));
}
if (C_d) {
HIPCHECK(hipFree(C_d));
HIP_CHECK(hipFree(C_d));
}
return freeArraysForHost(A_h, B_h, C_h, usePinnedHost);
catch/include/hip_test_common.hh
+32 -1
View File
@@ -1,10 +1,32 @@
/*
Copyright (c) 2021-present 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.
*/
#pragma once
#include "hip_test_context.hh"
#include <catch.hpp>
#define HIP_PRINT_STATUS(status) INFO(hipGetErrorName(status) << " at line: " << __LINE__);
#define HIPCHECK(error) \
#define HIP_CHECK(error) \
{ \
hipError_t localError = error; \
if ((localError != hipSuccess) && (localError != hipErrorPeerAccessAlreadyEnabled)) { \
@@ -14,3 +36,12 @@
} \
}
// Although its assert, it will be evaluated at runtime
#define HIP_ASSERT(x) \
{ REQUIRE((x)); }
// Utility Functions
namespace HipTest {
int getDeviceCount();
}
catch/include/hip_test_helper.hh
+32
View File
@@ -0,0 +1,32 @@
/*
Copyright (c) 2021-present 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.
*/
#pragma once
#include "hip_test_common.hh"
namespace HipTest {
static inline int getGeviceCount() {
int dev = 0;
HIPCHECK(hipGetDeviceCount(&dev));
return dev;
}
} // namespace HipTest
catch/include/hip_test_kernels.hh
+23 -1
View File
@@ -1,3 +1,25 @@
/*
Copyright (c) 2021-present 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.
*/
#pragma once
#include <hip/hip_runtime.h>
@@ -59,4 +81,4 @@ template <typename T> __global__ void memsetReverse(T* C_d, T val, int64_t NELEM
C_d[i] = val;
}
}
} // namespace HipTest
} // namespace HipTest
catch/multiproc/hipMallocConcurrency.cc
+10 -10
View File
@@ -20,9 +20,9 @@ unsigned threadsPerBlock = 256;
unsigned setNumBlocks(unsigned blocksPerCU, unsigned threadsPerBlock, size_t N) {
int device;
HIPCHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDevice(&device));
hipDeviceProp_t props;
HIPCHECK(hipGetDeviceProperties(&props, device));
HIP_CHECK(hipGetDeviceProperties(&props, device));
unsigned blocks = props.multiProcessorCount * blocksPerCU;
if (blocks * threadsPerBlock > N) {
@@ -43,20 +43,20 @@ bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) {
size_t prevAvl, prevTot, curAvl, curTot;
bool TestPassed = true;
HIPCHECK(hipSetDevice(gpu));
HIPCHECK(hipMemGetInfo(&prevAvl, &prevTot));
HIP_CHECK(hipSetDevice(gpu));
HIP_CHECK(hipMemGetInfo(&prevAvl, &prevTot));
printf("tgs allocating..\n");
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
unsigned blocks = setNumBlocks(blocksPerCU, threadsPerBlock, N);
HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
HIPCHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0,
static_cast<const int*>(A_d), static_cast<const int*>(B_d), C_d, N);
HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
if (!HipTest::checkVectorADD(A_h, B_h, C_h, N)) {
printf("Validation PASSED for gpu %d from pid %d\n", gpu, getpid());
@@ -66,7 +66,7 @@ bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) {
}
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIPCHECK(hipMemGetInfo(&curAvl, &curTot));
HIP_CHECK(hipMemGetInfo(&curAvl, &curTot));
if (!concurOnOneGPU && (prevAvl != curAvl || prevTot != curTot)) {
// In concurrent calls on one GPU, we cannot verify leaking in this way
@@ -116,7 +116,7 @@ void getDeviceCount1(int* pdevCnt) {
// writing only, no need for read-descriptor
close(fd[0]);
HIPCHECK(hipGetDeviceCount(&devCnt));
HIP_CHECK(hipGetDeviceCount(&devCnt));
// send the value on the write-descriptor:
write(fd[1], &devCnt, sizeof(devCnt));
@@ -129,7 +129,7 @@ void getDeviceCount1(int* pdevCnt) {
}
#else
HIPCHECK(hipGetDeviceCount(pdevCnt));
HIP_CHECK(hipGetDeviceCount(pdevCnt));
#endif
}
#endif
catch/unit/CMakeLists.txt
-1
View File
@@ -1,5 +1,4 @@
add_subdirectory(memory)
add_subdirectory(deviceLib)
add_subdirectory(kernels)
# Disable Saxpy test temporarily to see if CI Passes
# add_subdirectory(rtc)
catch/unit/deviceLib/CMakeLists.txt
+16 -2
View File
@@ -1,20 +1,34 @@
# Common Tests - Test independent of all platforms
set(TEST_SRC
floatMath.cc
anyAll.cc
ballot.cc
clz.cc
ffs.cc
funnelshift.cc
brev.cc
)
# AMD only tests
set(AMD_TEST_SRC
vectorTypesDevice.cc
bitExtract.cc
bitInsert.cc
floatTM.cc
)
if(HIP_PLATFORM MATCHES "amd")
set(TEST_SRC ${TEST_SRC} ${AMD_TEST_SRC})
set_source_files_properties(floatTM.cc PROPERTIES COMPILE_FLAGS -std=c++17)
endif()
# Create shared lib of all tests
add_library(DeviceLibs SHARED EXCLUDE_FROM_ALL ${TEST_SRC})
add_library(UnitDeviceTests SHARED EXCLUDE_FROM_ALL ${TEST_SRC})
if(HIP_PLATFORM MATCHES "nvidia")
target_compile_options(UnitDeviceTests PUBLIC --Wno-deprecated-declarations)
endif()
# Add dependency on build_tests to build it on this custom target
add_dependencies(build_tests DeviceLibs)
add_dependencies(build_tests UnitDeviceTests)
catch/unit/deviceLib/anyAll.cc
+83
View File
@@ -0,0 +1,83 @@
/*
Copyright (c) 2021 - present 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 <hip_test_common.hh>
__global__ void warpvote(int* device_any, int* device_all, int pshift) {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
device_any[threadIdx.x >> pshift] = __any(tid - 77);
device_all[threadIdx.x >> pshift] = __all(tid - 77);
}
TEST_CASE("Unit_AnyAll_CompileTest") {
int warpSize, pshift;
hipDeviceProp_t devProp;
hipGetDeviceProperties(&devProp, 0);
warpSize = devProp.warpSize;
int w = warpSize;
pshift = 0;
while (w >>= 1) ++pshift;
INFO("WarpSize:: " << warpSize << " pShift: " << pshift);
int anycount = 0;
int allcount = 0;
int Num_Threads_per_Block = 1024;
int Num_Blocks_per_Grid = 1;
int Num_Warps_per_Grid = (Num_Threads_per_Block * Num_Blocks_per_Grid) / warpSize;
int* host_any = (int*)malloc(Num_Warps_per_Grid * sizeof(int));
int* host_all = (int*)malloc(Num_Warps_per_Grid * sizeof(int));
int* device_any;
int* device_all;
HIP_CHECK(hipMalloc((void**)&device_any, Num_Warps_per_Grid * sizeof(int)));
HIP_CHECK(hipMalloc((void**)&device_all, Num_Warps_per_Grid * sizeof(int)));
for (int i = 0; i < Num_Warps_per_Grid; i++) {
host_any[i] = 0;
host_all[i] = 0;
}
HIP_CHECK(hipMemcpy(device_any, host_any, sizeof(int), hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(device_all, host_all, sizeof(int), hipMemcpyHostToDevice));
hipLaunchKernelGGL(warpvote, dim3(Num_Blocks_per_Grid), dim3(Num_Threads_per_Block), 0, 0,
device_any, device_all, pshift);
HIP_CHECK(
hipMemcpy(host_any, device_any, Num_Warps_per_Grid * sizeof(int), hipMemcpyDeviceToHost));
HIP_CHECK(
hipMemcpy(host_all, device_all, Num_Warps_per_Grid * sizeof(int), hipMemcpyDeviceToHost));
for (int i = 0; i < Num_Warps_per_Grid; i++) {
INFO("Warp Number: " << i << " __any: " << host_any[i] << " __all: " << host_all[i]);
if (host_all[i] != 1) ++allcount;
if (host_any[i] != 1) ++anycount;
}
HIP_CHECK(hipFree(device_any));
HIP_CHECK(hipFree(device_all));
REQUIRE(anycount == 0);
REQUIRE(allcount == 1);
}
catch/unit/deviceLib/ballot.cc
+85
View File
@@ -0,0 +1,85 @@
/*
Copyright (c) 2021 - present 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 <hip_test_common.hh>
#include <hip/device_functions.h>
__global__ void gpu_ballot(unsigned int* device_ballot, unsigned Num_Warps_per_Block,
unsigned pshift) {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
const unsigned int warp_num = threadIdx.x >> pshift;
// TODO try to remove this
#ifdef __HIP_PLATFORM_AMD__
atomicAdd(&device_ballot[warp_num + blockIdx.x * Num_Warps_per_Block],
__popcll(__ballot(tid - 245)));
#else
atomicAdd(&device_ballot[warp_num + blockIdx.x * Num_Warps_per_Block],
__popc(__ballot(tid - 245)));
#endif
}
TEST_CASE("Unit_ballot") {
unsigned warpSize, pshift;
hipDeviceProp_t devProp;
HIP_CHECK(hipGetDeviceProperties(&devProp, 0));
warpSize = devProp.warpSize;
int w = warpSize;
pshift = 0;
while (w >>= 1) ++pshift;
unsigned int Num_Threads_per_Block = 512;
unsigned int Num_Blocks_per_Grid = 1;
unsigned int Num_Warps_per_Block = Num_Threads_per_Block / warpSize;
unsigned int Num_Warps_per_Grid = (Num_Threads_per_Block * Num_Blocks_per_Grid) / warpSize;
unsigned int* host_ballot = (unsigned int*)malloc(Num_Warps_per_Grid * sizeof(unsigned int));
unsigned int* device_ballot;
HIP_CHECK(hipMalloc((void**)&device_ballot, Num_Warps_per_Grid * sizeof(unsigned int)));
int divergent_count = 0;
for (unsigned i = 0; i < Num_Warps_per_Grid; i++) host_ballot[i] = 0;
HIP_CHECK(hipMemcpy(device_ballot, host_ballot, Num_Warps_per_Grid * sizeof(unsigned int),
hipMemcpyHostToDevice));
hipLaunchKernelGGL(gpu_ballot, dim3(Num_Blocks_per_Grid), dim3(Num_Threads_per_Block), 0, 0,
device_ballot, Num_Warps_per_Block, pshift);
HIP_CHECK(hipMemcpy(host_ballot, device_ballot, Num_Warps_per_Grid * sizeof(unsigned int),
hipMemcpyDeviceToHost));
for (unsigned i = 0; i < Num_Warps_per_Grid; i++) {
if ((host_ballot[i] == 0) || (host_ballot[i] / warpSize == warpSize)) {
INFO("Warp: " << i << " is convergent - predicate true for " << (host_ballot[i] / warpSize)
<< " threads");
} else {
INFO("Warp: " << i << " is divergent - predicate true for " << (host_ballot[i] / warpSize)
<< " threads");
divergent_count++;
}
}
HIP_CHECK(hipFree(device_ballot));
REQUIRE(divergent_count == 1);
}
catch/unit/deviceLib/bitExtract.cc
+194
View File
@@ -0,0 +1,194 @@
/*
Copyright (c) 2021 - present 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 <hip_test_common.hh>
#include <hip/device_functions.h>
#include <assert.h>
#include <stdio.h>
#include <algorithm>
#include <stdlib.h>
#include <iostream>
#include <random>
// CPU implementation of bitextract
template <typename T> T bit_extract(T src0, unsigned int src1, unsigned int src2) {
unsigned int bits = sizeof(T) * 8;
T offset = src1 & (bits - 1);
T width = src2 & (bits - 1);
if (width == 0) {
return 0;
} else {
return (src0 << (bits - width - offset)) >> (bits - width);
}
}
__global__ void HIP_kernel(unsigned int* out32, unsigned int* in32_0, unsigned int* in32_1,
unsigned int* in32_2, unsigned long long int* out64,
unsigned long long int* in64_0, unsigned int* in64_1,
unsigned int* in64_2) {
int x = blockDim.x * blockIdx.x + threadIdx.x;
out32[x] = __bitextract_u32(in32_0[x], in32_1[x], in32_2[x]);
out64[x] = __bitextract_u64(in64_0[x], in64_1[x], in64_2[x]);
}
TEST_CASE("Unit_bitExtract") {
using namespace std;
unsigned int* hostOut32;
unsigned int* hostSrc032;
unsigned int* hostSrc132;
unsigned int* hostSrc232;
unsigned long long int* hostOut64;
unsigned long long int* hostSrc064;
unsigned int* hostSrc164;
unsigned int* hostSrc264;
unsigned int* deviceOut32;
unsigned int* deviceSrc032;
unsigned int* deviceSrc132;
unsigned int* deviceSrc232;
unsigned long long int* deviceOut64;
unsigned long long int* deviceSrc064;
unsigned int* deviceSrc164;
unsigned int* deviceSrc264;
hipDeviceProp_t devProp;
HIP_CHECK(hipGetDeviceProperties(&devProp, 0));
INFO("System minor : " << devProp.minor);
INFO("System major : " << devProp.major);
INFO("agent prop name : " << devProp.name);
INFO("hip Device prop succeeded");
unsigned int wave_size = devProp.warpSize;
unsigned int num_waves_per_block = 2;
unsigned int num_threads_per_block = wave_size * num_waves_per_block;
unsigned int num_blocks = 2;
unsigned int NUM = num_threads_per_block * num_blocks;
unsigned i;
int errors;
hostOut32 = (unsigned int*)malloc(NUM * sizeof(unsigned int));
hostSrc032 = (unsigned int*)malloc(NUM * sizeof(unsigned int));
hostSrc132 = (unsigned int*)malloc(NUM * sizeof(unsigned int));
hostSrc232 = (unsigned int*)malloc(NUM * sizeof(unsigned int));
hostOut64 = (unsigned long long int*)malloc(NUM * sizeof(unsigned long long int));
hostSrc064 = (unsigned long long int*)malloc(NUM * sizeof(unsigned long long int));
hostSrc164 = (unsigned int*)malloc(NUM * sizeof(unsigned int));
hostSrc264 = (unsigned int*)malloc(NUM * sizeof(unsigned int));
// initialize the input data
std::random_device rd;
std::uniform_int_distribution<uint32_t> uint32_src0_dist;
std::uniform_int_distribution<uint32_t> uint32_src12_dist(0, 31);
std::uniform_int_distribution<uint64_t> uint64_src0_dist;
std::uniform_int_distribution<uint32_t> uint64_src12_dist(0, 63);
for (i = 0; i < NUM; i++) {
hostOut32[i] = 0;
hostSrc032[i] = uint32_src0_dist(rd);
hostSrc132[i] = uint32_src12_dist(rd);
hostSrc232[i] = uint32_src12_dist(rd);
if (hostSrc132[i] + hostSrc232[i] > 32) hostSrc232[i] = 32 - hostSrc132[i];
hostOut64[i] = 0;
hostSrc064[i] = uint64_src0_dist(rd);
hostSrc164[i] = uint64_src12_dist(rd);
hostSrc264[i] = uint64_src12_dist(rd);
}
HIP_CHECK(hipMalloc((void**)&deviceOut32, NUM * sizeof(unsigned int)));
HIP_CHECK(hipMalloc((void**)&deviceSrc032, NUM * sizeof(unsigned int)));
HIP_CHECK(hipMalloc((void**)&deviceSrc132, NUM * sizeof(unsigned int)));
HIP_CHECK(hipMalloc((void**)&deviceSrc232, NUM * sizeof(unsigned int)));
HIP_CHECK(hipMalloc((void**)&deviceOut64, NUM * sizeof(unsigned long long int)));
HIP_CHECK(hipMalloc((void**)&deviceSrc064, NUM * sizeof(unsigned long long int)));
HIP_CHECK(hipMalloc((void**)&deviceSrc164, NUM * sizeof(unsigned int)));
HIP_CHECK(hipMalloc((void**)&deviceSrc264, NUM * sizeof(unsigned int)));
HIP_CHECK(
hipMemcpy(deviceSrc032, hostSrc032, NUM * sizeof(unsigned int), hipMemcpyHostToDevice));
HIP_CHECK(
hipMemcpy(deviceSrc132, hostSrc132, NUM * sizeof(unsigned int), hipMemcpyHostToDevice));
HIP_CHECK(
hipMemcpy(deviceSrc232, hostSrc232, NUM * sizeof(unsigned int), hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(deviceSrc064, hostSrc064, NUM * sizeof(unsigned long long int),
hipMemcpyHostToDevice));
HIP_CHECK(
hipMemcpy(deviceSrc164, hostSrc164, NUM * sizeof(unsigned int), hipMemcpyHostToDevice));
HIP_CHECK(
hipMemcpy(deviceSrc264, hostSrc264, NUM * sizeof(unsigned int), hipMemcpyHostToDevice));
hipLaunchKernelGGL(HIP_kernel, dim3(num_blocks), dim3(num_threads_per_block), 0, 0, deviceOut32,
deviceSrc032, deviceSrc132, deviceSrc232, deviceOut64, deviceSrc064,
deviceSrc164, deviceSrc264);
HIP_CHECK(hipMemcpy(hostOut32, deviceOut32, NUM * sizeof(unsigned int), hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpy(hostOut64, deviceOut64, NUM * sizeof(unsigned long long int),
hipMemcpyDeviceToHost));
// verify the results
errors = 0;
for (i = 0; i < NUM; i++) {
if (hostOut32[i] != bit_extract<uint32_t>(hostSrc032[i], hostSrc132[i], hostSrc232[i])) {
errors++;
INFO("device: " << hostOut32[i] << " host: "
<< bit_extract<uint32_t>(hostSrc032[i], hostSrc132[i], hostSrc232[i]) << " "
<< hostSrc032[i] << " " << hostSrc132[i] << " " << hostSrc232[i] << "\n");
}
}
for (i = 0; i < NUM; i++) {
if (hostOut64[i] != bit_extract<uint64_t>(hostSrc064[i], hostSrc164[i], hostSrc264[i])) {
errors++;
INFO("device: " << hostOut64[i] << " host: "
<< bit_extract<uint64_t>(hostSrc064[i], hostSrc164[i], hostSrc264[i]) << " "
<< hostSrc064[i] << " " << hostSrc164[i] << " " << hostSrc264[i] << "\n");
}
}
HIP_CHECK(hipFree(deviceOut32));
HIP_CHECK(hipFree(deviceSrc032));
HIP_CHECK(hipFree(deviceSrc132));
HIP_CHECK(hipFree(deviceSrc232));
HIP_CHECK(hipFree(deviceOut64));
HIP_CHECK(hipFree(deviceSrc064));
HIP_CHECK(hipFree(deviceSrc164));
HIP_CHECK(hipFree(deviceSrc264));
free(hostOut32);
free(hostSrc032);
free(hostSrc132);
free(hostSrc232);
free(hostOut64);
free(hostSrc064);
free(hostSrc164);
free(hostSrc264);
REQUIRE(errors == 0);
}
catch/unit/deviceLib/bitInsert.cc
+216
View File
@@ -0,0 +1,216 @@
/*
Copyright (c) 2021 - present 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 <hip_test_common.hh>
#include <hip/device_functions.h>
#include <assert.h>
#include <stdio.h>
#include <algorithm>
#include <stdlib.h>
#include <iostream>
#include <random>
// CPU implementation of bitinsert
template <typename T> T bit_insert(T src0, T src1, unsigned int src2, unsigned int src3) {
unsigned int bits = sizeof(T) * 8;
T offset = src2 & (bits - 1);
T width = src3 & (bits - 1);
T mask = (((T)1) << width) - 1;
return ((src0 & ~(mask << offset)) | ((src1 & mask) << offset));
}
__global__ void HIP_kernel(unsigned int* out32, unsigned int* in32_0, unsigned int* in32_1,
unsigned int* in32_2, unsigned int* in32_3,
unsigned long long int* out64, unsigned long long int* in64_0,
unsigned long long int* in64_1, unsigned int* in64_2,
unsigned int* in64_3) {
int x = blockDim.x * blockIdx.x + threadIdx.x;
out32[x] = __bitinsert_u32(in32_0[x], in32_1[x], in32_2[x], in32_3[x]);
out64[x] = __bitinsert_u64(in64_0[x], in64_1[x], in64_2[x], in64_3[x]);
}
TEST_CASE("Unit_bitInsert") {
using namespace std;
unsigned int* hostOut32;
unsigned int* hostSrc032;
unsigned int* hostSrc132;
unsigned int* hostSrc232;
unsigned int* hostSrc332;
unsigned long long int* hostOut64;
unsigned long long int* hostSrc064;
unsigned long long int* hostSrc164;
unsigned int* hostSrc264;
unsigned int* hostSrc364;
unsigned int* deviceOut32;
unsigned int* deviceSrc032;
unsigned int* deviceSrc132;
unsigned int* deviceSrc232;
unsigned int* deviceSrc332;
unsigned long long int* deviceOut64;
unsigned long long int* deviceSrc064;
unsigned long long int* deviceSrc164;
unsigned int* deviceSrc264;
unsigned int* deviceSrc364;
hipDeviceProp_t devProp;
HIP_CHECK(hipGetDeviceProperties(&devProp, 0));
INFO("System minor : " << devProp.minor);
INFO("System major : " << devProp.major);
INFO("agent prop name : " << devProp.name);
INFO("hip Device prop succeeded");
unsigned int wave_size = devProp.warpSize;
unsigned int num_waves_per_block = 2;
unsigned int num_threads_per_block = wave_size * num_waves_per_block;
unsigned int num_blocks = 2;
unsigned int NUM = num_threads_per_block * num_blocks;
unsigned i;
int errors;
hostOut32 = (unsigned int*)malloc(NUM * sizeof(unsigned int));
hostSrc032 = (unsigned int*)malloc(NUM * sizeof(unsigned int));
hostSrc132 = (unsigned int*)malloc(NUM * sizeof(unsigned int));
hostSrc232 = (unsigned int*)malloc(NUM * sizeof(unsigned int));
hostSrc332 = (unsigned int*)malloc(NUM * sizeof(unsigned int));
hostOut64 = (unsigned long long int*)malloc(NUM * sizeof(unsigned long long int));
hostSrc064 = (unsigned long long int*)malloc(NUM * sizeof(unsigned long long int));
hostSrc164 = (unsigned long long int*)malloc(NUM * sizeof(unsigned long long int));
hostSrc264 = (unsigned int*)malloc(NUM * sizeof(unsigned int));
hostSrc364 = (unsigned int*)malloc(NUM * sizeof(unsigned int));
// initialize the input data
std::random_device rd;
std::uniform_int_distribution<uint32_t> uint32_src01_dist;
std::uniform_int_distribution<uint32_t> uint32_src23_dist(0, 31);
std::uniform_int_distribution<uint64_t> uint64_src01_dist;
std::uniform_int_distribution<uint32_t> uint64_src23_dist(0, 63);
for (i = 0; i < NUM; i++) {
hostOut32[i] = 0;
hostSrc032[i] = uint32_src01_dist(rd);
hostSrc132[i] = uint32_src01_dist(rd);
hostSrc232[i] = uint32_src23_dist(rd);
hostSrc232[i] = uint32_src23_dist(rd);
hostOut64[i] = 0;
hostSrc064[i] = uint64_src01_dist(rd);
hostSrc164[i] = uint64_src01_dist(rd);
hostSrc264[i] = uint64_src23_dist(rd);
hostSrc264[i] = uint64_src23_dist(rd);
}
HIP_CHECK(hipMalloc((void**)&deviceOut32, NUM * sizeof(unsigned int)));
HIP_CHECK(hipMalloc((void**)&deviceSrc032, NUM * sizeof(unsigned int)));
HIP_CHECK(hipMalloc((void**)&deviceSrc132, NUM * sizeof(unsigned int)));
HIP_CHECK(hipMalloc((void**)&deviceSrc232, NUM * sizeof(unsigned int)));
HIP_CHECK(hipMalloc((void**)&deviceSrc332, NUM * sizeof(unsigned int)));
HIP_CHECK(hipMalloc((void**)&deviceOut64, NUM * sizeof(unsigned long long int)));
HIP_CHECK(hipMalloc((void**)&deviceSrc064, NUM * sizeof(unsigned long long int)));
HIP_CHECK(hipMalloc((void**)&deviceSrc164, NUM * sizeof(unsigned long long int)));
HIP_CHECK(hipMalloc((void**)&deviceSrc264, NUM * sizeof(unsigned int)));
HIP_CHECK(hipMalloc((void**)&deviceSrc364, NUM * sizeof(unsigned int)));
HIP_CHECK(
hipMemcpy(deviceSrc032, hostSrc032, NUM * sizeof(unsigned int), hipMemcpyHostToDevice));
HIP_CHECK(
hipMemcpy(deviceSrc132, hostSrc132, NUM * sizeof(unsigned int), hipMemcpyHostToDevice));
HIP_CHECK(
hipMemcpy(deviceSrc232, hostSrc232, NUM * sizeof(unsigned int), hipMemcpyHostToDevice));
HIP_CHECK(
hipMemcpy(deviceSrc332, hostSrc332, NUM * sizeof(unsigned int), hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(deviceSrc064, hostSrc064, NUM * sizeof(unsigned long long int),
hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(deviceSrc164, hostSrc164, NUM * sizeof(unsigned long long int),
hipMemcpyHostToDevice));
HIP_CHECK(
hipMemcpy(deviceSrc264, hostSrc264, NUM * sizeof(unsigned int), hipMemcpyHostToDevice));
HIP_CHECK(
hipMemcpy(deviceSrc364, hostSrc364, NUM * sizeof(unsigned int), hipMemcpyHostToDevice));
hipLaunchKernelGGL(HIP_kernel, dim3(num_blocks), dim3(num_threads_per_block), 0, 0, deviceOut32,
deviceSrc032, deviceSrc132, deviceSrc232, deviceSrc332, deviceOut64,
deviceSrc064, deviceSrc164, deviceSrc264, deviceSrc364);
HIP_CHECK(hipMemcpy(hostOut32, deviceOut32, NUM * sizeof(unsigned int), hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpy(hostOut64, deviceOut64, NUM * sizeof(unsigned long long int),
hipMemcpyDeviceToHost));
// verify the results
errors = 0;
for (i = 0; i < NUM; i++) {
if (hostOut32[i] !=
bit_insert<uint32_t>(hostSrc032[i], hostSrc132[i], hostSrc232[i], hostSrc332[i])) {
errors++;
INFO("device: " << hostOut32[i] << " host: "
<< bit_insert<uint32_t>(hostSrc032[i], hostSrc132[i], hostSrc232[i],
hostSrc332[i])
<< " " << hostSrc032[i] << " " << hostSrc132[i] << " " << hostSrc232[i] << " "
<< hostSrc332[i] << "\n");
}
}
for (i = 0; i < NUM; i++) {
if (hostOut64[i] !=
bit_insert<uint64_t>(hostSrc064[i], hostSrc164[i], hostSrc264[i], hostSrc364[i])) {
errors++;
INFO("device: " << hostOut64[i] << " host: "
<< bit_insert<uint64_t>(hostSrc064[i], hostSrc164[i], hostSrc264[i],
hostSrc364[i])
<< " " << hostSrc064[i] << " " << hostSrc164[i] << " " << hostSrc264[i] << " "
<< hostSrc364[i] << "\n");
}
}
HIP_CHECK(hipFree(deviceOut32));
HIP_CHECK(hipFree(deviceSrc032));
HIP_CHECK(hipFree(deviceSrc132));
HIP_CHECK(hipFree(deviceSrc232));
HIP_CHECK(hipFree(deviceSrc332));
HIP_CHECK(hipFree(deviceOut64));
HIP_CHECK(hipFree(deviceSrc064));
HIP_CHECK(hipFree(deviceSrc164));
HIP_CHECK(hipFree(deviceSrc264));
HIP_CHECK(hipFree(deviceSrc364));
free(hostOut32);
free(hostSrc032);
free(hostSrc132);
free(hostSrc232);
free(hostSrc332);
free(hostOut64);
free(hostSrc064);
free(hostSrc164);
free(hostSrc264);
free(hostSrc364);
REQUIRE(errors == 0);
}
catch/unit/deviceLib/brev.cc
+150
View File
@@ -0,0 +1,150 @@
/*
Copyright (c) 2021 - present 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 <hip_test_common.hh>
#include <hip/device_functions.h>
#include <assert.h>
#include <stdio.h>
#include <algorithm>
#include <stdlib.h>
#include <iostream>
#define WIDTH 32
#define HEIGHT 32
#define NUM (WIDTH * HEIGHT)
#define THREADS_PER_BLOCK_X 8
#define THREADS_PER_BLOCK_Y 8
#define THREADS_PER_BLOCK_Z 1
// CPU implementation of bitreverse
template <typename T> T bitreverse(T num) {
T count = sizeof(num) * 8 - 1;
T reverse_num = num;
num >>= 1;
while (num) {
reverse_num <<= 1;
reverse_num |= num & 1;
num >>= 1;
count--;
}
reverse_num <<= count;
return reverse_num;
}
__global__ void HIP_kernel(unsigned int* a, unsigned int* b, unsigned long long int* c,
unsigned long long int* d, int width, int height) {
int x = blockDim.x * blockIdx.x + threadIdx.x;
int y = blockDim.y * blockIdx.y + threadIdx.y;
int i = y * width + x;
if (i < (width * height)) {
a[i] = __brev(b[i]);
c[i] = __brevll(d[i]);
}
}
TEST_CASE("Unit_brev") {
using namespace std;
unsigned int* hostA;
unsigned int* hostB;
unsigned long long int* hostC;
unsigned long long int* hostD;
unsigned int* deviceA;
unsigned int* deviceB;
unsigned long long int* deviceC;
unsigned long long int* deviceD;
hipDeviceProp_t devProp;
HIP_CHECK(hipGetDeviceProperties(&devProp, 0));
INFO("System minor : " << devProp.minor);
INFO("System major : " << devProp.major);
INFO("agent prop name : " << devProp.name);
INFO("hip Device prop succeeded");
int i;
int errors;
hostA = (unsigned int*)malloc(NUM * sizeof(unsigned int));
hostB = (unsigned int*)malloc(NUM * sizeof(unsigned int));
hostC = (unsigned long long int*)malloc(NUM * sizeof(unsigned long long int));
hostD = (unsigned long long int*)malloc(NUM * sizeof(unsigned long long int));
// initialize the input data
for (i = 0; i < NUM; i++) {
hostB[i] = i;
hostD[i] = i;
}
HIP_CHECK(hipMalloc((void**)&deviceA, NUM * sizeof(unsigned int)));
HIP_CHECK(hipMalloc((void**)&deviceB, NUM * sizeof(unsigned int)));
HIP_CHECK(hipMalloc((void**)&deviceC, NUM * sizeof(unsigned long long int)));
HIP_CHECK(hipMalloc((void**)&deviceD, NUM * sizeof(unsigned long long int)));
HIP_CHECK(hipMemcpy(deviceB, hostB, NUM * sizeof(unsigned int), hipMemcpyHostToDevice));
HIP_CHECK(
hipMemcpy(deviceD, hostD, NUM * sizeof(unsigned long long int), hipMemcpyHostToDevice));
hipLaunchKernelGGL(HIP_kernel, dim3(WIDTH / THREADS_PER_BLOCK_X, HEIGHT / THREADS_PER_BLOCK_Y),
dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), 0, 0, deviceA, deviceB,
deviceC, deviceD, WIDTH, HEIGHT);
HIP_CHECK(hipMemcpy(hostA, deviceA, NUM * sizeof(unsigned int), hipMemcpyDeviceToHost));
HIP_CHECK(
hipMemcpy(hostC, deviceC, NUM * sizeof(unsigned long long int), hipMemcpyDeviceToHost));
// verify the results
errors = 0;
for (i = 0; i < NUM; i++) {
if (hostA[i] != bitreverse(hostB[i])) {
errors++;
}
}
for (i = 0; i < NUM; i++) {
if (hostC[i] != bitreverse(hostD[i])) {
errors++;
}
}
HIP_CHECK(hipFree(deviceA));
HIP_CHECK(hipFree(deviceB));
HIP_CHECK(hipFree(deviceC));
HIP_CHECK(hipFree(deviceD));
free(hostA);
free(hostB);
free(hostC);
free(hostD);
REQUIRE(errors == 0);
}
catch/unit/deviceLib/clz.cc
+172
View File
@@ -0,0 +1,172 @@
/*
Copyright (c) 2021 - present 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 <hip_test_common.hh>
#include <hip/device_functions.h>
#include <assert.h>
#include <stdio.h>
#include <algorithm>
#include <stdlib.h>
#include <iostream>
#define WIDTH 8
#define HEIGHT 8
#define NUM (WIDTH * HEIGHT)
#define THREADS_PER_BLOCK_X 8
#define THREADS_PER_BLOCK_Y 8
#define THREADS_PER_BLOCK_Z 1
unsigned int firstbit_u32(unsigned int a) {
if (a == 0) {
return 32;
}
unsigned int pos = 0;
while ((int)a > 0) {
a <<= 1;
pos++;
}
return pos;
}
unsigned int firstbit_u64(unsigned long long int a) {
if (a == 0) {
return 64;
}
unsigned int pos = 0;
while ((long long int)a > 0) {
a <<= 1;
pos++;
}
return pos;
}
// Check implicit conversion will not cause ambiguity.
__device__ void test_ambiguity() {
short s = 1;
unsigned short us = 2;
float f = 3;
int i = 4;
unsigned int ui = 5;
__clz(f);
__clz(s);
__clz(us);
__clzll(f);
__clzll(i);
__clzll(ui);
}
__global__ void clz_HIP_kernel(unsigned int* a, unsigned int* b, unsigned int* c,
unsigned long long int* d, int width, int height) {
int x = blockDim.x * blockIdx.x + threadIdx.x;
int y = blockDim.y * blockIdx.y + threadIdx.y;
int i = y * width + x;
if (i < (width * height)) {
a[i] = __clz(b[i]);
c[i] = __clzll(d[i]);
}
}
TEST_CASE("Unit_clz") {
using namespace std;
unsigned int* hostA;
unsigned int* hostB;
unsigned int* hostC;
unsigned long long int* hostD;
unsigned int* deviceA;
unsigned int* deviceB;
unsigned int* deviceC;
unsigned long long int* deviceD;
hipDeviceProp_t devProp;
HIP_CHECK(hipGetDeviceProperties(&devProp, 0));
INFO("System minor : " << devProp.minor);
INFO("System major : " << devProp.major);
INFO("agent prop name : " << devProp.name);
INFO("hip Device prop succeeded");
unsigned int i;
int errors;
hostA = (unsigned int*)malloc(NUM * sizeof(unsigned int));
hostB = (unsigned int*)malloc(NUM * sizeof(unsigned int));
hostC = (unsigned int*)malloc(NUM * sizeof(unsigned int));
hostD = (unsigned long long int*)malloc(NUM * sizeof(unsigned long long int));
// initialize the input data
for (i = 0; i < NUM; i++) {
hostB[i] = 419430 * i;
hostD[i] = i;
}
HIP_CHECK(hipMalloc((void**)&deviceA, NUM * sizeof(unsigned int)));
HIP_CHECK(hipMalloc((void**)&deviceB, NUM * sizeof(unsigned int)));
HIP_CHECK(hipMalloc((void**)&deviceC, NUM * sizeof(unsigned int)));
HIP_CHECK(hipMalloc((void**)&deviceD, NUM * sizeof(unsigned long long int)));
HIP_CHECK(hipMemcpy(deviceB, hostB, NUM * sizeof(unsigned int), hipMemcpyHostToDevice));
HIP_CHECK(
hipMemcpy(deviceD, hostD, NUM * sizeof(unsigned long long int), hipMemcpyHostToDevice));
hipLaunchKernelGGL(clz_HIP_kernel,
dim3(WIDTH / THREADS_PER_BLOCK_X, HEIGHT / THREADS_PER_BLOCK_Y),
dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), 0, 0, deviceA, deviceB,
deviceC, deviceD, WIDTH, HEIGHT);
HIP_CHECK(hipMemcpy(hostA, deviceA, NUM * sizeof(unsigned int), hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpy(hostC, deviceC, NUM * sizeof(unsigned int), hipMemcpyDeviceToHost));
// verify the results
errors = 0;
for (i = 0; i < NUM; i++) {
if (hostA[i] != firstbit_u32(hostB[i])) {
INFO("Match Failed: " << hostA[i] << " - " << firstbit_u32(hostB[i]));
errors++;
}
}
for (i = 0; i < NUM; i++) {
if (hostC[i] != firstbit_u64(hostD[i])) {
INFO("Match Failed: " << hostC[i] << " - " << firstbit_u32(hostD[i]));
errors++;
}
}
HIP_CHECK(hipFree(deviceA));
HIP_CHECK(hipFree(deviceB));
HIP_CHECK(hipFree(deviceC));
HIP_CHECK(hipFree(deviceD));
free(hostA);
free(hostB);
free(hostC);
free(hostD);
REQUIRE(errors == 0);
}
catch/unit/deviceLib/ffs.cc
+146
View File
@@ -0,0 +1,146 @@
/*
Copyright (c) 2021 - present 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 <hip_test_common.hh>
#include <hip/device_functions.h>
#include <assert.h>
#include <stdio.h>
#include <algorithm>
#include <stdlib.h>
#include <iostream>
#define WIDTH 8
#define HEIGHT 8
#define NUM (WIDTH * HEIGHT)
#define THREADS_PER_BLOCK_X 8
#define THREADS_PER_BLOCK_Y 8
#define THREADS_PER_BLOCK_Z 1
template <typename T> unsigned lastbit(T a) {
if (a == 0) return 0;
unsigned pos = 1;
while ((a & 1) != 1) {
a >>= 1;
pos++;
}
return pos;
}
__global__ void ffs_HIP_kernel(unsigned int* a, unsigned int* b, unsigned int* c,
unsigned long long int* d, int width, int height) {
int x = blockDim.x * blockIdx.x + threadIdx.x;
int y = blockDim.y * blockIdx.y + threadIdx.y;
int i = y * width + x;
if (i < (width * height)) {
a[i] = __ffs(b[i]);
c[i] = __ffsll(d[i]);
}
}
TEST_CASE("Unit_ffs") {
using namespace std;
unsigned int* hostA;
unsigned int* hostB;
unsigned int* hostC;
unsigned long long int* hostD;
unsigned int* deviceA;
unsigned int* deviceB;
unsigned int* deviceC;
unsigned long long int* deviceD;
hipDeviceProp_t devProp;
HIP_CHECK(hipGetDeviceProperties(&devProp, 0));
INFO("System minor : " << devProp.minor);
INFO("System major : " << devProp.major);
INFO("agent prop name : " << devProp.name);
INFO("hip Device prop succeeded");
int i;
int errors;
hostA = (unsigned int*)malloc(NUM * sizeof(unsigned int));
hostB = (unsigned int*)malloc(NUM * sizeof(unsigned int));
hostC = (unsigned int*)malloc(NUM * sizeof(unsigned int));
hostD = (unsigned long long int*)malloc(NUM * sizeof(unsigned long long int));
// initialize the input data
for (i = 0; i < NUM; i++) {
hostB[i] = i;
hostD[i] = 1099511627776 + i;
}
HIP_CHECK(hipMalloc((void**)&deviceA, NUM * sizeof(unsigned int)));
HIP_CHECK(hipMalloc((void**)&deviceB, NUM * sizeof(unsigned int)));
HIP_CHECK(hipMalloc((void**)&deviceC, NUM * sizeof(unsigned int)));
HIP_CHECK(hipMalloc((void**)&deviceD, NUM * sizeof(unsigned long long int)));
HIP_CHECK(hipMemcpy(deviceB, hostB, NUM * sizeof(unsigned int), hipMemcpyHostToDevice));
HIP_CHECK(
hipMemcpy(deviceD, hostD, NUM * sizeof(unsigned long long int), hipMemcpyHostToDevice));
hipLaunchKernelGGL(ffs_HIP_kernel,
dim3(WIDTH / THREADS_PER_BLOCK_X, HEIGHT / THREADS_PER_BLOCK_Y),
dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), 0, 0, deviceA, deviceB,
deviceC, deviceD, WIDTH, HEIGHT);
HIP_CHECK(hipMemcpy(hostA, deviceA, NUM * sizeof(unsigned int), hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpy(hostC, deviceC, NUM * sizeof(unsigned int), hipMemcpyDeviceToHost));
// verify the results
errors = 0;
for (i = 0; i < NUM; i++) {
if (hostA[i] != lastbit(hostB[i])) {
INFO("Mismatch: " << hostA[i] << " - " << lastbit(hostB[i]));
errors++;
}
}
for (i = 0; i < NUM; i++) {
if (hostC[i] != lastbit(hostD[i])) {
INFO("Mismatch: " << hostC[i] << " - " << lastbit(hostD[i]));
errors++;
}
}
HIP_CHECK(hipFree(deviceA));
HIP_CHECK(hipFree(deviceB));
HIP_CHECK(hipFree(deviceC));
HIP_CHECK(hipFree(deviceD));
free(hostA);
free(hostB);
free(hostC);
free(hostD);
REQUIRE(errors == 0);
}
catch/unit/deviceLib/floatMath.cc
+25 -1
View File
@@ -1,3 +1,25 @@
/*
Copyright (c) 2021 - present 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 <hip_test_common.hh>
#define LEN 512
@@ -27,7 +49,7 @@ __global__ void floatMath(float* In, float* Out) {
Out[tid] = __tanf(Out[tid]);
}
TEST_CASE("FloatMathTest") {
TEST_CASE("Unit_deviceFunctions_CompileTest") {
float *Ind, *Outd;
auto res = hipMalloc((void**)&Ind, SIZE);
REQUIRE(res == hipSuccess);
@@ -38,4 +60,6 @@ TEST_CASE("FloatMathTest") {
REQUIRE(res == hipSuccess);
res = hipGetLastError();
REQUIRE(res == hipSuccess);
HIP_CHECK(hipFree(Ind));
HIP_CHECK(hipFree(Outd));
}
catch/unit/deviceLib/floatTM.cc
+226
View File
@@ -0,0 +1,226 @@
/*
Copyright (c) 2021 - present 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 <hip_test_common.hh>
#include <type_traits>
#include <random>
static std::random_device dev;
static std::mt19937 rng(dev());
template <typename T, typename M> __host__ __device__ inline constexpr int count() {
return sizeof(T) / sizeof(M);
}
inline float getRandomFloat(float min = 10, float max = 100) {
std::uniform_real_distribution<float> gen(min, max);
return gen(rng);
}
template <typename T, typename B> void fillMatrix(T* a, int size) {
for (int i = 0; i < size; i++) {
T t;
t.x = getRandomFloat();
if constexpr (count<T, B>() >= 2) t.y = getRandomFloat();
if constexpr (count<T, B>() >= 3) t.z = getRandomFloat();
if constexpr (count<T, B>() >= 4) t.w = getRandomFloat();
a[i] = t;
}
}
// Test operations
template <typename T, typename B> __host__ __device__ void testOperations(T& a, T& b) {
a.x += b.x;
a.x++;
b.x++;
if constexpr (count<T, B>() >= 2) {
a.y = b.x;
a.x = b.y;
}
if constexpr (count<T, B>() >= 3) {
if (a.x > 0) b.x /= a.x;
a.x *= b.z;
a.y--;
}
if constexpr (count<T, B>() >= 4) {
b.w = a.x;
a.w += (-b.y);
}
}
template <typename T, typename B> __global__ void testOperationsGPU(T* d_a, T* d_b, int size) {
int id = threadIdx.x;
if (id > size) return;
T& a = d_a[id];
T& b = d_b[id];
testOperations<T, B>(a, b);
}
template <typename T> void dcopy(T* a, T* b, int size) {
for (int i = 0; i < size; i++) {
a[i] = b[i];
}
}
template <typename T> bool isEqual(T* a, T* b, int size) {
for (int i = 0; i < size; i++) {
if (a[i] != b[i]) {
return false;
}
}
return true;
}
// Main function that tests type
// T = what you want to test
// D = pack of 1 i.e. float1 int1
template <typename T, typename D> void testType(int msize) {
T *fa, *fb, *fc, *h_fa, *h_fb;
fa = new T[msize];
fb = new T[msize];
fc = new T[msize];
h_fa = new T[msize];
h_fb = new T[msize];
T *d_fa, *d_fb;
constexpr int c = count<T, D>();
if (c <= 0 || c >= 5) {
INFO("It should be between 1 to 5");
REQUIRE(false);
}
fillMatrix<T, D>(fa, msize);
dcopy(fb, fa, msize);
dcopy(h_fa, fa, msize);
dcopy(h_fb, fa, msize);
for (int i = 0; i < msize; i++) testOperations<T, D>(h_fa[i], h_fb[i]);
HIP_CHECK(hipMalloc(&d_fa, sizeof(T) * msize));
HIP_CHECK(hipMalloc(&d_fb, sizeof(T) * msize));
HIP_CHECK(hipMemcpy(d_fa, fa, sizeof(T) * msize, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(d_fb, fb, sizeof(T) * msize, hipMemcpyHostToDevice));
auto kernel = testOperationsGPU<T, D>;
hipLaunchKernelGGL(kernel, 1, msize, 0, 0, d_fa, d_fb, msize);
HIP_CHECK(hipMemcpy(fc, d_fa, sizeof(T) * msize, hipMemcpyDeviceToHost));
bool pass = true;
if (!isEqual<T>(h_fa, fc, msize)) {
pass = false;
}
delete[] fa;
delete[] fb;
delete[] fc;
delete[] h_fa;
delete[] h_fb;
HIP_CHECK(hipFree(d_fa));
HIP_CHECK(hipFree(d_fb));
REQUIRE(pass);
}
TEST_CASE("Unit_floatTM") {
constexpr int msize = 100;
// double
testType<double1, double1>(msize);
testType<double2, double1>(msize);
testType<double3, double1>(msize);
testType<double4, double1>(msize);
// floats
testType<float1, float1>(msize);
testType<float2, float1>(msize);
testType<float3, float1>(msize);
testType<float4, float1>(msize);
// ints
testType<int1, int1>(msize);
testType<int2, int1>(msize);
testType<int3, int1>(msize);
testType<int4, int1>(msize);
// chars
testType<char1, char1>(msize);
testType<char2, char1>(msize);
testType<char3, char1>(msize);
testType<char4, char1>(msize);
// long
testType<long1, long1>(msize);
testType<long2, long1>(msize);
testType<long3, long1>(msize);
testType<long4, long1>(msize);
// longlong
testType<longlong1, longlong1>(msize);
testType<longlong2, longlong1>(msize);
testType<longlong3, longlong1>(msize);
testType<longlong4, longlong1>(msize);
// short
testType<short1, short1>(msize);
testType<short2, short1>(msize);
testType<short3, short1>(msize);
testType<short4, short1>(msize);
// uints
testType<uint1, uint1>(msize);
testType<uint2, uint1>(msize);
testType<uint3, uint1>(msize);
testType<uint4, uint1>(msize);
// uchars
testType<uchar1, uchar1>(msize);
testType<uchar2, uchar1>(msize);
testType<uchar3, uchar1>(msize);
testType<uchar4, uchar1>(msize);
// ulong
testType<ulong1, ulong1>(msize);
testType<ulong2, ulong1>(msize);
testType<ulong3, ulong1>(msize);
testType<ulong4, ulong1>(msize);
// ulonglong
testType<ulonglong1, ulonglong1>(msize);
testType<ulonglong2, ulonglong1>(msize);
testType<ulonglong3, ulonglong1>(msize);
testType<ulonglong4, ulonglong1>(msize);
// ushort
testType<ushort1, ushort1>(msize);
testType<ushort2, ushort1>(msize);
testType<ushort3, ushort1>(msize);
testType<ushort4, ushort1>(msize);
}
catch/unit/deviceLib/funnelshift.cc
+208
View File
@@ -0,0 +1,208 @@
/*
Copyright (c) 2021 - present 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 <hip_test_common.hh>
#include <hip/device_functions.h>
#include <assert.h>
#include <stdio.h>
#include <stdint.h>
#include <stdlib.h>
#define NUM_TESTS 65
#define HI_INT 0xfacefeed
#define LO_INT 0xdeadbeef
__global__ void funnelshift_kernel(unsigned int* l_out, unsigned int* lc_out, unsigned int* r_out,
unsigned int* rc_out) {
for (int i = 0; i < NUM_TESTS; i++) {
l_out[i] = __funnelshift_l(LO_INT, HI_INT, i);
lc_out[i] = __funnelshift_lc(LO_INT, HI_INT, i);
r_out[i] = __funnelshift_r(LO_INT, HI_INT, i);
rc_out[i] = __funnelshift_rc(LO_INT, HI_INT, i);
}
}
static unsigned int cpu_funnelshift_l(unsigned int lo, unsigned int hi, unsigned int shift) {
// Concatenate hi:lo
uint64_t val = hi;
val <<= 32;
val |= lo;
// left shift by intput & 31
val <<= (shift & 31);
// pull out upper 32 bits and return them
val >>= 32;
return val & 0xffffffff;
}
static unsigned int cpu_funnelshift_lc(unsigned int lo, unsigned int hi, unsigned int shift) {
// Concatenate hi:lo
uint64_t val = hi;
val <<= 32;
val |= lo;
// left shift by min(input,32)
if (shift > 32) shift = 32;
val <<= shift;
// pull out upper 32 bits and return them
val >>= 32;
return val & 0xffffffff;
}
static unsigned int cpu_funnelshift_r(unsigned int lo, unsigned int hi, unsigned int shift) {
// Concatenate hi:lo
uint64_t val = hi;
val <<= 32;
val |= lo;
// right shift by intput & 31
val >>= (shift & 31);
// return lower 32 bits
return val & 0xffffffff;
}
static unsigned int cpu_funnelshift_rc(unsigned int lo, unsigned int hi, unsigned int shift) {
// Concatenate hi:lo
uint64_t val = hi;
val <<= 32;
val |= lo;
// left shift by min(input, 32)
if (shift > 32) shift = 32;
val >>= shift;
// return lower 32 bits
return val & 0xffffffff;
}
TEST_CASE("Unit_funnelshift") {
unsigned int* host_l_output;
unsigned int* host_lc_output;
unsigned int* host_r_output;
unsigned int* host_rc_output;
unsigned int* device_l_output;
unsigned int* device_lc_output;
unsigned int* device_r_output;
unsigned int* device_rc_output;
unsigned int* golden_l;
unsigned int* golden_lc;
unsigned int* golden_r;
unsigned int* golden_rc;
hipDeviceProp_t devProp;
HIP_CHECK(hipGetDeviceProperties(&devProp, 0));
INFO("System minor : " << devProp.minor);
INFO("System major : " << devProp.major);
INFO("agent prop name : " << devProp.name);
INFO("hip Device prop succeeded");
int i;
int errors;
host_l_output = (unsigned int*)calloc(NUM_TESTS, sizeof(unsigned int));
host_lc_output = (unsigned int*)calloc(NUM_TESTS, sizeof(unsigned int));
host_r_output = (unsigned int*)calloc(NUM_TESTS, sizeof(unsigned int));
host_rc_output = (unsigned int*)calloc(NUM_TESTS, sizeof(unsigned int));
golden_l = (unsigned int*)calloc(NUM_TESTS, sizeof(unsigned int));
golden_lc = (unsigned int*)calloc(NUM_TESTS, sizeof(unsigned int));
golden_r = (unsigned int*)calloc(NUM_TESTS, sizeof(unsigned int));
golden_rc = (unsigned int*)calloc(NUM_TESTS, sizeof(unsigned int));
for (int i = 0; i < NUM_TESTS; i++) {
golden_l[i] = cpu_funnelshift_l(LO_INT, HI_INT, i);
golden_lc[i] = cpu_funnelshift_lc(LO_INT, HI_INT, i);
golden_r[i] = cpu_funnelshift_r(LO_INT, HI_INT, i);
golden_rc[i] = cpu_funnelshift_rc(LO_INT, HI_INT, i);
}
HIP_CHECK(hipMalloc((void**)&device_l_output, NUM_TESTS * sizeof(unsigned int)));
HIP_CHECK(hipMalloc((void**)&device_lc_output, NUM_TESTS * sizeof(unsigned int)));
HIP_CHECK(hipMalloc((void**)&device_r_output, NUM_TESTS * sizeof(unsigned int)));
HIP_CHECK(hipMalloc((void**)&device_rc_output, NUM_TESTS * sizeof(unsigned int)));
hipLaunchKernelGGL(funnelshift_kernel, dim3(1), dim3(1), 0, 0, device_l_output, device_lc_output,
device_r_output, device_rc_output);
HIP_CHECK(hipMemcpy(host_l_output, device_l_output, NUM_TESTS * sizeof(unsigned int),
hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpy(host_lc_output, device_lc_output, NUM_TESTS * sizeof(unsigned int),
hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpy(host_r_output, device_r_output, NUM_TESTS * sizeof(unsigned int),
hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpy(host_rc_output, device_rc_output, NUM_TESTS * sizeof(unsigned int),
hipMemcpyDeviceToHost));
// verify the results
errors = 0;
printf("HI val: 0x%x\n", HI_INT);
printf("LO val: 0x%x\n", LO_INT);
for (i = 0; i < NUM_TESTS; i++) {
if (host_l_output[i] != golden_l[i]) {
errors++;
INFO("Mismatch: " << host_l_output[i] << " - " << golden_l[i]);
break;
}
}
for (i = 0; i < NUM_TESTS && errors == 0; i++) {
if (host_lc_output[i] != golden_lc[i]) {
errors++;
INFO("Mismatch: " << host_lc_output[i] << " - " << golden_lc[i]);
break;
}
}
errors = 0;
for (i = 0; i < NUM_TESTS && errors == 0; i++) {
if (host_r_output[i] != golden_r[i]) {
errors++;
INFO("Mismatch: " << host_r_output[i] << " - " << golden_r[i]);
break;
}
}
for (i = 0; i < NUM_TESTS && errors == 0; i++) {
if (host_rc_output[i] != golden_rc[i]) {
errors++;
INFO("Mismatch: " << host_rc_output[i] << " - " << golden_rc[i]);
break;
}
}
HIP_CHECK(hipFree(device_l_output));
HIP_CHECK(hipFree(device_lc_output));
HIP_CHECK(hipFree(device_r_output));
HIP_CHECK(hipFree(device_rc_output));
free(host_l_output);
free(host_lc_output);
free(host_r_output);
free(host_rc_output);
REQUIRE(errors == 0);
}
catch/unit/deviceLib/vectorTypesDevice.cc
+25 -1
View File
@@ -1,3 +1,25 @@
/*
Copyright (c) 2021 - present 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 <hip_test_common.hh>
#include <type_traits>
#include <memory>
@@ -218,7 +240,7 @@ template <typename V, typename... Vs> bool run_CheckSharedVectorTypes() {
return run_CheckSharedVectorType<V>() && run_CheckSharedVectorTypes<Vs...>();
}
TEST_CASE("VectorTypesTest") {
TEST_CASE("Unit_vectorTypes_CompileTest") {
static_assert(sizeof(float1) == 4, "");
static_assert(sizeof(float2) >= 8, "");
static_assert(sizeof(float3) >= 12, "");
@@ -242,5 +264,7 @@ TEST_CASE("VectorTypesTest") {
uint3, uint4, long1, long2, long3, long4, ulong1, ulong2, ulong3, ulong4, longlong1,
longlong2, longlong3, longlong4, ulonglong1, ulonglong2, ulonglong3, ulonglong4,
float1, float2, float3, float4, double1, double2, double3, double4>();
HIP_CHECK(hipFree(ptr));
REQUIRE(passed == true);
}
catch/unit/memory/malloc.cc
+2 -2
View File
@@ -1,6 +1,6 @@
#include <hip_test_common.hh>
TEST_CASE("HostAllocBasic") {
TEST_CASE("Unit_hipHostMalloc_4bytes") {
int* d_a;
auto res = hipHostMalloc(&d_a, sizeof(int), 0);
REQUIRE(res == hipSuccess);
@@ -8,7 +8,7 @@ TEST_CASE("HostAllocBasic") {
REQUIRE(res == hipSuccess);
}
TEST_CASE("AllocateBasic") {
TEST_CASE("Unit_hipMalloc_4bytes") {
int* d_a;
auto res = hipMalloc(&d_a, sizeof(int));
REQUIRE(res == hipSuccess);
catch/unit/memory/memset.cc
+2 -2
View File
@@ -1,6 +1,6 @@
#include <hip_test_common.hh>
TEST_CASE("MemsetBasic") {
TEST_CASE("Unit_hipMemset_4bytes") {
int* d_a;
auto res = hipMalloc(&d_a, sizeof(int));
REQUIRE(res == hipSuccess);
@@ -9,7 +9,7 @@ TEST_CASE("MemsetBasic") {
hipFree(d_a);
}
TEST_CASE("HostMemsetBasic") {
TEST_CASE("Unit_hipMemset_4bytes_hostMem") {
int* d_a;
auto res = hipHostMalloc(&d_a, sizeof(int), 0);
REQUIRE(res == hipSuccess);