gdb/testsuite/gdb.rocm/show-info.cpp

   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
  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 }