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

bit_extract reports PASSED when passed

Browse Source
This commit is contained in:
Maneesh Gupta
2016-05-03 14:19:25 +05:30
parent caadac1005
commit ab671b729d
1 changed files with 33 additions and 32 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/samples/0_Intro/bit_extract/bit_extract.cpp
+33 -32
View File
@@ -30,7 +30,7 @@ THE SOFTWARE.
if (error != hipSuccess) { \
fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error,__FILE__, __LINE__); \
exit(EXIT_FAILURE);\
}\
}\
}
void __global__
@@ -43,59 +43,60 @@ bit_extract_kernel(hipLaunchParm lp, uint32_t *C_d, const uint32_t *A_d, size_t
#ifdef __HIP_PLATFORM_HCC__
C_d[i] = hc::__bitextract_u32(A_d[i], 8, 4);
#else /* defined __HIP_PLATFORM_NVCC__ or other path */
C_d[i] = ((A_d[i] & 0xf00) >> 8);
C_d[i] = ((A_d[i] & 0xf00) >> 8);
#endif
}
}
}
int main(int argc, char *argv[])
{
uint32_t *A_d, *C_d;
uint32_t *A_h, *C_h;
size_t N = 1000000;
size_t Nbytes = N * sizeof(uint32_t);
uint32_t *A_d, *C_d;
uint32_t *A_h, *C_h;
size_t N = 1000000;
size_t Nbytes = N * sizeof(uint32_t);
int deviceId;
CHECK (hipGetDevice(&deviceId));
hipDeviceProp_t props;
CHECK(hipGetDeviceProperties(&props, deviceId));
printf ("info: running on device #%d %s\n", deviceId, props.name);
hipDeviceProp_t props;
CHECK(hipGetDeviceProperties(&props, deviceId));
printf ("info: running on device #%d %s\n", deviceId, props.name);
printf ("info: allocate host mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
A_h = (uint32_t*)malloc(Nbytes);
CHECK(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess );
C_h = (uint32_t*)malloc(Nbytes);
CHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess );
printf ("info: allocate host mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
A_h = (uint32_t*)malloc(Nbytes);
CHECK(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess );
C_h = (uint32_t*)malloc(Nbytes);
CHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess );
for (size_t i=0; i<N; i++)
{
A_h[i] = i;
}
{
A_h[i] = i;
}
printf ("info: allocate device mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
CHECK(hipMalloc(&A_d, Nbytes));
CHECK(hipMalloc(&C_d, Nbytes));
printf ("info: allocate device mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
CHECK(hipMalloc(&A_d, Nbytes));
CHECK(hipMalloc(&C_d, Nbytes));
printf ("info: copy Host2Device\n");
printf ("info: copy Host2Device\n");
CHECK ( hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
printf ("info: launch 'bit_extract_kernel' \n");
const unsigned blocks = 512;
const unsigned threadsPerBlock = 256;
hipLaunchKernel(bit_extract_kernel, dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, N);
printf ("info: launch 'bit_extract_kernel' \n");
const unsigned blocks = 512;
const unsigned threadsPerBlock = 256;
hipLaunchKernel(bit_extract_kernel, dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, N);
printf ("info: copy Device2Host\n");
printf ("info: copy Device2Host\n");
CHECK ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
printf ("info: check result\n");
printf ("info: check result\n");
for (size_t i=0; i<N; i++) {
unsigned Agold = ((A_h[i] & 0xf00) >> 8);
if (C_h[i] != Agold) {
if (C_h[i] != Agold) {
fprintf (stderr, "mismatch detected.\n");
printf ("%zu: %08x =? %08x (Ain=%08x)\n", i, C_h[i], Agold, A_h[i]);
CHECK(hipErrorUnknown);
}
}
CHECK(hipErrorUnknown);
}
}
printf ("PASSED!\n");
}