SWDEV-252381 - Added dtest for hipGetSymbolAddress api (#2594)

Change-Id: Ideb58497c6598fe6f492a904b58bdc125d6d41d8

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 <hip_test_common.hh>
 #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<void **>(&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<void **>(&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));
+    }
+}