From fda049fa5f056b846e90b2f3b95d408a18d23bd2 Mon Sep 17 00:00:00 2001 From: 7SK Date: Mon, 6 Jun 2016 13:20:45 +0530 Subject: [PATCH] fix_ldg Change-Id: I53de5fa91b4f57d496ffe46787d197ae84dde4a4 --- hipamd/include/hcc_detail/hip_ldg.h | 30 +-- hipamd/src/hip_ldg.cpp | 151 +---------- hipamd/tests/src/CMakeLists.txt | 67 ++--- hipamd/tests/src/hip_test_ldg.cpp | 401 +++++++++++++++++----------- 4 files changed, 277 insertions(+), 372 deletions(-) diff --git a/hipamd/include/hcc_detail/hip_ldg.h b/hipamd/include/hcc_detail/hip_ldg.h index 1b24060d44..8d16ee27ef 100644 --- a/hipamd/include/hcc_detail/hip_ldg.h +++ b/hipamd/include/hcc_detail/hip_ldg.h @@ -23,79 +23,51 @@ THE SOFTWARE. #if __HCC__ #include"hip_vector_types.h" #include"host_defines.h" -#if __hcc_workweek__ >= 16164 +#if __hcc_workweek__ >= 16164 #include"hip/hip_vector_types.h" #include"hip/hcc_detail/host_defines.h" __device__ char __ldg(const char* ); -__device__ char1 __ldg(const char1* ); __device__ char2 __ldg(const char2* ); -__device__ char3 __ldg(const char3* ); __device__ char4 __ldg(const char4* ); __device__ signed char __ldg(const signed char* ); __device__ unsigned char __ldg(const unsigned char* ); __device__ short __ldg(const short* ); -__device__ short1 __ldg(const short1* ); __device__ short2 __ldg(const short2* ); -__device__ short3 __ldg(const short3* ); __device__ short4 __ldg(const short4* ); __device__ unsigned short __ldg(const unsigned short* ); __device__ int __ldg(const int* ); -__device__ int1 __ldg(const int1* ); __device__ int2 __ldg(const int2* ); -__device__ int3 __ldg(const int3* ); __device__ int4 __ldg(const int4* ); __device__ unsigned int __ldg(const unsigned int* ); __device__ long __ldg(const long* ); -__device__ long1 __ldg(const long1* ); -__device__ long2 __ldg(const long2* ); -__device__ long3 __ldg(const long3* ); -__device__ long4 __ldg(const long4* ); __device__ unsigned long __ldg(const unsigned long* ); __device__ long long __ldg(const long long* ); -__device__ longlong1 __ldg(const longlong1* ); __device__ longlong2 __ldg(const longlong2* ); -__device__ longlong3 __ldg(const longlong3* ); -__device__ longlong4 __ldg(const longlong4* ); __device__ unsigned long long __ldg(const unsigned long long* ); -__device__ uchar1 __ldg(const uchar1* ); __device__ uchar2 __ldg(const uchar2* ); -__device__ uchar3 __ldg(const uchar3* ); __device__ uchar4 __ldg(const uchar4* ); -__device__ ushort1 __ldg(const ushort1* ); __device__ ushort2 __ldg(const ushort2* ); -__device__ ushort3 __ldg(const ushort3* ); -__device__ ushort4 __ldg(const ushort4* ); -__device__ uint1 __ldg(const uint1* ); __device__ uint2 __ldg(const uint2* ); -__device__ uint3 __ldg(const uint3* ); __device__ uint4 __ldg(const uint4* ); -__device__ ulonglong1 __ldg(const ulonglong1* ); __device__ ulonglong2 __ldg(const ulonglong2* ); -__device__ ulonglong3 __ldg(const ulonglong3* ); -__device__ ulonglong4 __ldg(const ulonglong4* ); __device__ float __ldg(const float* ); -__device__ float1 __ldg(const float1* ); __device__ float2 __ldg(const float2* ); -__device__ float3 __ldg(const float3* ); __device__ float4 __ldg(const float4* ); __device__ double __ldg(const double* ); -__device__ double1 __ldg(const double1* ); __device__ double2 __ldg(const double2* ); -__device__ double3 __ldg(const double3* ); -__device__ double4 __ldg(const double4* ); #endif // __hcc_workweek__ diff --git a/hipamd/src/hip_ldg.cpp b/hipamd/src/hip_ldg.cpp index 9c3189c322..b4b4e8f9e8 100644 --- a/hipamd/src/hip_ldg.cpp +++ b/hipamd/src/hip_ldg.cpp @@ -20,27 +20,17 @@ THE SOFTWARE. #include #include"hcc_detail/hip_ldg.h" -#if __hcc_workweek__ >= 16164 +#if __hcc_workweek__ >= 16164 __device__ char __ldg(const char* ptr) { return ptr[0]; } -__device__ char1 __ldg(const char1* ptr) -{ - return ptr[0]; -} - __device__ char2 __ldg(const char2* ptr) { return ptr[0]; } -__device__ char3 __ldg(const char3* ptr) -{ - return ptr[0]; -} - __device__ char4 __ldg(const char4* ptr) { return ptr[0]; @@ -56,27 +46,16 @@ __device__ unsigned char __ldg(const unsigned char* ptr) return ptr[0]; } - __device__ short __ldg(const short* ptr) { return ptr[0]; } -__device__ short1 __ldg(const short1* ptr) -{ - return ptr[0]; -} - __device__ short2 __ldg(const short2* ptr) { return ptr[0]; } -__device__ short3 __ldg(const short3* ptr) -{ - return ptr[0]; -} - __device__ short4 __ldg(const short4* ptr) { return ptr[0]; @@ -87,27 +66,16 @@ __device__ unsigned short __ldg(const unsigned short* ptr) return ptr[0]; } - __device__ int __ldg(const int* ptr) { return ptr[0]; } -__device__ int1 __ldg(const int1* ptr) -{ - return ptr[0]; -} - __device__ int2 __ldg(const int2* ptr) { return ptr[0]; } -__device__ int3 __ldg(const int3* ptr) -{ - return ptr[0]; -} - __device__ int4 __ldg(const int4* ptr) { return ptr[0]; @@ -124,198 +92,81 @@ __device__ long __ldg(const long* ptr) return ptr[0]; } -__device__ long1 __ldg(const long1* ptr) -{ - return ptr[0]; -} - -__device__ long2 __ldg(const long2* ptr) -{ - return ptr[0]; -} - -__device__ long3 __ldg(const long3* ptr) -{ - return ptr[0]; -} - -__device__ long4 __ldg(const long4* ptr) -{ - return ptr[0]; -} - __device__ unsigned long __ldg(const unsigned long* ptr) { return ptr[0]; } - __device__ long long __ldg(const long long* ptr) { return ptr[0]; } -__device__ longlong1 __ldg(const longlong1* ptr) -{ - return ptr[0]; -} - __device__ longlong2 __ldg(const longlong2* ptr) { return ptr[0]; } -__device__ longlong3 __ldg(const longlong3* ptr) -{ - return ptr[0]; -} - -__device__ longlong4 __ldg(const longlong4* ptr) -{ - return ptr[0]; -} - __device__ unsigned long long __ldg(const unsigned long long* ptr) { return ptr[0]; } - -__device__ uchar1 __ldg(const uchar1* ptr) -{ - return ptr[0]; -} - __device__ uchar2 __ldg(const uchar2* ptr) { return ptr[0]; } -__device__ uchar3 __ldg(const uchar3* ptr) -{ - return ptr[0]; -} - __device__ uchar4 __ldg(const uchar4* ptr) { return ptr[0]; } - -__device__ ushort1 __ldg(const ushort1* ptr) -{ - return ptr[0]; -} - __device__ ushort2 __ldg(const ushort2* ptr) { return ptr[0]; } -__device__ ushort3 __ldg(const ushort3* ptr) -{ - return ptr[0]; -} - -__device__ ushort4 __ldg(const ushort4* ptr) -{ - return ptr[0]; -} - - -__device__ uint1 __ldg(const uint1* ptr) -{ - return ptr[0]; -} - __device__ uint2 __ldg(const uint2* ptr) { return ptr[0]; } -__device__ uint3 __ldg(const uint3* ptr) -{ - return ptr[0]; -} - __device__ uint4 __ldg(const uint4* ptr) { return ptr[0]; } - -__device__ ulonglong1 __ldg(const ulonglong1* ptr) -{ - return ptr[0]; -} - __device__ ulonglong2 __ldg(const ulonglong2* ptr) { return ptr[0]; } -__device__ ulonglong3 __ldg(const ulonglong3* ptr) -{ - return ptr[0]; -} - -__device__ ulonglong4 __ldg(const ulonglong4* ptr) -{ - return ptr[0]; -} - - __device__ float __ldg(const float* ptr) { return ptr[0]; } -__device__ float1 __ldg(const float1* ptr) -{ - return ptr[0]; -} - __device__ float2 __ldg(const float2* ptr) { return ptr[0]; } -__device__ float3 __ldg(const float3* ptr) -{ - return ptr[0]; -} - __device__ float4 __ldg(const float4* ptr) { return ptr[0]; } - __device__ double __ldg(const double* ptr) { return ptr[0]; } -__device__ double1 __ldg(const double1* ptr) -{ - return ptr[0]; -} - __device__ double2 __ldg(const double2* ptr) { return ptr[0]; } -__device__ double3 __ldg(const double3* ptr) -{ - return ptr[0]; -} - -__device__ double4 __ldg(const double4* ptr) -{ - return ptr[0]; -} - #endif diff --git a/hipamd/tests/src/CMakeLists.txt b/hipamd/tests/src/CMakeLists.txt index e0894df5fa..c507ac2ba1 100644 --- a/hipamd/tests/src/CMakeLists.txt +++ b/hipamd/tests/src/CMakeLists.txt @@ -60,12 +60,12 @@ if (${HIP_PLATFORM} STREQUAL "hcc") set (CMAKE_CXX_FLAGS --hipcc_explicit_lib) endif() - + elseif (${HIP_PLATFORM} STREQUAL "nvcc") MESSAGE ("HIP_PLATFORM=nvcc") #Need C++11 for threads in some of the tests. - add_definitions(-std=c++11) + add_definitions(-std=c++11) # NVCC does not not support -rdynamic option set(CMAKE_SHARED_LIBRARY_LINK_CXX_FLAGS ) @@ -83,7 +83,7 @@ add_library(test_common OBJECT test_common.cpp ) # usage : make_hip_executable (exe_name CPP_FILES) -macro (make_hip_executable exe cpp) +macro (make_hip_executable exe cpp) if (${HIP_PLATFORM} STREQUAL "hcc") if (${HIP_BUILD_LOCAL}) #target_link_libraries(${exe} hip_hcc) @@ -93,12 +93,12 @@ macro (make_hip_executable exe cpp) endif() else() add_executable (${exe} ${cpp} ${ARGN} $ ) - endif() + endif() endmacro() # Make a hip executable, using libc++ -macro (make_hip_executable_libcpp exe cpp) +macro (make_hip_executable_libcpp exe cpp) make_hip_executable( ${exe} ${cpp} ${ARGN} ) if (${HIP_PLATFORM} STREQUAL "hcc") set_source_files_properties (${cpp} i${ARGN} PROPERTIES COMPILE_FLAGS --stdlib=libc++ ) @@ -106,7 +106,7 @@ macro (make_hip_executable_libcpp exe cpp) endmacro() macro (make_named_test exe testname ) - add_test (NAME ${testname} + add_test (NAME ${testname} COMMAND ${PROJECT_BINARY_DIR}/${exe} ${ARGN} ) set_tests_properties (${testname} @@ -126,7 +126,7 @@ macro (make_hipify_test sourceFile ) #string (REPLACE " " "" smush_args ${ARGN}) set (testname ${sourceFile}${smush_args}.tst) - add_test (NAME ${testname} + add_test (NAME ${testname} COMMAND ${HIP_PATH}/bin/hipify ${PROJECT_SOURCE_DIR}/${sourceFile} ${ARGN} ) endmacro() @@ -135,7 +135,7 @@ endmacro() macro (make_test_matches exe match_string) string (REPLACE " " "" smush_args ${ARGN}) set (testname ${exe}${smush_args}.tst) - add_test (NAME ${testname} + add_test (NAME ${testname} COMMAND ${PROJECT_BINARY_DIR}/${exe} ${ARGN} ) set_tests_properties (${testname} @@ -143,38 +143,44 @@ macro (make_test_matches exe match_string) ) endmacro() - +macro (make_hip_executable_ldg exe cpp) + if (${HIP_PLATFORM} STREQUAL "hcc") + make_hip_executable( ${exe} ${cpp} ${ARGN} ) +else () +make_hip_executable( ${exe} ${cpp} ${ARGN} ) + set_source_files_properties (${cpp} i${ARGN} PROPERTIES COMPILE_FLAGS --gpu-architecture=sm_35 ) +endif() +endmacro() #make_hip_executable (hipAPIStreamEnable hipAPIStreamEnable.cpp) #make_hip_executable (hipAPIStreamDisable hipAPIStreamDisable.cpp) make_hip_executable (hip_ballot hip_ballot.cpp) -make_hip_executable (hip_anyall hip_anyall.cpp) -make_hip_executable (hip_popc hip_popc.cpp) -make_hip_executable (hip_clz hip_clz.cpp) -make_hip_executable (hip_brev hip_brev.cpp) -make_hip_executable (hip_ffs hip_ffs.cpp) -if (${HIP_PLATFORM} STREQUAL "hcc") - make_hip_executable (hip_test_ldg hip_test_ldg.cpp) -endif() +make_hip_executable (hip_anyall hip_anyall.cpp) +make_hip_executable (hip_popc hip_popc.cpp) +make_hip_executable (hip_clz hip_clz.cpp) +make_hip_executable (hip_brev hip_brev.cpp) +make_hip_executable (hip_ffs hip_ffs.cpp) +make_hip_executable_ldg (hip_test_ldg hip_test_ldg.cpp) + make_hip_executable (hipGetDeviceAttribute hipGetDeviceAttribute.cpp) -make_hip_executable (hipEnvVar hipEnvVar.cpp) +make_hip_executable (hipEnvVar hipEnvVar.cpp) make_hip_executable (hipEnvVarDriver hipEnvVarDriver.cpp) make_hip_executable (hipMemcpy_simple hipMemcpy_simple.cpp) -make_hip_executable (hipMemcpy hipMemcpy.cpp) -make_hip_executable (hipMemcpyAsync hipMemcpyAsync.cpp) -make_hip_executable (hipMemset hipMemset.cpp) -make_hip_executable (hipEventRecord hipEventRecord.cpp) -make_hip_executable_libcpp (hipLanguageExtensions hipLanguageExtensions.cpp) -make_hip_executable (hipGridLaunch hipGridLaunch.cpp) +make_hip_executable (hipMemcpy hipMemcpy.cpp) +make_hip_executable (hipMemcpyAsync hipMemcpyAsync.cpp) +make_hip_executable (hipMemset hipMemset.cpp) +make_hip_executable (hipEventRecord hipEventRecord.cpp) +make_hip_executable_libcpp (hipLanguageExtensions hipLanguageExtensions.cpp) +make_hip_executable (hipGridLaunch hipGridLaunch.cpp) make_hip_executable_libcpp (hipHcc hipHcc.cpp) -#set_source_files_properties (hipHcc.cpp PROPERTIES COMPILE_FLAGS --stdlib=libc++ ) +#set_source_files_properties (hipHcc.cpp PROPERTIES COMPILE_FLAGS --stdlib=libc++ ) make_hip_executable (hipSimpleAtomicsTest hipSimpleAtomicsTest.cpp) make_hip_executable (hipMathFunctionsHost hipMathFunctions.cpp hipSinglePrecisionMathHost.cpp hipDoublePrecisionMathHost.cpp) make_hip_executable (hipMathFunctionsDevice hipMathFunctions.cpp hipSinglePrecisionMathDevice.cpp hipDoublePrecisionMathDevice.cpp) make_hip_executable (hipIntrinsics hipMathFunctions.cpp hipSinglePrecisionIntrinsics.cpp hipDoublePrecisionIntrinsics.cpp hipIntegerIntrinsics.cpp) -make_hip_executable_libcpp (hipPointerAttrib hipPointerAttrib.cpp) +make_hip_executable_libcpp (hipPointerAttrib hipPointerAttrib.cpp) make_hip_executable (hipMultiThreadStreams1 hipMultiThreadStreams1.cpp) make_hip_executable (hipMultiThreadStreams2 hipMultiThreadStreams2.cpp) make_hip_executable (hipHostAlloc hipHostAlloc.cpp) @@ -187,7 +193,7 @@ make_hip_executable (hipFuncSetDeviceFlags hipFuncSetDeviceFlags.cpp) make_hip_executable (hipFuncGetDevice hipFuncGetDevice.cpp) make_hip_executable (hipFuncSetDevice hipFuncSetDevice.cpp) make_hip_executable (hipFuncDeviceSynchronize hipFuncDeviceSynchronize.cpp) -make_hip_executable (hipPeerToPeer_simple hipPeerToPeer_simple.cpp) +make_hip_executable (hipPeerToPeer_simple hipPeerToPeer_simple.cpp) make_hip_executable (hipMemcpyAll hipMemcpyAll.cpp) make_hip_executable (hipMultiThreadDevice hipMultiThreadDevice.cpp) make_hip_executable (hipTestMemcpyPin hipTestMemcpyPin.cpp) @@ -216,7 +222,7 @@ make_test(hipMemcpy_simple " " ) make_named_test(hipMemcpy "hipMemcpy-modes" --tests 0x1 ) make_named_test(hipMemcpy "hipMemcpy-size" --tests 0x6 ) make_named_test(hipMemcpy "hipMemcpy-multithreaded" --tests 0x8 ) - +make_test(hip_test_ldg " " ) # Debug synchronization, then enable. #make_named_test(hipMemcpy_simple "hipMemcpyAsync-simple" --async) @@ -239,14 +245,13 @@ make_named_test (hipMultiThreadDevice "hipMultiThreadDevice-pyramid" --tests 0x4 make_named_test (hipMultiThreadDevice "hipMultiThreadDevice-nearzero" --tests 0x10) if (${HIP_MULTI_GPU}) - make_test(hipPeerToPeer_simple " ") # use current device for copy, this fails. - make_test(hipPeerToPeer_simple --memcpyWithPeer) + make_test(hipPeerToPeer_simple " ") # use current device for copy, this fails. + make_test(hipPeerToPeer_simple --memcpyWithPeer) make_test(hipPeerToPeer_simple --mirrorPeers) # mirror mapping: test to ensure mirror doesn't destroy orig mapping. endif() if (${HIP_PLATFORM} STREQUAL "hcc") - make_test(hip_test_ldg " " ) make_test(hipMemoryAllocate " ") make_test(hipFuncSetDevice " ") endif() diff --git a/hipamd/tests/src/hip_test_ldg.cpp b/hipamd/tests/src/hip_test_ldg.cpp index 31c7d4865a..184493783a 100644 --- a/hipamd/tests/src/hip_test_ldg.cpp +++ b/hipamd/tests/src/hip_test_ldg.cpp @@ -26,7 +26,9 @@ THE SOFTWARE. #include #include "hip_runtime.h" #include "test_common.h" -#if __hcc_workweek__ >= 16164 + +#if (__hcc_workweek__ >= 16164) || defined (__HIP_PLATFORM_NVCC__) + #define HIP_ASSERT(x) (assert((x)==hipSuccess)) @@ -39,10 +41,12 @@ THE SOFTWARE. #define THREADS_PER_BLOCK_Y 8 #define THREADS_PER_BLOCK_Z 1 +using namespace std; + template -__global__ void +__global__ void vectoradd_float(hipLaunchParm lp, - T* a, const T* bm, const T* cm, int width, int height) + T* a, const T* bm, int width, int height) { int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; @@ -50,65 +54,108 @@ vectoradd_float(hipLaunchParm lp, int i = y * width + x; if ( i < (width * height)) { - a[i] = __ldg(&bm[i]) + __ldg(&cm[i]); + a[i] = __ldg(&bm[i]) ; } } -#if 0 -__kernel__ void vectoradd_float(float* a, const float* b, const float* c, int width, int height) { - - - int x = blockDimX * blockIdx.x + threadIdx.x; - int y = blockDimY * blockIdy.y + threadIdx.y; - - int i = y * width + x; - if ( i < (width * height)) { - a[i] = b[i] + c[i]; - } +int2 make_vector2(int a){ + return make_int2(a,a); } -#endif -using namespace std; +char2 make_vector2(signed char a){ + return make_char2(a, a); +} -template +char4 make_vector4(signed char a){ + return make_char4(a, a, a ,a); +} + +short2 make_vector2(short a){ + return make_short2(a,a); +} + +ushort2 make_vector2(unsigned short a){ + return make_ushort2(a,a); +} + +short4 make_vector4(short a){ + return make_short4(a,a,a,a); +} + +int4 make_vector4(int a){ + return make_int4(a,a,a,a); +} + +uint2 make_vector2 (unsigned int a){ + return make_uint2 (a,a); +} + +uint4 make_vector4 (unsigned int a){ + return make_uint4 (a,a,a,a); +} + +float2 make_vector2 (float a){ + return make_float2 (a,a); +} + +float4 make_vector4 (float a){ + return make_float4 (a,a,a,a); +} + +uchar2 make_vector2 (unsigned char a){ + return make_uchar2 (a,a); +} + +uchar4 make_vector4 (unsigned char a){ + return make_uchar4 (a,a,a,a); +} + +double2 make_vector2 (double a){ + return make_double2 (a,a); +} + + + + + + + + +template bool dataTypesRun(){ T* hostA; T* hostB; - T* hostC; + T* deviceA; T* deviceB; - T* deviceC; + int i; int errors; hostA = (T*)malloc(NUM * sizeof(T)); hostB = (T*)malloc(NUM * sizeof(T)); - hostC = (T*)malloc(NUM * sizeof(T)); - + // initialize the input data for (i = 0; i < NUM; i++) { - hostB[i] = (T)i; - hostC[i] = (T)i; + hostB[i] = (U)i; } - + HIP_ASSERT(hipMalloc((void**)&deviceA, NUM * sizeof(T))); HIP_ASSERT(hipMalloc((void**)&deviceB, NUM * sizeof(T))); - HIP_ASSERT(hipMalloc((void**)&deviceC, NUM * sizeof(T))); - + HIP_ASSERT(hipMemcpy(deviceB, hostB, NUM*sizeof(T), hipMemcpyHostToDevice)); - HIP_ASSERT(hipMemcpy(deviceC, hostC, NUM*sizeof(T), hipMemcpyHostToDevice)); - hipLaunchKernel(vectoradd_float, + hipLaunchKernel(vectoradd_float, dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y), dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), 0, 0, - deviceA ,deviceB ,deviceC ,WIDTH ,HEIGHT); + deviceA ,deviceB ,WIDTH ,HEIGHT); HIP_ASSERT(hipMemcpy(hostA, deviceA, NUM*sizeof(T), hipMemcpyDeviceToHost)); @@ -117,12 +164,12 @@ bool dataTypesRun(){ // verify the results errors = 0; for (i = 0; i < NUM; i++) { - if (hostA[i] != (hostB[i] + hostC[i])) { + if (hostA[i] != (hostB[i])) { errors++; } } if (errors!=0) { - printf("FAILED: %d errors\n",errors); + std::cout << "FAILED\n"< +bool dataTypesRun2(){ + T* hostA; + T* hostB; + + + T* deviceA; + T* deviceB; + + + int i; + int errors; + + hostA = (T*)malloc(NUM * sizeof(T)); + hostB = (T*)malloc(NUM * sizeof(T)); + + // initialize the input data + for (i = 0; i < NUM; i++) { + hostB[i] = make_vector2((U)i); + + } + + HIP_ASSERT(hipMalloc((void**)&deviceA, NUM * sizeof(T))); + HIP_ASSERT(hipMalloc((void**)&deviceB, NUM * sizeof(T))); + + HIP_ASSERT(hipMemcpy(deviceB, hostB, NUM*sizeof(T), hipMemcpyHostToDevice)); + + hipLaunchKernel(vectoradd_float, + dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y), + dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), + 0, 0, + deviceA ,deviceB,WIDTH ,HEIGHT); + + + HIP_ASSERT(hipMemcpy(hostA, deviceA, NUM*sizeof(T), hipMemcpyDeviceToHost)); + + bool ret = false; + // verify the results + errors = 0; + for (i = 0; i < NUM; i++) { +if (hostA[i].x != (hostB[i].x) && hostA[i].y != (hostB[i].y)) { + errors++; + } + } + if (errors!=0) { + std::cout << "FAILED\n"< +bool dataTypesRun4(){ + T* hostA; + T* hostB; + + T* deviceA; + T* deviceB; + + int i; + int errors; + + hostA = (T*)malloc(NUM * sizeof(T)); + hostB = (T*)malloc(NUM * sizeof(T)); + + // initialize the input data + for (i = 0; i < NUM; i++) { + hostB[i] = make_vector4((U)i); + } + + HIP_ASSERT(hipMalloc((void**)&deviceA, NUM * sizeof(T))); + HIP_ASSERT(hipMalloc((void**)&deviceB, NUM * sizeof(T))); + + HIP_ASSERT(hipMemcpy(deviceB, hostB, NUM*sizeof(T), hipMemcpyHostToDevice)); + + + hipLaunchKernel(vectoradd_float, + dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y), + dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), + 0, 0, + deviceA ,deviceB ,WIDTH ,HEIGHT); + + + HIP_ASSERT(hipMemcpy(hostA, deviceA, NUM*sizeof(T), hipMemcpyDeviceToHost)); + + bool ret = false; + // verify the results + errors = 0; + for (i = 0; i < NUM; i++) { + if (hostA[i].x != (hostB[i].x ) && hostA[i].y != (hostB[i].y ) && hostA[i].z != (hostB[i].z ) && hostA[i].w != (hostB[i].w )) { + errors++; + } + } + if (errors!=0) { + std::cout << "FAILED\n"<() & - dataTypesRun() & - dataTypesRun() & - dataTypesRun() & - dataTypesRun() & - dataTypesRun() & - dataTypesRun(); + errors = dataTypesRun() & + dataTypesRun() & + dataTypesRun() & + dataTypesRun() & + dataTypesRun() & + dataTypesRun() & + dataTypesRun()& + dataTypesRun()& + dataTypesRun()& + dataTypesRun()& + dataTypesRun()& + dataTypesRun()& + dataTypesRun(); if(errors == 1){ errors = 0; + std::cout<<"ldg working for single element data types\n"<() & - dataTypesRun() & - dataTypesRun() & - dataTypesRun() & - dataTypesRun() & - dataTypesRun(); + errors = dataTypesRun2() & + dataTypesRun2() & + dataTypesRun2() & + dataTypesRun2() & + dataTypesRun2() & + dataTypesRun2() & + dataTypesRun2() & + dataTypesRun2(); if(errors == 1){ errors = 0; + std::cout<<"ldg working for two element data types\n"<() & - dataTypesRun() & - dataTypesRun() & - dataTypesRun() & - dataTypesRun() & - dataTypesRun(); + + + errors = dataTypesRun4() & + dataTypesRun4() & + dataTypesRun4() & + dataTypesRun4() & + dataTypesRun4() & + dataTypesRun4() ; if(errors == 1){ errors = 0; + std::cout<<"ldg working for four element data types\n"<() & - dataTypesRun() & - dataTypesRun() & - dataTypesRun() & - dataTypesRun() & - dataTypesRun(); - - if(errors == 1){ - errors = 0; - }else{ - std::cout<<"Failed Long"<() & - dataTypesRun() & - dataTypesRun() & - dataTypesRun() & - dataTypesRun() & - dataTypesRun(); - - if(errors == 1){ - errors = 0; - }else{ - std::cout<<"Failed Long Long"<() & - dataTypesRun() & - dataTypesRun() & - dataTypesRun(); - - if(errors == 1){ - errors = 0; - }else{ - std::cout<<"Failed Unsigned Char"<() & - dataTypesRun() & - dataTypesRun() & - dataTypesRun(); - - if(errors == 1){ - errors = 0; - }else{ - std::cout<<"Failed Unsigned Short"<() & - dataTypesRun() & - dataTypesRun() & - dataTypesRun(); - - if(errors == 1){ - errors = 0; - }else{ - std::cout<<"Failed Unsigned Int"<() & - dataTypesRun() & - dataTypesRun() & - dataTypesRun(); - - if(errors == 1){ - errors = 0; - }else{ - std::cout<<"Failed Unsigned Long Long"<() & - dataTypesRun() & - dataTypesRun() & - dataTypesRun() & - dataTypesRun(); - - if(errors == 1){ - errors = 0; - }else{ - std::cout<<"Failed Float"<() & - dataTypesRun() & - dataTypesRun() & - dataTypesRun() & - dataTypesRun(); - - - //hipResetDefaultAccelerator(); - if(errors == 1){ - passed(); - return 0; - }else{ - std::cout<<"Failed Float"<