Commit | Line | Data |
---|---|---|
5922befa LM |
1 | /* |
2 | Copyright (c) 2015-present Advanced Micro Devices, Inc. All rights reserved. | |
3 | Permission is hereby granted, free of charge, to any person obtaining a copy | |
4 | of this software and associated documentation files (the "Software"), to deal | |
5 | in the Software without restriction, including without limitation the rights | |
6 | to use, copy, modify, merge, publish, distribute, sublicense, and/or sell | |
7 | copies of the Software, and to permit persons to whom the Software is | |
8 | furnished to do so, subject to the following conditions: | |
9 | The above copyright notice and this permission notice shall be included in | |
10 | all copies or substantial portions of the Software. | |
11 | THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR | |
12 | IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, | |
13 | FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE | |
14 | AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER | |
15 | LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, | |
16 | OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN | |
17 | THE SOFTWARE. | |
18 | */ | |
19 | ||
20 | #include <stdio.h> | |
21 | #include <iostream> | |
22 | #include "hip/hip_runtime.h" | |
23 | ||
24 | #define CHECK(cmd) \ | |
25 | { \ | |
26 | hipError_t error = cmd; \ | |
27 | if (error != hipSuccess) { \ | |
28 | fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error, \ | |
29 | __FILE__, __LINE__); \ | |
30 | exit(EXIT_FAILURE); \ | |
31 | } \ | |
32 | } | |
33 | ||
34 | __global__ void bit_extract_kernel(uint32_t* C_d, const uint32_t* A_d, size_t N) { | |
35 | size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x); | |
36 | size_t stride = hipBlockDim_x * hipGridDim_x; | |
37 | ||
38 | for (size_t i = offset; i < N; i += stride) { | |
39 | C_d[i] = __bitextract_u32(A_d[i], 8, 4); | |
40 | } | |
41 | } | |
42 | ||
43 | ||
44 | int main(int argc, char* argv[]) { | |
45 | uint32_t *A_d, *C_d; | |
46 | uint32_t *A_h, *C_h; | |
47 | size_t N = 1000000; | |
48 | size_t Nbytes = N * sizeof(uint32_t); | |
49 | ||
50 | int deviceId; | |
51 | CHECK(hipGetDevice(&deviceId)); | |
52 | hipDeviceProp_t props; | |
53 | CHECK(hipGetDeviceProperties(&props, deviceId)); | |
54 | printf("info: running on device #%d %s\n", deviceId, props.name); | |
55 | ||
56 | ||
57 | printf("info: allocate host mem (%6.2f MB)\n", 2 * Nbytes / 1024.0 / 1024.0); | |
58 | A_h = (uint32_t*)malloc(Nbytes); | |
59 | CHECK(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess); | |
60 | C_h = (uint32_t*)malloc(Nbytes); | |
61 | CHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess); | |
62 | ||
63 | for (size_t i = 0; i < N; i++) { | |
64 | A_h[i] = i; | |
65 | } | |
66 | ||
67 | printf("info: allocate device mem (%6.2f MB)\n", 2 * Nbytes / 1024.0 / 1024.0); | |
68 | CHECK(hipMalloc(&A_d, Nbytes)); | |
69 | CHECK(hipMalloc(&C_d, Nbytes)); | |
70 | ||
71 | printf("info: copy Host2Device\n"); | |
72 | CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); | |
73 | ||
74 | printf("info: launch 'bit_extract_kernel' \n"); | |
75 | const unsigned blocks = 512; | |
76 | const unsigned threadsPerBlock = 256; | |
77 | hipLaunchKernelGGL(bit_extract_kernel, dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, N); | |
78 | ||
79 | printf("info: copy Device2Host\n"); | |
80 | CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); | |
81 | ||
82 | printf("info: check result\n"); | |
83 | for (size_t i = 0; i < N; i++) { | |
84 | unsigned Agold = ((A_h[i] & 0xf00) >> 8); | |
85 | if (C_h[i] != Agold) { | |
86 | fprintf(stderr, "mismatch detected.\n"); | |
87 | printf("%zu: %08x =? %08x (Ain=%08x)\n", i, C_h[i], Agold, A_h[i]); | |
88 | CHECK(hipErrorUnknown); | |
89 | } | |
90 | } | |
91 | printf("PASSED!\n"); | |
92 | } |