AI/rocm-systems
2
0
Rozštěpit 0
Zdrojový kód Úkoly Pull requesty Akce Balíčky Projekty Vydání Wiki Aktivita

[dtest] Tests for loading hip apis dynamically

Procházet zdrojové kódy 
SWDEV-238517 for enhancing hip unit tests

Change-Id: I9b6ef41db2f02ee3679360d0fae709a404a0ea4f
Tento commit je obsažen v:
Sudheer Kumar
2020-08-27 01:29:32 +05:30
odevzdal Mohan Kumar Mithur
rodič 3c4208aa7d
revize e112d1e13f
3 změnil soubory, kde provedl 481 přidání a 127 odebrání
Stáhněte soubor záplaty Stáhněte rozdílový soubor Rozbalit všechny soubory Sbalit všechny soubory
tests/src/dynamicLoading/complex_loading_behavior.cpp
+284 -127
Zobrazit soubor
@@ -1,6 +1,5 @@
/*
Copyright (c) 2019 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
@@ -20,10 +19,13 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/* Test for loading device kernels from a library created with extern "C" function
*/
/* HIT_START
* BUILD_CMD: libfoo_amd %hc %S/%s -o libfoo.so -Xcompiler -fPIC -lpthread -shared -DTEST_SHARED_LIBRARY EXCLUDE_HIP_PLATFORM nvcc
* BUILD_CMD: libfoo_nvidia %hc %S/%s -o libfoo.so -Xcompiler -fPIC -lpthread -shared -DTEST_SHARED_LIBRARY EXCLUDE_HIP_PLATFORM hcc rocclr
* BUILD_CMD: %t %hc %S/%s -o %T/%t -ldl
* BUILD_CMD: libLazyLoad_amd %hc %S/%s -o liblazyLoad.so -I%S/.. -fPIC -lpthread -shared -DTEST_SHARED_LIBRARY EXCLUDE_HIP_PLATFORM nvcc
* BUILD_CMD: libLazyLoad_nvidia %hc %S/%s --std=c++11 -o liblazyLoad.so -I%S/.. -Xcompiler -fPIC -lpthread -shared -DTEST_SHARED_LIBRARY EXCLUDE_HIP_PLATFORM hcc rocclr
* BUILD_CMD: %t %hc %S/%s --std=c++11 -o %T/%t -I%S/.. -ldl
* TEST: %t
* HIT_END
*/
@@ -32,17 +34,7 @@ THE SOFTWARE.
#include <dlfcn.h>
#include <iostream>
#include <hip/hip_runtime.h>
#define CHECK(cmd) \
{ \
hipError_t error = cmd; \
if (error != hipSuccess) { \
fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error, \
__FILE__, __LINE__); \
return (EXIT_FAILURE); \
} \
}
#include "test_common.h"
__global__ void vector_add(float* C, float* A, float* B, size_t N) {
size_t offset = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
@@ -52,160 +44,325 @@ __global__ void vector_add(float* C, float* A, float* B, size_t N) {
}
}
int launch_local_kernel() {
float *A_d, *B_d, *C_d;
float *A_h, *B_h, *C_h;
size_t N = 1000000;
size_t Nbytes = N * sizeof(float);
static int device = 0;
CHECK(hipSetDevice(device));
hipDeviceProp_t props;
CHECK(hipGetDeviceProperties(&props, device /*deviceID*/));
A_h = (float*)malloc(Nbytes);
CHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess);
B_h = (float*)malloc(Nbytes);
CHECK(B_h == 0 ? hipErrorOutOfMemory : hipSuccess);
C_h = (float*)malloc(Nbytes);
CHECK(C_h == 0 ? hipErrorOutOfMemory : hipSuccess);
// Fill with Phi + i
for (size_t i = 0; i < N; i++) {
A_h[i] = 1.618f + i;
B_h[i] = 1.618f + i;
bool launch_local_kernel() {
bool testResult = true;
float *A_d, *B_d, *C_d;
float *A_h, *B_h, *C_h;
size_t N = 1000000;
size_t Nbytes = N * sizeof(float);
static int device = 0;
HIPCHECK(hipSetDevice(device));
hipDeviceProp_t props;
HIPCHECK(hipGetDeviceProperties(&props, device /*deviceID*/));
A_h = reinterpret_cast<float*>(malloc(Nbytes));
HIPCHECK(A_h == nullptr ? hipErrorOutOfMemory : hipSuccess);
B_h = reinterpret_cast<float*>(malloc(Nbytes));
HIPCHECK(B_h == nullptr ? hipErrorOutOfMemory : hipSuccess);
C_h = reinterpret_cast<float*>(malloc(Nbytes));
HIPCHECK(C_h == nullptr ? hipErrorOutOfMemory : hipSuccess);
// Fill with Phi + i
for (size_t i = 0; i < N; i++) {
A_h[i] = 1.618f + i;
B_h[i] = 1.618f + i;
}
HIPCHECK(hipMalloc(&A_d, Nbytes));
HIPCHECK(hipMalloc(&B_d, Nbytes));
HIPCHECK(hipMalloc(&C_d, Nbytes));
HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
HIPCHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
const unsigned blocks = 512;
const unsigned threadsPerBlock = 256;
hipLaunchKernelGGL(vector_add, dim3(blocks), dim3(threadsPerBlock),
0, 0, C_d, A_d, B_d, N);
HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
for (size_t i=0; i < N ; i++) {
if (C_h[i] != (A_h[i] + B_h[i])) {
printf("data mismatch. Local kernel failed");
testResult = false;
break;
}
}
CHECK(hipMalloc(&A_d, Nbytes));
CHECK(hipMalloc(&B_d, Nbytes));
CHECK(hipMalloc(&C_d, Nbytes));
CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
HIPCHECK(hipFree(A_d));
HIPCHECK(hipFree(B_d));
HIPCHECK(hipFree(C_d));
const unsigned blocks = 512;
const unsigned threadsPerBlock = 256;
hipLaunchKernelGGL(vector_add, dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, B_d, N);
CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
free(A_h);
free(B_h);
free(C_h);
CHECK(hipFree(A_d));
CHECK(hipFree(B_d));
CHECK(hipFree(C_d));
free(A_h);
free(B_h);
free(C_h);
std::cout << "PASSED!\n";
return 0;
std::cout << "Local kernel executed successfully\n";
return testResult;
}
int launch_dynamically_loaded_kernel() {
void* handle = dlopen("./libfoo.so", RTLD_LAZY);
bool launch_dynamically_loaded_kernel() {
bool testResult = true;
int ret = 1;
void* handle = dlopen("./liblazyLoad.so", RTLD_LAZY);
if (!handle) {
std::cout << dlerror() << "\n";
return -1;
testResult = false;
return testResult;
}
std::cout << "loaded libfoo.so\n";
void* sym = dlsym(handle, "foo");
std::cout << "loaded liblazyLoad.so\n";
void* sym = dlsym(handle, "lazyLoad");
if (!sym) {
std::cout << "unable to locate foo within libfoo.so\n";
std::cout << "unable to locate lazyLoad within lazyLoad.so\n";
std::cout << dlerror() << "\n";
dlclose(handle);
return -1;
testResult = false;
return testResult;
}
int(*fp)() = reinterpret_cast<int(*)()>(sym);
int ret = fp();
if (ret) {
ret = fp();
if (ret == 0) {
std::cout << "dynamic launch failed\n";
testResult = false;
} else {
std::cout << "dynamic launch succeeded\n";
}
dlclose(handle);
return ret;
return testResult;
}
int main() {
int ret = 0;
ret = launch_local_kernel();
if (ret) {
return ret;
}
bool testResult = true;
ret = launch_dynamically_loaded_kernel();
if (ret) {
return ret;
}
testResult &= launch_local_kernel();
testResult &= launch_dynamically_loaded_kernel();
return 0;
if (testResult == true) {
passed();
} else {
failed("One or more tests failed");
}
}
#else // !defined(TEST_SHARED_LIBRARY)
#else // !defined(TEST_SHARED_LIBRARY)
#include <dlfcn.h>
#include <iostream>
#include <hip/hip_runtime.h>
#include "test_common.h"
#define CHECK(cmd) \
{ \
hipError_t error = cmd; \
if (error != hipSuccess) { \
fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error, \
__FILE__, __LINE__); \
return (EXIT_FAILURE); \
} \
}
__global__ void vadd(float* C, float* A, float* B, size_t N) {
__global__ void vAdd(float* C, float* A, float* B, size_t N) {
size_t offset = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
size_t stride = hipBlockDim_x * hipGridDim_x;
for (size_t i = offset; i < N; i += stride) {
C[i] = A[i] + B[i];
}
}
extern "C" int foo() {
float *A_d, *B_d, *C_d;
float *A_h, *B_h, *C_h;
size_t N = 1000000;
size_t Nbytes = N * sizeof(float);
static int device = 0;
CHECK(hipSetDevice(device));
hipDeviceProp_t props;
CHECK(hipGetDeviceProperties(&props, device /*deviceID*/));
A_h = (float*)malloc(Nbytes);
CHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess);
B_h = (float*)malloc(Nbytes);
CHECK(B_h == 0 ? hipErrorOutOfMemory : hipSuccess);
C_h = (float*)malloc(Nbytes);
CHECK(C_h == 0 ? hipErrorOutOfMemory : hipSuccess);
// Fill with Phi + i
for (size_t i = 0; i < N; i++) {
A_h[i] = 1.618f + i;
B_h[i] = 1.618f + i;
int vectorAddKernelTest() {
int testResult = 1;
float *A_d, *B_d, *C_d;
float *A_h, *B_h, *C_h;
size_t N = 1000000;
size_t Nbytes = N * sizeof(float);
static int device = 0;
HIPCHECK(hipSetDevice(device));
hipDeviceProp_t props;
HIPCHECK(hipGetDeviceProperties(&props, device /*deviceID*/));
A_h = reinterpret_cast<float*>(malloc(Nbytes));
HIPCHECK(A_h == nullptr ? hipErrorOutOfMemory : hipSuccess);
B_h = reinterpret_cast<float*>(malloc(Nbytes));
HIPCHECK(B_h == nullptr ? hipErrorOutOfMemory : hipSuccess);
C_h = reinterpret_cast<float*>(malloc(Nbytes));
HIPCHECK(C_h == nullptr ? hipErrorOutOfMemory : hipSuccess);
// Fill with Phi + i
for (size_t i = 0; i < N; i++) {
A_h[i] = 1.618f + i;
B_h[i] = 1.618f + i;
}
HIPCHECK(hipMalloc(&A_d, Nbytes));
HIPCHECK(hipMalloc(&B_d, Nbytes));
HIPCHECK(hipMalloc(&C_d, Nbytes));
HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
HIPCHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
const unsigned blocks = 512;
const unsigned threadsPerBlock = 256;
std::cout << "info: Launching vAdd kernel\n";
hipLaunchKernelGGL(vAdd, dim3(blocks), dim3(threadsPerBlock),
0, 0, C_d, A_d, B_d, N);
HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
for (size_t i=0; i < N ; i++) {
if (C_h[i] != (A_h[i] + B_h[i])) {
printf("info: data mismatch. vAdd kernel failed");
testResult = 0;
break;
}
}
CHECK(hipMalloc(&A_d, Nbytes));
CHECK(hipMalloc(&B_d, Nbytes));
CHECK(hipMalloc(&C_d, Nbytes));
CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
if (testResult) {
std::cout << "info: vAdd kernel executed fine\n";
}
const unsigned blocks = 512;
const unsigned threadsPerBlock = 256;
std::cout << "Launch vadd\n";
hipLaunchKernelGGL(vadd, dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, B_d, N);
CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
HIPCHECK(hipFree(A_d));
HIPCHECK(hipFree(B_d));
HIPCHECK(hipFree(C_d));
CHECK(hipFree(A_d));
CHECK(hipFree(B_d));
CHECK(hipFree(C_d));
free(A_h);
free(B_h);
free(C_h);
return 0;
free(A_h);
free(B_h);
free(C_h);
return testResult;
}
#endif // !defined(TEST_SHARED_LIBRARY)
#include "hip/hip_runtime.h"
#include "hip/hip_runtime_api.h"
#include "hip/hip_cooperative_groups.h"
namespace cg = cooperative_groups;
static const uint BufferSizeInDwords = 448 * 1024 * 1024;
__global__ void test_gws(uint* buf, uint bufSize,
long* tmpBuf, long* result) {
extern __shared__ long tmp[];
uint offset = blockIdx.x * blockDim.x + threadIdx.x;
uint stride = gridDim.x * blockDim.x;
cg::grid_group gg = cg::this_grid();
long sum = 0;
for (uint i = offset; i < bufSize; i += stride) {
sum += buf[i];
}
tmp[threadIdx.x] = sum;
__syncthreads();
if (threadIdx.x == 0) {
sum = 0;
for (uint i = 0; i < blockDim.x; i++) {
sum += tmp[i];
}
tmpBuf[blockIdx.x] = sum;
}
gg.sync();
if (offset == 0) {
for (uint i = 1; i < gridDim.x; ++i) {
sum += tmpBuf[i];
}
*result = sum;
}
}
int cooperativeKernelTest() {
int testResult = 1;
uint* dA;
long* dB;
long* dC;
long* Ah;
hipDeviceProp_t deviceProp;
hipGetDeviceProperties(&deviceProp, 0);
if (!deviceProp.cooperativeLaunch) {
std::cout << "info: Device doesn't support cooperative launch!"
"skipping the test!\n";
return testResult;
}
uint32_t* init = new uint32_t[BufferSizeInDwords];
for (uint32_t i = 0; i < BufferSizeInDwords; ++i) {
init[i] = i;
}
std::cout << "info: Launch kernel to test hipLaunchCooperativeKernel api\n";
std::cout << "info: running on bus 0x" << deviceProp.pciBusID << " " <<
deviceProp.name << "\n";
size_t SIZE = BufferSizeInDwords * sizeof(uint);
HIPCHECK(hipMalloc(reinterpret_cast<void**>(&dA), SIZE));
HIPCHECK(hipMalloc(reinterpret_cast<void**>(&dC), sizeof(long)));
HIPCHECK(hipMemcpy(dA, init, SIZE, hipMemcpyHostToDevice));
Ah = reinterpret_cast<long*>(malloc(sizeof(long)));
hipStream_t stream;
HIPCHECK(hipStreamCreate(&stream));
dim3 dimBlock = dim3(1);
dim3 dimGrid = dim3(1);
int numBlocks = 0;
uint workgroups[4] = {32, 64, 128, 256};
for (uint i = 0; i < 4; ++i) {
dimBlock.x = workgroups[i];
/* Calculate the device occupancy to know how many blocks can be
run concurrently */
hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocks,
test_gws, dimBlock.x * dimBlock.y * dimBlock.z, dimBlock.x * sizeof(long));
dimGrid.x = deviceProp.multiProcessorCount * std::min(numBlocks, 32);
HIPCHECK(hipMalloc(reinterpret_cast<void**>(&dB),
dimGrid.x * sizeof(long)));
void *params[4];
params[0] = reinterpret_cast<void*>(&dA);
params[1] = (void*)&BufferSizeInDwords;
params[2] = reinterpret_cast<void*>(&dB);
params[3] = reinterpret_cast<void*>(&dC);
std::cout << "Testing with grid size = " << dimGrid.x <<
" and block size = " << dimBlock.x << "\n";
HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(test_gws),
dimGrid, dimBlock, params,
dimBlock.x * sizeof(long), stream));
HIPCHECK(hipMemcpy(Ah, dC, sizeof(long), hipMemcpyDeviceToHost));
if (*Ah != (((long)(BufferSizeInDwords) * (BufferSizeInDwords - 1)) / 2)) {
std::cout << "Data validation failed for grid size = " << dimGrid.x <<
" and block size = " << dimBlock.x << "\n";
HIPCHECK(hipFree(dB));
std::cout << "Test failed! \n";
testResult = 0;
break;
} else {
std::cout << "info: data validated!\n";
HIPCHECK(hipFree(dB));
}
}
if (testResult) {
std::cout <<"info: hipLaunchCooperativeKernel api executed fine\n";
}
HIPCHECK(hipStreamDestroy(stream));
HIPCHECK(hipFree(dC));
HIPCHECK(hipFree(dA));
delete [] init;
free(Ah);
return testResult;
}
extern "C" int lazyLoad() {
return vectorAddKernelTest() & cooperativeKernelTest();
}
#endif // !defined(TEST_SHARED_LIBRARY)