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

Resolved test issues

Browse Source 
[ROCm/clr commit: 2f8d56e903]
This commit is contained in:
streamhsa
2016-02-10 20:01:16 +08:00
parent 91a5b0aa77
commit 5d857b2bc3
4 changed files with 26 additions and 17 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/tests/src/CMakeLists.txt
+1 -1
View File
@@ -126,7 +126,7 @@ make_test(hipEventRecord --iterations 10)
make_test(hipMemset " " ) make_test(hipMemset " " )
make_test(hipMemset --N 10 --memsetval 0x42 ) # small copy, just 10 bytes. make_test(hipMemset --N 10 --memsetval 0x42 ) # small copy, just 10 bytes.
make_test(hipMemset --N 10013 --memsetval 0x5a ) # oddball size. make_test(hipMemset --N 10013 --memsetval 0x5a ) # oddball size.
make_test(hipMemset --N 500M --memsetval 0xa6 ) # big copy make_test(hipMemset --N 256M --memsetval 0xa6 ) # big copy
make_test(hipGridLaunch " " ) make_test(hipGridLaunch " " )
make_test(hipMemcpy " " ) make_test(hipMemcpy " " )
projects/clr/hipamd/tests/src/hipGridLaunch.cpp
+1 -1
View File
@@ -37,7 +37,7 @@ __device__ int foo(int i)
//Syntax we would like to support with GRID_LAUNCH enabled: //Syntax we would like to support with GRID_LAUNCH enabled:
template <typename T> template <typename T>
__global__ void __global__ void
vectorADD2( grid_launch_parm lp, vectorADD2( hipLaunchParm lp,
T *A_d, T *A_d,
T *B_d, T *B_d,
T *C_d, T *C_d,
projects/clr/hipamd/tests/src/hip_anyall.cpp
+12 -6
View File
@@ -39,12 +39,17 @@ __global__ void
int main(int argc, char *argv[]) int main(int argc, char *argv[])
{ { int warpSize;
hipDeviceProp_t devProp;
hipDeviceGetProperties(&devProp, 0);
if(strncmp(devProp.name,"Fiji",1)==0) warpSize =64;
else warpSize =32;
int anycount =0;
int allcount =0;
int Num_Threads_per_Block = 1024; int Num_Threads_per_Block = 1024;
int Num_Blocks_per_Grid = 1; int Num_Blocks_per_Grid = 1;
int Num_Warps_per_Block = Num_Threads_per_Block/64; int Num_Warps_per_Block = Num_Threads_per_Block/warpSize;
int Num_Warps_per_Grid = (Num_Threads_per_Block*Num_Blocks_per_Grid)/64; int Num_Warps_per_Grid = (Num_Threads_per_Block*Num_Blocks_per_Grid)/warpSize;
int * host_any = ( int*)malloc(Num_Warps_per_Grid*sizeof(int)); int * host_any = ( int*)malloc(Num_Warps_per_Grid*sizeof(int));
int * host_all = ( int*)malloc(Num_Warps_per_Grid*sizeof(int)); int * host_all = ( int*)malloc(Num_Warps_per_Grid*sizeof(int));
@@ -69,10 +74,11 @@ for (int i=0; i<Num_Warps_per_Grid; i++)
printf("warp no. %d __any = %d \n",i,host_any[i]); printf("warp no. %d __any = %d \n",i,host_any[i]);
printf("warp no. %d __all = %d \n",i,host_all[i]); printf("warp no. %d __all = %d \n",i,host_all[i]);
if (host_any[i]!=1) ++anycount;
if (host_all[i]!=1) ++allcount;
} }
if (anycount == 0 && allcount ==1) printf("PASSED"); else printf("FAILED");
return EXIT_SUCCESS; return EXIT_SUCCESS;
projects/clr/hipamd/tests/src/hip_ballot.cpp
+12 -9
View File
@@ -15,16 +15,19 @@ __global__ void
int main(int argc, char *argv[]) int main(int argc, char *argv[])
{ { int warpSize;
hipDeviceProp_t devProp;
hipDeviceGetProperties(&devProp, 0);
if(strncmp(devProp.name,"Fiji",1)==0) warpSize =64;
else warpSize =32;
unsigned int Num_Threads_per_Block = 512; unsigned int Num_Threads_per_Block = 512;
unsigned int Num_Blocks_per_Grid = 1; unsigned int Num_Blocks_per_Grid = 1;
unsigned int Num_Warps_per_Block = Num_Threads_per_Block/64; unsigned int Num_Warps_per_Block = Num_Threads_per_Block/warpSize;
unsigned int Num_Warps_per_Grid = (Num_Threads_per_Block*Num_Blocks_per_Grid)/64; unsigned int Num_Warps_per_Grid = (Num_Threads_per_Block*Num_Blocks_per_Grid)/warpSize;
unsigned int* host_ballot = (unsigned int*)malloc(Num_Warps_per_Grid*sizeof(unsigned int)); unsigned int* host_ballot = (unsigned int*)malloc(Num_Warps_per_Grid*sizeof(unsigned int));
unsigned int* device_ballot; unsigned int* device_ballot;
HIP_ASSERT(hipMalloc((void**)&device_ballot, Num_Warps_per_Grid*sizeof(unsigned int))); HIP_ASSERT(hipMalloc((void**)&device_ballot, Num_Warps_per_Grid*sizeof(unsigned int)));
int divergent_count =0;
for (int i=0; i<Num_Warps_per_Grid; i++) host_ballot[i] = 0; for (int i=0; i<Num_Warps_per_Grid; i++) host_ballot[i] = 0;
@@ -36,13 +39,13 @@ int main(int argc, char *argv[])
HIP_ASSERT(hipMemcpy(host_ballot, device_ballot, Num_Warps_per_Grid*sizeof(unsigned int), hipMemcpyDeviceToHost)); HIP_ASSERT(hipMemcpy(host_ballot, device_ballot, Num_Warps_per_Grid*sizeof(unsigned int), hipMemcpyDeviceToHost));
for (int i=0; i<Num_Warps_per_Grid; i++) { for (int i=0; i<Num_Warps_per_Grid; i++) {
if ((host_ballot[i] == 0)||(host_ballot[i]/64 == 64)) std::cout << "Warp " << i << " IS convergent- Predicate true for " << host_ballot[i]/64 << " threads\n"; if ((host_ballot[i] == 0)||(host_ballot[i]/warpSize == warpSize)) std::cout << "Warp " << i << " IS convergent- Predicate true for " << host_ballot[i]/warpSize << " threads\n";
else std::cout << "Warp " << i << " IS divergent - Predicate true for " << host_ballot[i]/64<< " threads\n";
else {std::cout << "Warp " << i << " IS divergent - Predicate true for " << host_ballot[i]/warpSize<< " threads\n";
divergent_count++;}
} }
if (divergent_count==1) printf("PASSED"); else printf("FAILED");
return EXIT_SUCCESS; return EXIT_SUCCESS;
} }