SWDEV-366138 - Add surface tests (#309)
Add surface tests Optimize some codes Change-Id: I1691b81d597bfd722d561778f919ebf8ca530e89
Esse commit está contido em:
ROCm CI Service Account
commit de
GitHub
GitHub
pai
f4feab33bc
commit
2337705b55
@@ -81,4 +81,136 @@ struct vector_info<ushort4>
|
||||
: type_and_size_and_format<unsigned short, 4, HIP_AD_FORMAT_UNSIGNED_INT16> {};
|
||||
template <>
|
||||
struct vector_info<uchar4>
|
||||
: type_and_size_and_format<unsigned char, 4, HIP_AD_FORMAT_UNSIGNED_INT8> {};
|
||||
: type_and_size_and_format<unsigned char, 4, HIP_AD_FORMAT_UNSIGNED_INT8> {};
|
||||
|
||||
template <
|
||||
typename T,
|
||||
typename std::enable_if<std::is_scalar<T>::value == false>::type* = nullptr>
|
||||
static inline __host__ __device__ constexpr int rank() {
|
||||
return sizeof(T) / sizeof(decltype(T::x));
|
||||
}
|
||||
|
||||
template<
|
||||
typename T,
|
||||
typename std::enable_if<rank<T>() == 1>::type* = nullptr>
|
||||
static inline bool isEqual(const T &val0, const T &val1) {
|
||||
return val0.x == val1.x;
|
||||
}
|
||||
|
||||
template<
|
||||
typename T,
|
||||
typename std::enable_if<rank<T>() == 2>::type* = nullptr>
|
||||
static inline bool isEqual(const T &val0, const T &val1) {
|
||||
return val0.x == val1.x &&
|
||||
val0.y == val1.y;
|
||||
}
|
||||
|
||||
template<
|
||||
typename T,
|
||||
typename std::enable_if<rank<T>() == 4>::type* = nullptr>
|
||||
static inline bool isEqual(const T &val0, const T &val1) {
|
||||
return val0.x == val1.x &&
|
||||
val0.y == val1.y &&
|
||||
val0.z == val1.z &&
|
||||
val0.w == val1.w;
|
||||
}
|
||||
|
||||
template<
|
||||
typename T,
|
||||
typename std::enable_if<std::is_scalar<T>::value>::type* = nullptr>
|
||||
static inline bool isEqual(const T &val0, const T &val1) {
|
||||
return val0 == val1;
|
||||
}
|
||||
|
||||
template<
|
||||
typename T,
|
||||
typename std::enable_if<rank<T>() == 1>::type* = nullptr>
|
||||
const std::string getString(const T& t)
|
||||
{
|
||||
std::ostringstream os;
|
||||
os<< "(" << t.x << ")";
|
||||
return os.str();
|
||||
}
|
||||
|
||||
template<
|
||||
typename T,
|
||||
typename std::enable_if<rank<T>() == 2>::type* = nullptr>
|
||||
const std::string getString(const T& t)
|
||||
{
|
||||
std::ostringstream os;
|
||||
os<< "(" << t.x << ", " << t.y << ")";
|
||||
return os.str();
|
||||
}
|
||||
|
||||
template<
|
||||
typename T,
|
||||
typename std::enable_if<rank<T>() == 3>::type* = nullptr>
|
||||
const std::string getString(const T& t)
|
||||
{
|
||||
std::ostringstream os;
|
||||
os<< "(" << t.x << ", " << t.y << ", " << t.z << ")";
|
||||
return os.str();
|
||||
}
|
||||
|
||||
template<
|
||||
typename T,
|
||||
typename std::enable_if<rank<T>() == 4>::type* = nullptr>
|
||||
const std::string getString(const T& t)
|
||||
{
|
||||
std::ostringstream os;
|
||||
os<< "(" << t.x << ", " << t.y << ", " << t.z << ", " << t.w << ")";
|
||||
return os.str();
|
||||
}
|
||||
|
||||
template<
|
||||
typename T,
|
||||
typename std::enable_if<std::is_scalar<T>::value>::type* = nullptr>
|
||||
std::string getString(const T& t)
|
||||
{
|
||||
std::ostringstream os;
|
||||
os << t;
|
||||
return os.str();
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
static inline T getRandom() {
|
||||
double r = 0;
|
||||
if (std::is_signed<T>::value) {
|
||||
r = (std::rand() - RAND_MAX / 2.0) / (RAND_MAX / 2.0 + 1.);
|
||||
} else {
|
||||
r = std::rand() / (RAND_MAX + 1.);
|
||||
}
|
||||
return static_cast<T>(std::numeric_limits < T > ::max() * r);
|
||||
}
|
||||
|
||||
template<
|
||||
typename T,
|
||||
typename std::enable_if<rank<T>() == 1>::type* = nullptr>
|
||||
static inline void initVal(T &val) {
|
||||
val.x = getRandom<decltype(T::x)>();
|
||||
}
|
||||
|
||||
template<
|
||||
typename T,
|
||||
typename std::enable_if<rank<T>() == 2>::type* = nullptr>
|
||||
static inline void initVal(T &val) {
|
||||
val.x = getRandom<decltype(T::x)>();
|
||||
val.y = getRandom<decltype(T::x)>();
|
||||
}
|
||||
|
||||
template<
|
||||
typename T,
|
||||
typename std::enable_if<rank<T>() == 4>::type* = nullptr>
|
||||
static inline void initVal(T &val) {
|
||||
val.x = getRandom<decltype(T::x)>();
|
||||
val.y = getRandom<decltype(T::x)>();
|
||||
val.z = getRandom<decltype(T::x)>();
|
||||
val.w = getRandom<decltype(T::x)>();
|
||||
}
|
||||
|
||||
template<
|
||||
typename T,
|
||||
typename std::enable_if<std::is_scalar<T>::value>::type* = nullptr>
|
||||
static inline void initVal(T &val) {
|
||||
val = getRandom<T>();
|
||||
}
|
||||
|
||||
@@ -8,6 +8,8 @@
|
||||
#define HIP_SAMPLING_VERIFY_ABSOLUTE_THRESHOLD 0.1
|
||||
|
||||
#if HT_NVIDIA
|
||||
typedef unsigned char uchar;
|
||||
|
||||
template<typename T>
|
||||
typename std::enable_if<sizeof(T) / sizeof(decltype(T::x)) == 4, T>::type
|
||||
inline __host__ __device__ operator+(const T &a, const T &b)
|
||||
|
||||
@@ -28,6 +28,7 @@ add_subdirectory(occupancy)
|
||||
add_subdirectory(device)
|
||||
add_subdirectory(printf)
|
||||
add_subdirectory(texture)
|
||||
add_subdirectory(surface)
|
||||
add_subdirectory(streamperthread)
|
||||
add_subdirectory(kernel)
|
||||
add_subdirectory(multiThread)
|
||||
|
||||
@@ -0,0 +1,30 @@
|
||||
# Copyright (c) 2023 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.
|
||||
|
||||
# Common Tests - Test independent of all platforms
|
||||
set(TEST_SRC
|
||||
hipSurfaceObj1D.cc
|
||||
hipSurfaceObj2D.cc
|
||||
hipSurfaceObj3D.cc
|
||||
)
|
||||
|
||||
hip_add_exe_to_target(NAME SurfaceTest
|
||||
TEST_SRC ${TEST_SRC}
|
||||
TEST_TARGET_NAME build_tests)
|
||||
@@ -0,0 +1,317 @@
|
||||
/*
|
||||
Copyright (c) 2023 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_array_common.hh>
|
||||
#include <hip_texture_helper.hh>
|
||||
|
||||
template <typename T>
|
||||
__global__ void
|
||||
surf1DKernelR(hipSurfaceObject_t surfaceObject,
|
||||
T* outputData, int width)
|
||||
{
|
||||
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (x < width) {
|
||||
surf1Dread(outputData + x, surfaceObject, x * sizeof(T));
|
||||
}
|
||||
#endif
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
__global__ void
|
||||
surf1DKernelW(hipSurfaceObject_t surfaceObject,
|
||||
T* inputData, int width)
|
||||
{
|
||||
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (x < width) {
|
||||
surf1Dwrite(inputData[x], surfaceObject, x * sizeof(T));
|
||||
}
|
||||
#endif
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
__global__ void
|
||||
surf1DKernelRW(hipSurfaceObject_t surfaceObject,
|
||||
hipSurfaceObject_t outputSurfObj, int width)
|
||||
{
|
||||
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (x < width) {
|
||||
T data;
|
||||
surf1Dread(&data, surfaceObject, x * sizeof(T));
|
||||
surf1Dwrite(data, outputSurfObj, x * sizeof(T));
|
||||
}
|
||||
#endif
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
static void runTestR(const int width)
|
||||
{
|
||||
unsigned int size = width * sizeof(T);
|
||||
T *hData = (T*) malloc (size);
|
||||
memset(hData, 0, size);
|
||||
for (int j = 0; j < width; j++)
|
||||
{
|
||||
initVal(hData[j]);
|
||||
}
|
||||
|
||||
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<T>();
|
||||
|
||||
hipArray *hipArray = nullptr;
|
||||
HIP_CHECK(hipMallocArray(&hipArray, &channelDesc, width, 0, hipArraySurfaceLoadStore));
|
||||
|
||||
HIP_CHECK(hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice));
|
||||
|
||||
hipResourceDesc resDesc;
|
||||
memset(&resDesc, 0, sizeof(resDesc));
|
||||
resDesc.resType = hipResourceTypeArray;
|
||||
resDesc.res.array.array = hipArray;
|
||||
|
||||
// Create surface object
|
||||
hipSurfaceObject_t surfaceObject = 0;
|
||||
HIP_CHECK(hipCreateSurfaceObject(&surfaceObject, &resDesc));
|
||||
|
||||
T *hOutputData = nullptr;
|
||||
HIP_CHECK(hipHostMalloc((void**)&hOutputData, size));
|
||||
memset(hOutputData, 0, size);
|
||||
|
||||
dim3 dimBlock (16, 1, 1);
|
||||
dim3 dimGrid ((width + dimBlock.x - 1) / dimBlock.x, 1, 1);
|
||||
|
||||
surf1DKernelR<T><<<dimGrid, dimBlock>>>(surfaceObject, hOutputData, width);
|
||||
|
||||
HIP_CHECK(hipGetLastError());
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
for (int j = 0; j < width; j++) {
|
||||
if (!isEqual(hData[j], hOutputData[j])) {
|
||||
printf("Difference [ %d ]:%s ----%s\n", j,
|
||||
getString(hData[j]).c_str(), getString(hOutputData[j]).c_str());
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
|
||||
HIP_CHECK(hipDestroySurfaceObject(surfaceObject));
|
||||
HIP_CHECK(hipFreeArray(hipArray));
|
||||
free(hData);
|
||||
HIP_CHECK(hipHostFree(hOutputData));
|
||||
REQUIRE(true);
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
static void runTestW(const int width)
|
||||
{
|
||||
unsigned int size = width * sizeof(T);
|
||||
T *hData = nullptr;
|
||||
HIP_CHECK(hipHostMalloc((void**)&hData, size));
|
||||
memset(hData, 0, size);
|
||||
|
||||
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<T>();
|
||||
|
||||
hipArray *hipArray = nullptr;
|
||||
HIP_CHECK(hipMallocArray(&hipArray, &channelDesc, width, 0, hipArraySurfaceLoadStore));
|
||||
|
||||
HIP_CHECK(hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice));
|
||||
|
||||
hipResourceDesc resDesc;
|
||||
memset(&resDesc, 0, sizeof(resDesc));
|
||||
resDesc.resType = hipResourceTypeArray;
|
||||
resDesc.res.array.array = hipArray;
|
||||
|
||||
// Create surface object
|
||||
hipSurfaceObject_t surfaceObject = 0;
|
||||
HIP_CHECK(hipCreateSurfaceObject(&surfaceObject, &resDesc));
|
||||
|
||||
for (int j = 0; j < width; j++)
|
||||
{
|
||||
initVal(hData[j]);
|
||||
}
|
||||
|
||||
dim3 dimBlock (16, 1, 1);
|
||||
dim3 dimGrid ((width + dimBlock.x - 1) / dimBlock.x, 1, 1);
|
||||
|
||||
surf1DKernelW<T><<<dimGrid, dimBlock>>>(surfaceObject, hData, width);
|
||||
|
||||
HIP_CHECK(hipGetLastError());
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
T *hOutputData = (T*) malloc (size);
|
||||
memset(hOutputData, 0, size);
|
||||
HIP_CHECK(hipMemcpyFromArray(hOutputData, hipArray, 0, 0, size, hipMemcpyDeviceToHost));
|
||||
|
||||
for (int j = 0; j < width; j++) {
|
||||
if (!isEqual(hData[j], hOutputData[j])) {
|
||||
printf("Difference [ %d ]:%s ----%s\n", j,
|
||||
getString(hData[j]).c_str(), getString(hOutputData[j]).c_str());
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
|
||||
HIP_CHECK(hipDestroySurfaceObject(surfaceObject));
|
||||
HIP_CHECK(hipFreeArray(hipArray));
|
||||
HIP_CHECK(hipHostFree(hData));
|
||||
free(hOutputData);
|
||||
REQUIRE(true);
|
||||
}
|
||||
|
||||
|
||||
template <typename T>
|
||||
static void runTestRW(const int width)
|
||||
{
|
||||
unsigned int size = width * sizeof(T);
|
||||
T *hData = (T*) malloc (size);
|
||||
memset(hData, 0, size);
|
||||
for (int j = 0; j < width; j++)
|
||||
{
|
||||
initVal(hData[j]);
|
||||
}
|
||||
|
||||
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<T>();
|
||||
|
||||
hipArray *hipArray = nullptr, *hipOutArray = nullptr;
|
||||
HIP_CHECK(hipMallocArray(&hipArray, &channelDesc, width, 0, hipArraySurfaceLoadStore));
|
||||
|
||||
HIP_CHECK(hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice));
|
||||
|
||||
hipResourceDesc resDesc;
|
||||
memset(&resDesc, 0, sizeof(resDesc));
|
||||
resDesc.resType = hipResourceTypeArray;
|
||||
resDesc.res.array.array = hipArray;
|
||||
|
||||
// Create surface object
|
||||
hipSurfaceObject_t surfaceObject = 0;
|
||||
HIP_CHECK(hipCreateSurfaceObject(&surfaceObject, &resDesc));
|
||||
|
||||
HIP_CHECK(hipMallocArray(&hipOutArray, &channelDesc, width, 0, hipArraySurfaceLoadStore));
|
||||
|
||||
hipResourceDesc resOutDesc;
|
||||
memset(&resOutDesc, 0, sizeof(resOutDesc));
|
||||
resOutDesc.resType = hipResourceTypeArray;
|
||||
resOutDesc.res.array.array = hipOutArray;
|
||||
|
||||
hipSurfaceObject_t outSurfaceObject = 0;
|
||||
HIP_CHECK(hipCreateSurfaceObject (&outSurfaceObject, &resOutDesc));
|
||||
|
||||
dim3 dimBlock (16, 1, 1);
|
||||
dim3 dimGrid ((width + dimBlock.x - 1) / dimBlock.x, 1, 1);
|
||||
|
||||
surf1DKernelRW<T><<<dimGrid, dimBlock>>>(surfaceObject, outSurfaceObject, width);
|
||||
|
||||
HIP_CHECK(hipGetLastError());
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
T *hOutputData = (T*) malloc (size);
|
||||
memset(hOutputData, 0, size);
|
||||
HIP_CHECK(hipMemcpyFromArray(hOutputData, hipOutArray, 0, 0, size, hipMemcpyDeviceToHost));
|
||||
|
||||
for (int j = 0; j < width; j++) {
|
||||
if (!isEqual(hData[j], hOutputData[j])) {
|
||||
printf("Difference [ %d ]:%s ----%s\n", j,
|
||||
getString(hData[j]).c_str(), getString(hOutputData[j]).c_str());
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
|
||||
HIP_CHECK(hipDestroySurfaceObject(surfaceObject));
|
||||
HIP_CHECK(hipDestroySurfaceObject(outSurfaceObject));
|
||||
HIP_CHECK(hipFreeArray(hipArray));
|
||||
HIP_CHECK(hipFreeArray(hipOutArray));
|
||||
free(hData);
|
||||
free(hOutputData);
|
||||
REQUIRE(true);
|
||||
}
|
||||
|
||||
TEMPLATE_TEST_CASE("Unit_hipSurfaceObj1D_type_R", "",
|
||||
char, uchar, short, ushort, int, uint, float,
|
||||
char1, uchar1, short1, ushort1, int1, uint1, float1,
|
||||
char2, uchar2, short2, ushort2, int2, uint2, float2,
|
||||
char4, uchar4, short4, ushort4, int4, uint4, float4)
|
||||
{
|
||||
CHECK_IMAGE_SUPPORT
|
||||
auto err = hipGetLastError(); // reset last err due to previous negative tests
|
||||
|
||||
SECTION("Unit_hipSurfaceObj1D_type_R - 31") {
|
||||
runTestR<TestType>(31);
|
||||
}
|
||||
|
||||
SECTION("Unit_hipSurfaceObj1D_type_R - 67") {
|
||||
runTestR<TestType>(67);
|
||||
}
|
||||
|
||||
SECTION("Unit_hipSurfaceObj1D_type_R - 131") {
|
||||
runTestR<TestType>(131);
|
||||
}
|
||||
|
||||
SECTION("Unit_hipSurfaceObj1D_type_R - 263") {
|
||||
runTestR<TestType>(263);
|
||||
}
|
||||
}
|
||||
|
||||
TEMPLATE_TEST_CASE("Unit_hipSurfaceObj1D_type_W", "",
|
||||
char, uchar, short, ushort, int, uint, float,
|
||||
char1, uchar1, short1, ushort1, int1, uint1, float1,
|
||||
char2, uchar2, short2, ushort2, int2, uint2, float2,
|
||||
char4, uchar4, short4, ushort4, int4, uint4, float4)
|
||||
{
|
||||
CHECK_IMAGE_SUPPORT
|
||||
auto err = hipGetLastError(); // reset last err due to previous negative tests
|
||||
|
||||
SECTION("Unit_hipSurfaceObj1D_type_W - 31") {
|
||||
runTestW<TestType>(31);
|
||||
}
|
||||
|
||||
SECTION("Unit_hipSurfaceObj1D_type_W - 63") {
|
||||
runTestW<TestType>(63);
|
||||
}
|
||||
|
||||
SECTION("Unit_hipSurfaceObj1D_type_W - 131") {
|
||||
runTestW<TestType>(131);
|
||||
}
|
||||
|
||||
SECTION("Unit_hipSurfaceObj1D_type_W - 263") {
|
||||
runTestW<TestType>(263);
|
||||
}
|
||||
}
|
||||
|
||||
TEMPLATE_TEST_CASE("Unit_hipSurfaceObj1D_type_RW", "",
|
||||
char, uchar, short, ushort, int, uint, float,
|
||||
char1, uchar1, short1, ushort1, int1, uint1, float1,
|
||||
char2, uchar2, short2, ushort2, int2, uint2, float2,
|
||||
char4, uchar4, short4, ushort4, int4, uint4, float4)
|
||||
{
|
||||
CHECK_IMAGE_SUPPORT
|
||||
auto err = hipGetLastError(); // reset last err due to previous negative tests
|
||||
|
||||
SECTION("Unit_hipSurfaceObj1D_type_RW - 23") {
|
||||
runTestRW<TestType>(23);
|
||||
}
|
||||
|
||||
SECTION("Unit_hipSurfaceObj1D_type_RW - 67") {
|
||||
runTestRW<TestType>(67);
|
||||
}
|
||||
|
||||
SECTION("Unit_hipSurfaceObj1D_type_RW - 131") {
|
||||
runTestRW<TestType>(131);
|
||||
}
|
||||
|
||||
SECTION("Unit_hipSurfaceObj1D_type_RW - 263") {
|
||||
runTestRW<TestType>(263);
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,364 @@
|
||||
/*
|
||||
Copyright (c) 2023 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_array_common.hh>
|
||||
#include <hip_texture_helper.hh>
|
||||
|
||||
#define LOG_DATA 0
|
||||
|
||||
template <typename T>
|
||||
__global__ void
|
||||
surf2DKernelR(hipSurfaceObject_t surfaceObject,
|
||||
T* outputData, int width, int height)
|
||||
{
|
||||
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
if (x < width && y < height) {
|
||||
surf2Dread<T>(outputData + y * width + x, surfaceObject, x * sizeof(T), y);
|
||||
}
|
||||
#endif
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
__global__ void
|
||||
surf2DKernelW(hipSurfaceObject_t surfaceObject,
|
||||
T* inputData, int width, int height)
|
||||
{
|
||||
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
if (x < width && y < height) {
|
||||
surf2Dwrite<T>(inputData[y * width + x], surfaceObject, x * sizeof(T), y);
|
||||
}
|
||||
#endif
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
__global__ void
|
||||
surf2DKernelRW(hipSurfaceObject_t surfaceObject,
|
||||
hipSurfaceObject_t outputSurfObj, int width, int height)
|
||||
{
|
||||
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
if (x < width && y < height) {
|
||||
T data;
|
||||
surf2Dread<T>(&data, surfaceObject, x * sizeof(T), y);
|
||||
surf2Dwrite<T>(data, outputSurfObj, x * sizeof(T), y);
|
||||
}
|
||||
#endif
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
static void runTestR(const int width, const int height)
|
||||
{
|
||||
unsigned int size = width * height * sizeof(T);
|
||||
T* hData = (T*) malloc(size);
|
||||
memset(hData, 0, size);
|
||||
for (int i = 0; i < height; i++)
|
||||
{
|
||||
for (int j = 0; j < width; j++)
|
||||
{
|
||||
initVal(hData[i * width + j]);
|
||||
}
|
||||
}
|
||||
|
||||
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<T>();
|
||||
hipArray *hipArray = nullptr;
|
||||
HIP_CHECK(hipMallocArray (&hipArray, &channelDesc, width, height,
|
||||
hipArraySurfaceLoadStore));
|
||||
|
||||
// Need set source pitch, but we don't have any padding here
|
||||
const size_t spitch = width * sizeof(T);
|
||||
HIP_CHECK(hipMemcpy2DToArray(hipArray, 0, 0, hData, spitch, spitch, height,
|
||||
hipMemcpyHostToDevice));
|
||||
|
||||
hipResourceDesc resDesc;
|
||||
memset(&resDesc, 0, sizeof(resDesc));
|
||||
resDesc.resType = hipResourceTypeArray;
|
||||
resDesc.res.array.array = hipArray;
|
||||
|
||||
// Create surface object
|
||||
hipSurfaceObject_t surfaceObject = 0;
|
||||
HIP_CHECK(hipCreateSurfaceObject(&surfaceObject, &resDesc));
|
||||
|
||||
T* hOutputData = nullptr;
|
||||
HIP_CHECK(hipHostMalloc((void**)&hOutputData, size));
|
||||
memset(hOutputData, 0, size);
|
||||
|
||||
dim3 dimBlock (16, 16, 1);
|
||||
dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y -1)/ dimBlock.y, 1);
|
||||
surf2DKernelR<T><<<dimGrid, dimBlock>>>(surfaceObject, hOutputData, width, height);
|
||||
HIP_CHECK(hipGetLastError());
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
for (int i = 0; i < height; i++) {
|
||||
for (int j = 0; j < width; j++) {
|
||||
int index = i * width + j;
|
||||
if (!isEqual(hData[index], hOutputData[index])) {
|
||||
printf("Difference [ %d %d ]:%s ----%s\n", i, j,
|
||||
getString(hData[index]).c_str(), getString(hOutputData[index]).c_str());
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
HIP_CHECK(hipDestroySurfaceObject(surfaceObject));
|
||||
HIP_CHECK(hipFreeArray(hipArray));
|
||||
free(hData);
|
||||
HIP_CHECK(hipHostFree(hOutputData));
|
||||
REQUIRE(true);
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
static void runTestW(const int width, const int height)
|
||||
{
|
||||
unsigned int size = width * height * sizeof(T);
|
||||
T* hData = nullptr;
|
||||
HIP_CHECK(hipHostMalloc((void**)&hData, size));
|
||||
memset(hData, 0, size);
|
||||
|
||||
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<T>();
|
||||
hipArray *hipArray = nullptr;
|
||||
HIP_CHECK(hipMallocArray (&hipArray, &channelDesc, width, height,
|
||||
hipArraySurfaceLoadStore));
|
||||
|
||||
// Need set source pitch, but we don't have any padding here
|
||||
const size_t spitch = width * sizeof(T);
|
||||
HIP_CHECK(hipMemcpy2DToArray(hipArray, 0, 0, hData, spitch, spitch, height,
|
||||
hipMemcpyHostToDevice));
|
||||
|
||||
hipResourceDesc resDesc;
|
||||
memset(&resDesc, 0, sizeof(resDesc));
|
||||
resDesc.resType = hipResourceTypeArray;
|
||||
resDesc.res.array.array = hipArray;
|
||||
|
||||
// Create surface object
|
||||
hipSurfaceObject_t surfaceObject = 0;
|
||||
HIP_CHECK(hipCreateSurfaceObject(&surfaceObject, &resDesc));
|
||||
|
||||
for (int i = 0; i < height; i++)
|
||||
{
|
||||
for (int j = 0; j < width; j++)
|
||||
{
|
||||
initVal(hData[i * width + j]);
|
||||
}
|
||||
}
|
||||
|
||||
dim3 dimBlock (16, 16, 1);
|
||||
dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y -1)/ dimBlock.y, 1);
|
||||
surf2DKernelW<T><<<dimGrid, dimBlock>>>(surfaceObject, hData, width, height);
|
||||
HIP_CHECK(hipGetLastError());
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
T* hOutputData = (T*) malloc(size);
|
||||
|
||||
memset(hOutputData, 0, size);
|
||||
HIP_CHECK(hipMemcpy2DFromArray(hOutputData, spitch, hipArray, 0, 0, spitch,
|
||||
height, hipMemcpyDeviceToHost));
|
||||
|
||||
for (int i = 0; i < height; i++) {
|
||||
for (int j = 0; j < width; j++) {
|
||||
int index = i * width + j;
|
||||
if (!isEqual(hData[index], hOutputData[index])) {
|
||||
printf("Difference [ %d %d ]:%s ----%s\n", i, j,
|
||||
getString(hData[index]).c_str(), getString(hOutputData[index]).c_str());
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
HIP_CHECK(hipDestroySurfaceObject(surfaceObject));
|
||||
HIP_CHECK(hipFreeArray(hipArray));
|
||||
HIP_CHECK(hipHostFree(hData));
|
||||
free(hOutputData);
|
||||
REQUIRE(true);
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
static void runTestRW(const int width, const int height)
|
||||
{
|
||||
unsigned int size = width * height * sizeof(T);
|
||||
T* hData = (T*) malloc(size);
|
||||
memset(hData, 0, size);
|
||||
for (int i = 0; i < height; i++)
|
||||
{
|
||||
for (int j = 0; j < width; j++)
|
||||
{
|
||||
initVal(hData[i * width + j]);
|
||||
}
|
||||
}
|
||||
#if LOG_DATA
|
||||
printf ("hData: ");
|
||||
for (int i = 0; i < 32; i++)
|
||||
{
|
||||
printf ("%s ", getString(hData[i]).c_str());
|
||||
}
|
||||
printf ("\n");
|
||||
#endif
|
||||
|
||||
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<T>();
|
||||
hipArray *hipArray = nullptr, *hipOutArray = nullptr;
|
||||
HIP_CHECK(hipMallocArray (&hipArray, &channelDesc, width, height,
|
||||
hipArraySurfaceLoadStore));
|
||||
|
||||
// Need set source pitch, but we don't have any padding here
|
||||
const size_t spitch = width * sizeof(T);
|
||||
HIP_CHECK(hipMemcpy2DToArray(hipArray, 0, 0, hData, spitch, spitch, height,
|
||||
hipMemcpyHostToDevice));
|
||||
|
||||
hipResourceDesc resDesc;
|
||||
memset(&resDesc, 0, sizeof(resDesc));
|
||||
resDesc.resType = hipResourceTypeArray;
|
||||
resDesc.res.array.array = hipArray;
|
||||
|
||||
// Create surface object
|
||||
hipSurfaceObject_t surfaceObject = 0;
|
||||
HIP_CHECK(hipCreateSurfaceObject(&surfaceObject, &resDesc));
|
||||
|
||||
HIP_CHECK(hipMallocArray(&hipOutArray, &channelDesc, width, height,
|
||||
hipArraySurfaceLoadStore));
|
||||
|
||||
hipResourceDesc resOutDesc;
|
||||
memset(&resOutDesc, 0, sizeof(resOutDesc));
|
||||
resOutDesc.resType = hipResourceTypeArray;
|
||||
resOutDesc.res.array.array = hipOutArray;
|
||||
|
||||
hipSurfaceObject_t outSurfaceObject = 0;
|
||||
HIP_CHECK(hipCreateSurfaceObject (&outSurfaceObject, &resOutDesc));
|
||||
|
||||
dim3 dimBlock (16, 16, 1);
|
||||
dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y -1)/ dimBlock.y, 1);
|
||||
surf2DKernelRW<T><<<dimGrid, dimBlock>>>(surfaceObject, outSurfaceObject, width, height);
|
||||
HIP_CHECK(hipGetLastError());
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
T* hOutputData = (T*) malloc(size);
|
||||
|
||||
memset(hOutputData, 0, size);
|
||||
HIP_CHECK(hipMemcpy2DFromArray(hOutputData, spitch, hipOutArray, 0, 0, spitch,
|
||||
height, hipMemcpyDeviceToHost));
|
||||
|
||||
#if LOG_DATA
|
||||
printf ("dData: ");
|
||||
for (int i = 0; i < 32; i++)
|
||||
{
|
||||
printf ("%s ", getString(hOutputData[i]).c_str());
|
||||
}
|
||||
printf ("\n");
|
||||
#endif
|
||||
|
||||
for (int i = 0; i < height; i++) {
|
||||
for (int j = 0; j < width; j++) {
|
||||
int index = i * width + j;
|
||||
if (!isEqual(hData[index], hOutputData[index])) {
|
||||
printf("Difference [ %d %d ]:%s ----%s\n", i, j,
|
||||
getString(hData[index]).c_str(), getString(hOutputData[index]).c_str());
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
HIP_CHECK(hipDestroySurfaceObject(surfaceObject));
|
||||
HIP_CHECK(hipDestroySurfaceObject(outSurfaceObject));
|
||||
HIP_CHECK(hipFreeArray(hipArray));
|
||||
HIP_CHECK(hipFreeArray(hipOutArray));
|
||||
free(hData);
|
||||
free(hOutputData);
|
||||
REQUIRE(true);
|
||||
}
|
||||
|
||||
TEMPLATE_TEST_CASE("Unit_hipSurfaceObj2D_type_R", "",
|
||||
char, uchar, short, ushort, int, uint, float,
|
||||
char1, uchar1, short1, ushort1, int1, uint1, float1,
|
||||
char2, uchar2, short2, ushort2, int2, uint2, float2,
|
||||
char4, uchar4, short4, ushort4, int4, uint4, float4)
|
||||
{
|
||||
CHECK_IMAGE_SUPPORT
|
||||
auto err = hipGetLastError(); // reset last err due to previous negative tests
|
||||
|
||||
SECTION("Unit_hipSurfaceObj2D_type_R - 23, 67") {
|
||||
runTestR<TestType>(23, 67);
|
||||
}
|
||||
|
||||
SECTION("Unit_hipSurfaceObj2D_type_R - 67, 23") {
|
||||
runTestR<TestType>(67, 23);
|
||||
}
|
||||
|
||||
SECTION("Unit_hipSurfaceObj2D_type_R - 131, 67") {
|
||||
runTestR<TestType>(131, 67);
|
||||
}
|
||||
|
||||
SECTION("Unit_hipSurfaceObj2D_type_R - 263, 131") {
|
||||
runTestR<TestType>(263, 131);
|
||||
}
|
||||
}
|
||||
|
||||
TEMPLATE_TEST_CASE("Unit_hipSurfaceObj2D_type_W", "",
|
||||
char, uchar, short, ushort, int, uint, float,
|
||||
char1, uchar1, short1, ushort1, int1, uint1, float1,
|
||||
char2, uchar2, short2, ushort2, int2, uint2, float2,
|
||||
char4, uchar4, short4, ushort4, int4, uint4, float4)
|
||||
{
|
||||
CHECK_IMAGE_SUPPORT
|
||||
auto err = hipGetLastError(); // reset last err due to previous negative tests
|
||||
|
||||
SECTION("Unit_hipSurfaceObj2D_type_W - 23, 67") {
|
||||
runTestW<TestType>(23, 67);
|
||||
}
|
||||
|
||||
SECTION("Unit_hipSurfaceObj2D_type_W - 67, 23") {
|
||||
runTestW<TestType>(67, 23);
|
||||
}
|
||||
|
||||
SECTION("Unit_hipSurfaceObj2D_type_W - 131, 67") {
|
||||
runTestW<TestType>(131, 67);
|
||||
}
|
||||
|
||||
SECTION("Unit_hipSurfaceObj2D_type_W - 263, 23") {
|
||||
runTestW<TestType>(263, 23);
|
||||
}
|
||||
}
|
||||
|
||||
TEMPLATE_TEST_CASE("Unit_hipSurfaceObj2D_type_RW", "",
|
||||
char, uchar, short, ushort, int, uint, float,
|
||||
char1, uchar1, short1, ushort1, int1, uint1, float1,
|
||||
char2, uchar2, short2, ushort2, int2, uint2, float2,
|
||||
char4, uchar4, short4, ushort4, int4, uint4, float4)
|
||||
{
|
||||
CHECK_IMAGE_SUPPORT
|
||||
auto err = hipGetLastError(); // reset last err due to previous negative tests
|
||||
|
||||
SECTION("Unit_hipSurfaceObj2D_type_RW - 23, 67") {
|
||||
runTestRW<TestType>(23, 67);
|
||||
}
|
||||
|
||||
SECTION("Unit_hipSurfaceObj2D_type_RW - 67, 131") {
|
||||
runTestRW<TestType>(67, 131);
|
||||
}
|
||||
|
||||
SECTION("Unit_hipSurfaceObj2D_type_RW - 131, 263") {
|
||||
runTestRW<TestType>(131, 263);
|
||||
}
|
||||
|
||||
SECTION("Unit_hipSurfaceObj2D_type_RW - 263, 67") {
|
||||
runTestRW<TestType>(263, 67);
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,400 @@
|
||||
/*
|
||||
Copyright (c) 2023 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_array_common.hh>
|
||||
#include <hip_texture_helper.hh>
|
||||
|
||||
template <typename T>
|
||||
__global__ void
|
||||
surf3DKernelR(hipSurfaceObject_t surfaceObject,
|
||||
T* outputData, int width, int height, int depth)
|
||||
{
|
||||
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
int z = blockIdx.z * blockDim.z + threadIdx.z;
|
||||
if (x < width && y < height && z < depth) {
|
||||
surf3Dread(outputData + z * width * height + y * width + x,
|
||||
surfaceObject, x * sizeof(T), y, z);
|
||||
}
|
||||
#endif
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
__global__ void
|
||||
surf3DKernelW(hipSurfaceObject_t surfaceObject,
|
||||
T* inputData, int width, int height, int depth)
|
||||
{
|
||||
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
int z = blockIdx.z * blockDim.z + threadIdx.z;
|
||||
if (x < width && y < height && z < depth) {
|
||||
surf3Dwrite(inputData[z * width * height + y * width + x],
|
||||
surfaceObject, x * sizeof(T), y, z);
|
||||
}
|
||||
#endif
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
__global__ void
|
||||
surf3DKernelRW(hipSurfaceObject_t surfaceObject,
|
||||
hipSurfaceObject_t outputSurfObj, int width, int height, int depth)
|
||||
{
|
||||
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
int z = blockIdx.z * blockDim.z + threadIdx.z;
|
||||
if (x < width && y < height && z < depth) {
|
||||
T data;
|
||||
surf3Dread(&data, surfaceObject, x * sizeof(T), y, z);
|
||||
surf3Dwrite(data, outputSurfObj, x * sizeof(T), y, z);
|
||||
}
|
||||
#endif
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
static void runTestR(const int width, const int height, const int depth)
|
||||
{
|
||||
unsigned int size = width * height * depth * sizeof(T);
|
||||
T *hData = (T*) malloc(size);
|
||||
memset(hData, 0, size);
|
||||
for (int i = 0; i < depth; i++) {
|
||||
for (int j = 0; j < height; j++) {
|
||||
for (int k = 0; k < width; k++) {
|
||||
initVal(hData[i * width * height + j * width + k]);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
// Allocate array and copy image data
|
||||
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<T>();
|
||||
hipArray *hipArray = nullptr;
|
||||
HIP_CHECK(hipMalloc3DArray(&hipArray, &channelDesc, make_hipExtent(width, height, depth),
|
||||
hipArraySurfaceLoadStore));
|
||||
|
||||
hipMemcpy3DParms myparms;
|
||||
memset(&myparms, 0, sizeof(myparms));
|
||||
myparms.srcPos = make_hipPos(0,0,0);
|
||||
myparms.dstPos = make_hipPos(0,0,0);
|
||||
myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height);
|
||||
myparms.dstArray = hipArray;
|
||||
myparms.extent = make_hipExtent(width, height, depth);
|
||||
myparms.kind = hipMemcpyHostToDevice;
|
||||
|
||||
HIP_CHECK(hipMemcpy3D(&myparms));
|
||||
|
||||
hipResourceDesc resDesc;
|
||||
memset(&resDesc, 0, sizeof(resDesc));
|
||||
resDesc.resType = hipResourceTypeArray;
|
||||
resDesc.res.array.array = hipArray;
|
||||
|
||||
// Create surface object
|
||||
hipSurfaceObject_t surfaceObject = 0;
|
||||
HIP_CHECK(hipCreateSurfaceObject(&surfaceObject, &resDesc));
|
||||
|
||||
T *hOutputData = nullptr;
|
||||
HIP_CHECK(hipHostMalloc((void**)&hOutputData, size));
|
||||
memset(hOutputData, 0, size);
|
||||
|
||||
dim3 dimBlock(8, 8, 8); // 512 threads
|
||||
dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y -1)/ dimBlock.y,
|
||||
(depth + dimBlock.z - 1) / dimBlock.z);
|
||||
|
||||
surf3DKernelR<T><<<dimGrid, dimBlock>>>(surfaceObject, hOutputData, width, height, depth);
|
||||
|
||||
HIP_CHECK(hipGetLastError());
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
for (int i = 0; i < depth; i++) {
|
||||
for (int j = 0; j < height; j++) {
|
||||
for (int k = 0; k < width; k++) {
|
||||
int index = i * width * height + j * width + k;
|
||||
if (!isEqual(hData[index], hOutputData[index])) {
|
||||
printf("Difference [ %d %d %d]:%s ----%s\n", i, j, k,
|
||||
getString(hData[index]).c_str(), getString(hOutputData[index]).c_str());
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
HIP_CHECK(hipDestroySurfaceObject (surfaceObject));
|
||||
HIP_CHECK(hipFreeArray(hipArray));
|
||||
free(hData);
|
||||
HIP_CHECK(hipHostFree(hOutputData));
|
||||
REQUIRE(true);
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
static void runTestW(const int width, const int height, const int depth)
|
||||
{
|
||||
unsigned int size = width * height * depth * sizeof(T);
|
||||
T *hData = nullptr;
|
||||
HIP_CHECK(hipHostMalloc((void**)&hData, size));
|
||||
memset(hData, 0, size);
|
||||
|
||||
// Allocate array and copy image data
|
||||
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<T>();
|
||||
hipArray *hipArray = nullptr;
|
||||
HIP_CHECK(hipMalloc3DArray(&hipArray, &channelDesc, make_hipExtent(width, height, depth),
|
||||
hipArraySurfaceLoadStore));
|
||||
|
||||
hipMemcpy3DParms myparms;
|
||||
memset(&myparms, 0, sizeof(myparms));
|
||||
myparms.srcPos = make_hipPos(0,0,0);
|
||||
myparms.dstPos = make_hipPos(0,0,0);
|
||||
myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height);
|
||||
myparms.dstArray = hipArray;
|
||||
myparms.extent = make_hipExtent(width, height, depth);
|
||||
myparms.kind = hipMemcpyHostToDevice;
|
||||
|
||||
HIP_CHECK(hipMemcpy3D(&myparms));
|
||||
|
||||
hipResourceDesc resDesc;
|
||||
memset(&resDesc, 0, sizeof(resDesc));
|
||||
resDesc.resType = hipResourceTypeArray;
|
||||
resDesc.res.array.array = hipArray;
|
||||
|
||||
// Create surface object
|
||||
hipSurfaceObject_t surfaceObject = 0;
|
||||
HIP_CHECK(hipCreateSurfaceObject(&surfaceObject, &resDesc));
|
||||
|
||||
for (int i = 0; i < depth; i++) {
|
||||
for (int j = 0; j < height; j++) {
|
||||
for (int k = 0; k < width; k++) {
|
||||
initVal(hData[i * width * height + j * width + k]);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
dim3 dimBlock(8, 8, 8); // 512 threads
|
||||
dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y -1)/ dimBlock.y,
|
||||
(depth + dimBlock.z - 1) / dimBlock.z);
|
||||
|
||||
surf3DKernelW<T><<<dimGrid, dimBlock>>>(surfaceObject, hData, width, height, depth);
|
||||
|
||||
HIP_CHECK(hipGetLastError());
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
T *hOutputData = (T*) malloc (size);
|
||||
memset(hOutputData, 0, size);
|
||||
|
||||
memset(&myparms, 0, sizeof(myparms));
|
||||
myparms.srcPos = make_hipPos(0,0,0);
|
||||
myparms.dstPos = make_hipPos(0,0,0);
|
||||
myparms.srcArray= hipArray;
|
||||
myparms.dstPtr = make_hipPitchedPtr(hOutputData, width * sizeof(T), width, height);
|
||||
myparms.extent = make_hipExtent(width, height, depth);
|
||||
myparms.kind = hipMemcpyDeviceToHost;
|
||||
|
||||
HIP_CHECK(hipMemcpy3D(&myparms));
|
||||
|
||||
for (int i = 0; i < depth; i++) {
|
||||
for (int j = 0; j < height; j++) {
|
||||
for (int k = 0; k < width; k++) {
|
||||
int index = i * width * height + j * width + k;
|
||||
if (!isEqual(hData[index], hOutputData[index])) {
|
||||
printf("Difference [ %d %d %d]:%s ----%s\n", i, j, k,
|
||||
getString(hData[index]).c_str(), getString(hOutputData[index]).c_str());
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
HIP_CHECK(hipDestroySurfaceObject (surfaceObject));
|
||||
HIP_CHECK(hipFreeArray(hipArray));
|
||||
HIP_CHECK(hipHostFree(hData));
|
||||
free(hOutputData);
|
||||
REQUIRE(true);
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
static void runTestRW(const int width, const int height, const int depth)
|
||||
{
|
||||
unsigned int size = width * height * depth * sizeof(T);
|
||||
T *hData = (T*) malloc(size);
|
||||
memset(hData, 0, size);
|
||||
for (int i = 0; i < depth; i++) {
|
||||
for (int j = 0; j < height; j++) {
|
||||
for (int k = 0; k < width; k++) {
|
||||
initVal(hData[i * width * height + j * width + k]);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
// Allocate array and copy image data
|
||||
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<T>();
|
||||
hipArray *hipArray = nullptr, *hipOutArray = nullptr;
|
||||
HIP_CHECK(hipMalloc3DArray(&hipArray, &channelDesc, make_hipExtent(width, height, depth),
|
||||
hipArraySurfaceLoadStore));
|
||||
|
||||
hipMemcpy3DParms myparms;
|
||||
memset(&myparms, 0, sizeof(myparms));
|
||||
myparms.srcPos = make_hipPos(0,0,0);
|
||||
myparms.dstPos = make_hipPos(0,0,0);
|
||||
myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height);
|
||||
myparms.dstArray = hipArray;
|
||||
myparms.extent = make_hipExtent(width, height, depth);
|
||||
myparms.kind = hipMemcpyHostToDevice;
|
||||
|
||||
HIP_CHECK(hipMemcpy3D(&myparms));
|
||||
|
||||
hipResourceDesc resDesc;
|
||||
memset(&resDesc, 0, sizeof(resDesc));
|
||||
resDesc.resType = hipResourceTypeArray;
|
||||
resDesc.res.array.array = hipArray;
|
||||
|
||||
// Create surface object
|
||||
hipSurfaceObject_t surfaceObject = 0;
|
||||
HIP_CHECK(hipCreateSurfaceObject(&surfaceObject, &resDesc));
|
||||
|
||||
HIP_CHECK(hipMalloc3DArray(&hipOutArray, &channelDesc, make_hipExtent(width, height, depth),
|
||||
hipArraySurfaceLoadStore));
|
||||
|
||||
hipResourceDesc resOutDesc;
|
||||
memset(&resOutDesc, 0, sizeof(resOutDesc));
|
||||
resOutDesc.resType = hipResourceTypeArray;
|
||||
resOutDesc.res.array.array = hipOutArray;
|
||||
|
||||
hipSurfaceObject_t outSurfaceObject = 0;
|
||||
HIP_CHECK(hipCreateSurfaceObject(&outSurfaceObject, &resOutDesc));
|
||||
|
||||
dim3 dimBlock(8, 8, 8); // 512 threads
|
||||
dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y -1)/ dimBlock.y,
|
||||
(depth + dimBlock.z - 1) / dimBlock.z);
|
||||
|
||||
surf3DKernelRW<T><<<dimGrid, dimBlock>>>(surfaceObject, outSurfaceObject, width, height, depth);
|
||||
|
||||
HIP_CHECK(hipGetLastError());
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
T *hOutputData = (T*) malloc (size);
|
||||
memset(hOutputData, 0, size);
|
||||
|
||||
memset(&myparms, 0, sizeof(myparms));
|
||||
myparms.srcPos = make_hipPos(0,0,0);
|
||||
myparms.dstPos = make_hipPos(0,0,0);
|
||||
myparms.srcArray= hipOutArray;
|
||||
myparms.dstPtr = make_hipPitchedPtr(hOutputData, width * sizeof(T), width, height);
|
||||
myparms.extent = make_hipExtent(width, height, depth);
|
||||
myparms.kind = hipMemcpyDeviceToHost;
|
||||
|
||||
HIP_CHECK(hipMemcpy3D(&myparms));
|
||||
|
||||
for (int i = 0; i < depth; i++) {
|
||||
for (int j = 0; j < height; j++) {
|
||||
for (int k = 0; k < width; k++) {
|
||||
int index = i * width * height + j * width + k;
|
||||
if (!isEqual(hData[index], hOutputData[index])) {
|
||||
printf("Difference [ %d %d %d]:%s ----%s\n", i, j, k,
|
||||
getString(hData[index]).c_str(), getString(hOutputData[index]).c_str());
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
HIP_CHECK(hipDestroySurfaceObject (surfaceObject));
|
||||
HIP_CHECK(hipDestroySurfaceObject (outSurfaceObject));
|
||||
HIP_CHECK(hipFreeArray(hipArray));
|
||||
HIP_CHECK(hipFreeArray(hipOutArray));
|
||||
free(hData);
|
||||
free(hOutputData);
|
||||
REQUIRE(true);
|
||||
}
|
||||
|
||||
TEMPLATE_TEST_CASE("Unit_hipSurfaceObj3D_type_R", "",
|
||||
char, uchar, short, ushort, int, uint, float,
|
||||
char1, uchar1, short1, ushort1, int1, uint1, float1,
|
||||
char2, uchar2, short2, ushort2, int2, uint2, float2,
|
||||
char4, uchar4, short4, ushort4, int4, uint4, float4)
|
||||
{
|
||||
CHECK_IMAGE_SUPPORT
|
||||
auto err = hipGetLastError(); // reset last err due to previous negative tests
|
||||
|
||||
SECTION("Unit_hipSurfaceObj3D_type_R - 31, 67, 131") {
|
||||
runTestR<TestType>(31, 67, 131);
|
||||
}
|
||||
|
||||
SECTION("Unit_hipSurfaceObj3D_type_R - 67, 31, 263") {
|
||||
runTestR<TestType>(67, 31, 263);
|
||||
}
|
||||
|
||||
SECTION("Unit_hipSurfaceObj3D_type_R - 131, 131, 67") {
|
||||
runTestR<TestType>(131, 131, 67);
|
||||
}
|
||||
|
||||
SECTION("Unit_hipSurfaceObj3D_type_R - 263, 131, 263") {
|
||||
runTestR<TestType>(263, 131, 263);
|
||||
}
|
||||
}
|
||||
|
||||
TEMPLATE_TEST_CASE("Unit_hipSurfaceObj3D_type_W", "",
|
||||
char, uchar, short, ushort, int, uint, float,
|
||||
char1, uchar1, short1, ushort1, int1, uint1, float1,
|
||||
char2, uchar2, short2, ushort2, int2, uint2, float2,
|
||||
char4, uchar4, short4, ushort4, int4, uint4, float4)
|
||||
{
|
||||
CHECK_IMAGE_SUPPORT
|
||||
auto err = hipGetLastError(); // reset last err due to previous negative tests
|
||||
|
||||
SECTION("Unit_hipSurfaceObj3D_type_W - 31, 67, 131") {
|
||||
runTestW<TestType>(31, 67, 131);
|
||||
}
|
||||
|
||||
SECTION("Unit_hipSurfaceObj3D_type_W - 67, 67, 31") {
|
||||
runTestW<TestType>(67, 67, 31);
|
||||
}
|
||||
|
||||
SECTION("Unit_hipSurfaceObj3D_type_W - 131, 131, 67") {
|
||||
runTestW<TestType>(131, 131, 67);
|
||||
}
|
||||
|
||||
SECTION("Unit_hipSurfaceObj3D_type_W - 263, 131, 263") {
|
||||
runTestW<TestType>(263, 131, 263);
|
||||
}
|
||||
}
|
||||
|
||||
TEMPLATE_TEST_CASE("Unit_hipSurfaceObj3D_type_RW", "",
|
||||
char, uchar, short, ushort, int, uint, float,
|
||||
char1, uchar1, short1, ushort1, int1, uint1, float1,
|
||||
char2, uchar2, short2, ushort2, int2, uint2, float2,
|
||||
char4, uchar4, short4, ushort4, int4, uint4, float4)
|
||||
{
|
||||
CHECK_IMAGE_SUPPORT
|
||||
auto err = hipGetLastError(); // reset last err due to previous negative tests
|
||||
|
||||
SECTION("Unit_hipSurfaceObj3D_type_RW - 31, 31, 67") {
|
||||
runTestRW<TestType>(31, 31, 67);
|
||||
}
|
||||
|
||||
SECTION("Unit_hipSurfaceObj3D_type_RW - 67, 67, 31") {
|
||||
runTestRW<TestType>(67, 67, 31);
|
||||
}
|
||||
|
||||
SECTION("Unit_hipSurfaceObj3D_type_RW - 131, 67, 263") {
|
||||
runTestRW<TestType>(131, 67, 263);
|
||||
}
|
||||
|
||||
SECTION("Unit_hipSurfaceObj3D_type_RW - 263, 131, 263") {
|
||||
runTestRW<TestType>(263, 131, 263);
|
||||
}
|
||||
}
|
||||
@@ -18,6 +18,7 @@ THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip_array_common.hh>
|
||||
#include <vector>
|
||||
#include <iostream>
|
||||
|
||||
@@ -31,47 +32,6 @@ __global__ void tex1dKernelFetch(T *val, hipTextureObject_t obj, int N) {
|
||||
#endif
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
static inline __host__ __device__ constexpr int rank() {
|
||||
return sizeof(T) / sizeof(decltype(T::x));
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
static inline T getRandom() {
|
||||
double r = 0;
|
||||
if (std::is_signed < T > ::value) {
|
||||
r = (std::rand() - RAND_MAX / 2.0) / (RAND_MAX / 2.0 + 1.);
|
||||
} else {
|
||||
r = std::rand() / (RAND_MAX + 1.);
|
||||
}
|
||||
return static_cast<T>(std::numeric_limits < T > ::max() * r);
|
||||
}
|
||||
|
||||
template<
|
||||
typename T,
|
||||
typename std::enable_if<rank<T>() == 1>::type* = nullptr>
|
||||
static inline void initVal(T &val) {
|
||||
val.x = getRandom<decltype(T::x)>();
|
||||
}
|
||||
|
||||
template<
|
||||
typename T,
|
||||
typename std::enable_if<rank<T>() == 2>::type* = nullptr>
|
||||
static inline void initVal(T &val) {
|
||||
val.x = getRandom<decltype(T::x)>();
|
||||
val.y = getRandom<decltype(T::x)>();
|
||||
}
|
||||
|
||||
template<
|
||||
typename T,
|
||||
typename std::enable_if<rank<T>() == 4>::type* = nullptr>
|
||||
static inline void initVal(T &val) {
|
||||
val.x = getRandom<decltype(T::x)>();
|
||||
val.y = getRandom<decltype(T::x)>();
|
||||
val.z = getRandom<decltype(T::x)>();
|
||||
val.w = getRandom<decltype(T::x)>();
|
||||
}
|
||||
|
||||
template<
|
||||
typename T,
|
||||
typename std::enable_if<rank<T>() == 1>::type* = nullptr>
|
||||
@@ -112,31 +72,6 @@ static inline void printVector(T &val) {
|
||||
std::cout << ")";
|
||||
}
|
||||
|
||||
template<
|
||||
typename T,
|
||||
typename std::enable_if<rank<T>() == 1>::type* = nullptr>
|
||||
static inline bool isEqual(const T &val0, const T &val1) {
|
||||
return val0.x == val1.x;
|
||||
}
|
||||
|
||||
template<
|
||||
typename T,
|
||||
typename std::enable_if<rank<T>() == 2>::type* = nullptr>
|
||||
static inline bool isEqual(const T &val0, const T &val1) {
|
||||
return val0.x == val1.x &&
|
||||
val0.y == val1.y;
|
||||
}
|
||||
|
||||
template<
|
||||
typename T,
|
||||
typename std::enable_if<rank<T>() == 4>::type* = nullptr>
|
||||
static inline bool isEqual(const T &val0, const T &val1) {
|
||||
return val0.x == val1.x &&
|
||||
val0.y == val1.y &&
|
||||
val0.z == val1.z &&
|
||||
val0.w == val1.w;
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
bool runTest() {
|
||||
const int N = 1024;
|
||||
@@ -144,7 +79,7 @@ bool runTest() {
|
||||
// Allocating the required buffer on gpu device
|
||||
T *texBuf, *texBufOut;
|
||||
T val[N], output[N];
|
||||
|
||||
hipGetLastError(); // Clear err due to negative tests
|
||||
memset(output, 0, sizeof(output));
|
||||
std::srand(std::time(nullptr)); // use current time as seed for random generator
|
||||
|
||||
|
||||
Referência em uma Nova Issue
Bloquear um usuário