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/>. */ |
18 | |
04202a50 |
19 | #include "hip/hip_runtime.h" |
20 | #include "stdio.h" |
21 | #include <iostream> |
22 | #include <hip/hip_runtime.h> |
23 | #include <hip/hip_runtime.h> |
24 | |
25 | // Defining number of elements in Array |
26 | #define N 64 |
27 | #define NO_OF_BLOCKS 4 |
28 | #define NO_OF_THREADS 256 |
29 | |
30 | |
31 | |
32 | #define HIPCHECK(cmd) \ |
33 | do { \ |
34 | hipError_t error = (cmd); \ |
35 | if (error != hipSuccess) \ |
36 | { \ |
37 | std::cerr << "Encountered HIP error (" << error << ") at line " \ |
38 | << __LINE__ << " in file " << __FILE__ << "\n"; \ |
39 | exit(-1); \ |
40 | } \ |
41 | } while (0) |
42 | |
43 | #define MAX_GPU 8 |
44 | |
45 | |
46 | |
47 | // Defining Kernel function for vector addition |
48 | __global__ void VectorAdd(int *d_a, int *d_b, int *d_c) |
49 | { |
50 | // Getting block index of current kernel |
51 | int tid = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; |
52 | if (tid < N) |
53 | d_c[tid] = d_a[tid] + d_b[tid]; |
54 | } |
55 | |
56 | |
57 | int main(void) |
58 | { |
59 | // Defining host arrays |
60 | int h_a[N], h_b[N], h_c[N]; |
61 | // Defining device pointers |
62 | int *d_a[N], *d_b[N], *d_c[N]; |
63 | // allocate the memory |
64 | |
65 | hipStream_t stream[MAX_GPU]; |
66 | |
67 | int nGpu = 1; |
68 | HIPCHECK(hipGetDeviceCount(&nGpu)); |
69 | for (int i = 0; i < nGpu; i ++) { |
70 | HIPCHECK(hipSetDevice(i)); |
71 | hipDeviceProp_t prop; |
72 | HIPCHECK(hipGetDeviceProperties(&prop, i)); |
73 | printf("# device %d [0x%02x] %s\n", |
74 | i, prop.pciBusID, prop.name); |
75 | //create stream |
76 | HIPCHECK(hipStreamCreate(&stream[i])); |
77 | |
78 | hipMalloc((void**)&d_a[i], N * sizeof(int)); |
79 | hipMalloc((void**)&d_b[i], N * sizeof(int)); |
80 | hipMalloc((void**)&d_c[i], N * sizeof(int)); |
81 | // Initializing Arrays |
82 | for (int i = 0; i < N; i++) { |
83 | h_a[i] = 2*i; |
84 | h_b[i] = i ; |
85 | } |
86 | |
87 | // Copy input arrays from host to device memory |
88 | hipMemcpyAsync(d_a[i], h_a, N * sizeof(int), hipMemcpyHostToDevice, stream[i]); |
89 | hipMemcpyAsync(d_b[i], h_b, N * sizeof(int), hipMemcpyHostToDevice, stream[i]); |
90 | } |
91 | |
92 | for (int i = 0; i < nGpu; i ++) { |
93 | HIPCHECK(hipSetDevice(i)); |
94 | |
95 | // Calling VectorAdd kernels with NO_OF_BLOCKS and NO_OF_THREADS per block, passing |
96 | // device pointers as parameters |
97 | hipLaunchKernelGGL(VectorAdd, dim3(NO_OF_BLOCKS), dim3(NO_OF_THREADS), 0, stream[i], d_a[i], d_b[i], d_c[i]); |
98 | } |
99 | |
100 | for (int i = 0; i < nGpu; i ++) { |
101 | HIPCHECK(hipSetDevice(i)); |
102 | // Copy result back to host memory from device memory |
103 | hipMemcpyAsync(h_c, d_c[i], N * sizeof(int), hipMemcpyDeviceToHost, stream[i]); |
104 | HIPCHECK(hipStreamSynchronize(stream[i])); |
105 | //printf("Vector addition on GPU \n"); |
106 | // Printing result on console |
107 | for (int i = 0; i < N; i++) { |
108 | /*printf("Operation result of %d element is %d + %d = %d\n", |
109 | i, h_a[i], h_b[i],h_c[i]);*/ |
110 | if(h_a[i]+h_b[i] !=h_c[i]) { |
111 | HIPCHECK(hipErrorUnknown); |
112 | } |
113 | } |
114 | // Free up memory |
115 | HIPCHECK(hipStreamDestroy(stream[i])); |
116 | hipFree(d_a[i]); |
117 | hipFree(d_b[i]); |
118 | hipFree(d_c[i]); |
119 | } |
120 | return 0; |
121 | } |