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

fix_testcase

Browse Source 
Change-Id: I878e79ac29b685672662db6ba8cd8d94d725d24f
This commit is contained in:
7SK
2016-06-08 16:19:43 +05:30
committed by Kent Russell
parent fda049fa5f
commit 0201bf8191
4 changed files with 111 additions and 155 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/tests/src/hip_popc.cpp
+19 -29
View File
@@ -52,12 +52,12 @@ unsigned int popcountCPU( T value) {
}
return ret;
}
__global__ void
HIP_kernel(hipLaunchParm lp,
unsigned int* a, unsigned int* b,unsigned int* c, unsigned long long int* d, int width, int height)
__global__ void
HIP_kernel(hipLaunchParm lp,
unsigned int* a, unsigned int* b,unsigned int* c, unsigned long long int* d, int width, int height)
{
int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
@@ -74,10 +74,10 @@ HIP_kernel(hipLaunchParm lp,
using namespace std;
int main() {
unsigned int* hostA;
unsigned int* hostB;
unsigned int* hostC;
unsigned int* hostC;
unsigned long long int* hostD;
unsigned int* deviceA;
@@ -107,17 +107,17 @@ int main() {
hostB[i] = i;
hostD[i] = 1099511627776-i;
}
HIP_ASSERT(hipMalloc((void**)&deviceA, NUM * sizeof(unsigned int)));
HIP_ASSERT(hipMalloc((void**)&deviceB, NUM * sizeof(unsigned int)));
HIP_ASSERT(hipMalloc((void**)&deviceC, NUM * sizeof(unsigned int)));
HIP_ASSERT(hipMalloc((void**)&deviceD, NUM * sizeof(unsigned long long int)));
HIP_ASSERT(hipMemcpy(deviceB, hostB, NUM*sizeof(unsigned int), hipMemcpyHostToDevice));
HIP_ASSERT(hipMemcpy(deviceD, hostD, NUM*sizeof(unsigned long long int), hipMemcpyHostToDevice));
hipLaunchKernel(HIP_kernel,
hipLaunchKernel(HIP_kernel,
dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y),
dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
0, 0,
@@ -129,28 +129,31 @@ int main() {
// verify the results
errors = 0;
for (i = 0; i < NUM; i++) {
printf("gpu_popc =%d, cpu_popc =%d \n",hostA[i],popcountCPU(hostB[i]));
if (hostA[i] != popcountCPU(hostB[i])) {
errors++;
}
}
if (errors!=0) {
printf("FAILED: %d errors\n",errors);
cout << "FAILED: popc" << endl;
return -1;
} else {
printf ("__popc() PASSED!\n");
cout << "__popc() checked!" << endl;
}
errors = 0;
for (i = 0; i < NUM; i++) {
printf("gpu_popcll =%d, cpu_popcll =%d \n",hostC[i],popcountCPU(hostD[i]));
if (hostC[i] != popcountCPU(hostD[i])) {
errors++;
}
}
if (errors!=0) {
printf("FAILED: %d errors\n",errors);
cout << "FAILED:popc" << endl;
return -1;
} else {
printf ("__popcll() PASSED!\n");
cout << "__popcll() checked!" << endl;
}
cout << "popc test PASSED!" << endl;
HIP_ASSERT(hipFree(deviceA));
HIP_ASSERT(hipFree(deviceB));
HIP_ASSERT(hipFree(deviceC));
@@ -161,19 +164,6 @@ int main() {
free(hostC);
free(hostD);
//hipResetDefaultAccelerator();
return errors;
}