Add tests for hipGetSymbolAddress/hipGetSymbolSize (#2662)
Этот коммит содержится в:
Jatin Chaudhary
коммит произвёл
GitHub
GitHub
родитель
574f65f03d
Коммит
a4e1ca3bdc
@@ -26,117 +26,131 @@ THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#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<int>(i);
|
||||
B[i] = 0;
|
||||
C[i] = 0;
|
||||
Am[i] = -1 * static_cast<int>(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<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 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));
|
||||
}
|
||||
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));
|
||||
}
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipGetSymbolAddress_Negative") {
|
||||
SECTION("Invalid symbol") {
|
||||
int notADeviceSymbol{0};
|
||||
int* addr{nullptr};
|
||||
HIP_CHECK_ERROR(
|
||||
hipGetSymbolAddress(reinterpret_cast<void**>(&addr), HIP_SYMBOL(notADeviceSymbol)),
|
||||
hipErrorInvalidSymbol);
|
||||
}
|
||||
|
||||
SECTION("Nullptr symbol") {
|
||||
int* addr{nullptr};
|
||||
HIP_CHECK_ERROR(hipGetSymbolAddress(reinterpret_cast<void**>(&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);
|
||||
}
|
||||
}
|
||||
|
||||
Ссылка в новой задаче
Block a user