diff --git a/hip-config-clang.cmake.in b/hip-config-clang.cmake.in index de7858a63c..7344458ffd 100644 --- a/hip-config-clang.cmake.in +++ b/hip-config-clang.cmake.in @@ -50,14 +50,24 @@ set_and_check( hip_BIN_INSTALL_DIR "@PACKAGE_BIN_INSTALL_DIR@" ) set_and_check(hip_HIPCC_EXECUTABLE "${hip_BIN_INSTALL_DIR}/hipcc") set_and_check(hip_HIPCONFIG_EXECUTABLE "${hip_BIN_INSTALL_DIR}/hipconfig") -get_filename_component(HIP_CLANG_ROOT "${CMAKE_CXX_COMPILER}" PATH) -get_filename_component(HIP_CLANG_ROOT "${HIP_CLANG_ROOT}" PATH) +if(CMAKE_CXX_COMPILER MATCHES ".*hipcc") + execute_process(COMMAND ${CMAKE_CXX_COMPILER} --version + OUTPUT_STRIP_TRAILING_WHITESPACE + OUTPUT_VARIABLE HIP_CLANG_CXX_COMPILER_VERSION_OUTPUT) + if(HIP_CLANG_CXX_COMPILER_VERSION_OUTPUT MATCHES "InstalledDir:[\t\r\n][\t\r\n]*([^\t\r\n])") + set(HIP_CLANG_ROOT ${CMAKE_MATCH_1}) + else() + set(HIP_CLANG_ROOT /opt/rocm/llvm) + endif() +else() + get_filename_component(HIP_CLANG_ROOT "${CMAKE_CXX_COMPILER}" PATH) + get_filename_component(HIP_CLANG_ROOT "${HIP_CLANG_ROOT}" PATH) +endif() file(GLOB HIP_CLANG_INCLUDE_SEARCH_PATHS ${HIP_CLANG_ROOT}/lib/clang/*/include) find_path(HIP_CLANG_INCLUDE_PATH stddef.h HINTS ${HIP_CLANG_INCLUDE_SEARCH_PATHS} NO_DEFAULT_PATH) - find_dependency(amd_comgr) find_dependency(AMDDeviceLibs) set(AMDGPU_TARGETS "gfx900;gfx906" CACHE STRING "AMD GPU targets to compile for") diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index 025688e98c..cf6a64ad65 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -128,6 +128,14 @@ typedef struct hipDeviceProp_t { int kernelExecTimeoutEnabled; /// + +// Global string constants don't work inside device functions, so we +// use a macro to repeat the declaration in host and device contexts. +DECLARE_DATA(); + +__global__ void kernel_uniform0(int *retval) { + uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + retval[tid] = printf("Hello World\n"); +} + +static void test_uniform0(int *retval, uint num_blocks, + uint threads_per_block) { + CaptureStream captured(stdout); + + uint num_threads = num_blocks * threads_per_block; + for (uint i = 0; i != num_threads; ++i) { + retval[i] = 0x23232323; + } + + hipLaunchKernelGGL(kernel_uniform0, dim3(num_blocks), dim3(threads_per_block), + 0, 0, retval); + hipStreamSynchronize(0); + auto CapturedData = captured.getCapturedData(); + + for (uint ii = 0; ii != num_threads; ++ii) { + HIPASSERT(retval[ii] == strlen("Hello World\n")); + } + + std::map linecount; + for (std::string line; std::getline(CapturedData, line);) { + linecount[line]++; + } + + HIPASSERT(linecount.size() == 1); + HIPASSERT(linecount["Hello World"] == num_threads); +} + +__global__ void kernel_uniform1(int *retval) { + uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + retval[tid] = printf("Six times Eight is %d\n", 42); +} + +static void test_uniform1(int *retval, uint num_blocks, + uint threads_per_block) { + CaptureStream captured(stdout); + + uint num_threads = num_blocks * threads_per_block; + for (uint i = 0; i != num_threads; ++i) { + retval[i] = 0x23232323; + } + + hipLaunchKernelGGL(kernel_uniform1, dim3(num_blocks), dim3(threads_per_block), + 0, 0, retval); + hipStreamSynchronize(0); + auto CapturedData = captured.getCapturedData(); + + for (uint ii = 0; ii != num_threads; ++ii) { + HIPASSERT(retval[ii] == strlen("Six times Eight is 42") + 1); + } + + std::map linecount; + for (std::string line; std::getline(CapturedData, line);) { + linecount[line]++; + } + + HIPASSERT(linecount.size() == 1); + HIPASSERT(linecount["Six times Eight is 42"] == num_threads); +} + +__global__ void kernel_divergent0(int *retval) { + uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + retval[tid] = printf("Thread ID: %d\n", tid); +} + +static void test_divergent0(int *retval, uint num_blocks, + uint threads_per_block) { + CaptureStream captured(stdout); + + uint num_threads = num_blocks * threads_per_block; + for (uint i = 0; i != num_threads; ++i) { + retval[i] = 0x23232323; + } + + hipLaunchKernelGGL(kernel_divergent0, dim3(num_blocks), + dim3(threads_per_block), 0, 0, retval); + hipStreamSynchronize(0); + auto CapturedData = captured.getCapturedData(); + + for (uint ii = 0; ii != 10; ++ii) { + HIPASSERT(retval[ii] == 13); + } + + for (uint ii = 10; ii != num_threads; ++ii) { + HIPASSERT(retval[ii] == 14); + } + + std::vector threadIds; + for (std::string line; std::getline(CapturedData, line);) { + auto pos = line.find(':'); + HIPASSERT(line.substr(0, pos) == "Thread ID"); + threadIds.push_back(std::stoul(line.substr(pos + 2))); + } + + std::sort(threadIds.begin(), threadIds.end()); + HIPASSERT(threadIds.size() == num_threads); + HIPASSERT(threadIds.back() == num_threads - 1); +} + +__global__ void kernel_divergent1(int *retval) { + uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + if (tid % 2) { + retval[tid] = printf("Hello World\n"); + } else { + retval[tid] = -1; + } +} + +static void test_divergent1(int *retval, uint num_blocks, + uint threads_per_block) { + CaptureStream captured(stdout); + + uint num_threads = num_blocks * threads_per_block; + for (uint i = 0; i != num_threads; ++i) { + retval[i] = 0x23232323; + } + + hipLaunchKernelGGL(kernel_divergent1, dim3(num_blocks), + dim3(threads_per_block), 0, 0, retval); + hipStreamSynchronize(0); + auto CapturedData = captured.getCapturedData(); + + for (uint ii = 0; ii != num_threads; ++ii) { + if (ii % 2) { + HIPASSERT(retval[ii] == strlen("Hello World\n")); + } else { + HIPASSERT(retval[ii] == -1); + } + } + + std::map linecount; + for (std::string line; std::getline(CapturedData, line);) { + linecount[line]++; + } + + HIPASSERT(linecount.size() == 1); + HIPASSERT(linecount["Hello World"] == num_threads / 2); +} + +__global__ void kernel_series(int *retval) { + DECLARE_DATA(); + + const uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + int result = 0; + + result += printf("%s\n", msg_long1); + result += printf("%s\n", msg_short); + result += printf("%s\n", msg_long2); + + retval[tid] = result; +} + +static void test_series(int *retval, uint num_blocks, uint threads_per_block) { + CaptureStream captured(stdout); + + uint num_threads = num_blocks * threads_per_block; + for (uint i = 0; i != num_threads; ++i) { + retval[i] = 0x23232323; + } + + hipLaunchKernelGGL(kernel_series, dim3(num_blocks), dim3(threads_per_block), + 0, 0, retval); + hipStreamSynchronize(0); + auto CapturedData = captured.getCapturedData(); + + for (uint ii = 0; ii != num_threads; ++ii) { + HIPASSERT(retval[ii] == + strlen(msg_long1) + strlen(msg_short) + strlen(msg_long2) + 3); + } + + std::map linecount; + for (std::string line; std::getline(CapturedData, line);) { + linecount[line]++; + } + + HIPASSERT(linecount.size() == 3); + HIPASSERT(linecount[msg_long1] == num_threads); + HIPASSERT(linecount[msg_long2] == num_threads); + HIPASSERT(linecount[msg_short] == num_threads); +} + +int main() { + uint num_blocks = 1; + uint threads_per_block = 64; + uint num_threads = num_blocks * threads_per_block; + + void *retval_void; + HIPCHECK(hipHostMalloc(&retval_void, 4 * num_threads)); + auto retval = reinterpret_cast(retval_void); + + test_uniform0(retval, num_blocks, threads_per_block); + test_uniform1(retval, num_blocks, threads_per_block); + test_divergent0(retval, num_blocks, threads_per_block); + test_divergent1(retval, num_blocks, threads_per_block); + test_series(retval, num_blocks, threads_per_block); + + passed(); +} diff --git a/tests/src/printf/hipPrintfFlags.cpp b/tests/src/printf/hipPrintfFlags.cpp new file mode 100644 index 0000000000..5093c5e1c9 --- /dev/null +++ b/tests/src/printf/hipPrintfFlags.cpp @@ -0,0 +1,68 @@ +/* +Copyright (c) 2020 - present 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. +*/ + +/* HIT_START + * BUILD: %t %s EXCLUDE_HIP_PLATFORM NVCC EXCLUDE_HIP_RUNTIME HCC EXCLUDE_HIP_COMPILER hcc + * TEST: %t EXCLUDE_HIP_PLATFORM NVCC EXCLUDE_HIP_RUNTIME HCC EXCLUDE_HIP_COMPILER hcc + * HIT_END + */ + +#include "test_common.h" +#include "printf_common.h" + +__global__ void test_kernel() { + printf("%08d\n", 42); + printf("%08i\n", -42); + printf("%08u\n", 42); + printf("%08g\n", 123.456); + printf("%0+8d\n", 42); + printf("%+d\n", -42); + printf("%+08d\n", 42); + printf("%-8s\n", "xyzzy"); + printf("% i\n", -42); + printf("%-16.8d\n", 42); + printf("%16.8d\n", 42); +} + +int main(int argc, char **argv) { + std::string reference(R"here(00000042 +-0000042 +00000042 +0123.456 ++0000042 +-42 ++0000042 +xyzzy +-42 +00000042 + 00000042 +)here"); + + CaptureStream captured(stdout); + hipLaunchKernelGGL(test_kernel, dim3(1), dim3(1), 0, 0); + hipStreamSynchronize(0); + auto CapturedData = captured.getCapturedData(); + std::string device_output = gulp(CapturedData); + + HIPASSERT(device_output == reference); + passed(); +} diff --git a/tests/src/printf/hipPrintfManyDevices.cpp b/tests/src/printf/hipPrintfManyDevices.cpp new file mode 100644 index 0000000000..23751e26ea --- /dev/null +++ b/tests/src/printf/hipPrintfManyDevices.cpp @@ -0,0 +1,77 @@ +/* +Copyright (c) 2020 - present 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. +*/ + +/* HIT_START + * BUILD: %t %s EXCLUDE_HIP_PLATFORM NVCC EXCLUDE_HIP_RUNTIME HCC EXCLUDE_HIP_COMPILER hcc + * TEST: %t EXCLUDE_HIP_PLATFORM NVCC EXCLUDE_HIP_RUNTIME HCC EXCLUDE_HIP_COMPILER hcc + * HIT_END + */ + +#include "test_common.h" +#include "printf_common.h" + +DECLARE_DATA(); + +__global__ void print_things() { + DECLARE_DATA(); + + uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + const char *msg[] = {msg_short, msg_long1, msg_long2}; + + printf("%s\n", msg[tid % 3]); + if (tid % 3 == 0) + printf("%s\n", msg_short); + printf("%s\n", msg[(tid + 1) % 3]); + printf("%s\n", msg[(tid + 2) % 3]); +} + +int main() { + uint num_blocks = 14; + uint threads_per_block = 250; + uint threads_per_device = num_blocks * threads_per_block; + + int num_devices = 0; + hipGetDeviceCount(&num_devices); + + CaptureStream captured(stdout); + for (int i = 0; i != num_devices; ++i) { + hipSetDevice(i); + hipLaunchKernelGGL(print_things, dim3(num_blocks), dim3(threads_per_block), + 0, 0); + hipDeviceSynchronize(); + } + auto CapturedData = captured.getCapturedData(); + + std::map linecount; + for (std::string line; std::getline(CapturedData, line);) { + linecount[line]++; + } + + uint num_threads = threads_per_device * num_devices; + HIPASSERT(linecount.size() == 3); + HIPASSERT(linecount[msg_long1] == num_threads); + HIPASSERT(linecount[msg_long2] == num_threads); + HIPASSERT(linecount[msg_short] == + num_threads + ((threads_per_device + 2) / 3) * num_devices); + + passed(); +} diff --git a/tests/src/printf/hipPrintfManyWaves.cpp b/tests/src/printf/hipPrintfManyWaves.cpp new file mode 100644 index 0000000000..c15e695424 --- /dev/null +++ b/tests/src/printf/hipPrintfManyWaves.cpp @@ -0,0 +1,301 @@ +/* +Copyright (c) 2020 - present 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. +*/ + +/* HIT_START + * BUILD: %t %s EXCLUDE_HIP_PLATFORM NVCC EXCLUDE_HIP_RUNTIME HCC EXCLUDE_HIP_COMPILER hcc + * TEST: %t EXCLUDE_HIP_PLATFORM NVCC EXCLUDE_HIP_RUNTIME HCC EXCLUDE_HIP_COMPILER hcc + * HIT_END + */ + +#include "test_common.h" +#include "printf_common.h" +#include + +// Global string constants don't work inside device functions, so we +// use a macro to repeat the declaration in host and device contexts. +DECLARE_DATA(); + +__global__ void kernel_mixed0(int *retval) { + DECLARE_DATA(); + + uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + ulong result = 0; + + // Three strings passed as divergent values to the same hostcall. + const char *msg; + switch (tid % 3) { + case 0: + msg = msg_short; + break; + case 1: + msg = msg_long1; + break; + case 2: + msg = msg_long2; + break; + } + + retval[tid] = printf("%s\n", msg); +} + +static void test_mixed0(int *retval, uint num_blocks, uint threads_per_block) { + CaptureStream captured(stdout); + + uint num_threads = num_blocks * threads_per_block; + for (uint i = 0; i != num_threads; ++i) { + retval[i] = 0x23232323; + } + + hipLaunchKernelGGL(kernel_mixed0, dim3(num_blocks), dim3(threads_per_block), + 0, 0, retval); + hipStreamSynchronize(0); + auto CapturedData = captured.getCapturedData(); + + for (uint ii = 0; ii != num_threads; ++ii) { + switch (ii % 3) { + case 0: + HIPASSERT(retval[ii] == strlen(msg_short) + 1); + break; + case 1: + HIPASSERT(retval[ii] == strlen(msg_long1) + 1); + break; + case 2: + HIPASSERT(retval[ii] == strlen(msg_long2) + 1); + break; + } + } + + std::map linecount; + for (std::string line; std::getline(CapturedData, line);) { + linecount[line]++; + } + + HIPASSERT(linecount.size() == 3); + HIPASSERT(linecount[msg_short] == (num_threads + 2) / 3); + HIPASSERT(linecount[msg_long1] == (num_threads + 1) / 3); + HIPASSERT(linecount[msg_long2] == (num_threads + 0) / 3); +} + +__global__ void kernel_mixed1(int *retval) { + DECLARE_DATA(); + + const uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + + // Three strings passed to divergent hostcalls. + switch (tid % 3) { + case 0: + retval[tid] = printf("%s\n", msg_short); + break; + case 1: + retval[tid] = printf("%s\n", msg_long1); + break; + case 2: + retval[tid] = printf("%s\n", msg_long2); + break; + } +} + +static void test_mixed1(int *retval, uint num_blocks, uint threads_per_block) { + CaptureStream captured(stdout); + + uint num_threads = num_blocks * threads_per_block; + for (uint i = 0; i != num_threads; ++i) { + retval[i] = 0x23232323; + } + + hipLaunchKernelGGL(kernel_mixed1, dim3(num_blocks), dim3(threads_per_block), + 0, 0, retval); + hipStreamSynchronize(0); + auto CapturedData = captured.getCapturedData(); + + for (uint ii = 0; ii != num_threads; ++ii) { + switch (ii % 3) { + case 0: + HIPASSERT(retval[ii] == strlen(msg_short) + 1); + break; + case 1: + HIPASSERT(retval[ii] == strlen(msg_long1) + 1); + break; + case 2: + HIPASSERT(retval[ii] == strlen(msg_long2) + 1); + break; + } + } + + std::map linecount; + for (std::string line; std::getline(CapturedData, line);) { + linecount[line]++; + } + + HIPASSERT(linecount.size() == 3); + HIPASSERT(linecount[msg_short] == (num_threads + 2) / 3); + HIPASSERT(linecount[msg_long1] == (num_threads + 1) / 3); + HIPASSERT(linecount[msg_long2] == (num_threads + 0) / 3); +} + +__global__ void kernel_mixed2(int *retval) { + DECLARE_DATA(); + + const uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + + // Three different strings. All workitems print all three, but + // in different orders. + const char *msg[] = {msg_short, msg_long1, msg_long2}; + retval[tid] = + printf("%s%s%s\n", msg[tid % 3], msg[(tid + 1) % 3], msg[(tid + 2) % 3]); +} + +static void test_mixed2(int *retval, uint num_blocks, uint threads_per_block) { + CaptureStream captured(stdout); + + uint num_threads = num_blocks * threads_per_block; + for (uint i = 0; i != num_threads; ++i) { + retval[i] = 0x23232323; + } + + hipLaunchKernelGGL(kernel_mixed2, dim3(num_blocks), dim3(threads_per_block), + 0, 0, retval); + hipStreamSynchronize(0); + auto CapturedData = captured.getCapturedData(); + + for (uint ii = 0; ii != num_threads; ++ii) { + HIPASSERT(retval[ii] == + strlen(msg_short) + strlen(msg_long1) + strlen(msg_long2) + 1); + } + + std::map linecount; + for (std::string line; std::getline(CapturedData, line);) { + linecount[line]++; + } + + std::string str1 = + std::string(msg_short) + std::string(msg_long1) + std::string(msg_long2); + std::string str2 = + std::string(msg_long1) + std::string(msg_long2) + std::string(msg_short); + std::string str3 = + std::string(msg_long2) + std::string(msg_short) + std::string(msg_long1); + + HIPASSERT(linecount.size() == 3); + HIPASSERT(linecount[str1] == (num_threads + 2) / 3); + HIPASSERT(linecount[str2] == (num_threads + 1) / 3); + HIPASSERT(linecount[str3] == (num_threads + 0) / 3); +} + +__global__ void kernel_mixed3(int *retval) { + DECLARE_DATA(); + + const uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + int result = 0; + + result += printf("%s\n", msg_long1); + if (tid % 3 == 0) { + result += printf("%s\n", msg_short); + } + result += printf("%s\n", msg_long2); + + retval[tid] = result; +} + +static void test_mixed3(int *retval, uint num_blocks, uint threads_per_block) { + CaptureStream captured(stdout); + + uint num_threads = num_blocks * threads_per_block; + for (uint i = 0; i != num_threads; ++i) { + retval[i] = 0x23232323; + } + + hipLaunchKernelGGL(kernel_mixed3, dim3(num_blocks), dim3(threads_per_block), + 0, 0, retval); + hipStreamSynchronize(0); + auto CapturedData = captured.getCapturedData(); + + for (uint ii = 0; ii != num_threads; ++ii) { + if (ii % 3 == 0) { + HIPASSERT(retval[ii] == + strlen(msg_long1) + strlen(msg_short) + strlen(msg_long2) + 3); + } else { + HIPASSERT(retval[ii] == strlen(msg_long1) + strlen(msg_long2) + 2); + } + } + + std::map linecount; + for (std::string line; std::getline(CapturedData, line);) { + linecount[line]++; + } + + HIPASSERT(linecount.size() == 3); + HIPASSERT(linecount[msg_long1] == num_threads); + HIPASSERT(linecount[msg_long2] == num_threads); + HIPASSERT(linecount[msg_short] == (num_threads + 2) / 3); +} + +__global__ void kernel_numbers() { + uint tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + for (uint i = 0; i != 7; ++i) { + uint base = tid * 21 + i * 3; + printf("%d %d %d\n", base, base + 1, base + 2); + } +} + +static void test_numbers(uint num_blocks, uint threads_per_block) { + CaptureStream captured(stdout); + uint num_threads = num_blocks * threads_per_block; + + hipLaunchKernelGGL(kernel_numbers, dim3(num_blocks), dim3(threads_per_block), + 0, 0); + hipStreamSynchronize(0); + auto CapturedData = captured.getCapturedData(); + + std::vector points; + while (true) { + uint i; + CapturedData >> i; + if (CapturedData.fail()) + break; + points.push_back(i); + } + + std::sort(points.begin(), points.end()); + points.erase(std::unique(points.begin(), points.end()), points.end()); + HIPASSERT(points.size() == 21 * num_threads); + HIPASSERT(points.back() == 21 * num_threads - 1); + + passed(); +} + +int main(int argc, char **argv) { + uint num_blocks = 150; + uint threads_per_block = 250; + uint num_threads = num_blocks * threads_per_block; + + void *retval_void; + HIPCHECK(hipHostMalloc(&retval_void, 4 * num_threads)); + auto retval = reinterpret_cast(retval_void); + + test_mixed0(retval, num_blocks, threads_per_block); + test_mixed1(retval, num_blocks, threads_per_block); + test_mixed2(retval, num_blocks, threads_per_block); + test_mixed3(retval, num_blocks, threads_per_block); + test_numbers(num_blocks, threads_per_block); + + passed(); +} diff --git a/tests/src/printf/hipPrintfSpecifiers.cpp b/tests/src/printf/hipPrintfSpecifiers.cpp new file mode 100644 index 0000000000..03507658b9 --- /dev/null +++ b/tests/src/printf/hipPrintfSpecifiers.cpp @@ -0,0 +1,90 @@ +/* +Copyright (c) 2020 - present 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. +*/ + +/* HIT_START + * BUILD: %t %s EXCLUDE_HIP_PLATFORM NVCC EXCLUDE_HIP_RUNTIME HCC EXCLUDE_HIP_COMPILER hcc + * TEST: %t EXCLUDE_HIP_PLATFORM NVCC EXCLUDE_HIP_RUNTIME HCC EXCLUDE_HIP_COMPILER hcc + * HIT_END + */ + +#include "test_common.h" +#include "printf_common.h" + +__global__ void test_kernel() { + const char *N = nullptr; + const char *s = "hello world"; + + printf("xyzzy\n"); + printf("%%\n"); + printf("hello %% world\n"); + printf("%%s\n"); + // Two special tests to make sure that the compiler pass correctly + // skips over a '%%' without affecting the logic for locating + // string arguments. + printf("%%s%p\n", (void *)0xf01dab1eca55e77e); + printf("%%c%s\n", "xyzzy"); + printf("%c%c%c\n", 's', 'e', 'p'); + printf("%d\n", -42); + printf("%u\n", 42); + printf("%f\n", 123.456); + printf("%F\n", -123.456); + printf("%e\n", -123.456); + printf("%E\n", 123.456); + printf("%g\n", 123.456); + printf("%G\n", -123.456); + printf("%c\n", 'x'); + printf("%s\n", N); + printf("%p\n", N); + printf("%.*f %*.*s %p\n", 8, 3.14159, 8, 5, s, (void *)0xf01dab1eca55e77e); +} + +int main(int argc, char **argv) { + std::string reference(R"here(xyzzy +% +hello % world +%s +%s0xf01dab1eca55e77e +%cxyzzy +sep +-42 +42 +123.456000 +-123.456000 +-1.234560e+02 +1.234560E+02 +123.456 +-123.456 +x + +(nil) +3.14159000 hello 0xf01dab1eca55e77e +)here"); + + CaptureStream captured(stdout); + hipLaunchKernelGGL(test_kernel, dim3(1), dim3(1), 0, 0); + hipStreamSynchronize(0); + auto CapturedData = captured.getCapturedData(); + std::string device_output = gulp(CapturedData); + + HIPASSERT(device_output == reference); + passed(); +} diff --git a/tests/src/printf/hipPrintfStar.cpp b/tests/src/printf/hipPrintfStar.cpp new file mode 100644 index 0000000000..5e97d6eae5 --- /dev/null +++ b/tests/src/printf/hipPrintfStar.cpp @@ -0,0 +1,54 @@ +/* +Copyright (c) 2020 - present 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. +*/ + +/* HIT_START + * BUILD: %t %s EXCLUDE_HIP_PLATFORM NVCC EXCLUDE_HIP_RUNTIME HCC EXCLUDE_HIP_COMPILER hcc + * TEST: %t EXCLUDE_HIP_PLATFORM NVCC EXCLUDE_HIP_RUNTIME HCC EXCLUDE_HIP_COMPILER hcc + * HIT_END + */ + +#include "test_common.h" +#include "printf_common.h" + +__global__ void test_kernel() { + printf("%*d\n", 16, 42); + printf("%.*d\n", 8, 42); + printf("%*.*d\n", -16, 8, 42); + printf("%*.*f %s * %.*s\n", 16, 8, 123.456, "hello", 5, "worldxyz"); +} + +int main(int argc, char **argv) { + std::string reference(R"here( 42 +00000042 +00000042 + 123.45600000 hello * world +)here"); + + CaptureStream captured(stdout); + hipLaunchKernelGGL(test_kernel, dim3(1), dim3(1), 0, 0); + hipStreamSynchronize(0); + auto CapturedData = captured.getCapturedData(); + std::string device_output = gulp(CapturedData); + + HIPASSERT(device_output == reference); + passed(); +} diff --git a/tests/src/printf/hipPrintfWidthPrecision.cpp b/tests/src/printf/hipPrintfWidthPrecision.cpp new file mode 100644 index 0000000000..db003377a9 --- /dev/null +++ b/tests/src/printf/hipPrintfWidthPrecision.cpp @@ -0,0 +1,74 @@ +/* +Copyright (c) 2020 - present 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. +*/ + +/* HIT_START + * BUILD: %t %s EXCLUDE_HIP_PLATFORM NVCC EXCLUDE_HIP_RUNTIME HCC EXCLUDE_HIP_COMPILER hcc + * TEST: %t EXCLUDE_HIP_PLATFORM NVCC EXCLUDE_HIP_RUNTIME HCC EXCLUDE_HIP_COMPILER hcc + * HIT_END + */ + +#include "test_common.h" +#include "printf_common.h" + +__global__ void test_kernel() { + printf("%16d\n", 42); + printf("%.8d\n", 42); + printf("%16.5d\n", -42); + printf("%.8x\n", 0x42); + printf("%.8o\n", 042); + printf("%16.8e\n", 12345.67891); + printf("%16.8f\n", -12345.67891); + printf("%16.8g\n", 12345.67891); + printf("%8.4e\n", -12345.67891); + printf("%8.4f\n", 12345.67891); + printf("%8.4g\n", 12345.67891); + printf("%4.2f\n", 12345.67891); + printf("%.1f\n", 12345.67891); + printf("%.5s\n", "helloxyz"); +} + +int main(int argc, char **argv) { + std::string reference(R"here( 42 +00000042 + -00042 +00000042 +00000042 + 1.23456789e+04 + -12345.67891000 + 12345.679 +-1.2346e+04 +12345.6789 +1.235e+04 +12345.68 +12345.7 +hello +)here"); + + CaptureStream captured(stdout); + hipLaunchKernelGGL(test_kernel, dim3(1), dim3(1), 0, 0); + hipStreamSynchronize(0); + auto CapturedData = captured.getCapturedData(); + std::string device_output = gulp(CapturedData); + + HIPASSERT(device_output == reference); + passed(); +} diff --git a/tests/src/printf/printf_common.h b/tests/src/printf/printf_common.h new file mode 100644 index 0000000000..a2df88db9f --- /dev/null +++ b/tests/src/printf/printf_common.h @@ -0,0 +1,94 @@ +#ifndef COMMON_H +#define COMMON_H + +#include +#include +#include +#include +#include +#include +#include +#include + +struct CaptureStream { + int saved_fd; + int orig_fd; + int temp_fd; + + char tempname[13] = "mytestXXXXXX"; + + CaptureStream(FILE *original) { + orig_fd = fileno(original); + saved_fd = dup(orig_fd); + + temp_fd = mkstemp(tempname); + if (errno) { + error(0, errno, "Error"); + assert(false); + } + + fflush(nullptr); + dup2(temp_fd, orig_fd); + if (errno) { + error(0, errno, "Error"); + assert(false); + } + close(temp_fd); + if (errno) { + error(0, errno, "Error"); + assert(false); + } + } + + void restoreStream() { + if (saved_fd == -1) + return; + fflush(nullptr); + dup2(saved_fd, orig_fd); + if (errno) { + error(0, errno, "Error"); + assert(false); + } + close(saved_fd); + if (errno) { + error(0, errno, "Error"); + assert(false); + } + saved_fd = -1; + } + + std::ifstream getCapturedData() { + restoreStream(); + std::ifstream temp(tempname); + return temp; + } + + ~CaptureStream() { + restoreStream(); + remove(tempname); + if (errno) { + error(0, errno, "Error"); + assert(false); + } + } +}; + +static std::string gulp(std::ifstream &input) { + std::string retval; + input.seekg(0, std::ios_base::end); + retval.resize(input.tellg()); + input.seekg(0, std::ios_base::beg); + input.read(&retval[0], retval.size()); + input.close(); + return retval; +} + +#define DECLARE_DATA() \ + const char *msg_short = "Carpe diem."; \ + const char *msg_long1 = "Lorem ipsum dolor sit amet, consectetur nullam. " \ + "In mollis imperdiet nibh nec ullamcorper."; \ + const char *msg_long2 = "Curabitur nec metus sit amet augue vehicula " \ + "ultrices ut id leo. Lorem ipsum dolor sit amet, " \ + "consectetur adipiscing elit amet."; + +#endif diff --git a/tests/src/runtimeApi/module/hipLaunchCoopMultiKernel.cpp b/tests/src/runtimeApi/module/hipLaunchCoopMultiKernel.cpp index b4d57a7693..31b9d94ac2 100644 --- a/tests/src/runtimeApi/module/hipLaunchCoopMultiKernel.cpp +++ b/tests/src/runtimeApi/module/hipLaunchCoopMultiKernel.cpp @@ -20,7 +20,7 @@ THE SOFTWARE. // Simple test for hipLaunchCooperativeKernelMultiDevice API. /* HIT_START - * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM all + * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc * TEST: %t * HIT_END */ @@ -179,8 +179,6 @@ int main() { hipLaunchCooperativeKernelMultiDevice(launchParamsList, nGpu, 0); - HIPCHECK(hipMemcpy(init, dC, sizeof(long), hipMemcpyDeviceToHost)); - if (*dC != (((long)(BufferSizeInDwords) * (BufferSizeInDwords - 1)) / 2)) { std::cout << "Data validation failed for grid size = " << dimGrid.x << " and block size = " << dimBlock.x << "\n"; std::cout << "Test failed! \n"; diff --git a/tests/src/runtimeApi/module/hipLaunchCooperativeKernel.cpp b/tests/src/runtimeApi/module/hipLaunchCooperativeKernel.cpp index c76685fa89..6862c61b1e 100644 --- a/tests/src/runtimeApi/module/hipLaunchCooperativeKernel.cpp +++ b/tests/src/runtimeApi/module/hipLaunchCooperativeKernel.cpp @@ -22,7 +22,7 @@ THE SOFTWARE. // Simple test for hipLaunchCooperativeKernel API. /* HIT_START - * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM all + * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc * TEST: %t * HIT_END */ diff --git a/vdi/hip_conversions.hpp b/vdi/hip_conversions.hpp index 0515c5c56b..d79a39c321 100644 --- a/vdi/hip_conversions.hpp +++ b/vdi/hip_conversions.hpp @@ -169,19 +169,19 @@ cl_mem_object_type getCLMemObjectType(const hipResourceType hipResType) { } inline -size_t getElementSize(const hipArray_Format arrayFormat) { - switch (arrayFormat) { +size_t getElementSize(const hipArray_const_t array) { + switch (array->Format) { case HIP_AD_FORMAT_UNSIGNED_INT8: case HIP_AD_FORMAT_SIGNED_INT8: - return 1; + return 1 * array->NumChannels; case HIP_AD_FORMAT_UNSIGNED_INT16: case HIP_AD_FORMAT_SIGNED_INT16: case HIP_AD_FORMAT_HALF: - return 2; + return 2 * array->NumChannels; case HIP_AD_FORMAT_UNSIGNED_INT32: case HIP_AD_FORMAT_SIGNED_INT32: case HIP_AD_FORMAT_FLOAT: - return 4; + return 4 * array->NumChannels; } ShouldNotReachHere(); @@ -617,7 +617,7 @@ HIP_MEMCPY3D getDrvMemcpy3DDesc(const hipMemcpy3DParms& desc) { descDrv.srcMemoryType = hipMemoryTypeArray; descDrv.srcArray = desc.srcArray; // When reffering to array memory, hipPos::x is in elements. - descDrv.srcXInBytes *= getElementSize(desc.srcArray->Format); + descDrv.srcXInBytes *= getElementSize(desc.srcArray); } if (desc.srcPtr.ptr != nullptr) { @@ -632,7 +632,7 @@ HIP_MEMCPY3D getDrvMemcpy3DDesc(const hipMemcpy3DParms& desc) { descDrv.dstMemoryType = hipMemoryTypeArray; descDrv.dstArray = desc.dstArray; // When reffering to array memory, hipPos::x is in elements. - descDrv.dstXInBytes *= getElementSize(desc.dstArray->Format); + descDrv.dstXInBytes *= getElementSize(desc.dstArray); } if (desc.dstPtr.ptr != nullptr) { @@ -645,11 +645,11 @@ HIP_MEMCPY3D getDrvMemcpy3DDesc(const hipMemcpy3DParms& desc) { // If a HIP array is participating in the copy, the extent is defined in terms of that array's elements. if ((desc.srcArray != nullptr) && (desc.dstArray == nullptr)) { - descDrv.WidthInBytes *= getElementSize(desc.srcArray->Format); + descDrv.WidthInBytes *= getElementSize(desc.srcArray); } else if ((desc.srcArray == nullptr) && (desc.dstArray != nullptr)) { - descDrv.WidthInBytes *= getElementSize(desc.dstArray->Format); + descDrv.WidthInBytes *= getElementSize(desc.dstArray); } else if ((desc.srcArray != nullptr) && (desc.dstArray != nullptr)) { - descDrv.WidthInBytes *= getElementSize(desc.dstArray->Format); + descDrv.WidthInBytes *= getElementSize(desc.dstArray); } return descDrv; diff --git a/vdi/hip_device.cpp b/vdi/hip_device.cpp index d1df525877..30b2292271 100644 --- a/vdi/hip_device.cpp +++ b/vdi/hip_device.cpp @@ -197,6 +197,11 @@ hipError_t hipGetDeviceProperties ( hipDeviceProp_t* props, hipDevice_t device ) deviceProps.cooperativeLaunch = info.cooperativeGroups_; deviceProps.cooperativeMultiDeviceLaunch = info.cooperativeMultiDeviceGroups_; + deviceProps.cooperativeMultiDeviceUnmatchedFunc = info.cooperativeMultiDeviceGroups_; + deviceProps.cooperativeMultiDeviceUnmatchedGridDim = info.cooperativeMultiDeviceGroups_; + deviceProps.cooperativeMultiDeviceUnmatchedBlockDim = info.cooperativeMultiDeviceGroups_; + deviceProps.cooperativeMultiDeviceUnmatchedSharedMem = info.cooperativeMultiDeviceGroups_; + deviceProps.maxTexture1D = info.imageMaxBufferSize_; deviceProps.maxTexture2D[0] = info.image2DMaxWidth_; deviceProps.maxTexture2D[1] = info.image2DMaxHeight_; diff --git a/vdi/hip_device_runtime.cpp b/vdi/hip_device_runtime.cpp index acf73fa656..febf64d116 100644 --- a/vdi/hip_device_runtime.cpp +++ b/vdi/hip_device_runtime.cpp @@ -281,6 +281,18 @@ hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device) case hipDeviceAttributeEccEnabled: *pi = prop.ECCEnabled; break; + case hipDeviceAttributeCooperativeMultiDeviceUnmatchedFunc: + *pi = prop.cooperativeMultiDeviceUnmatchedFunc; + break; + case hipDeviceAttributeCooperativeMultiDeviceUnmatchedGridDim: + *pi = prop.cooperativeMultiDeviceUnmatchedGridDim; + break; + case hipDeviceAttributeCooperativeMultiDeviceUnmatchedBlockDim: + *pi = prop.cooperativeMultiDeviceUnmatchedBlockDim; + break; + case hipDeviceAttributeCooperativeMultiDeviceUnmatchedSharedMem: + *pi = prop.cooperativeMultiDeviceUnmatchedSharedMem; + break; default: HIP_RETURN(hipErrorInvalidValue); } diff --git a/vdi/hip_hcc.def.in b/vdi/hip_hcc.def.in index d444e34147..ef511ee43a 100644 --- a/vdi/hip_hcc.def.in +++ b/vdi/hip_hcc.def.in @@ -238,6 +238,7 @@ hipFreeMipmappedArray hipMipmappedArrayGetLevel hipGetMipmappedArrayLevel hipMallocHost +hipFreeHost hipTexObjectCreate hipTexObjectDestroy hipTexObjectGetResourceDesc diff --git a/vdi/hip_hcc.map.in b/vdi/hip_hcc.map.in index 46c59f8123..2139f45ab8 100644 --- a/vdi/hip_hcc.map.in +++ b/vdi/hip_hcc.map.in @@ -232,6 +232,7 @@ global: hipMipmappedArrayGetLevel; hipGetMipmappedArrayLevel; hipMallocHost; + hipFreeHost; hipTexObjectCreate; hipTexObjectDestroy; hipTexObjectGetResourceDesc; diff --git a/vdi/hip_internal.hpp b/vdi/hip_internal.hpp index 92397d1255..a2851fd978 100755 --- a/vdi/hip_internal.hpp +++ b/vdi/hip_internal.hpp @@ -31,7 +31,11 @@ #include #include #include - +#ifdef _WIN32 +#include +#else +#include +#endif /*! IHIP IPC MEMORY Structure */ #define IHIP_IPC_MEM_HANDLE_SIZE 32 @@ -43,6 +47,10 @@ typedef struct ihipIpcMemHandle_st { char reserved[IHIP_IPC_MEM_RESERVED_SIZE]; } ihipIpcMemHandle_t; +#ifdef _WIN32 + int getpid() { return _getpid();} +#endif + #define HIP_INIT() \ std::call_once(hip::g_ihipInitialized, hip::init); \ if (hip::g_device == nullptr && g_devices.size() > 0) { \ @@ -51,7 +59,7 @@ typedef struct ihipIpcMemHandle_st { // This macro should be called at the beginning of every HIP API. #define HIP_INIT_API(cid, ...) \ - ClPrint(amd::LOG_INFO, amd::LOG_API, "[%zx] %s ( %s )", std::this_thread::get_id(), __func__, ToString( __VA_ARGS__ ).c_str()); \ + ClPrint(amd::LOG_INFO, amd::LOG_API, "%-5d: [%zx] %s ( %s )", getpid(), std::this_thread::get_id(), __func__, ToString( __VA_ARGS__ ).c_str()); \ amd::Thread* thread = amd::Thread::current(); \ if (!VDI_CHECK_THREAD(thread)) { \ HIP_RETURN(hipErrorOutOfMemory); \ @@ -61,7 +69,7 @@ typedef struct ihipIpcMemHandle_st { #define HIP_RETURN(ret) \ hip::g_lastError = ret; \ - ClPrint(amd::LOG_INFO, amd::LOG_API, "[%zx] %s: Returned %s", std::this_thread::get_id(), __func__, hipGetErrorName(hip::g_lastError)); \ + ClPrint(amd::LOG_INFO, amd::LOG_API, "%-5d: [%zx] %s: Returned %s", getpid(), std::this_thread::get_id(), __func__, hipGetErrorName(hip::g_lastError)); \ return hip::g_lastError; namespace hc { @@ -274,6 +282,7 @@ public: void configureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem, hipStream_t stream); void popExec(ihipExec_t& exec); + }; extern std::vector g_devices; diff --git a/vdi/hip_memory.cpp b/vdi/hip_memory.cpp index 5438064e9b..1625f7703a 100644 --- a/vdi/hip_memory.cpp +++ b/vdi/hip_memory.cpp @@ -41,6 +41,25 @@ amd::Memory* getMemoryObject(const void* ptr, size_t& offset) { return memObj; } +hipError_t ihipFree(void *ptr) +{ + if (ptr == nullptr) { + return hipSuccess; + } + if (amd::SvmBuffer::malloced(ptr)) { + for (auto& dev : g_devices) { + amd::HostQueue* queue = hip::getNullStream(*dev->asContext()); + if (queue != nullptr) { + queue->finish(); + } + hip::syncStreams(dev->deviceId()); + } + amd::SvmBuffer::free(*hip::getCurrentDevice()->asContext(), ptr); + return hipSuccess; + } + return hipErrorInvalidValue; +} + hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { if (sizeBytes == 0) { @@ -66,7 +85,7 @@ hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags) if (*ptr == nullptr) { return hipErrorOutOfMemory; } - ClPrint(amd::LOG_INFO, amd::LOG_API, "ihipMalloc ptr=0x%zx", *ptr); + ClPrint(amd::LOG_INFO, amd::LOG_API, "%-5d: [%zx] ihipMalloc ptr=0x%zx", getpid(),std::this_thread::get_id(), *ptr); return hipSuccess; } @@ -190,7 +209,8 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { } unsigned int ihipFlags = CL_MEM_SVM_FINE_GRAIN_BUFFER | (flags << 16); - if (flags & (hipHostMallocCoherent | hipHostMallocMapped) || + if (flags == 0 || + flags & (hipHostMallocCoherent | hipHostMallocMapped) || (!(flags & hipHostMallocNonCoherent) && HIP_HOST_COHERENT)) { ihipFlags |= CL_MEM_SVM_ATOMICS; } @@ -212,21 +232,7 @@ hipError_t hipMallocManaged(void** devPtr, size_t size, hipError_t hipFree(void* ptr) { HIP_INIT_API(hipFree, ptr); - if (ptr == nullptr) { - HIP_RETURN(hipSuccess); - } - if (amd::SvmBuffer::malloced(ptr)) { - for (auto& dev : g_devices) { - amd::HostQueue* queue = hip::getNullStream(*dev->asContext()); - if (queue != nullptr) { - queue->finish(); - } - hip::syncStreams(dev->deviceId()); - } - amd::SvmBuffer::free(*hip::getCurrentDevice()->asContext(), ptr); - HIP_RETURN(hipSuccess); - } - HIP_RETURN(hipErrorInvalidValue); + HIP_RETURN(ihipFree(ptr)); } hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind) { @@ -264,18 +270,7 @@ hipError_t hipMemPtrGetInfo(void *ptr, size_t *size) { hipError_t hipHostFree(void* ptr) { HIP_INIT_API(hipHostFree, ptr); - if (amd::SvmBuffer::malloced(ptr)) { - for (auto& dev : g_devices) { - amd::HostQueue* queue = hip::getNullStream(*dev->asContext()); - if (queue != nullptr) { - queue->finish(); - } - hip::syncStreams(dev->deviceId()); - } - amd::SvmBuffer::free(*hip::getCurrentDevice()->asContext(), ptr); - HIP_RETURN(hipSuccess); - } - HIP_RETURN(hipErrorInvalidValue); + HIP_RETURN(ihipFree(ptr)); } hipError_t ihipArrayDestroy(hipArray* array) { @@ -889,19 +884,6 @@ hipError_t hipMemcpyDtoHAsync(void* dstHost, HIP_RETURN(ihipMemcpy(dstHost, srcDevice, ByteCount, hipMemcpyDeviceToHost, *hip::getQueue(stream), true)); } -inline void adjustOrigin(amd::Coord3D &origin, - size_t offset, - size_t rowPitch, - size_t slicePitch) { - size_t zOffset = offset / (slicePitch ? slicePitch : 1); - size_t yOffset = (offset - slicePitch * zOffset) / (rowPitch ? rowPitch : 1); - size_t xOffset = (offset - slicePitch * zOffset - rowPitch * yOffset); - - static_cast(origin)[0] += xOffset; - static_cast(origin)[1] += yOffset; - static_cast(origin)[2] += zOffset; -} - hipError_t ihipMemcpyAtoD(hipArray* srcArray, void* dstDevice, amd::Coord3D srcOrigin, @@ -919,7 +901,6 @@ hipError_t ihipMemcpyAtoD(hipArray* srcArray, amd::Image* srcImage = as_amd(srcMemObj)->asImage(); size_t dstOffset = 0; amd::Memory* dstMemory = getMemoryObject(dstDevice, dstOffset); - adjustOrigin(dstOrigin, dstOffset, dstRowPitch, dstSlicePitch); amd::BufferRect srcRect; if (!srcRect.create(static_cast(srcOrigin), static_cast(copyRegion), srcImage->getRowPitch(), srcImage->getSlicePitch())) { @@ -930,6 +911,8 @@ hipError_t ihipMemcpyAtoD(hipArray* srcArray, if (!dstRect.create(static_cast(dstOrigin), static_cast(copyRegion), dstRowPitch, dstSlicePitch)) { return hipErrorInvalidValue; } + dstRect.start_ += dstOffset; + dstRect.end_ += dstOffset; const size_t copySizeInBytes = copyRegion[0] * copyRegion[1] * copyRegion[2] * srcImage->getImageFormat().getElementSize(); if (!srcImage->validateRegion(srcOrigin, copyRegion) || @@ -977,13 +960,14 @@ hipError_t ihipMemcpyDtoA(void* srcDevice, size_t srcOffset = 0; amd::Memory* srcMemory = getMemoryObject(srcDevice, srcOffset); - adjustOrigin(srcOrigin, srcOffset, srcRowPitch, srcSlicePitch); amd::Image* dstImage = as_amd(dstMemObj)->asImage(); amd::BufferRect srcRect; if (!srcRect.create(static_cast(srcOrigin), static_cast(copyRegion), srcRowPitch, srcSlicePitch)) { return hipErrorInvalidValue; } + srcRect.start_ += srcOffset; + srcRect.end_ += srcOffset; amd::BufferRect dstRect; if (!dstRect.create(static_cast(dstOrigin), static_cast(copyRegion), dstImage->getRowPitch(), dstImage->getSlicePitch())) { @@ -1033,15 +1017,15 @@ hipError_t ihipMemcpyDtoD(void* srcDevice, bool isAsync = false) { size_t srcOffset = 0; amd::Memory *srcMemory = getMemoryObject(srcDevice, srcOffset); - adjustOrigin(srcOrigin, srcOffset, srcRowPitch, srcSlicePitch); size_t dstOffset = 0; amd::Memory *dstMemory = getMemoryObject(dstDevice, dstOffset); - adjustOrigin(dstOrigin, dstOffset, dstRowPitch, dstSlicePitch); amd::BufferRect srcRect; if (!srcRect.create(static_cast(srcOrigin), static_cast(copyRegion), srcRowPitch, srcSlicePitch)) { return hipErrorInvalidValue; } + srcRect.start_ += srcOffset; + srcRect.end_ += srcOffset; amd::Coord3D srcStart(srcRect.start_, 0, 0); amd::Coord3D srcEnd(srcRect.end_, 1, 1); @@ -1053,6 +1037,8 @@ hipError_t ihipMemcpyDtoD(void* srcDevice, if (!dstRect.create(static_cast(dstOrigin), static_cast(copyRegion), dstRowPitch, dstSlicePitch)) { return hipErrorInvalidValue; } + dstRect.start_ += dstOffset; + dstRect.end_ += dstOffset; amd::Coord3D dstStart(dstRect.start_, 0, 0); amd::Coord3D dstEnd(dstRect.end_, 1, 1); @@ -1097,12 +1083,13 @@ hipError_t ihipMemcpyDtoH(void* srcDevice, bool isAsync = false) { size_t srcOffset = 0; amd::Memory *srcMemory = getMemoryObject(srcDevice, srcOffset); - adjustOrigin(srcOrigin, srcOffset, srcRowPitch, srcSlicePitch); amd::BufferRect srcRect; if (!srcRect.create(static_cast(srcOrigin), static_cast(copyRegion), srcRowPitch, srcSlicePitch)) { return hipErrorInvalidValue; } + srcRect.start_ += srcOffset; + srcRect.end_ += srcOffset; amd::Coord3D srcStart(srcRect.start_, 0, 0); amd::Coord3D srcEnd(srcRect.end_, 1, 1); @@ -1151,7 +1138,6 @@ hipError_t ihipMemcpyHtoD(const void* srcHost, bool isAsync = false) { size_t dstOffset = 0; amd::Memory *dstMemory = getMemoryObject(dstDevice, dstOffset); - adjustOrigin(dstOrigin, dstOffset, dstRowPitch, dstSlicePitch); amd::BufferRect srcRect; if (!srcRect.create(static_cast(srcOrigin), static_cast(copyRegion), srcRowPitch, srcSlicePitch)) { @@ -1162,6 +1148,8 @@ hipError_t ihipMemcpyHtoD(const void* srcHost, if (!dstRect.create(static_cast(dstOrigin), static_cast(copyRegion), dstRowPitch, dstSlicePitch)) { return hipErrorInvalidValue; } + dstRect.start_ += dstOffset; + dstRect.end_ += dstOffset; amd::Coord3D dstStart(dstRect.start_, 0, 0); amd::Coord3D dstEnd(dstRect.end_, 1, 1); @@ -1534,7 +1522,7 @@ hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, const const size_t arrayHeight = (dst->height != 0) ? dst->height : 1; const size_t witdthInBytes = count / arrayHeight; - const size_t height = (count / dst->width) / (hip::getElementSize(dst->Format) * dst->NumChannels); + const size_t height = (count / dst->width) / hip::getElementSize(dst); HIP_RETURN(ihipMemcpy2DToArray(dst, wOffset, hOffset, src, 0 /* spitch */, witdthInBytes, height, kind, nullptr)); } @@ -1574,7 +1562,7 @@ hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t src, size_t wOffsetSrc const size_t arrayHeight = (src->height != 0) ? src->height : 1; const size_t witdthInBytes = count / arrayHeight; - const size_t height = (count / src->width) / (hip::getElementSize(src->Format) * src->NumChannels); + const size_t height = (count / src->width) / hip::getElementSize(src); HIP_RETURN(ihipMemcpy2DFromArray(dst, 0 /* dpitch */, src, wOffsetSrc, hOffset, witdthInBytes, height, kind, nullptr)); } @@ -1609,7 +1597,7 @@ hipError_t ihipMemcpy3D(const hipMemcpy3DParms* p, // If the source and destination are both arrays, hipMemcpy3D() will return an error if they do not have the same element size. if (((p->srcArray != nullptr) && (p->dstArray != nullptr)) && - (hip::getElementSize(p->dstArray->Format) != hip::getElementSize(p->dstArray->Format))) { + (hip::getElementSize(p->dstArray) != hip::getElementSize(p->dstArray))) { return hipErrorInvalidValue; } @@ -2062,7 +2050,7 @@ hipError_t hipMemcpyFromArrayAsync(void* dst, hipArray_const_t src, size_t wOffs const size_t arrayHeight = (src->height != 0) ? src->height : 1; const size_t widthInBytes = count / arrayHeight; - const size_t height = (count / src->width) / (hip::getElementSize(src->Format) * src->NumChannels); + const size_t height = (count / src->width) / hip::getElementSize(src); HIP_RETURN(ihipMemcpy2DFromArray(dst, 0 /* dpitch */, src, wOffsetSrc, hOffsetSrc, widthInBytes, height, kind, stream, true)); } @@ -2083,7 +2071,7 @@ hipError_t hipMemcpyToArrayAsync(hipArray_t dst, size_t wOffset, size_t hOffset, const size_t arrayHeight = (dst->height != 0) ? dst->height : 1; const size_t widthInBytes = count / arrayHeight; - const size_t height = (count / dst->width) / (hip::getElementSize(dst->Format) * dst->NumChannels); + const size_t height = (count / dst->width) / hip::getElementSize(dst); HIP_RETURN(ihipMemcpy2DToArray(dst, wOffset, hOffset, src, 0 /* spitch */, widthInBytes, height, kind, stream, true)); } @@ -2192,3 +2180,9 @@ hipError_t hipMallocHost(void** ptr, HIP_RETURN(ihipMalloc(ptr, size, CL_MEM_SVM_FINE_GRAIN_BUFFER)); } + +hipError_t hipFreeHost(void *ptr) { + HIP_INIT_API(hipFreeHost, ptr); + + HIP_RETURN(ihipFree(ptr)); +} diff --git a/vdi/hip_module.cpp b/vdi/hip_module.cpp index a7822b315e..1e5f7d8b31 100755 --- a/vdi/hip_module.cpp +++ b/vdi/hip_module.cpp @@ -25,6 +25,7 @@ #include "hip_internal.hpp" #include "platform/program.hpp" #include "hip_event.hpp" +#include "hip_platform.hpp" hipError_t ihipModuleLoadData(hipModule_t *module, const void *image); @@ -272,7 +273,6 @@ hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func) HIP_RETURN(hipSuccess); } - hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t gridDimX, uint32_t gridDimY, uint32_t gridDimZ, uint32_t blockDimX, uint32_t blockDimY, uint32_t blockDimZ, @@ -294,13 +294,28 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, amd::HostQueue* queue = hip::getQueue(hStream); const amd::Device& device = queue->vdev()->device(); - if ((params & amd::NDRangeKernelCommand::CooperativeGroups) && - !device.info().cooperativeGroups_) { - return hipErrorLaunchFailure; + // Make sure dispatch doesn't exceed max workgroup size limit + if (blockDimX * blockDimY * blockDimZ > device.info().maxWorkGroupSize_) { + return hipErrorInvalidConfiguration; } - if ((params & amd::NDRangeKernelCommand::CooperativeMultiDeviceGroups) && - !device.info().cooperativeMultiDeviceGroups_) { - return hipErrorLaunchFailure; + + if (params & amd::NDRangeKernelCommand::CooperativeGroups) { + if (!device.info().cooperativeGroups_) { + return hipErrorLaunchFailure; + } + int num_blocks = 0; + int num_grids = 0; + int block_size = blockDimX * blockDimY * blockDimZ; + hip_impl::ihipOccupancyMaxActiveBlocksPerMultiprocessor( + &num_blocks, &num_grids, device, f, block_size, sharedMemBytes, true); + if (((gridDimX * gridDimY * gridDimZ) / block_size) > unsigned(num_grids)) { + return hipErrorCooperativeLaunchTooLarge; + } + } + if (params & amd::NDRangeKernelCommand::CooperativeMultiDeviceGroups) { + if (!device.info().cooperativeMultiDeviceGroups_) { + return hipErrorLaunchFailure; + } } if (!queue) { return hipErrorOutOfMemory; @@ -466,12 +481,42 @@ hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsL hipError_t result = hipErrorUnknown; uint64_t allGridSize = 0; + std::vector mgpu_list(numDevices); + for (int i = 0; i < numDevices; ++i) { const hipLaunchParams& launch = launchParamsList[i]; allGridSize += launch.gridDim.x * launch.gridDim.y * launch.gridDim.z; + + // Make sure block dimensions are valid + if (0 == launch.blockDim.x * launch.blockDim.y * launch.blockDim.z) { + return hipErrorInvalidConfiguration; + } + if (launch.stream != nullptr) { + // Validate devices to make sure it dosn't have duplicates + amd::HostQueue* queue = reinterpret_cast(launch.stream)->asHostQueue(); + auto device = &queue->vdev()->device(); + for (int j = 0; j < numDevices; ++j) { + if (mgpu_list[j] == device) { + return hipErrorInvalidDevice; + } + } + mgpu_list[i] = device; + } else { + return hipErrorInvalidResourceHandle; + } } uint64_t prevGridSize = 0; uint32_t firstDevice = 0; + + // Sync the execution streams on all devices + if ((flags & hipCooperativeLaunchMultiDeviceNoPreSync) == 0) { + for (int i = 0; i < numDevices; ++i) { + amd::HostQueue* queue = + reinterpret_cast(launchParamsList[i].stream)->asHostQueue(); + queue->finish(); + } + } + for (int i = 0; i < numDevices; ++i) { const hipLaunchParams& launch = launchParamsList[i]; amd::HostQueue* queue = reinterpret_cast(launch.stream)->asHostQueue(); @@ -506,6 +551,15 @@ hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsL prevGridSize += launch.gridDim.x * launch.gridDim.y * launch.gridDim.z; } + // Sync the execution streams on all devices + if ((flags & hipCooperativeLaunchMultiDeviceNoPostSync) == 0) { + for (int i = 0; i < numDevices; ++i) { + amd::HostQueue* queue = + reinterpret_cast(launchParamsList[i].stream)->asHostQueue(); + queue->finish(); + } + } + return result; } diff --git a/vdi/hip_platform.cpp b/vdi/hip_platform.cpp index 7e668b16cc..5ece473e06 100755 --- a/vdi/hip_platform.cpp +++ b/vdi/hip_platform.cpp @@ -544,15 +544,15 @@ extern "C" void __hipRegisterFunction( // executions. extern "C" void __hipRegisterVar( std::vector >* modules, // The device modules containing code object - char* var, // The shadow variable in host code + void* var, // The shadow variable in host code char* hostVar, // Variable name in host code char* deviceVar, // Variable name in device code int ext, // Whether this variable is external - int size, // Size of the variable + size_t size, // Size of the variable int constant, // Whether this variable is constant int global) // Unknown, always 0 { - PlatformState::DeviceVar dvar{var, std::string{ hostVar }, static_cast(size), modules, + PlatformState::DeviceVar dvar{var, std::string{ hostVar }, size, modules, std::vector{g_devices.size()}, false }; PlatformState::instance().registerVar(hostVar, dvar); @@ -704,7 +704,7 @@ hipError_t ihipCreateGlobalVarObj(const char* name, hipModule_t hmod, amd::Memor HIP_RETURN(hipErrorInvalidDeviceFunction); } /* Find the global Symbols */ - if(!dev_program->createGlobalVarObj(amd_mem_obj, dptr, bytes, name)) { + if (!dev_program->createGlobalVarObj(amd_mem_obj, dptr, bytes, name)) { HIP_RETURN(hipErrorInvalidSymbol); } @@ -713,52 +713,46 @@ hipError_t ihipCreateGlobalVarObj(const char* name, hipModule_t hmod, amd::Memor namespace hip_impl { -hipError_t ihipOccupancyMaxActiveBlocksPerMultiprocessor(int* numBlocks, int* numGrids, - hipFunction_t f, int blockSize, - size_t dynamicSMemSize, bool bCalcPotentialBlkSz) +hipError_t ihipOccupancyMaxActiveBlocksPerMultiprocessor( + int* numBlocks, int* numGrids, + const amd::Device& device, hipFunction_t func, int blockSize, + size_t dynamicSMemSize, bool bCalcPotentialBlkSz) { - HIP_INIT_API(NONE, f, blockSize, dynamicSMemSize, bCalcPotentialBlkSz); - if(numBlocks == nullptr){HIP_RETURN(hipErrorInvalidValue);} - int deviceId = ihipGetDevice(); - // FIXME: Function may not be a device function and may have been obtaiend via - // hipModuleGetFunction and thus not in the functions_ map. Check the map - // else interpret as a hip::Function for now. - hipFunction_t func = PlatformState::instance().getFunc(f, deviceId); - if (func == nullptr) { - func = f; - } hip::Function* function = hip::Function::asFunction(func); - if (function == nullptr) { - HIP_RETURN(hipErrorInvalidDeviceFunction); - } - amd::Kernel* kernel = function->function_; - if (!kernel) { - HIP_RETURN(hipErrorOutOfMemory); - } - amd::Device* device = hip::getCurrentDevice()->devices()[0]; - const device::Kernel::WorkGroupInfo* wrkGrpInfo = kernel->getDeviceKernel(*device)->workGroupInfo(); + const amd::Kernel& kernel = *function->function_; + + const device::Kernel::WorkGroupInfo* wrkGrpInfo = kernel.getDeviceKernel(device)->workGroupInfo(); if (blockSize == 0) { if (bCalcPotentialBlkSz == false){ - HIP_RETURN(hipErrorInvalidValue); + return hipErrorInvalidValue; } else { - blockSize = device->info().maxWorkGroupSize_; // maxwavefrontperblock + blockSize = device.info().maxWorkGroupSize_; // maxwavefrontperblock } } + + // Make sure the requested block size is smaller than max supported + if (blockSize > int(device.info().maxWorkGroupSize_)) { + numBlocks = 0; + numGrids = 0; + return hipSuccess; + } + // Find threads accupancy per CU => simd_per_cu * GPR usage constexpr size_t MaxWavesPerSimd = 8; // Limited by SPI 32 per CU, hence 8 per SIMD size_t VgprWaves = MaxWavesPerSimd; if (wrkGrpInfo->usedVGPRs_ > 0) { VgprWaves = wrkGrpInfo->availableVGPRs_ / amd::alignUp(wrkGrpInfo->usedVGPRs_, 4); } + size_t GprWaves = VgprWaves; if (wrkGrpInfo->usedSGPRs_ > 0) { - const size_t maxSGPRs = (device->info().gfxipVersion_ < 800) ? 512 : 800; + const size_t maxSGPRs = (device.info().gfxipVersion_ < 800) ? 512 : 800; size_t SgprWaves = maxSGPRs / amd::alignUp(wrkGrpInfo->usedSGPRs_, 16); GprWaves = std::min(VgprWaves, SgprWaves); } - size_t alu_accupancy = device->info().simdPerCU_ * std::min(MaxWavesPerSimd, GprWaves); + size_t alu_accupancy = device.info().simdPerCU_ * std::min(MaxWavesPerSimd, GprWaves); alu_accupancy *= wrkGrpInfo->wavefrontSize_; // Calculate blocks occupancy per CU *numBlocks = alu_accupancy / amd::alignUp(blockSize, wrkGrpInfo->wavefrontSize_); @@ -766,17 +760,15 @@ hipError_t ihipOccupancyMaxActiveBlocksPerMultiprocessor(int* numBlocks, int* nu size_t total_used_lds = wrkGrpInfo->usedLDSSize_ + dynamicSMemSize; if (total_used_lds != 0) { // Calculate LDS occupancy per CU. lds_per_cu / (static_lsd + dynamic_lds) - int lds_occupancy = static_cast(device->info().localMemSize_ / total_used_lds); + int lds_occupancy = static_cast(device.info().localMemSize_ / total_used_lds); *numBlocks = std::min(*numBlocks, lds_occupancy); } - if (bCalcPotentialBlkSz){ - if (numGrids == nullptr){ - HIP_RETURN(hipErrorInvalidValue); - } - *numGrids = *numBlocks * device->info().numRTCUs_; + + if (bCalcPotentialBlkSz) { + *numGrids = *numBlocks * device.info().numRTCUs_; } - HIP_RETURN(hipSuccess); + return hipSuccess; } } @@ -786,14 +778,28 @@ hipError_t hipOccupancyMaxPotentialBlockSize(uint32_t* gridSize, uint32_t* block hipFunction_t f, size_t dynSharedMemPerBlk, uint32_t blockSizeLimit) { - int numGrids = 0; - int numBlocks = 0; - hipError_t Ret = hip_impl::ihipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocks, &numGrids, f, 0, dynSharedMemPerBlk,true); - if (Ret == hipSuccess){ - *blockSize = numBlocks; - *gridSize = numGrids; + HIP_INIT_API(hipOccupancyMaxPotentialBlockSize, f, dynSharedMemPerBlk, blockSizeLimit); + + if ((gridSize == nullptr) || (blockSize == nullptr)) { + return HIP_RETURN(hipErrorInvalidValue); } - HIP_RETURN(Ret); + + hipFunction_t func = PlatformState::instance().getFunc(f, ihipGetDevice()); + if (func == nullptr) { + func = f; + } + + const amd::Device& device = *hip::getCurrentDevice()->devices()[0]; + + int num_grids = 0; + int num_blocks = 0; + hipError_t ret = hip_impl::ihipOccupancyMaxActiveBlocksPerMultiprocessor( + &num_blocks, &num_grids, device, func, 0, dynSharedMemPerBlk,true); + if (ret == hipSuccess) { + *blockSize = num_blocks; + *gridSize = num_grids; + } + HIP_RETURN(ret); } // FIXME: Need to replace `uint32_t` with `int` finally. @@ -802,10 +808,23 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor(uint32_t* numBlocks, uint32_t blockSize, size_t dynamicSMemSize) { - int NB; - hipError_t Ret = hip_impl::ihipOccupancyMaxActiveBlocksPerMultiprocessor(&NB, nullptr, f, blockSize, dynamicSMemSize, false); - *numBlocks = NB; - HIP_RETURN(Ret); + HIP_INIT_API(hipOccupancyMaxActiveBlocksPerMultiprocessor, f, blockSize, dynamicSMemSize); + if (numBlocks == nullptr) { + return HIP_RETURN(hipErrorInvalidValue); + } + + hipFunction_t func = PlatformState::instance().getFunc(f, ihipGetDevice()); + if (func == nullptr) { + func = f; + } + + const amd::Device& device = *hip::getCurrentDevice()->devices()[0]; + + int num_blocks = 0; + hipError_t ret = hip_impl::ihipOccupancyMaxActiveBlocksPerMultiprocessor( + &num_blocks, nullptr, device, func, blockSize, dynamicSMemSize, false); + *numBlocks = num_blocks; + HIP_RETURN(ret); } // FIXME: Need to replace `uint32_t` with `int` finally. @@ -815,10 +834,23 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(uint32_t* numBl size_t dynamicSMemSize, unsigned int flags) { - int NB; - hipError_t Ret = hip_impl::ihipOccupancyMaxActiveBlocksPerMultiprocessor(&NB, nullptr, f, blockSize, dynamicSMemSize, false); - *numBlocks = NB; - HIP_RETURN(Ret); + HIP_INIT_API(hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags, f, blockSize, dynamicSMemSize, flags); + if (numBlocks == nullptr) { + return HIP_RETURN(hipErrorInvalidValue); + } + + hipFunction_t func = PlatformState::instance().getFunc(f, ihipGetDevice()); + if (func == nullptr) { + func = f; + } + + const amd::Device& device = *hip::getCurrentDevice()->devices()[0]; + + int num_blocks = 0; + hipError_t ret = hip_impl::ihipOccupancyMaxActiveBlocksPerMultiprocessor( + &num_blocks, nullptr, device, func, blockSize, dynamicSMemSize, false); + *numBlocks = num_blocks; + HIP_RETURN(ret); } } diff --git a/vdi/hip_platform.hpp b/vdi/hip_platform.hpp new file mode 100644 index 0000000000..8e5eaa191f --- /dev/null +++ b/vdi/hip_platform.hpp @@ -0,0 +1,29 @@ +/* Copyright (c) 2015-present Advanced Micro Devices, Inc. + + 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. */ +#pragma once + +#include "device/device.hpp" + +namespace hip_impl { +hipError_t ihipOccupancyMaxActiveBlocksPerMultiprocessor( + int* numBlocks, int* numGrids, + const amd::Device& device, hipFunction_t func, int blockSize, + size_t dynamicSMemSize, bool bCalcPotentialBlkSz); +} \ No newline at end of file diff --git a/vdi/hip_texture.cpp b/vdi/hip_texture.cpp index fcea0bbb47..0f8e818948 100644 --- a/vdi/hip_texture.cpp +++ b/vdi/hip_texture.cpp @@ -207,8 +207,8 @@ hipError_t ihipCreateTextureObject(hipTextureObject_t* pTexObject, hipTextureReadMode readMode = pTexDesc->readMode; // 32-bit integer format will not be promoted, regardless of whether or not // this hipTextureDesc::readMode is set hipReadModeNormalizedFloat is specified. - if ((hip::getElementSize(pResDesc->res.array.array->Format) == 4) && - (pResDesc->res.array.array->Format != HIP_AD_FORMAT_FLOAT)) { + if ((pResDesc->res.array.array->Format == HIP_AD_FORMAT_SIGNED_INT32) || + (pResDesc->res.array.array->Format == HIP_AD_FORMAT_UNSIGNED_INT32)) { readMode = hipReadModeElementType; }