AI/rocm-systems
2
0
انشعاب 0
کد مسائل تقاضاهای واکشی Actions Packages پروژه‌ها انتشارها دانشنامه فعالیت

[dtest] Tests for hipModuleGetTexRef() API

فهرست منبع 
1. Negative Test Cases
2. API functionality test using multiple streams.
3. API functionality validation using multiple GPUs.
4. Setting hipTexRefSetFlags() with HIP_TRSF_READ_AS_INTEGER flag instead of 0 as suggested in SWDEV-256096.
5. Enabling test cases 0x03 (short) and 0x04 (char).
6. Implemented external code review comments.

SWDEV-238517 for enhancing hip unit tests

Change-Id: If42796047ec1cf2e3695dc2b7f40a2d9dd50f5bd


[ROCm/hip commit: 9d92c90b55]
This commit is contained in:
Rupam Chetia
2020-10-04 20:45:09 +05:30
کامیت شده توسط Mohan Kumar Mithur
والد 47b1005aac
کامیت 3e2dc0d39f
2فایلهای تغییر یافته به همراه634 افزوده شده و 100 حذف شده
دانلود پرونده وصله دانلود فایل تغییرات diff Expand all files Collapse all files
projects/hip/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp
+594 -89
مشاهده پرونده
@@ -22,123 +22,628 @@ THE SOFTWARE.
/* HIT_START
* BUILD_CMD: tex2d_kernel.code %hc --genco %S/tex2d_kernel.cpp -o tex2d_kernel.code
* BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvidia EXCLUDE_HIP_RUNTIME rocclr
* TEST: %t
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
* TEST: %t --tests 0x01
* TEST: %t --tests 0x02
* TEST: %t --tests 0x03
* TEST: %t --tests 0x04
* TEST: %t --tests 0x05
* TEST: %t --tests 0x06
* TEST: %t --tests 0x07
* TEST: %t --tests 0x10
* TEST: %t --tests 0x11
* TEST: %t --tests 0x12
* TEST: %t --tests 0x13
* TEST: %t --tests 0x14 EXCLUDE_HIP_PLATFORM amd
* TEST: %t --tests 0x15
* HIT_END
*/
#include "hip/hip_runtime.h"
#include <iostream>
#include <fstream>
#include <vector>
#include <type_traits>
#include <limits>
#include <atomic>
#include "test_common.h"
#define fileName "tex2d_kernel.code"
#define CODEOBJ_FILE "tex2d_kernel.code"
#define NON_EXISTING_TEX_NAME "xyz"
#define EMPTY_TEX_NAME ""
#define GLOBAL_KERNEL_VAR "deviceGlobalFloat"
#define TEX_REF "ftex"
#define WIDTH 256
#define HEIGHT 256
#define MAX_STREAMS 4
#define GRIDDIMX 16
#define GRIDDIMY 16
#define GRIDDIMZ 1
#define BLOCKDIMZ 1
texture<float, 2, hipReadModeElementType> tex;
bool testResult = false;
#ifdef __HIP_PLATFORM_NVCC__
#define HIP_CHECK(cmd) \
{ \
hipError_t status = cmd; \
if (status != hipSuccess) { \
std::cout << "error: #" << status << " (" << hipGetErrorString(status) \
<< ") at line:" << __LINE__ << ": " << #cmd << std::endl; \
abort(); \
} \
}
#define CTX_CREATE() \
hipCtx_t context;\
initHipCtx(&context);
bool runTest(int argc, char** argv) {
unsigned int width = 256;
unsigned int height = 256;
unsigned int size = width * height * sizeof(float);
float* hData = (float*)malloc(size);
memset(hData, 0, size);
#define CTX_DESTROY() HIPCHECK(hipCtxDestroy(context));
#define ARRAY_DESTROY(array) HIPCHECK(hipArrayDestroy(array));
#define HIP_TEX_REFERENCE hipTexRef
#define HIP_ARRAY hiparray
/**
* Internal Function
*/
void initHipCtx(hipCtx_t *pcontext) {
HIPCHECK(hipInit(0));
hipDevice_t device;
HIPCHECK(hipDeviceGet(&device, 0));
HIPCHECK(hipCtxCreate(pcontext, 0, device));
}
#else // __HIP_PLATFORM_NVCC__
#define CTX_CREATE()
#define CTX_DESTROY()
#define ARRAY_DESTROY(array) HIPCHECK(hipFreeArray(array));
#define HIP_TEX_REFERENCE textureReference*
#define HIP_ARRAY hipArray*
#endif // __HIP_PLATFORM_NVCC__
std::atomic<int> g_thTestPassed(1);
/**
* Validates negative scenarios for hipModuleGetTexRef
* texRef = nullptr
*/
bool testTexRefEqNullPtr() {
bool TestPassed = false;
hipModule_t Module;
CTX_CREATE()
HIPCHECK(hipModuleLoad(&Module, CODEOBJ_FILE));
if (hipSuccess != hipModuleGetTexRef(nullptr, Module, "tex")) {
TestPassed = true;
} else {
printf("Test Failed as texRef = nullptr returns hipSuccess \n");
}
CTX_DESTROY()
return TestPassed;
}
/**
* Validates negative scenarios for hipModuleGetTexRef
* name = nullptr
*/
bool testNameEqNullPtr() {
bool TestPassed = false;
hipModule_t Module;
HIP_TEX_REFERENCE texref;
CTX_CREATE()
HIPCHECK(hipModuleLoad(&Module, CODEOBJ_FILE));
if (hipSuccess != hipModuleGetTexRef(&texref, Module, nullptr)) {
TestPassed = true;
} else {
printf("Test Failed as name = nullptr returns hipSuccess \n");
}
CTX_DESTROY()
return TestPassed;
}
/**
* Validates negative scenarios for hipModuleGetTexRef
* name = Non Existing Tex Name
*/
bool testInvalidTexName() {
bool TestPassed = false;
hipModule_t Module;
HIP_TEX_REFERENCE texref;
CTX_CREATE()
HIPCHECK(hipModuleLoad(&Module, CODEOBJ_FILE));
if (hipSuccess != hipModuleGetTexRef(&texref, Module,
NON_EXISTING_TEX_NAME)) {
TestPassed = true;
} else {
printf("Test Failed as invalid tex ref returns hipSuccess \n");
}
CTX_DESTROY()
return TestPassed;
}
/**
* Validates negative scenarios for hipModuleGetTexRef
* name = Empty Tex Name
*/
bool testEmptyTexName() {
bool TestPassed = false;
hipModule_t Module;
HIP_TEX_REFERENCE texref;
CTX_CREATE()
HIPCHECK(hipModuleLoad(&Module, CODEOBJ_FILE));
if (hipSuccess != hipModuleGetTexRef(&texref, Module, EMPTY_TEX_NAME)) {
TestPassed = true;
} else {
printf("Test Failed as empty tex ref returns hipSuccess \n");
}
CTX_DESTROY()
return TestPassed;
}
/**
* Validates negative scenarios for hipModuleGetTexRef
* name = Global Kernel Variable
*/
bool testWrongTexRef() {
bool TestPassed = false;
hipModule_t Module;
HIP_TEX_REFERENCE texref;
CTX_CREATE()
HIPCHECK(hipModuleLoad(&Module, CODEOBJ_FILE));
if (hipSuccess != hipModuleGetTexRef(&texref, Module, GLOBAL_KERNEL_VAR)) {
TestPassed = true;
} else {
printf("Test Failed as global tex ref returns hipSuccess \n");
}
CTX_DESTROY()
return TestPassed;
}
/**
* Validates negative scenarios for hipModuleGetTexRef
* module = unloaded module
*/
bool testUnloadedMod() {
bool TestPassed = false;
hipModule_t Module;
HIP_TEX_REFERENCE texref;
CTX_CREATE()
HIPCHECK(hipModuleLoad(&Module, CODEOBJ_FILE));
HIPCHECK(hipModuleUnload(Module));
if (hipSuccess != hipModuleGetTexRef(&texref, Module, TEX_REF)) {
TestPassed = true;
} else {
printf("Test Failed as unloaded module returns hipSuccess \n");
}
CTX_DESTROY()
return TestPassed;
}
/**
* Internal Functions
*
*/
std::vector<char> load_file() {
std::ifstream file(CODEOBJ_FILE, std::ios::binary | std::ios::ate);
std::streamsize fsize = file.tellg();
file.seekg(0, std::ios::beg);
std::vector<char> buffer(fsize);
if (!file.read(buffer.data(), fsize)) {
failed("could not open code object '%s'\n", CODEOBJ_FILE);
}
return buffer;
}
template <class T> void fillTestBuffer(unsigned int width,
unsigned int height,
T* hData) {
if (std::is_same<T, float>::value) {
for (int i = 0; i < height; i++) {
for (int j = 0; j < width; j++) {
hData[i * width + j] = i * width + j;
}
for (int j = 0; j < width; j++) {
hData[i * width + j] = i * width + j + 0.5;
}
}
hipModule_t Module;
HIP_CHECK(hipModuleLoad(&Module, fileName));
} else if (std::is_same<T, int>::value) {
for (int i = 0; i < height; i++) {
for (int j = 0; j < width; j++) {
hData[i * width + j] = i * width + j;
}
}
} else if (std::is_same<T, short>::value) {
for (int i = 0; i < height; i++) {
for (int j = 0; j < width; j++) {
hData[i * width + j] = (i * width + j)%
(std::numeric_limits<short>::max());
}
}
} else if (std::is_same<T, char>::value) {
for (int i = 0; i < height; i++) {
for (int j = 0; j < width; j++) {
hData[i * width + j] = (i * width + j)%
(std::numeric_limits<char>::max());
}
}
}
}
hipArray* array;
HIP_ARRAY_DESCRIPTOR desc;
desc.Format = HIP_AD_FORMAT_FLOAT;
desc.NumChannels = 1;
desc.Width = width;
desc.Height = height;
hipArrayCreate(&array, &desc);
void allocInitArray(unsigned int width,
unsigned int height,
hipArray_Format format,
HIP_ARRAY* array
) {
HIP_ARRAY_DESCRIPTOR desc;
desc.Format = format;
desc.NumChannels = 1;
desc.Width = width;
desc.Height = height;
HIPCHECK(hipArrayCreate(array, &desc));
}
hip_Memcpy2D copyParam;
memset(&copyParam, 0, sizeof(copyParam));
copyParam.dstMemoryType = hipMemoryTypeArray;
copyParam.dstArray = array;
copyParam.srcMemoryType = hipMemoryTypeHost;
copyParam.srcHost = hData;
copyParam.srcPitch = width * sizeof(float);
copyParam.WidthInBytes = copyParam.srcPitch;
copyParam.Height = height;
hipMemcpyParam2D(&copyParam);
template <class T, class T1> void copyBuffer2Array(unsigned int width,
unsigned int height,
T* hData,
T1 array
) {
hip_Memcpy2D copyParam;
memset(&copyParam, 0, sizeof(copyParam));
#ifdef __HIP_PLATFORM_NVCC__
copyParam.dstMemoryType = CU_MEMORYTYPE_ARRAY;
copyParam.srcMemoryType = CU_MEMORYTYPE_HOST;
copyParam.dstArray = *array;
#else
copyParam.dstMemoryType = hipMemoryTypeArray;
copyParam.srcMemoryType = hipMemoryTypeHost;
copyParam.dstArray = array;
#endif
copyParam.srcHost = hData;
copyParam.srcPitch = width * sizeof(T);
copyParam.WidthInBytes = copyParam.srcPitch;
copyParam.Height = height;
HIPCHECK(hipMemcpyParam2D(&copyParam));
}
textureReference* texref;
hipModuleGetTexRef(&texref, Module, "tex");
hipTexRefSetAddressMode(texref, 0, hipAddressModeWrap);
hipTexRefSetAddressMode(texref, 1, hipAddressModeWrap);
hipTexRefSetFilterMode(texref, hipFilterModePoint);
hipTexRefSetFlags(texref, 0);
hipTexRefSetFormat(texref, HIP_AD_FORMAT_FLOAT, 1);
hipTexRefSetArray(texref, array, HIP_TRSA_OVERRIDE_FORMAT);
template <class T> void assignArray2TexRef(hipArray_Format format,
const char* texRefName,
hipModule_t Module,
T array
) {
HIP_TEX_REFERENCE texref;
#ifdef __HIP_PLATFORM_NVCC__
HIPCHECK(hipModuleGetTexRef(&texref, Module, texRefName));
HIPCHECK(hipTexRefSetAddressMode(texref, 0, CU_TR_ADDRESS_MODE_WRAP));
HIPCHECK(hipTexRefSetAddressMode(texref, 1, CU_TR_ADDRESS_MODE_WRAP));
HIPCHECK(hipTexRefSetFilterMode(texref, HIP_TR_FILTER_MODE_POINT));
HIPCHECK(hipTexRefSetFlags(texref, CU_TRSF_READ_AS_INTEGER));
HIPCHECK(hipTexRefSetFormat(texref, format, 1));
HIPCHECK(hipTexRefSetArray(texref, *array, CU_TRSA_OVERRIDE_FORMAT));
#else
HIPCHECK(hipModuleGetTexRef(&texref, Module, texRefName));
HIPCHECK(hipTexRefSetAddressMode(texref, 0, hipAddressModeWrap));
HIPCHECK(hipTexRefSetAddressMode(texref, 1, hipAddressModeWrap));
HIPCHECK(hipTexRefSetFilterMode(texref, hipFilterModePoint));
HIPCHECK(hipTexRefSetFlags(texref, HIP_TRSF_READ_AS_INTEGER));
HIPCHECK(hipTexRefSetFormat(texref, format, 1));
HIPCHECK(hipTexRefSetArray(texref, array, HIP_TRSA_OVERRIDE_FORMAT));
#endif
}
float* dData = NULL;
hipMalloc((void**)&dData, size);
template <class T> bool validateOutput(unsigned int width,
unsigned int height,
T* hData,
T* hOutputData) {
for (int i = 0; i < height; i++) {
for (int j = 0; j < width; j++) {
if (hData[i * width + j] != hOutputData[i * width + j]) {
std::cout << "Difference [ " << i << " " << j << "]:" <<
(int)hData[i * width + j] << "---" << (int)hOutputData[i * width + j]
<< std::endl;
return false;
}
}
}
return true;
}
/**
* Validates texture type data functionality for hipModuleGetTexRef
*
*/
template <class T> bool testTexType(hipArray_Format format,
const char* texRefName,
const char* kerFuncName) {
bool TestPassed = true;
unsigned int width = WIDTH;
unsigned int height = HEIGHT;
unsigned int size = width * height * sizeof(T);
T* hData = reinterpret_cast<T*>(malloc(size));
if (NULL == hData) {
printf("Failed to allocate using malloc in testTexType.\n");
return false;
}
CTX_CREATE()
fillTestBuffer<T>(width, height, hData);
// Load Kernel File and create hipArray
hipModule_t Module;
HIPCHECK(hipModuleLoad(&Module, CODEOBJ_FILE));
HIP_ARRAY array;
allocInitArray(width, height, format, &array);
#ifdef __HIP_PLATFORM_NVCC__
// Copy from hData to array using hipMemcpyParam2D
copyBuffer2Array<T, HIP_ARRAY*>(width, height, hData, &array);
// Get tex reference from the loaded kernel file
// Assign array to the tex reference
assignArray2TexRef<HIP_ARRAY*>(format, texRefName, Module, &array);
#else
// Copy from hData to array using hipMemcpyParam2D
copyBuffer2Array<T, HIP_ARRAY>(width, height, hData, array);
// Get tex reference from the loaded kernel file
// Assign array to the tex reference
assignArray2TexRef<HIP_ARRAY>(format, texRefName, Module, array);
#endif
hipFunction_t Function;
HIPCHECK(hipModuleGetFunction(&Function, Module, kerFuncName));
T* dData = NULL;
HIPCHECK(hipMalloc(reinterpret_cast<void**>(&dData), size));
struct {
void* _Ad;
unsigned int _Bd;
unsigned int _Cd;
} args;
args._Ad = reinterpret_cast<void*>(dData);
args._Bd = width;
args._Cd = height;
size_t sizeTemp = sizeof(args);
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER,
&args,
HIP_LAUNCH_PARAM_BUFFER_SIZE,
&sizeTemp,
HIP_LAUNCH_PARAM_END};
int temp1 = width / GRIDDIMX;
int temp2 = height / GRIDDIMY;
HIPCHECK(
hipModuleLaunchKernel(Function, GRIDDIMX, GRIDDIMY, GRIDDIMZ,
temp1, temp2, BLOCKDIMZ, 0, 0,
NULL, reinterpret_cast<void**>(&config)));
HIPCHECK(hipDeviceSynchronize());
T* hOutputData = reinterpret_cast<T*>(malloc(size));
if (NULL == hOutputData) {
printf("Failed to allocate using malloc in testTexType.\n");
TestPassed = false;
} else {
memset(hOutputData, 0, size);
HIPCHECK(hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost));
TestPassed = validateOutput<T>(width, height, hData, hOutputData);
}
free(hOutputData);
HIPCHECK(hipFree(dData));
ARRAY_DESTROY(array)
HIPCHECK(hipModuleUnload(Module));
free(hData);
CTX_DESTROY()
return TestPassed;
}
/**
* Validates texture functionality with multiple streams for hipModuleGetTexRef
*
*/
template <class T> bool testTexMultStream(const std::vector<char>& buffer,
hipArray_Format format,
const char* texRefName,
const char* kerFuncName,
unsigned int numOfStreams) {
bool TestPassed = true;
unsigned int width = WIDTH;
unsigned int height = HEIGHT;
unsigned int size = width * height * sizeof(T);
T* hData = reinterpret_cast<T*>(malloc(size));
if (NULL == hData) {
printf("Failed to allocate using malloc in testTexMultStream.\n");
return false;
}
CTX_CREATE()
fillTestBuffer<T>(width, height, hData);
// Load Kernel File and create hipArray
hipModule_t Module;
HIPCHECK(hipModuleLoadData(&Module, &buffer[0]));
HIP_ARRAY array;
allocInitArray(width, height, format, &array);
#ifdef __HIP_PLATFORM_NVCC__
// Copy from hData to array using hipMemcpyParam2D
copyBuffer2Array<T, HIP_ARRAY*>(width, height, hData, &array);
// Get tex reference from the loaded kernel file
// Assign array to the tex reference
assignArray2TexRef<HIP_ARRAY*>(format, texRefName, Module, &array);
#else
// Copy from hData to array using hipMemcpyParam2D
copyBuffer2Array<T, HIP_ARRAY>(width, height, hData, array);
// Get tex reference from the loaded kernel file
// Assign array to the tex reference
assignArray2TexRef<HIP_ARRAY>(format, texRefName, Module, array);
#endif
hipFunction_t Function;
HIPCHECK(hipModuleGetFunction(&Function, Module, kerFuncName));
// Create Multiple Strings
hipStream_t streams[MAX_STREAMS]={0};
T* dData[MAX_STREAMS] = {NULL};
T* hOutputData[MAX_STREAMS] = {NULL};
if (numOfStreams > MAX_STREAMS) {
numOfStreams = MAX_STREAMS;
}
unsigned int totalStreamsCreated = 0;
for (int stream_num = 0; stream_num < numOfStreams; stream_num++) {
hOutputData[stream_num] = reinterpret_cast<T*>(malloc(size));
if (NULL == hOutputData[stream_num]) {
printf("Failed to allocate using malloc in testTexMultStream.\n");
TestPassed &= false;
break;
}
HIPCHECK(hipStreamCreate(&streams[stream_num]));
HIPCHECK(hipMalloc(reinterpret_cast<void**>(&dData[stream_num]), size));
memset(hOutputData[stream_num], 0, size);
struct {
void* _Ad;
unsigned int _Bd;
unsigned int _Cd;
void* _Ad;
unsigned int _Bd;
unsigned int _Cd;
} args;
args._Ad = (void*) dData;
args._Ad = reinterpret_cast<void*>(dData[stream_num]);
args._Bd = width;
args._Cd = height;
size_t sizeTemp = sizeof(args);
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE,
&sizeTemp, HIP_LAUNCH_PARAM_END};
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER,
&args,
HIP_LAUNCH_PARAM_BUFFER_SIZE,
&sizeTemp,
HIP_LAUNCH_PARAM_END};
hipFunction_t Function;
HIP_CHECK(hipModuleGetFunction(&Function, Module, "tex2dKernel"));
int temp1 = width / GRIDDIMX;
int temp2 = height / GRIDDIMY;
HIPCHECK(
hipModuleLaunchKernel(Function, GRIDDIMX, GRIDDIMY, GRIDDIMZ,
temp1, temp2, BLOCKDIMZ, 0, streams[stream_num],
NULL, reinterpret_cast<void**>(&config)));
totalStreamsCreated++;
}
// Check the kernel results separately
for (int stream_num = 0; stream_num < totalStreamsCreated; stream_num++) {
HIPCHECK(hipStreamSynchronize(streams[stream_num]));
HIPCHECK(hipMemcpy(hOutputData[stream_num], dData[stream_num], size,
hipMemcpyDeviceToHost));
TestPassed &= validateOutput<T>(width, height, hData,
hOutputData[stream_num]);
}
for (int i = 0; i < totalStreamsCreated; i++) {
HIPCHECK(hipFree(dData[i]));
HIPCHECK(hipStreamDestroy(streams[i]));
free(hOutputData[i]);
}
ARRAY_DESTROY(array)
HIPCHECK(hipModuleUnload(Module));
free(hData);
CTX_DESTROY()
return TestPassed;
}
int temp1 = width / 16;
int temp2 = height / 16;
HIP_CHECK(
hipModuleLaunchKernel(Function, 16, 16, 1, temp1, temp2, 1, 0, 0, NULL, (void**)&config));
hipDeviceSynchronize();
/**
* Internal Thread Functions
*
*/
void launchSingleStreamMultGPU(int gpu, const std::vector<char>& buffer) {
bool TestPassed = true;
HIPCHECK(hipSetDevice(gpu));
TestPassed = testTexMultStream<float>(buffer,
HIP_AD_FORMAT_FLOAT,
"ftex",
"tex2dKernelFloat", 1);
g_thTestPassed &= static_cast<int>(TestPassed);
}
float* hOutputData = (float*)malloc(size);
memset(hOutputData, 0, size);
hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost);
void launchMultStreamMultGPU(int gpu, const std::vector<char>& buffer) {
bool TestPassed = true;
HIPCHECK(hipSetDevice(gpu));
TestPassed = testTexMultStream<float>(buffer,
HIP_AD_FORMAT_FLOAT,
"ftex",
"tex2dKernelFloat", 3);
g_thTestPassed &= static_cast<int>(TestPassed);
}
/**
* Validates texture functionality with Multiple Streams on multuple GPU
* for hipModuleGetTexRef
*
*/
bool testTexMultStreamMultGPU(unsigned int numOfGPUs,
const std::vector<char>& buffer) {
bool TestPassed = true;
std::thread T[numOfGPUs];
for (int i = 0; i < height; i++) {
for (int j = 0; j < width; j++) {
if (hData[i * width + j] != hOutputData[i * width + j]) {
printf("Difference [ %d %d ]:%f ----%f\n", i, j, hData[i * width + j],
hOutputData[i * width + j]);
testResult = false;
break;
}
}
}
hipFree(dData);
hipFreeArray(array);
HIP_CHECK(hipModuleUnload(Module));
return true;
for (int gpu = 0; gpu < numOfGPUs; gpu++) {
T[gpu] = std::thread(launchMultStreamMultGPU, gpu, buffer);
}
for (int gpu = 0; gpu < numOfGPUs; gpu++) {
T[gpu].join();
}
if (g_thTestPassed) {
TestPassed = true;
} else {
TestPassed = false;
}
return TestPassed;
}
/**
* Validates texture functionality with Single Stream on multuple GPU
* for hipModuleGetTexRef
*
*/
bool testTexSingleStreamMultGPU(unsigned int numOfGPUs,
const std::vector<char>& buffer) {
bool TestPassed = true;
std::thread T[numOfGPUs];
for (int gpu = 0; gpu < numOfGPUs; gpu++) {
T[gpu] = std::thread(launchSingleStreamMultGPU, gpu, buffer);
}
for (int gpu = 0; gpu < numOfGPUs; gpu++) {
T[gpu].join();
}
if (g_thTestPassed) {
TestPassed = true;
} else {
TestPassed = false;
}
return TestPassed;
}
int main(int argc, char** argv) {
hipInit(0);
testResult = runTest(argc, argv);
printf("%s ...\n", testResult ? "PASSED" : "FAILED");
exit(testResult ? EXIT_SUCCESS : EXIT_FAILURE);
return 0;
HipTest::parseStandardArguments(argc, argv, true);
bool TestPassed = true;
if (p_tests == 0x01) {
TestPassed = testTexType<float>(HIP_AD_FORMAT_FLOAT,
"ftex",
"tex2dKernelFloat");
} else if (p_tests == 0x02) {
TestPassed = testTexType<int>(HIP_AD_FORMAT_SIGNED_INT32,
"itex",
"tex2dKernelInt");
} else if (p_tests == 0x03) {
TestPassed = testTexType<short>(HIP_AD_FORMAT_SIGNED_INT16,
"stex",
"tex2dKernelInt16");
} else if (p_tests == 0x04) {
TestPassed = testTexType<char>(HIP_AD_FORMAT_SIGNED_INT8,
"ctex",
"tex2dKernelInt8");
} else if (p_tests == 0x05) {
auto buffer = load_file();
TestPassed = testTexMultStream<float>(buffer,
HIP_AD_FORMAT_FLOAT,
"ftex",
"tex2dKernelFloat",
MAX_STREAMS);
} else if (p_tests == 0x06) {
int gpu_cnt = 0;
auto buffer = load_file();
HIPCHECK(hipGetDeviceCount(&gpu_cnt));
TestPassed = testTexSingleStreamMultGPU(gpu_cnt, buffer);
} else if (p_tests == 0x07) {
int gpu_cnt = 0;
auto buffer = load_file();
HIPCHECK(hipGetDeviceCount(&gpu_cnt));
TestPassed = testTexMultStreamMultGPU(gpu_cnt, buffer);
} else if (p_tests == 0x10) {
TestPassed = testTexRefEqNullPtr();
} else if (p_tests == 0x11) {
TestPassed = testNameEqNullPtr();
} else if (p_tests == 0x12) {
TestPassed = testInvalidTexName();
} else if (p_tests == 0x13) {
TestPassed = testEmptyTexName();
} else if (p_tests == 0x14) {
TestPassed = testWrongTexRef();
} else if (p_tests == 0x15) {
TestPassed = testUnloadedMod();
} else {
printf("Invalid Test Case \n");
exit(1);
}
if (TestPassed) {
passed();
} else {
failed("Test Case %x Failed!", p_tests);
}
}
projects/hip/tests/src/runtimeApi/module/tex2d_kernel.cpp
+40 -11
مشاهده پرونده
@@ -19,18 +19,47 @@ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/* HIT_START
* BUILD_CMD: tex2d_kernel.code %hc --genco %S/tex2d_kernel.cpp -o tex2d_kernel.code EXCLUDE_HIP_PLATFORM nvidia EXCLUDE_HIP_RUNTIME rocclr
* HIT_END
*/
#include "hip/hip_runtime.h"
texture<float, 2, hipReadModeElementType> tex;
texture<float, 2, hipReadModeElementType> ftex;
texture<int, 2, hipReadModeElementType> itex;
texture<short, 2, hipReadModeElementType> stex;
texture<char, 2, hipReadModeElementType> ctex;
extern "C" __global__ void tex2dKernel(float* outputData, int width, int height) {
int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
outputData[y * width + x] = tex2D(tex, x, y);
__device__ float deviceGlobalFloat;
extern "C" __global__ void tex2dKernelFloat(float* outputData,
int width, int height) {
int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
if ((x < width) && (y < width)) {
outputData[y * width + x] = tex2D(ftex, x, y);
}
}
extern "C" __global__ void tex2dKernelInt(int* outputData,
int width, int height) {
int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
if ((x < width) && (y < width)) {
outputData[y * width + x] = tex2D(itex, x, y);
}
}
extern "C" __global__ void tex2dKernelInt16(short* outputData,
int width, int height) {
int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
if ((x < width) && (y < width)) {
outputData[y * width + x] = tex2D(stex, x, y);
}
}
extern "C" __global__ void tex2dKernelInt8(char* outputData,
int width, int height) {
int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
if ((x < width) && (y < width)) {
outputData[y * width + x] = tex2D(ctex, x, y);
}
}