diff --git a/projects/hip-tests/catch/unit/deviceLib/hipTestDeviceSymbol.cc b/projects/hip-tests/catch/unit/deviceLib/hipTestDeviceSymbol.cc index 501a0805cb..d34d88c846 100644 --- a/projects/hip-tests/catch/unit/deviceLib/hipTestDeviceSymbol.cc +++ b/projects/hip-tests/catch/unit/deviceLib/hipTestDeviceSymbol.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2022 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 @@ -18,8 +18,12 @@ THE SOFTWARE. */ -/* Test Case Description: Calling hipMemcpyTo/FromSymbolAsync() using user - declared stream obj and hipStreamPerThread*/ +/* + Test Scenarios : + 1) Calling hipMemcpyTo/FromSymbolAsync() using user declared stream obj and hipStreamPerThread. + 2) Validate get symbol address/size for global const array. + 3) Validate get symbol address/size for static const variable. +*/ #include #define NUM 1024 @@ -34,12 +38,20 @@ __global__ void Assign(int* Out) { globalOut[tid] = globalIn[tid]; } -__device__ __constant__ int globalConst[NUM]; +__device__ __constant__ int globalConstArr[NUM]; +__device__ __constant__ static float statConstVar = 1.0f; -__global__ void checkAddress(int* addr, bool* out) { - *out = (globalConst == addr); +__global__ void checkGlobalConstAddress(int* addr, bool* out) { + *out = (globalConstArr == addr); } +__global__ void checkStaticConstVarAddress(float* addr, bool* out) { + *out = (&statConstVar == addr); +} + +/** + Calling hipMemcpyTo/FromSymbolAsync() using user declared stream obj and hipStreamPerThread. + */ TEST_CASE("Unit_hipMemcpyToSymbolAsync_ToNFrom") { int *A, *Am, *B, *Ad, *C, *Cm; A = new int[NUM]; @@ -51,9 +63,9 @@ TEST_CASE("Unit_hipMemcpyToSymbolAsync_ToNFrom") { C[i] = 0; } - HIP_CHECK(hipMalloc((void**)&Ad, SIZE)); - HIP_CHECK(hipHostMalloc((void**)&Am, SIZE)); - HIP_CHECK(hipHostMalloc((void**)&Cm, SIZE)); + HIP_CHECK(hipMalloc(&Ad, SIZE)); + HIP_CHECK(hipHostMalloc(&Am, SIZE)); + HIP_CHECK(hipHostMalloc(&Cm, SIZE)); for (int i = 0; i < NUM; i++) { Am[i] = -1 * i; Cm[i] = 0; @@ -70,8 +82,8 @@ TEST_CASE("Unit_hipMemcpyToSymbolAsync_ToNFrom") { hipMemcpyDeviceToHost, stream)); HIP_CHECK(hipStreamSynchronize(stream)); for (int i = 0; i < NUM; i++) { - assert(Am[i] == B[i]); - assert(Am[i] == Cm[i]); + REQUIRE(Am[i] == B[i]); + REQUIRE(Am[i] == Cm[i]); } for (int i = 0; i < NUM; i++) { @@ -86,8 +98,8 @@ TEST_CASE("Unit_hipMemcpyToSymbolAsync_ToNFrom") { HIP_CHECK(hipMemcpyFromSymbol(C, HIP_SYMBOL(globalOut), SIZE, 0, hipMemcpyDeviceToHost)); for (int i = 0; i < NUM; i++) { - assert(A[i] == B[i]); - assert(A[i] == C[i]); + REQUIRE(A[i] == B[i]); + REQUIRE(A[i] == C[i]); } for (int i = 0; i < NUM; i++) { @@ -106,32 +118,19 @@ TEST_CASE("Unit_hipMemcpyToSymbolAsync_ToNFrom") { } SECTION("Calling hipMemcpyTo/FromSymbol using hipStreamPerThread") { HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), A, SIZE, 0, - hipMemcpyHostToDevice, hipStreamPerThread)); + hipMemcpyHostToDevice, hipStreamPerThread)); HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad); HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost)); HIP_CHECK(hipMemcpyFromSymbolAsync(C, HIP_SYMBOL(globalOut), SIZE, 0, - hipMemcpyDeviceToHost, hipStreamPerThread)); + hipMemcpyDeviceToHost, hipStreamPerThread)); HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); } for (int i = 0; i < NUM; i++) { - assert(A[i] == B[i]); - assert(A[i] == C[i]); + REQUIRE(A[i] == B[i]); + REQUIRE(A[i] == C[i]); } - - bool *checkOkD; - bool checkOk = false; - size_t symbolSize = 0; - int *symbolAddress; - HIP_CHECK(hipGetSymbolSize(&symbolSize, HIP_SYMBOL(globalConst))); - HIP_CHECK(hipGetSymbolAddress((void**) &symbolAddress, HIP_SYMBOL(globalConst))); - HIP_CHECK(hipMalloc((void**)&checkOkD, sizeof(bool))); - hipLaunchKernelGGL(checkAddress, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, symbolAddress, checkOkD); - HIP_CHECK(hipMemcpy(&checkOk, checkOkD, sizeof(bool), hipMemcpyDeviceToHost)); - HIP_CHECK(hipFree(checkOkD)); - HIP_ASSERT(checkOk); - HIP_ASSERT((symbolSize == SIZE)); - + HIP_CHECK(hipStreamDestroy(stream)); HIP_CHECK(hipHostFree(Am)); HIP_CHECK(hipHostFree(Cm)); HIP_CHECK(hipFree(Ad)); @@ -139,3 +138,45 @@ TEST_CASE("Unit_hipMemcpyToSymbolAsync_ToNFrom") { delete[] B; delete[] C; } + +/** + 1) Validate get symbol address/size for global const array. + 2) Validate get symbol address/size for static const variable. + */ +TEST_CASE("Unit_hipGetSymbolAddressAndSize_Validation") { + bool *checkOkD; + bool checkOk = false; + size_t symbolSize{}; + int *symbolArrAddress{}; + float *symbolVarAddress{}; + + SECTION("Validate symbol size/address of global const array") { + HIP_CHECK(hipGetSymbolSize(&symbolSize, HIP_SYMBOL(globalConstArr))); + HIP_CHECK(hipGetSymbolAddress( + reinterpret_cast(&symbolArrAddress), + HIP_SYMBOL(globalConstArr))); + HIP_CHECK(hipMalloc(&checkOkD, sizeof(bool))); + hipLaunchKernelGGL(checkGlobalConstAddress, dim3(1, 1, 1), dim3(1, 1, 1), + 0, 0, symbolArrAddress, checkOkD); + HIP_CHECK(hipMemcpy(&checkOk, checkOkD, sizeof(bool), + hipMemcpyDeviceToHost)); + HIP_CHECK(hipFree(checkOkD)); + HIP_ASSERT(checkOk); + HIP_ASSERT(symbolSize == SIZE); + } + + SECTION("Validate symbol size/address of static const variable") { + HIP_CHECK(hipGetSymbolSize(&symbolSize, HIP_SYMBOL(statConstVar))); + HIP_CHECK(hipGetSymbolAddress( + reinterpret_cast(&symbolVarAddress), + HIP_SYMBOL(statConstVar))); + HIP_CHECK(hipMalloc(&checkOkD, sizeof(bool))); + hipLaunchKernelGGL(checkStaticConstVarAddress, dim3(1, 1, 1), + dim3(1, 1, 1), 0, 0, symbolVarAddress, checkOkD); + HIP_CHECK(hipMemcpy(&checkOk, checkOkD, sizeof(bool), + hipMemcpyDeviceToHost)); + HIP_CHECK(hipFree(checkOkD)); + HIP_ASSERT(checkOk); + HIP_ASSERT(symbolSize == sizeof(float)); + } +}