From a4e1ca3bdc80937dc5b18559c442b911796cb0ae Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Wed, 25 May 2022 11:05:55 +0100 Subject: [PATCH] Add tests for hipGetSymbolAddress/hipGetSymbolSize (#2662) --- catch/unit/deviceLib/hipTestDeviceSymbol.cc | 291 +++++++++++--------- 1 file changed, 165 insertions(+), 126 deletions(-) diff --git a/catch/unit/deviceLib/hipTestDeviceSymbol.cc b/catch/unit/deviceLib/hipTestDeviceSymbol.cc index d34d88c846..d1c39b600a 100644 --- a/catch/unit/deviceLib/hipTestDeviceSymbol.cc +++ b/catch/unit/deviceLib/hipTestDeviceSymbol.cc @@ -26,117 +26,131 @@ THE SOFTWARE. */ #include -#define NUM 1024 -#define SIZE 1024 * 4 + +constexpr size_t NUM = 1024; +constexpr size_t SIZE = 1024 * 4; __device__ int globalIn[NUM]; __device__ int globalOut[NUM]; __global__ void Assign(int* Out) { - int tid = threadIdx.x + blockIdx.x * blockDim.x; - Out[tid] = globalIn[tid]; - globalOut[tid] = globalIn[tid]; + int tid = threadIdx.x + blockIdx.x * blockDim.x; + Out[tid] = globalIn[tid]; + globalOut[tid] = globalIn[tid]; } -__device__ __constant__ int globalConstArr[NUM]; -__device__ __constant__ static float statConstVar = 1.0f; +__device__ __constant__ int globalConst[NUM]; -__global__ void checkGlobalConstAddress(int* addr, bool* out) { - *out = (globalConstArr == addr); -} +__global__ void checkAddress(int* addr, bool* out) { *out = (globalConst == 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]; - B = new int[NUM]; - C = new int[NUM]; - for (int i = 0; i < NUM; i++) { - A[i] = -1 * i; - B[i] = 0; - C[i] = 0; - } + int *A{nullptr}, *Am{nullptr}, *B{nullptr}, *Ad{nullptr}, *C{nullptr}, *Cm{nullptr}; + A = new int[NUM]; + B = new int[NUM]; + C = new int[NUM]; - 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; - } + HIP_CHECK(hipMalloc((void**)&Ad, SIZE)); + HIP_CHECK(hipHostMalloc((void**)&Am, SIZE)); + HIP_CHECK(hipHostMalloc((void**)&Cm, SIZE)); - hipStream_t stream; + for (size_t i = 0; i < NUM; i++) { + A[i] = -1 * static_cast(i); + B[i] = 0; + C[i] = 0; + Am[i] = -1 * static_cast(i); + Cm[i] = 0; + } + + + SECTION("Calling hipMemcpyTo/FromSymbol using stream") { + hipStream_t stream{}; HIP_CHECK(hipStreamCreate(&stream)); - HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), Am, SIZE, 0, - hipMemcpyHostToDevice, stream)); + HIP_CHECK( + hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), Am, SIZE, 0, hipMemcpyHostToDevice, stream)); HIP_CHECK(hipStreamSynchronize(stream)); hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad); HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost)); - HIP_CHECK(hipMemcpyFromSymbolAsync(Cm, HIP_SYMBOL(globalOut), SIZE, 0, - hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipMemcpyFromSymbolAsync(Cm, HIP_SYMBOL(globalOut), SIZE, 0, hipMemcpyDeviceToHost, + stream)); HIP_CHECK(hipStreamSynchronize(stream)); - for (int i = 0; i < NUM; i++) { - REQUIRE(Am[i] == B[i]); - REQUIRE(Am[i] == Cm[i]); - } - - for (int i = 0; i < NUM; i++) { - A[i] = -2 * i; - B[i] = 0; - } - - HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(globalIn), A, SIZE, 0, - hipMemcpyHostToDevice)); - hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad); - HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost)); - HIP_CHECK(hipMemcpyFromSymbol(C, HIP_SYMBOL(globalOut), SIZE, 0, - hipMemcpyDeviceToHost)); - for (int i = 0; i < NUM; i++) { - REQUIRE(A[i] == B[i]); - REQUIRE(A[i] == C[i]); - } - - for (int i = 0; i < NUM; i++) { - A[i] = -3 * i; - B[i] = 0; - } - SECTION("Calling hipMemcpyTo/FromSymbol using user declared stream obj") { - HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), A, SIZE, 0, - hipMemcpyHostToDevice, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - 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, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - } - SECTION("Calling hipMemcpyTo/FromSymbol using hipStreamPerThread") { - HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), A, SIZE, 0, - 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)); - HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); - } - for (int i = 0; i < NUM; i++) { - REQUIRE(A[i] == B[i]); - REQUIRE(A[i] == C[i]); - } HIP_CHECK(hipStreamDestroy(stream)); - HIP_CHECK(hipHostFree(Am)); - HIP_CHECK(hipHostFree(Cm)); - HIP_CHECK(hipFree(Ad)); - delete[] A; - delete[] B; - delete[] C; + for (size_t i = 0; i < NUM; i++) { + REQUIRE(Am[i] == B[i]); + REQUIRE(Am[i] == Cm[i]); + } + } + + SECTION("Calling hipMemcpyTo/FromSymbol - validate value in host memory") { + HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(globalIn), A, SIZE, 0, hipMemcpyHostToDevice)); + hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad); + HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpyFromSymbol(C, HIP_SYMBOL(globalOut), SIZE, 0, hipMemcpyDeviceToHost)); + + for (size_t i = 0; i < NUM; i++) { + REQUIRE(A[i] == B[i]); + REQUIRE(A[i] == C[i]); + } + } + + SECTION("Calling hipMemcpyTo/FromSymbol using user declared stream obj") { + hipStream_t stream{}; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK( + hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), A, SIZE, 0, hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + 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, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipStreamDestroy(stream)); + + for (size_t i = 0; i < NUM; i++) { + REQUIRE(A[i] == B[i]); + REQUIRE(A[i] == C[i]); + } + } + + SECTION("Calling hipMemcpyTo/FromSymbol using hipStreamPerThread") { + HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), A, SIZE, 0, 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)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + + for (size_t i = 0; i < NUM; i++) { + REQUIRE(A[i] == B[i]); + REQUIRE(A[i] == C[i]); + } + } + + // Check for address on GPU and CPU side and compare it + // If address mismatch report error + // Validate size of symbol as well, compare it with output of hipGetSymbolSize + SECTION("Validate address on GPU") { + bool* checkOkD{nullptr}; + bool checkOk = false; + size_t symbolSize = 0; + int* symbolAddress{nullptr}; + 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(hipHostFree(Am)); + HIP_CHECK(hipHostFree(Cm)); + HIP_CHECK(hipFree(Ad)); + delete[] A; + delete[] B; + delete[] C; } /** @@ -144,39 +158,64 @@ TEST_CASE("Unit_hipMemcpyToSymbolAsync_ToNFrom") { 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{}; + bool* checkOkD{nullptr}; + 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 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)); - } + 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)); + } +} + +TEST_CASE("Unit_hipGetSymbolAddress_Negative") { + SECTION("Invalid symbol") { + int notADeviceSymbol{0}; + int* addr{nullptr}; + HIP_CHECK_ERROR( + hipGetSymbolAddress(reinterpret_cast(&addr), HIP_SYMBOL(notADeviceSymbol)), + hipErrorInvalidSymbol); + } + + SECTION("Nullptr symbol") { + int* addr{nullptr}; + HIP_CHECK_ERROR(hipGetSymbolAddress(reinterpret_cast(&addr), nullptr), + hipErrorInvalidSymbol); + } +} + +TEST_CASE("Unit_hipGetSymbolSize_Negative") { + SECTION("Invalid symbol") { + int notADeviceSymbol{0}; + size_t dsize{0}; + HIP_CHECK_ERROR(hipGetSymbolSize(&dsize, HIP_SYMBOL(notADeviceSymbol)), hipErrorInvalidSymbol); + } + + SECTION("Nullptr symbol") { + size_t size{0}; + HIP_CHECK_ERROR(hipGetSymbolSize(&size, nullptr), hipErrorInvalidSymbol); + } }