From de3a7dfe434cc015badf32c6df450aa8a919d649 Mon Sep 17 00:00:00 2001 From: ansurya <50609411+ansurya@users.noreply.github.com> Date: Fri, 19 Jul 2019 10:15:20 +0530 Subject: [PATCH] [HIP][Tests] Added new testcases for Module API (#1150) * [HIP][tests] New testcases for module api * [HIP][Tests]Support for CUDA devices * Updated tests as per latest master & test GetGlobal to work on all platforms [ROCm/hip commit: fa4d6b353ab31b70354f51c5ddcd24fb33264457] --- .../src/runtimeApi/module/global_kernel.cpp | 39 +++++ .../runtimeApi/module/hipModuleGetGlobal.cpp | 159 ++++++++++++++++++ .../runtimeApi/module/hipModuleLoadData.cpp | 98 +++++++++++ .../module/hipModuleTexture2dDrv.cpp | 144 ++++++++++++++++ .../src/runtimeApi/module/tex2d_kernel.cpp | 35 ++++ 5 files changed, 475 insertions(+) create mode 100644 projects/hip/tests/src/runtimeApi/module/global_kernel.cpp create mode 100644 projects/hip/tests/src/runtimeApi/module/hipModuleGetGlobal.cpp create mode 100644 projects/hip/tests/src/runtimeApi/module/hipModuleLoadData.cpp create mode 100644 projects/hip/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp create mode 100644 projects/hip/tests/src/runtimeApi/module/tex2d_kernel.cpp diff --git a/projects/hip/tests/src/runtimeApi/module/global_kernel.cpp b/projects/hip/tests/src/runtimeApi/module/global_kernel.cpp new file mode 100644 index 0000000000..d256d7ca10 --- /dev/null +++ b/projects/hip/tests/src/runtimeApi/module/global_kernel.cpp @@ -0,0 +1,39 @@ +/* +Copyright (c) 2017-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/hip_runtime.h" + +#define ARRAY_SIZE (16) + +__device__ float myDeviceGlobal; +__device__ float myDeviceGlobalArray[16]; + + +extern "C" __global__ void hello_world(const float* a, float* b) { + int tx = hipThreadIdx_x; + b[tx] = a[tx]; +} + +extern "C" __global__ void test_globals(const float* a, float* b) { + int tx = hipThreadIdx_x; + b[tx] = a[tx] + myDeviceGlobal + myDeviceGlobalArray[tx % ARRAY_SIZE]; +} diff --git a/projects/hip/tests/src/runtimeApi/module/hipModuleGetGlobal.cpp b/projects/hip/tests/src/runtimeApi/module/hipModuleGetGlobal.cpp new file mode 100644 index 0000000000..956ba4e50a --- /dev/null +++ b/projects/hip/tests/src/runtimeApi/module/hipModuleGetGlobal.cpp @@ -0,0 +1,159 @@ +/* +Copyright (c) 2017-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. +*/ + +/* HIT_START + * BUILD_CMD: global_kernel.code %hc --genco %S/global_kernel.cpp -o global_kernel.code + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 + * TEST: %t + * HIT_END + */ + +#include "hip/hip_runtime.h" +#include "hip/hip_runtime_api.h" +#include +#include +#include +#include + +#define LEN 64 +#define SIZE LEN * sizeof(float) + +#define fileName "global_kernel.code" +#define HIP_CHECK(cmd) \ + { \ + hipError_t status = cmd; \ + if (status != hipSuccess) { \ + std::cout << "error: #" << status << " (" << hipGetErrorString(status) \ + << ") at line:" << __LINE__ << ": " << #cmd << std::endl; \ + abort(); \ + } \ + } + +int main() { + float *A, *B; + float *Ad, *Bd; + A = new float[LEN]; + B = new float[LEN]; + + for (uint32_t i = 0; i < LEN; i++) { + A[i] = i * 1.0f; + B[i] = 0.0f; + } + + hipInit(0); + hipDevice_t device; + hipCtx_t context; + hipDeviceGet(&device, 0); + hipCtxCreate(&context, 0, device); + + hipMalloc((void**)&Ad, SIZE); + hipMalloc((void**)&Bd, SIZE); + + hipMemcpyHtoD(hipDeviceptr_t(Ad), A, SIZE); + hipMemcpyHtoD((hipDeviceptr_t)(Bd), B, SIZE); + hipModule_t Module; + HIP_CHECK(hipModuleLoad(&Module, fileName)); + + float myDeviceGlobal_h = 42.0; + hipDeviceptr_t deviceGlobal; + size_t deviceGlobalSize; + HIP_CHECK(hipModuleGetGlobal(&deviceGlobal, &deviceGlobalSize, Module, "myDeviceGlobal")); + HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(deviceGlobal), &myDeviceGlobal_h, deviceGlobalSize)); +#define ARRAY_SIZE 16 + float myDeviceGlobalArray_h[ARRAY_SIZE]; + hipDeviceptr_t myDeviceGlobalArray; + size_t myDeviceGlobalArraySize; + + HIP_CHECK(hipModuleGetGlobal((hipDeviceptr_t*)&myDeviceGlobalArray, &myDeviceGlobalArraySize, Module, "myDeviceGlobalArray")); + + for (int i = 0; i < ARRAY_SIZE; i++) { + myDeviceGlobalArray_h[i] = i * 1000.0f; + HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(myDeviceGlobalArray), &myDeviceGlobalArray_h, myDeviceGlobalArraySize)); + } + + struct { + void* _Ad; + void* _Bd; + } args; + + args._Ad = (void*) Ad; + args._Bd = (void*) Bd; + + size_t size = sizeof(args); + + void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, + HIP_LAUNCH_PARAM_END}; + + { + hipFunction_t Function; + HIP_CHECK(hipModuleGetFunction(&Function, Module, "hello_world")); + HIP_CHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config)); + + hipMemcpyDtoH(B, hipDeviceptr_t(Bd), SIZE); + + int mismatchCount = 0; + for (uint32_t i = 0; i < LEN; i++) { + if (A[i] != B[i]) { + mismatchCount++; + std::cout << "error: mismatch " << A[i] << " != " << B[i] << std::endl; + if (mismatchCount >= 10) { + break; + } + } + } + + if (mismatchCount == 0) { + std::cout << "PASSED!\n"; + } else { + std::cout << "FAILED!\n"; + }; + } + + { + hipFunction_t Function; + HIP_CHECK(hipModuleGetFunction(&Function, Module, "test_globals")); + HIP_CHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config)); + + hipMemcpyDtoH(B, hipDeviceptr_t(Bd), SIZE); + + int mismatchCount = 0; + for (uint32_t i = 0; i < LEN; i++) { + float expected = A[i] + myDeviceGlobal_h + + myDeviceGlobalArray_h[i % 16]; + if (expected != B[i]) { + mismatchCount++; + std::cout << "error: mismatch " << expected << " != " << B[i] << std::endl; + if (mismatchCount >= 10) { + break; + } + } + } + + if (mismatchCount == 0) { + std::cout << "PASSED!\n"; + } else { + std::cout << "FAILED!\n"; + }; + } + + hipCtxDestroy(context); + return 0; +} diff --git a/projects/hip/tests/src/runtimeApi/module/hipModuleLoadData.cpp b/projects/hip/tests/src/runtimeApi/module/hipModuleLoadData.cpp new file mode 100644 index 0000000000..30a0352f26 --- /dev/null +++ b/projects/hip/tests/src/runtimeApi/module/hipModuleLoadData.cpp @@ -0,0 +1,98 @@ +/* +Copyright (c) 2015-2016 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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 + * TEST: %t + * HIT_END + */ + +#include "hip/hip_runtime.h" +#include "hip/hip_runtime_api.h" +#include +#include +#include +#include +#include + +#include "test_common.h" + +#define LEN 64 +#define SIZE LEN << 2 + +#define FILENAME "vcpy_kernel.code" +#define kernel_name "hello_world" + +int main() { + float *A, *B, *Ad, *Bd; + A = new float[LEN]; + B = new float[LEN]; + + for (uint32_t i = 0; i < LEN; i++) { + A[i] = i * 1.0f; + B[i] = 0.0f; + } + + HIPCHECK(hipInit(0)); + HIPCHECK(hipMalloc((void**)&Ad, SIZE)); + HIPCHECK(hipMalloc((void**)&Bd, SIZE)); + + HIPCHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice)); + HIPCHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice)); + + hipModule_t Module; + hipFunction_t Function; + std::ifstream file(FILENAME, std::ios::binary | std::ios::ate); + std::streamsize fsize = file.tellg(); + file.seekg(0, std::ios::beg); + + std::vector buffer(fsize); + if (file.read(buffer.data(), fsize)) { + HIPCHECK(hipModuleLoadData(&Module, &buffer[0])); + HIPCHECK(hipModuleGetFunction(&Function, Module, kernel_name)); + } + else { + failed("could not open code object '%s'\n", FILENAME); + } + + hipStream_t stream; + HIPCHECK(hipStreamCreate(&stream)); + + struct { + void* _Ad; + void* _Bd; + } args; + args._Ad = (void*) Ad; + args._Bd = (void*) Bd; + size_t size = sizeof(args); + + void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, + HIP_LAUNCH_PARAM_END}; + HIPCHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, stream, NULL, (void**)&config)); + + HIPCHECK(hipStreamDestroy(stream)); + + HIPCHECK(hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost)); + + for (uint32_t i = 0; i < LEN; i++) { + assert(A[i] == B[i]); + } + + passed(); +} diff --git a/projects/hip/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp b/projects/hip/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp new file mode 100644 index 0000000000..e678c35d6e --- /dev/null +++ b/projects/hip/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp @@ -0,0 +1,144 @@ +/* +Copyright (c) 2015-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. +*/ + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc + * TEST: %t + * HIT_END + */ + +#include "hip/hip_runtime.h" +//#include "hip/hip_runtime_api.h" +#include +#include +#include +//#include + +#define fileName "tex2d_kernel.code" + +texture tex; +bool testResult = false; + +#define HIP_CHECK(cmd) \ + { \ + hipError_t status = cmd; \ + if (status != hipSuccess) { \ + std::cout << "error: #" << status << " (" << hipGetErrorString(status) \ + << ") at line:" << __LINE__ << ": " << #cmd << std::endl; \ + abort(); \ + } \ + } + +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; + } + } + 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; + 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; + hipMemcpyParam2D(©Param); + + textureReference* texref; + hipModuleGetTexRef(&texref, Module, "tex"); + hipTexRefSetAddressMode(texref, 0, hipAddressModeWrap); + hipTexRefSetAddressMode(texref, 1, hipAddressModeWrap); + hipTexRefSetFilterMode(texref, hipFilterModePoint); + hipTexRefSetFlags(texref, 0); + hipTexRefSetFormat(texref, HIP_AD_FORMAT_FLOAT, 1); + hipTexRefSetArray(texref, array, HIP_TRSA_OVERRIDE_FORMAT); + + float* dData = NULL; + hipMalloc((void**)&dData, size); + + 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); + + 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")); + + 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); + 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; + } + } + } + hipFree(dData); + hipFreeArray(array); + return true; +} + +int main(int argc, char** argv) { + hipInit(0); + testResult = runTest(argc, argv); + printf("%s ...\n", testResult ? "PASSED" : "FAILED"); + exit(testResult ? EXIT_SUCCESS : EXIT_FAILURE); + return 0; +} diff --git a/projects/hip/tests/src/runtimeApi/module/tex2d_kernel.cpp b/projects/hip/tests/src/runtimeApi/module/tex2d_kernel.cpp new file mode 100644 index 0000000000..b12dd1815d --- /dev/null +++ b/projects/hip/tests/src/runtimeApi/module/tex2d_kernel.cpp @@ -0,0 +1,35 @@ +/* +Copyright (c) 2015-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. +*/ + +/* HIT_START + * BUILD_CMD: tex2d_kernel.code %hc --genco %S/tex2d_kernel.cpp -o tex2d_kernel.code + * HIT_END + */ + +#include "hip/hip_runtime.h" +extern texture tex; + +extern "C" __global__ void tex2dKernel(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(tex, x, y); +}