From 620a07102d4e9dc74d4176aa71bd615bb30752e9 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Tue, 7 May 2019 05:09:02 +0530 Subject: [PATCH 1/4] Maintain HIP_VISIBLE_DEVICES for kernel launch --- src/hip_hcc.cpp | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index ac285e6a83..6f2e41cd49 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -2504,14 +2504,16 @@ hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view** a namespace hip_impl { std::vector all_hsa_agents() { std::vector r{}; - for (auto&& acc : hc::accelerator::get_all()) { + std::vector visible_accelerators; + for (int i=0; i < g_deviceCnt; i++) + visible_accelerators.push_back(g_deviceArray[i]->_acc); + for (auto&& acc : visible_accelerators) { const auto agent = acc.get_hsa_agent(); if (!agent || !acc.is_hsa_accelerator()) continue; r.emplace_back(*static_cast(agent)); } - return r; } From 9b9257f9b0366de567f7f6e4cd2d710c0d0351cc Mon Sep 17 00:00:00 2001 From: "Wen-Heng (Jack) Chung" Date: Mon, 13 May 2019 04:12:05 -0500 Subject: [PATCH 2/4] Revert "HACK for SWDEV-173477" (#1004) * Revert "HACK for SWDEV-173477" This reverts commit d941f19399e9b4f040aec5d41123e7073913cbc3. --- include/hip/hcc_detail/hip_runtime_api.h | 15 +-------------- 1 file changed, 1 insertion(+), 14 deletions(-) diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index 3eb59e485e..5b598b54a8 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -2684,20 +2684,7 @@ public: std::tie(*dptr, *bytes) = read_global_description(it0->second.cbegin(), it0->second.cend(), name); - // HACK for SWDEV-173477 - // - // For code objects with global symbols of length 0, ROCR runtime would - // ignore them even though they exist in the symbol table. Therefore the - // result from read_agent_globals() can't be trusted entirely. - // - // As a workaround to tame applications which depend on the existence of - // global symbols with length 0, always return hipSuccess here. - // - // This behavior shall be reverted once ROCR runtime has been fixed to - // address SWDEV-173477 - - //return *dptr ? hipSuccess : hipErrorNotFound; - return hipSuccess; + return *dptr ? hipSuccess : hipErrorNotFound; } hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes, From 4b861bca3947bc1772c7b2b1274cf62579220856 Mon Sep 17 00:00:00 2001 From: emankov Date: Mon, 13 May 2019 19:37:18 +0300 Subject: [PATCH 3/4] [HIPIFY][tests] Add reverse engineered HIP sample MatrixTranspose --- .../0_MatrixTranspose/MatrixTranspose.cpp | 130 ++++++++++++++++++ 1 file changed, 130 insertions(+) create mode 100644 tests/hipify-clang/unit_tests/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp diff --git a/tests/hipify-clang/unit_tests/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp b/tests/hipify-clang/unit_tests/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp new file mode 100644 index 0000000000..2898587143 --- /dev/null +++ b/tests/hipify-clang/unit_tests/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp @@ -0,0 +1,130 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +/* +Copyright (c) 2015-2019 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. +*/ + +#include + +// CHECK: #include +#include + +#define WIDTH 1024 + +#define NUM (WIDTH * WIDTH) + +#define THREADS_PER_BLOCK_X 4 +#define THREADS_PER_BLOCK_Y 4 +#define THREADS_PER_BLOCK_Z 1 + +// Device (Kernel) function, it must be void +__global__ void matrixTranspose(float* out, float* in, const int width) { + int x = blockDim.x * blockIdx.x + threadIdx.x; + int y = blockDim.y * blockIdx.y + threadIdx.y; + out[y * width + x] = in[x * width + y]; +} + +// CPU implementation of matrix transpose +void matrixTransposeCPUReference(float* output, float* input, const unsigned int width) { + for (unsigned int j = 0; j < width; j++) { + for (unsigned int i = 0; i < width; i++) { + output[i * width + j] = input[j * width + i]; + } + } +} + +int main() { + float* Matrix; + float* TransposeMatrix; + float* cpuTransposeMatrix; + + float* gpuMatrix; + float* gpuTransposeMatrix; + + // CHECK: hipDeviceProp_t devProp; + cudaDeviceProp devProp; + // CHECK: hipGetDeviceProperties(&devProp, 0); + cudaGetDeviceProperties(&devProp, 0); + + std::cout << "Device name " << devProp.name << std::endl; + + int i; + int errors; + + Matrix = (float*)malloc(NUM * sizeof(float)); + TransposeMatrix = (float*)malloc(NUM * sizeof(float)); + cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float)); + + // initialize the input data + for (i = 0; i < NUM; i++) { + Matrix[i] = (float)i * 10.0f; + } + + // allocate the memory on the device side + // CHECK: hipMalloc((void**)&gpuMatrix, NUM * sizeof(float)); + cudaMalloc((void**)&gpuMatrix, NUM * sizeof(float)); + // CHECK: hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float)); + cudaMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float)); + + // Memory transfer from host to device + // CHECK: hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice); + cudaMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), cudaMemcpyHostToDevice); + + // Lauching kernel from host + + dim3 dimGrid(WIDTH / THREADS_PER_BLOCK_X, WIDTH / THREADS_PER_BLOCK_Y); + dim3 dimBlock(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y); + // CHECK: hipLaunchKernelGGL(matrixTranspose, dim3(dimGrid), dim3(dimBlock), 0, 0, gpuTransposeMatrix, gpuMatrix, WIDTH); + matrixTranspose <<>>(gpuTransposeMatrix, gpuMatrix, WIDTH); + + // Memory transfer from device to host + // CHECK: hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost); + cudaMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), cudaMemcpyDeviceToHost); + + // CPU MatrixTranspose computation + matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH); + + // verify the results + errors = 0; + double eps = 1.0E-6; + for (i = 0; i < NUM; i++) { + if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps) { + errors++; + } + } + if (errors != 0) { + printf("FAILED: %d errors\n", errors); + } else { + printf("PASSED!\n"); + } + + // free the resources on device side + // CHECK: hipFree(gpuMatrix); + cudaFree(gpuMatrix); + // CHECK: hipFree(gpuTransposeMatrix); + cudaFree(gpuTransposeMatrix); + + // free the resources on host side + free(Matrix); + free(TransposeMatrix); + free(cpuTransposeMatrix); + + return errors; +} From 3bc3b61fb4dfb0eaf7102394e404f183fae8c599 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 13 May 2019 22:12:43 +0300 Subject: [PATCH 4/4] [HIPIFY][tests] Add reverse engineered HIP sample hipEvent --- .../samples/1_hipEvent/hipEvent.cpp | 181 ++++++++++++++++++ 1 file changed, 181 insertions(+) create mode 100644 tests/hipify-clang/unit_tests/samples/1_hipEvent/hipEvent.cpp diff --git a/tests/hipify-clang/unit_tests/samples/1_hipEvent/hipEvent.cpp b/tests/hipify-clang/unit_tests/samples/1_hipEvent/hipEvent.cpp new file mode 100644 index 0000000000..1f5e90dd3a --- /dev/null +++ b/tests/hipify-clang/unit_tests/samples/1_hipEvent/hipEvent.cpp @@ -0,0 +1,181 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +/* +Copyright (c) 2015-2019 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. +*/ + +#include + +// CHECK: #include +#include + +#define WIDTH 1024 + +#define NUM (WIDTH * WIDTH) + +#define THREADS_PER_BLOCK_X 4 +#define THREADS_PER_BLOCK_Y 4 +#define THREADS_PER_BLOCK_Z 1 + +// Device (Kernel) function, it must be void +__global__ void matrixTranspose(float* out, float* in, const int width) { + int x = blockDim.x * blockIdx.x + threadIdx.x; + int y = blockDim.y * blockIdx.y + threadIdx.y; + out[y * width + x] = in[x * width + y]; +} + +// CPU implementation of matrix transpose +void matrixTransposeCPUReference(float* output, float* input, const unsigned int width) { + for (unsigned int j = 0; j < width; j++) { + for (unsigned int i = 0; i < width; i++) { + output[i * width + j] = input[j * width + i]; + } + } +} + +int main() { + float* Matrix; + float* TransposeMatrix; + float* cpuTransposeMatrix; + + float* gpuMatrix; + float* gpuTransposeMatrix; + + // CHECK: hipDeviceProp_t devProp; + cudaDeviceProp devProp; + // CHECK: hipGetDeviceProperties(&devProp, 0); + cudaGetDeviceProperties(&devProp, 0); + + std::cout << "Device name " << devProp.name << std::endl; + + // CHECK: hipEvent_t start, stop; + cudaEvent_t start, stop; + // CHECK: hipEventCreate(&start); + cudaEventCreate(&start); + // CHECK: hipEventCreate(&stop); + cudaEventCreate(&stop); + float eventMs = 1.0f; + + int i; + int errors; + + Matrix = (float*)malloc(NUM * sizeof(float)); + TransposeMatrix = (float*)malloc(NUM * sizeof(float)); + cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float)); + + // initialize the input data + for (i = 0; i < NUM; i++) { + Matrix[i] = (float)i * 10.0f; + } + + // allocate the memory on the device side + // CHECK: hipMalloc((void**)&gpuMatrix, NUM * sizeof(float)); + cudaMalloc((void**)&gpuMatrix, NUM * sizeof(float)); + // CHECK: hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float)); + cudaMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float)); + + // Record the start event + // CHECK: hipEventRecord(start, NULL); + cudaEventRecord(start, NULL); + + // Memory transfer from host to device + // CHECK: hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice); + cudaMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), cudaMemcpyHostToDevice); + + // Record the stop event + // CHECK: hipEventRecord(stop, NULL); + cudaEventRecord(stop, NULL); + // CHECK: hipEventSynchronize(stop); + cudaEventSynchronize(stop); + // CHECK: hipEventElapsedTime(&eventMs, start, stop); + cudaEventElapsedTime(&eventMs, start, stop); + + printf("hipMemcpyHostToDevice time taken = %6.3fms\n", eventMs); + + // Record the start event + // CHECK: hipEventRecord(start, NULL); + cudaEventRecord(start, NULL); + + // Lauching kernel from host + + dim3 dimGrid(WIDTH / THREADS_PER_BLOCK_X, WIDTH / THREADS_PER_BLOCK_Y); + dim3 dimBlock(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y); + // CHECK: hipLaunchKernelGGL(matrixTranspose, dim3(dimGrid), dim3(dimBlock), 0, 0, gpuTransposeMatrix, gpuMatrix, WIDTH); + matrixTranspose <<>>(gpuTransposeMatrix, gpuMatrix, WIDTH); + + // Record the stop event + // CHECK: hipEventRecord(stop, NULL); + cudaEventRecord(stop, NULL); + // CHECK: hipEventSynchronize(stop); + cudaEventSynchronize(stop); + + // CHECK: hipEventElapsedTime(&eventMs, start, stop); + cudaEventElapsedTime(&eventMs, start, stop); + + printf("kernel Execution time = %6.3fms\n", eventMs); + + // Record the start event + // CHECK: hipEventRecord(start, NULL); + cudaEventRecord(start, NULL); + + // Memory transfer from device to host + // CHECK: hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost); + cudaMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), cudaMemcpyDeviceToHost); + + // Record the stop event + // CHECK: hipEventRecord(stop, NULL); + cudaEventRecord(stop, NULL); + // CHECK: hipEventSynchronize(stop); + cudaEventSynchronize(stop); + // CHECK: hipEventElapsedTime(&eventMs, start, stop); + cudaEventElapsedTime(&eventMs, start, stop); + + printf("hipMemcpyDeviceToHost time taken = %6.3fms\n", eventMs); + + // CPU MatrixTranspose computation + matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH); + + // verify the results + errors = 0; + double eps = 1.0E-6; + for (i = 0; i < NUM; i++) { + if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps) { + errors++; + } + } + if (errors != 0) { + printf("FAILED: %d errors\n", errors); + } else { + printf("PASSED!\n"); + } + + // free the resources on device side + // CHECK: hipFree(gpuMatrix); + cudaFree(gpuMatrix); + // CHECK: hipFree(gpuTransposeMatrix); + cudaFree(gpuTransposeMatrix); + + // free the resources on host side + free(Matrix); + free(TransposeMatrix); + free(cpuTransposeMatrix); + + return errors; +}