diff --git a/tests/featuretests/profiler/apps/hello_world_hip.cpp b/tests/featuretests/profiler/apps/hello_world_hip.cpp index a3842dba31..8aa6fd0751 100755 --- a/tests/featuretests/profiler/apps/hello_world_hip.cpp +++ b/tests/featuretests/profiler/apps/hello_world_hip.cpp @@ -25,6 +25,7 @@ THE SOFTWARE. #include #include +#include "utils/test_helper.h" #include #include @@ -33,44 +34,39 @@ THE SOFTWARE. #define SUCCESS 0 #define FAILURE 1 -__global__ void helloworld(char *in, char *out) { +__global__ void helloworld(char* in, char* out) { int num = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x; out[num] = in[num] + 1; } -int main(int argc, char *argv[]) { +int main(int argc, char* argv[]) { hipDeviceProp_t devProp; - hipGetDeviceProperties(&devProp, 0); + HIP_RC(hipGetDeviceProperties(&devProp, 0)); std::cout << " System minor " << devProp.minor << std::endl; std::cout << " System major " << devProp.major << std::endl; std::cout << " agent prop name " << devProp.name << std::endl; /* Initial input,output for the host and create memory objects for the * kernel*/ - const char *input = "GdkknVnqkc"; + const char* input = "GdkknVnqkc"; size_t strlength = strlen(input); std::cout << "input string:" << std::endl; std::cout << input << std::endl; - char *output = reinterpret_cast(malloc(strlength + 1)); + char* output = reinterpret_cast(malloc(strlength + 1)); - char *inputBuffer; - char *outputBuffer; - hipMalloc(reinterpret_cast(&inputBuffer), - (strlength + 1) * sizeof(char)); - hipMalloc(reinterpret_cast(&outputBuffer), - (strlength + 1) * sizeof(char)); + char* inputBuffer; + char* outputBuffer; + HIP_RC(hipMalloc(reinterpret_cast(&inputBuffer), (strlength + 1) * sizeof(char))); + HIP_RC(hipMalloc(reinterpret_cast(&outputBuffer), (strlength + 1) * sizeof(char))); - hipMemcpy(inputBuffer, input, (strlength + 1) * sizeof(char), - hipMemcpyHostToDevice); + HIP_RC(hipMemcpy(inputBuffer, input, (strlength + 1) * sizeof(char), hipMemcpyHostToDevice)); - hipLaunchKernelGGL(helloworld, dim3(1), dim3(strlength), 0, 0, inputBuffer, - outputBuffer); + HIP_KL(hipLaunchKernelGGL(helloworld, dim3(1), dim3(strlength), 0, 0, inputBuffer, outputBuffer)); - hipMemcpy(output, outputBuffer, (strlength + 1) * sizeof(char), - hipMemcpyDeviceToHost); + HIP_RC(hipMemcpy(output, outputBuffer, (strlength + 1) * sizeof(char), hipMemcpyDeviceToHost)); - hipFree(inputBuffer); - hipFree(outputBuffer); + HIP_RC(hipFree(inputBuffer)); + HIP_RC(hipFree(outputBuffer)); output[strlength] = '\0'; // Add the terminal character to the end of output. std::cout << "\noutput string:" << std::endl; diff --git a/tests/featuretests/profiler/apps/hello_world_omp.cpp b/tests/featuretests/profiler/apps/hello_world_omp.cpp index d1e7ec6246..0b240d4ac7 100755 --- a/tests/featuretests/profiler/apps/hello_world_omp.cpp +++ b/tests/featuretests/profiler/apps/hello_world_omp.cpp @@ -31,39 +31,31 @@ THE SOFTWARE. // HIP header #include +#include "utils/test_helper.h" #define NUM_THREADS 16 -#define CHECK(cmd) \ - { \ - hipError_t error = cmd; \ - if (error != hipSuccess) { \ - fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), \ - error, __FILE__, __LINE__); \ - exit(EXIT_FAILURE); \ - } \ - } -__global__ void hip_helloworld(unsigned omp_id, int *A_d) { +__global__ void hip_helloworld(unsigned omp_id, int* A_d) { // Note: the printf command will only work if printf is enabled in your build. // printf("Hello World... from HIP thread = %u\n", omp_id); A_d[omp_id] = omp_id; } -int main(int argc, char *argv[]) { +int main(int argc, char* argv[]) { int *A_h, *A_d; size_t Nbytes = NUM_THREADS * sizeof(int); hipDeviceProp_t props; - CHECK(hipGetDeviceProperties(&props, 0 /*deviceID*/)); + HIP_RC(hipGetDeviceProperties(&props, 0 /*deviceID*/)); // printf("info: running on device %s\n", props.name); - A_h = reinterpret_cast(malloc(Nbytes)); - CHECK(hipMalloc(&A_d, Nbytes)); + A_h = reinterpret_cast(malloc(Nbytes)); + HIP_RC(hipMalloc(&A_d, Nbytes)); for (int i = 0; i < NUM_THREADS; i++) { A_h[i] = 0; } - CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIP_RC(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); // Beginning of parallel region #pragma omp parallel num_threads(NUM_THREADS) @@ -71,13 +63,12 @@ int main(int argc, char *argv[]) { // fprintf(stderr, "Hello World... from OMP thread = %d\n", // omp_get_thread_num()); - hipLaunchKernelGGL(hip_helloworld, dim3(1), dim3(1), 0, 0, - omp_get_thread_num(), A_d); + HIP_KL(hipLaunchKernelGGL(hip_helloworld, dim3(1), dim3(1), 0, 0, omp_get_thread_num(), A_d)); } // Ending of parallel region - hipStreamSynchronize(0); - CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost)); + HIP_RC(hipStreamSynchronize(0)); + HIP_RC(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost)); // printf("Device Results:\n"); for (int i = 0; i < NUM_THREADS; i++) { // printf(" A_d[%d] = %d\n", i, A_h[i]); @@ -86,6 +77,6 @@ int main(int argc, char *argv[]) { printf("PASSED!\n"); free(A_h); - CHECK(hipFree(A_d)); + HIP_RC(hipFree(A_d)); return 0; } diff --git a/tests/featuretests/profiler/apps/vector_add_hip.cpp b/tests/featuretests/profiler/apps/vector_add_hip.cpp index a12a641e0b..c43223a16b 100755 --- a/tests/featuretests/profiler/apps/vector_add_hip.cpp +++ b/tests/featuretests/profiler/apps/vector_add_hip.cpp @@ -27,8 +27,7 @@ THE SOFTWARE. #include #include "hip/hip_runtime.h" - -#define HIP_ASSERT(x) (assert((x) == hipSuccess)) +#include "utils/test_helper.h" #define WIDTH 1024 #define HEIGHT 1024 @@ -39,10 +38,8 @@ THE SOFTWARE. #define THREADS_PER_BLOCK_Y 16 #define THREADS_PER_BLOCK_Z 1 -__global__ void vectoradd_float(float *__restrict__ a, - const float *__restrict__ b, - const float *__restrict__ c, int width, - int height) { +__global__ void vectoradd_float(float* __restrict__ a, const float* __restrict__ b, + const float* __restrict__ c, int width, int height) { int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; @@ -53,16 +50,16 @@ __global__ void vectoradd_float(float *__restrict__ a, } int main() { - float *hostA; - float *hostB; - float *hostC; + float* hostA; + float* hostB; + float* hostC; - float *deviceA; - float *deviceB; - float *deviceC; + float* deviceA; + float* deviceB; + float* deviceC; hipDeviceProp_t devProp; - hipGetDeviceProperties(&devProp, 0); + HIP_RC(hipGetDeviceProperties(&devProp, 0)); std::cout << " System minor " << devProp.minor << std::endl; std::cout << " System major " << devProp.major << std::endl; std::cout << " agent prop name " << devProp.name << std::endl; @@ -72,9 +69,9 @@ int main() { int i; int errors; - hostA = reinterpret_cast(malloc(NUM * sizeof(float))); - hostB = reinterpret_cast(malloc(NUM * sizeof(float))); - hostC = reinterpret_cast(malloc(NUM * sizeof(float))); + hostA = reinterpret_cast(malloc(NUM * sizeof(float))); + hostB = reinterpret_cast(malloc(NUM * sizeof(float))); + hostC = reinterpret_cast(malloc(NUM * sizeof(float))); // initialize the input data for (i = 0; i < NUM; i++) { @@ -82,26 +79,19 @@ int main() { hostC[i] = static_cast(i) * 100.0f; } - HIP_ASSERT( - hipMalloc(reinterpret_cast(&deviceA), NUM * sizeof(float))); - HIP_ASSERT( - hipMalloc(reinterpret_cast(&deviceB), NUM * sizeof(float))); - HIP_ASSERT( - hipMalloc(reinterpret_cast(&deviceC), NUM * sizeof(float))); + HIP_RC(hipMalloc(reinterpret_cast(&deviceA), NUM * sizeof(float))); + HIP_RC(hipMalloc(reinterpret_cast(&deviceB), NUM * sizeof(float))); + HIP_RC(hipMalloc(reinterpret_cast(&deviceC), NUM * sizeof(float))); - HIP_ASSERT( - hipMemcpy(deviceB, hostB, NUM * sizeof(float), hipMemcpyHostToDevice)); - HIP_ASSERT( - hipMemcpy(deviceC, hostC, NUM * sizeof(float), hipMemcpyHostToDevice)); + HIP_RC(hipMemcpy(deviceB, hostB, NUM * sizeof(float), hipMemcpyHostToDevice)); + HIP_RC(hipMemcpy(deviceC, hostC, NUM * sizeof(float), hipMemcpyHostToDevice)); - hipLaunchKernelGGL( - vectoradd_float, - dim3(WIDTH / THREADS_PER_BLOCK_X, HEIGHT / THREADS_PER_BLOCK_Y), - dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), 0, 0, deviceA, deviceB, - deviceC, WIDTH, HEIGHT); + HIP_KL(hipLaunchKernelGGL(vectoradd_float, + dim3(WIDTH / THREADS_PER_BLOCK_X, HEIGHT / THREADS_PER_BLOCK_Y), + dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), 0, 0, deviceA, deviceB, + deviceC, WIDTH, HEIGHT)); - HIP_ASSERT( - hipMemcpy(hostA, deviceA, NUM * sizeof(float), hipMemcpyDeviceToHost)); + HIP_RC(hipMemcpy(hostA, deviceA, NUM * sizeof(float), hipMemcpyDeviceToHost)); // verify the results errors = 0; @@ -116,9 +106,9 @@ int main() { printf("PASSED!\n"); } - HIP_ASSERT(hipFree(deviceA)); - HIP_ASSERT(hipFree(deviceB)); - HIP_ASSERT(hipFree(deviceC)); + HIP_RC(hipFree(deviceA)); + HIP_RC(hipFree(deviceB)); + HIP_RC(hipFree(deviceC)); free(hostA); free(hostB); diff --git a/tests/featuretests/profiler/apps/vector_add_mpi.cpp b/tests/featuretests/profiler/apps/vector_add_mpi.cpp index f0d406a6ee..8c2dbc3192 100755 --- a/tests/featuretests/profiler/apps/vector_add_mpi.cpp +++ b/tests/featuretests/profiler/apps/vector_add_mpi.cpp @@ -27,44 +27,21 @@ THE SOFTWARE. #include #include "hip/hip_runtime.h" - -#define HIP_RC(call) \ - do { \ - hipError_t err = call; \ - if (hipSuccess != err) { \ - printf("HIP ERROR (code = %d, %s) at %s:%d\n", err, \ - hipGetErrorString(err), __FILE__, __LINE__); \ - assert(0); \ - exit(1); \ - } \ - } while (0) - -#define HIP_KL(call) \ - do { \ - call; \ - hipError_t err = hipGetLastError(); \ - if (hipSuccess != err) { \ - printf("HIP ERROR (code = %d, %s) at %s:%d\n", err, \ - hipGetErrorString(err), __FILE__, __LINE__); \ - assert(0); \ - exit(1); \ - } \ - } while (0) +#include "utils/test_helper.h" // CUDA kernel to add elements of two arrays -__global__ void add(int n, float *x, float *y) { +__global__ void add(int n, float* x, float* y) { int index = blockIdx.x * blockDim.x + threadIdx.x; int stride = blockDim.x * gridDim.x; - for (int i = index; i < n; i += stride) - y[i] = x[i] + y[i]; + for (int i = index; i < n; i += stride) y[i] = x[i] + y[i]; } -int main(int argc, char *argv[]) { +int main(int argc, char* argv[]) { int N = 1 << 20; - float *x = new float[N]; - float *y = new float[N]; - float *d_x; - float *d_y; + float* x = new float[N]; + float* y = new float[N]; + float* d_x; + float* d_y; int myId; int devId; @@ -75,24 +52,22 @@ int main(int argc, char *argv[]) { MPI_Init(&argc, &argv); MPI_Comm_rank(MPI_COMM_WORLD, &myId); MPI_Comm_size(MPI_COMM_WORLD, &numRank); - hipGetDeviceCount(&deviceCount); + HIP_RC(hipGetDeviceCount(&deviceCount)); - std::cout << "device count and rank is" << deviceCount << ": " << numRank - << std::endl; + std::cout << "device count and rank is" << deviceCount << ": " << numRank << std::endl; // set the device ID to the rank ID mod deviceCount (in this case 4 since // there are 4 devices on a node) devId = myId % deviceCount; // set the device ID - hipSetDevice(devId); + HIP_RC(hipSetDevice(devId)); - printf("Rank Id: %d | Device Id : %d | Num Devices: %d\n", myId, devId, - deviceCount); + printf("Rank Id: %d | Device Id : %d | Num Devices: %d\n", myId, devId, deviceCount); fflush(stdout); // Allocate Unified Memory -- accessible from CPU or GPU - hipMallocManaged(&d_x, N * sizeof(float)); - hipMallocManaged(&d_y, N * sizeof(float)); + HIP_RC(hipMallocManaged(&d_x, N * sizeof(float))); + HIP_RC(hipMallocManaged(&d_y, N * sizeof(float))); // initialize x and y arrays on the host for (int i = 0; i < N; i++) { @@ -116,8 +91,7 @@ int main(int argc, char *argv[]) { // Check for errors (all values should be 3.0f) float maxError = 0.0f; - for (int i = 0; i < N; i++) - maxError = fmax(maxError, fabs(y[i] - 3.0f)); + for (int i = 0; i < N; i++) maxError = fmax(maxError, fabs(y[i] - 3.0f)); printf("Max error: %f\n", maxError); // Free memory diff --git a/tests/featuretests/utils/test_helper.h b/tests/featuretests/utils/test_helper.h new file mode 100644 index 0000000000..099296f081 --- /dev/null +++ b/tests/featuretests/utils/test_helper.h @@ -0,0 +1,49 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#ifndef TESTS_FEATURETESTS_PROFILER_UTILS_TEST_HELPER_H_ +#define TESTS_FEATURETESTS_PROFILER_UTILS_TEST_HELPER_H_ + +#include +#include + +#define HIP_RC(call) \ + do { \ + hipError_t err = call; \ + if (hipSuccess != err) { \ + fprintf(stderr, "HIP ERROR (%s) at %s:%d\n", hipGetErrorString(err), __FILE__, __LINE__); \ + assert(0); \ + exit(1); \ + } \ + } while (0) + +#define HIP_KL(call) \ + do { \ + call; \ + hipError_t err = hipGetLastError(); \ + if (hipSuccess != err) { \ + fprintf(stderr, "HIP ERROR (%s) at %s:%d\n", hipGetErrorString(err), __FILE__, __LINE__); \ + assert(0); \ + exit(1); \ + } \ + } while (0) + +#endif // TESTS_FEATURETESTS_PROFILER_UTILS_TEST_HELPER_H_ \ No newline at end of file