AI/rocm-systems
2
0
Forkuj 0
Kod Zgłoszenia Oczekujące zmiany Actions Packages Projekty Wydania Wiki Aktywność

fix_ldg

Przeglądaj źródła 
Change-Id: I53de5fa91b4f57d496ffe46787d197ae84dde4a4
This commit is contained in:
7SK
2016-06-06 13:20:45 +05:30
zatwierdzone przez Andres Rodriguez
rodzic e797253165
commit fda049fa5f
4 zmienionych plików z 277 dodań i 372 usunięć
Ściągnij plik aktualizacji Ściągnij plik porównania Expand all files Collapse all files
hipamd/include/hcc_detail/hip_ldg.h
+1 -29
Wyświetl plik
@@ -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__
hipamd/src/hip_ldg.cpp
+1 -150
Wyświetl plik
@@ -20,27 +20,17 @@ THE SOFTWARE.
#include <hc.hpp>
#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
hipamd/tests/src/CMakeLists.txt
+36 -31
Wyświetl plik
@@ -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} $<TARGET_OBJECTS:test_common> )
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()
hipamd/tests/src/hip_test_ldg.cpp
+239 -162
Wyświetl plik
@@ -26,7 +26,9 @@ THE SOFTWARE.
#include<iostream>
#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<typename T>
__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<typename T>
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<typename T, typename U>
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"<<std::endl;
ret = false;
} else {
ret = true;
@@ -130,175 +177,205 @@ bool dataTypesRun(){
HIP_ASSERT(hipFree(deviceA));
HIP_ASSERT(hipFree(deviceB));
HIP_ASSERT(hipFree(deviceC));
free(hostA);
free(hostB);
free(hostC);
return ret;
}
template<typename T, typename U>
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"<<std::endl;
ret = false;
} else {
ret = true;
}
HIP_ASSERT(hipFree(deviceA));
HIP_ASSERT(hipFree(deviceB));
free(hostA);
free(hostB);
return ret;
}
template<typename T, typename U>
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"<<std::endl;
ret = false;
} else {
ret = true;
}
HIP_ASSERT(hipFree(deviceA));
HIP_ASSERT(hipFree(deviceB));
free(hostA);
free(hostB);
return ret;
}
int main() {
hipDeviceProp_t devProp;
hipGetDeviceProperties(&devProp, 0);
cout << " System minor " << devProp.minor << endl;
cout << " System major " << devProp.major << endl;
cout << " agent prop name " << devProp.name << endl;
int errors;
errors = dataTypesRun<char>() &
dataTypesRun<char1>() &
dataTypesRun<char2>() &
dataTypesRun<char3>() &
dataTypesRun<char4>() &
dataTypesRun<signed char>() &
dataTypesRun<unsigned char>();
errors = dataTypesRun<char,char>() &
dataTypesRun<short, short>() &
dataTypesRun<int,int>() &
dataTypesRun<long, long>() &
dataTypesRun<long long, long long>() &
dataTypesRun<signed char,signed char>() &
dataTypesRun<unsigned char, unsigned char>()&
dataTypesRun<unsigned short, unsigned short>()&
dataTypesRun<unsigned int, unsigned int>()&
dataTypesRun<unsigned long, unsigned long>()&
dataTypesRun<unsigned long long,unsigned long long>()&
dataTypesRun<float, float>()&
dataTypesRun<double, double>();
if(errors == 1){
errors = 0;
std::cout<<"ldg working for single element data types\n"<<std::endl;
}else{
std::cout<<"Failed Char"<<std::endl;
std::cout<<"Failed single element data types"<<std::endl;
return -1;
}
errors = dataTypesRun<short>() &
dataTypesRun<short1>() &
dataTypesRun<short2>() &
dataTypesRun<short3>() &
dataTypesRun<short4>() &
dataTypesRun<unsigned short>();
errors = dataTypesRun2<int2,int>() &
dataTypesRun2<short2,short>() &
dataTypesRun2<ushort2,unsigned short>() &
dataTypesRun2<char2,signed char>() &
dataTypesRun2<uchar2,unsigned char>() &
dataTypesRun2<uint2,unsigned int>() &
dataTypesRun2<float2,float>() &
dataTypesRun2<double2,double>();
if(errors == 1){
errors = 0;
std::cout<<"ldg working for two element data types\n"<<std::endl;
}else{
std::cout<<"Failed Short"<<std::endl;
std::cout<<"Failed two element vector data types"<<std::endl;
return -1;
}
errors = dataTypesRun<int>() &
dataTypesRun<int1>() &
dataTypesRun<int2>() &
dataTypesRun<int3>() &
dataTypesRun<int4>() &
dataTypesRun<unsigned int>();
errors = dataTypesRun4<int4,int>() &
dataTypesRun4<char4,signed char>() &
dataTypesRun4<uchar4,unsigned char>() &
dataTypesRun4<short4, short>() &
dataTypesRun4<uint4,unsigned int>() &
dataTypesRun4<float4,float>() ;
if(errors == 1){
errors = 0;
std::cout<<"ldg working for four element data types\n"<<std::endl;
}else{
std::cout<<"Failed Int"<<std::endl;
std::cout<<"Failed four element vector data types"<<std::endl;
return -1;
}
errors = dataTypesRun<long>() &
dataTypesRun<long1>() &
dataTypesRun<long2>() &
dataTypesRun<long3>() &
dataTypesRun<long4>() &
dataTypesRun<unsigned long>();
if(errors == 1){
errors = 0;
}else{
std::cout<<"Failed Long"<<std::endl;
return -1;
}
errors = dataTypesRun<long long>() &
dataTypesRun<longlong1>() &
dataTypesRun<longlong2>() &
dataTypesRun<longlong3>() &
dataTypesRun<longlong4>() &
dataTypesRun<unsigned long long>();
if(errors == 1){
errors = 0;
}else{
std::cout<<"Failed Long Long"<<std::endl;
return -1;
}
errors = dataTypesRun<uchar1>() &
dataTypesRun<uchar2>() &
dataTypesRun<uchar3>() &
dataTypesRun<uchar4>();
if(errors == 1){
errors = 0;
}else{
std::cout<<"Failed Unsigned Char"<<std::endl;
return -1;
}
errors = dataTypesRun<ushort1>() &
dataTypesRun<ushort2>() &
dataTypesRun<ushort3>() &
dataTypesRun<ushort4>();
if(errors == 1){
errors = 0;
}else{
std::cout<<"Failed Unsigned Short"<<std::endl;
return -1;
}
errors = dataTypesRun<uint1>() &
dataTypesRun<uint2>() &
dataTypesRun<uint3>() &
dataTypesRun<uint4>();
if(errors == 1){
errors = 0;
}else{
std::cout<<"Failed Unsigned Int"<<std::endl;
return -1;
}
errors = dataTypesRun<ulonglong1>() &
dataTypesRun<ulonglong2>() &
dataTypesRun<ulonglong3>() &
dataTypesRun<ulonglong4>();
if(errors == 1){
errors = 0;
}else{
std::cout<<"Failed Unsigned Long Long"<<std::endl;
return -1;
}
errors = dataTypesRun<float>() &
dataTypesRun<float1>() &
dataTypesRun<float2>() &
dataTypesRun<float3>() &
dataTypesRun<float4>();
if(errors == 1){
errors = 0;
}else{
std::cout<<"Failed Float"<<std::endl;
return -1;
}
errors = dataTypesRun<double>() &
dataTypesRun<double1>() &
dataTypesRun<double2>() &
dataTypesRun<double3>() &
dataTypesRun<double4>();
//hipResetDefaultAccelerator();
if(errors == 1){
passed();
return 0;
}else{
std::cout<<"Failed Float"<<std::endl;
return -1;
}
std::cout<<"ldg test PASSED \n"<<std::endl;
}
#endif