From fb0bef74c1ab8fdfd1118d116d39c4c37bbfa50b Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 2 Jun 2022 17:01:14 +0530 Subject: [PATCH] SWDEV-331248 - Add more image tests in sample (#2709) In samples/2_Cookbook/11_texture_driver, add Vector data types(char4, short4, int4, float4); More arithmetic data types(char, short, int); Change-Id: I54aa482213d340d32cf912601adead0812c2323a --- .../11_texture_driver/tex2dKernel.cpp | 59 +++- .../11_texture_driver/texture2dDrv.cpp | 261 ++++++++++++------ 2 files changed, 233 insertions(+), 87 deletions(-) diff --git a/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp b/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp index 5f2ded8518..120f31c610 100644 --- a/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp +++ b/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp @@ -19,15 +19,62 @@ 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/hip_runtime.h" -texture tex; +texture texChar; +texture texShort; +texture texInt; +texture texFloat; -extern "C" __global__ void tex2dKernel(float* outputData, int width, int height) { -#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT +texture texChar4; +texture texShort4; +texture texInt4; +texture texFloat4; + +extern "C" __global__ void tex2dKernelChar(char* outputData, int width, int height) { int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y; - outputData[y * width + x] = tex2D(tex, x, y); -#endif + outputData[y * width + x] = tex2D(texChar, x, y); +} + +extern "C" __global__ void tex2dKernelShort(short* outputData, int width, int height) { + int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; + int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y; + outputData[y * width + x] = tex2D(texShort, x, y); +} + +extern "C" __global__ void tex2dKernelInt(int* outputData, int width, int height) { + int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; + int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y; + outputData[y * width + x] = tex2D(texInt, x, y); +} + +extern "C" __global__ void tex2dKernelFloat(float* outputData, int width, int height) { + int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; + int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y; + outputData[y * width + x] = tex2D(texFloat, x, y); +} + +extern "C" __global__ void tex2dKernelChar4(char4* outputData, int width, int height) { + int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; + int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y; + outputData[y * width + x] = tex2D(texChar4, x, y); +} + +extern "C" __global__ void tex2dKernelShort4(short4* outputData, int width, int height) { + int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; + int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y; + outputData[y * width + x] = tex2D(texShort4, x, y); +} + +extern "C" __global__ void tex2dKernelInt4(int4* outputData, int width, int height) { + int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; + int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y; + outputData[y * width + x] = tex2D(texInt4, x, y); +} + +extern "C" __global__ void tex2dKernelFloat4(float4* outputData, int width, int height) { + int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; + int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y; + outputData[y * width + x] = tex2D(texFloat4, x, y); } diff --git a/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp b/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp index 744ca2d965..fbeac3d41a 100644 --- a/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp +++ b/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp @@ -39,92 +39,181 @@ bool testResult = true; } \ } -bool runTest(int argc, char** argv) { - unsigned int width = 256; - unsigned int height = 256; - unsigned int size = width * height * sizeof(float); - float* hData = (float*)malloc(size); - memset(hData, 0, size); - for (int i = 0; i < height; i++) { - for (int j = 0; j < width; j++) { - hData[i * width + j] = i * width + j; - } +template::value>::type *t = nullptr> +static inline hipArray_Format getArrayFormat() { + if (std::is_same::value) { + return HIP_AD_FORMAT_SIGNED_INT8; + } else if (std::is_same::value) { + return HIP_AD_FORMAT_SIGNED_INT16; + } else if (std::is_same::value) { + return HIP_AD_FORMAT_SIGNED_INT32; + } else if (std::is_same::value) { + return HIP_AD_FORMAT_FLOAT; + } + return HIP_AD_FORMAT_HALF; +} + +template::value>::type *t = nullptr> +static inline hipArray_Format getArrayFormat() { + return getArrayFormat(); +} + +template +static inline constexpr int rank() { + return sizeof(T) / sizeof(decltype(T::x)); +} + +template +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(std::numeric_limits < T > ::max() * r); +} + +template::value>::type* = nullptr> +static inline constexpr int getChannels() { + return 1; +} + +template::value>::type *t = nullptr, + typename std::enable_if() != 0>::type *r = nullptr> +static inline constexpr int getChannels() { + return rank(); +} + +template::value>::type* = nullptr> +static inline void printDiff(const int &i, const int &j, const T &expected, + const T &output) { + std::cout << "Difference [" << i << " " << j << "]: " << expected << " - " + << output << "\n"; +} + +template::value>::type* = nullptr, + typename std::enable_if() == 4>::type* = nullptr> +static inline void printDiff(const int &i, const int &j, const T &expected, + const T &output) { + std::cout << "Difference [" << i << " " << j << "]: " << expected.x << "," + << expected.y << "," << expected.z << "," << expected.w << " - " + << output.x << "," << output.y << "," << output.z << "," << output.w + << "\n"; +} + +template::value>::type* = nullptr> +static inline void initVal(T &val) { + val = getRandom(); +} + +template::value>::type* = nullptr, + typename std::enable_if() == 4>::type* = nullptr> +static inline void initVal(T &val) { + val.x = getRandom(); + val.y = getRandom(); + val.z = getRandom(); + val.w = getRandom(); +} + +template +bool runTest(hipModule_t &module, const char *refName, const char *funcName) { + hipArray_Format format = getArrayFormat(); + int channels = getChannels(); + unsigned int width = 256; + unsigned int height = 256; + 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]); } - hipModule_t Module; - HIP_CHECK(hipModuleLoad(&Module, fileName)); + } - hipArray* array; - HIP_ARRAY_DESCRIPTOR desc; - desc.Format = HIP_AD_FORMAT_FLOAT; - desc.NumChannels = 1; - desc.Width = width; - desc.Height = height; - HIP_CHECK(hipArrayCreate(&array, &desc)); + hipArray *array; + HIP_ARRAY_DESCRIPTOR desc; + desc.Format = format; + desc.NumChannels = channels; + desc.Width = width; + desc.Height = height; + HIP_CHECK(hipArrayCreate(&array, &desc)); - hip_Memcpy2D copyParam; - memset(©Param, 0, sizeof(copyParam)); - copyParam.dstMemoryType = hipMemoryTypeArray; - copyParam.dstArray = array; - copyParam.srcMemoryType = hipMemoryTypeHost; - copyParam.srcHost = hData; - copyParam.srcPitch = width * sizeof(float); - copyParam.WidthInBytes = copyParam.srcPitch; - copyParam.Height = height; - HIP_CHECK(hipMemcpyParam2D(©Param)); + hip_Memcpy2D copyParam; + memset(©Param, 0, sizeof(copyParam)); + copyParam.dstMemoryType = hipMemoryTypeArray; + copyParam.dstArray = array; + copyParam.srcMemoryType = hipMemoryTypeHost; + copyParam.srcHost = hData; + copyParam.srcPitch = width * sizeof(T); + copyParam.WidthInBytes = copyParam.srcPitch; + copyParam.Height = height; + HIP_CHECK(hipMemcpyParam2D(©Param)); - textureReference* texref; - HIP_CHECK(hipModuleGetTexRef(&texref, Module, "tex")); - HIP_CHECK(hipTexRefSetAddressMode(texref, 0, hipAddressModeWrap)); - HIP_CHECK(hipTexRefSetAddressMode(texref, 1, hipAddressModeWrap)); - HIP_CHECK(hipTexRefSetFilterMode(texref, hipFilterModePoint)); - HIP_CHECK(hipTexRefSetFlags(texref, 0)); - HIP_CHECK(hipTexRefSetFormat(texref, HIP_AD_FORMAT_FLOAT, 1)); - HIP_CHECK(hipTexRefSetArray(texref, array, HIP_TRSA_OVERRIDE_FORMAT)); + textureReference *texref; + HIP_CHECK(hipModuleGetTexRef(&texref, module, refName)); + HIP_CHECK(hipTexRefSetAddressMode(texref, 0, hipAddressModeClamp)); + HIP_CHECK(hipTexRefSetAddressMode(texref, 1, hipAddressModeClamp)); + HIP_CHECK(hipTexRefSetFilterMode(texref, hipFilterModePoint)); + HIP_CHECK(hipTexRefSetFlags(texref, HIP_TRSF_READ_AS_INTEGER)); + HIP_CHECK(hipTexRefSetFormat(texref, format, channels)); + HIP_CHECK(hipTexRefSetArray(texref, array, HIP_TRSA_OVERRIDE_FORMAT)); - float* dData = NULL; - HIP_CHECK(hipMalloc((void**)&dData, size)); + T *dData = NULL; + HIP_CHECK(hipMalloc((void** )&dData, size)); - struct { - void* _Ad; - unsigned int _Bd; - unsigned int _Cd; - } args; - args._Ad = (void*) dData; - args._Bd = width; - args._Cd = height; + struct { + void *_Ad; + unsigned int _Bd; + unsigned int _Cd; + } args; + args._Ad = (void*) dData; + args._Bd = width; + args._Cd = height; - size_t sizeTemp = sizeof(args); + size_t sizeTemp = sizeof(args); - void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, - &sizeTemp, HIP_LAUNCH_PARAM_END}; + void *config[] = { HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &sizeTemp, HIP_LAUNCH_PARAM_END }; - hipFunction_t Function; - HIP_CHECK(hipModuleGetFunction(&Function, Module, "tex2dKernel")); + hipFunction_t Function; + HIP_CHECK(hipModuleGetFunction(&Function, module, funcName)); - int temp1 = width / 16; - int temp2 = height / 16; - HIP_CHECK( - hipModuleLaunchKernel(Function, 16, 16, 1, temp1, temp2, 1, 0, 0, NULL, (void**)&config)); - hipDeviceSynchronize(); + int temp1 = width / 16; + int temp2 = height / 16; + HIP_CHECK( + hipModuleLaunchKernel(Function, 16, 16, 1, temp1, temp2, 1, 0, 0, NULL, + (void** )&config)); + hipDeviceSynchronize(); - float* hOutputData = (float*)malloc(size); - memset(hOutputData, 0, size); - HIP_CHECK(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost)); + T *hOutputData = (T*) malloc(size); + memset(hOutputData, 0, size); + HIP_CHECK(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost)); - for (int i = 0; i < height; i++) { - for (int j = 0; j < width; j++) { - if (hData[i * width + j] != hOutputData[i * width + j]) { - printf("Difference [ %d %d ]:%f ----%f\n", i, j, hData[i * width + j], - hOutputData[i * width + j]); - testResult = false; - break; - } - } + for (int i = 0; i < height; i++) { + for (int j = 0; j < width; j++) { + if (hData[i * width + j] != hOutputData[i * width + j]) { + printDiff(i, j, hData[i * width + j], hOutputData[i * width + j]); + testResult = false; + break; + } } - HIP_CHECK(hipUnbindTexture(texref)); - HIP_CHECK(hipFree(dData)); - HIP_CHECK(hipFreeArray(array)); - return testResult; + } + HIP_CHECK(hipUnbindTexture(texref)); + HIP_CHECK(hipFree(dData)); + HIP_CHECK(hipFreeArray(array)); + free(hOutputData); + free(hData); + printf("%s test %s ...\n", funcName, testResult ? "PASSED" : "FAILED"); + return testResult; } inline bool isImageSupported() { @@ -137,13 +226,23 @@ inline bool isImageSupported() { } int main(int argc, char** argv) { - if (!isImageSupported()) { - printf("Texture is not support on the device. Skipped.\n"); - return 0; - } - hipInit(0); - testResult = runTest(argc, argv); - printf("%s ...\n", testResult ? "PASSED" : "FAILED"); - exit(testResult ? EXIT_SUCCESS : EXIT_FAILURE); + if (!isImageSupported()) { + printf("Texture is not support on the device. Skipped.\n"); return 0; + } + hipInit(0); + hipModule_t module; + HIP_CHECK(hipModuleLoad(&module, fileName)); + testResult = testResult && runTest(module, "texChar", "tex2dKernelChar"); + testResult = testResult && runTest(module, "texShort", "tex2dKernelShort"); + testResult = testResult && runTest(module, "texInt", "tex2dKernelInt"); + testResult = testResult && runTest(module, "texFloat", "tex2dKernelFloat"); + testResult = testResult && runTest(module, "texChar4", "tex2dKernelChar4"); + testResult = testResult && runTest(module, "texShort4", "tex2dKernelShort4"); + testResult = testResult && runTest(module, "texInt4", "tex2dKernelInt4"); + testResult = testResult && runTest(module, "texFloat4", "tex2dKernelFloat4"); + + HIP_CHECK(hipModuleUnload(module)); + printf("texture2dDrv %s ...\n", testResult ? "PASSED" : "FAILED"); + return testResult ? EXIT_SUCCESS : EXIT_FAILURE; }