AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

fix_ldg

Browse Source 
Change-Id: I53de5fa91b4f57d496ffe46787d197ae84dde4a4


[ROCm/clr commit: fda049fa5f]
This commit is contained in:
7SK
2016-06-06 13:20:45 +05:30
committed by Andres Rodriguez
parent d7e6475869
commit 2c5530e9f2
4 changed files with 277 additions and 372 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/tests/src/hip_test_ldg.cpp
+239 -162
View File
@@ -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