Commit | Line | Data |
---|---|---|
ec980634 | 1 | /* Copyright (C) 2019-2020 Free Software Foundation, Inc. |
2 | Copyright (C) 2019-2020 Advanced Micro Devices, Inc. All rights reserved. | |
3 | ||
4 | This file is part of GDB. | |
5 | ||
6 | This program is free software; you can redistribute it and/or modify | |
7 | it under the terms of the GNU General Public License as published by | |
8 | the Free Software Foundation; either version 3 of the License, or | |
9 | (at your option) any later version. | |
10 | ||
11 | This program is distributed in the hope that it will be useful, | |
12 | but WITHOUT ANY WARRANTY; without even the implied warranty of | |
13 | MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the | |
14 | GNU General Public License for more details. | |
15 | ||
16 | You should have received a copy of the GNU General Public License | |
17 | along with this program. If not, see <http://www.gnu.org/licenses/>. */ | |
abeeff98 LM |
18 | |
19 | #include <stdio.h> | |
20 | #include <iostream> | |
21 | #include "hip/hip_runtime.h" | |
22 | ||
23 | #define CHECK(cmd) \ | |
24 | { \ | |
25 | hipError_t error = cmd; \ | |
26 | if (error != hipSuccess) { \ | |
27 | fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error, \ | |
28 | __FILE__, __LINE__); \ | |
29 | exit(EXIT_FAILURE); \ | |
30 | } \ | |
31 | } | |
32 | ||
33 | __global__ void bit_extract_kernel(uint32_t* C_d, const uint32_t* A_d, size_t N) { | |
34 | size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x); | |
35 | size_t stride = hipBlockDim_x * hipGridDim_x; | |
36 | ||
37 | for (size_t i = offset; i < N; i += stride) { | |
38 | C_d[i] = __bitextract_u32(A_d[i], 8, 4); | |
39 | } | |
40 | } | |
41 | ||
42 | ||
43 | int main(int argc, char* argv[]) { | |
44 | uint32_t *A_d, *C_d; | |
45 | uint32_t *A_h, *C_h; | |
46 | size_t N = 1000000; | |
47 | size_t Nbytes = N * sizeof(uint32_t); | |
48 | ||
49 | int deviceId; | |
50 | CHECK(hipGetDevice(&deviceId)); | |
51 | hipDeviceProp_t props; | |
52 | CHECK(hipGetDeviceProperties(&props, deviceId)); | |
53 | printf("info: running on device #%d %s\n", deviceId, props.name); | |
54 | ||
55 | ||
56 | printf("info: allocate host mem (%6.2f MB)\n", 2 * Nbytes / 1024.0 / 1024.0); | |
57 | A_h = (uint32_t*)malloc(Nbytes); | |
58 | CHECK(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess); | |
59 | C_h = (uint32_t*)malloc(Nbytes); | |
60 | CHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess); | |
61 | ||
62 | for (size_t i = 0; i < N; i++) { | |
63 | A_h[i] = i; | |
64 | } | |
65 | ||
66 | printf("info: allocate device mem (%6.2f MB)\n", 2 * Nbytes / 1024.0 / 1024.0); | |
67 | CHECK(hipMalloc(&A_d, Nbytes)); | |
68 | CHECK(hipMalloc(&C_d, Nbytes)); | |
69 | ||
70 | printf("info: copy Host2Device\n"); | |
71 | CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); | |
72 | ||
73 | printf("info: launch 'bit_extract_kernel' \n"); | |
74 | const unsigned blocks = 512; | |
75 | const unsigned threadsPerBlock = 256; | |
76 | hipLaunchKernelGGL(bit_extract_kernel, dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, N); | |
77 | ||
78 | printf("info: copy Device2Host\n"); | |
79 | CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); | |
80 | ||
81 | printf("info: check result\n"); | |
82 | for (size_t i = 0; i < N; i++) { | |
83 | unsigned Agold = ((A_h[i] & 0xf00) >> 8); | |
84 | if (C_h[i] != Agold) { | |
85 | fprintf(stderr, "mismatch detected.\n"); | |
86 | printf("%zu: %08x =? %08x (Ain=%08x)\n", i, C_h[i], Agold, A_h[i]); | |
87 | CHECK(hipErrorUnknown); | |
88 | } | |
89 | } | |
90 | printf("PASSED!\n"); | |
91 | } |