From 548ec20ffdc139d230076768f49795daf450910b Mon Sep 17 00:00:00 2001 From: Lakhan singh Thakur Date: Thu, 1 Oct 2020 21:33:32 +0530 Subject: [PATCH] [dtest] Adding Test case for hipFuncSetSharedMemConfig() and hipFuncSetAttribute() SWDEV-251504 - QA task for validation Change-Id: I11ed98f648352cb32161c41290516c77231236bc --- .../runtimeApi/module/hipFuncSetAttribute.cpp | 46 ++++++++ .../module/hipFuncSetSharedMemConfig.cpp | 101 ++++++++++++++++-- 2 files changed, 141 insertions(+), 6 deletions(-) create mode 100644 tests/src/runtimeApi/module/hipFuncSetAttribute.cpp diff --git a/tests/src/runtimeApi/module/hipFuncSetAttribute.cpp b/tests/src/runtimeApi/module/hipFuncSetAttribute.cpp new file mode 100644 index 0000000000..df86f34deb --- /dev/null +++ b/tests/src/runtimeApi/module/hipFuncSetAttribute.cpp @@ -0,0 +1,46 @@ +/* +Copyright (c) 2020 - 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 "test_common.h" + +__global__ void fn(float* px, float* py) { + bool a[42]; + __shared__ double b[69]; + + for (auto&& x : b) x = *py++; + for (auto&& x : a) x = *px++ > 0.0; + for (auto&& x : a) if (x) *--py = *--px; +} + +int main() { + HIPCHECK(hipFuncSetAttribute(reinterpret_cast(&fn), + hipFuncAttributeMaxDynamicSharedMemorySize, + 0)); + HIPCHECK(hipFuncSetAttribute(reinterpret_cast(&fn), + hipFuncAttributePreferredSharedMemoryCarveout, + 0)); + passed(); +} diff --git a/tests/src/runtimeApi/module/hipFuncSetSharedMemConfig.cpp b/tests/src/runtimeApi/module/hipFuncSetSharedMemConfig.cpp index 3f04563c71..46114af865 100644 --- a/tests/src/runtimeApi/module/hipFuncSetSharedMemConfig.cpp +++ b/tests/src/runtimeApi/module/hipFuncSetSharedMemConfig.cpp @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2020 - 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 @@ -17,11 +17,100 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#include +/* HIT_START + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 + * TEST: %t + * HIT_END + */ + + + +// Test Description: +// This test case verifies the working of hipFuncSetSharedMemConfig() api and +// the flag parameter + #include "test_common.h" -int main() { - hipSharedMemConfig_t config; - HIP_PRINT_STATUS(hipFuncSetSharedMemConfig(NULL)); - HIP_PRINT_STATUS(hipFuncSetSharedMemConfig(&config)); + +__global__ void ReverseSeq(int *A, int *B, int N) { + extern __shared__ int SMem[]; + int offset = threadIdx.x; + int MirrorVal = N - offset - 1; + SMem[offset] = A[offset]; + __syncthreads(); + B[offset] = SMem[MirrorVal]; +} + +int main() { + bool IfTestPassed = true; + int *Ah = NULL, *RAh = NULL, NELMTS = 128; + int *Ad = NULL, *RAd = NULL; + Ah = reinterpret_cast(malloc(NELMTS * sizeof(int))); + RAh = reinterpret_cast(malloc(NELMTS * sizeof(int))); + HIPCHECK(hipMalloc(&Ad, NELMTS * sizeof(int))); + HIPCHECK(hipMalloc(&RAd, NELMTS * sizeof(int))); + for (int i = 0; i < NELMTS; ++i) { + Ah[i] = i; + RAh[i] = NELMTS - i - 1; + } + HIPCHECK(hipMemcpy(Ad, Ah, NELMTS * sizeof(int), hipMemcpyHostToDevice)); + HIPCHECK(hipMemset(RAd, 0, NELMTS * sizeof(int))); + // Testing hipFuncSetSharedMemConfig() with hipSharedMemBankSizeDefault flag + HIPCHECK(hipFuncSetSharedMemConfig(reinterpret_cast(&ReverseSeq), + hipSharedMemBankSizeDefault)); + // Kernel Launch with shared mem size of = NELMTS * sizeof(int) + ReverseSeq<<<1, NELMTS, NELMTS * sizeof(int)>>>(Ad, RAd, NELMTS); + memset(Ah, 0, NELMTS * sizeof(int)); + // Verifying the results + HIPCHECK(hipMemcpy(Ah, RAd, NELMTS * sizeof(int), hipMemcpyDeviceToHost)); + for (int i = 0; i < NELMTS; ++i) { + if (Ah[i] != RAh[i]) { + printf("Mismatch found at %d value of array\n", i); + printf(" after setting the flag hipSharedMemBankSizeDefault\n"); + IfTestPassed = false; + } + } + // Testing hipFuncSetSharedMemConfig() with hipSharedMemBankSizeFourBytes flg + HIPCHECK(hipFuncSetSharedMemConfig(reinterpret_cast(&ReverseSeq), + hipSharedMemBankSizeFourByte)); + HIPCHECK(hipMemset(RAd, 0, NELMTS * sizeof(int))); + // Kernel Launch with shared mem size of = NELMTS * sizeof(int) + ReverseSeq<<<1, NELMTS, NELMTS * sizeof(int)>>>(Ad, RAd, NELMTS); + memset(Ah, 0, NELMTS * sizeof(int)); + // Verifying the results + HIPCHECK(hipMemcpy(Ah, RAd, NELMTS * sizeof(int), hipMemcpyDeviceToHost)); + for (int i = 0; i < NELMTS; ++i) { + if (Ah[i] != RAh[i]) { + printf("Mismatch found at %d value of array\n", i); + printf(" after setting the flag hipSharedMemBankSizeFourByte\n"); + IfTestPassed = false; + } + } + // Testing hipFuncSetSharedMemConfig() with hipSharedMemBankSizeEightBytes flg + HIPCHECK(hipFuncSetSharedMemConfig(reinterpret_cast(&ReverseSeq), + hipSharedMemBankSizeEightByte)); + HIPCHECK(hipMemset(RAd, 0, NELMTS * sizeof(int))); + // Kernel Launch with shared mem size of = NELMTS * sizeof(int) + ReverseSeq<<<1, NELMTS, NELMTS * sizeof(int)>>>(Ad, RAd, NELMTS); + memset(Ah, 0, NELMTS * sizeof(int)); + // Verifying the results + HIPCHECK(hipMemcpy(Ah, RAd, NELMTS * sizeof(int), hipMemcpyDeviceToHost)); + for (int i = 0; i < NELMTS; ++i) { + if (Ah[i] != RAh[i]) { + printf("Mismatch found at %d value of array\n", i); + printf(" after setting the flag hipSharedMemBankSizeEightByte\n"); + IfTestPassed = false; + } + } + + free(Ah); + free(RAh); + HIPCHECK(hipFree(Ad)); + HIPCHECK(hipFree(RAd)); + + if (IfTestPassed) { + passed(); + } else { + failed(""); + } }