AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge branch 'amd-master-next' into amd-npi-next

Browse Source 
Change-Id: Ia4495849a02f23f2a2e87abae12c0e9db7c44b08
This commit is contained in:
Vladislav Sytchenko
2020-04-07 19:02:56 -04:00
parent 6912807674 f99db12c82
commit 9b5bd1bf41
24 changed files with 1368 additions and 135 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hip-config-clang.cmake.in
+13 -3
View File
@@ -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")
include/hip/hip_runtime_api.h
+17 -1
View File
@@ -128,6 +128,14 @@ typedef struct hipDeviceProp_t {
int kernelExecTimeoutEnabled; ///<Run time limit for kernels executed on the device
int ECCEnabled; ///<Device has ECC support enabled
int tccDriver; ///< 1:If device is Tesla device using TCC driver, else 0
int cooperativeMultiDeviceUnmatchedFunc; ///< HIP device supports cooperative launch on multiple
///devices with unmatched functions
int cooperativeMultiDeviceUnmatchedGridDim; ///< HIP device supports cooperative launch on multiple
///devices with unmatched grid dimensions
int cooperativeMultiDeviceUnmatchedBlockDim; ///< HIP device supports cooperative launch on multiple
///devices with unmatched block dimensions
int cooperativeMultiDeviceUnmatchedSharedMem; ///< HIP device supports cooperative launch on multiple
///devices with unmatched shared memories
} hipDeviceProp_t;
@@ -329,8 +337,16 @@ typedef enum hipDeviceAttribute_t {
hipDeviceAttributeTexturePitchAlignment, ///<Pitch alignment requirement for 2D texture references bound to pitched memory;
hipDeviceAttributeKernelExecTimeout, ///<Run time limit for kernels executed on the device
hipDeviceAttributeCanMapHostMemory, ///<Device can map host memory into device address space
hipDeviceAttributeEccEnabled ///<Device has ECC support enabled
hipDeviceAttributeEccEnabled, ///<Device has ECC support enabled
hipDeviceAttributeCooperativeMultiDeviceUnmatchedFunc, ///< Supports cooperative launch on multiple
///devices with unmatched functions
hipDeviceAttributeCooperativeMultiDeviceUnmatchedGridDim, ///< Supports cooperative launch on multiple
///devices with unmatched grid dimensions
hipDeviceAttributeCooperativeMultiDeviceUnmatchedBlockDim, ///< Supports cooperative launch on multiple
///devices with unmatched block dimensions
hipDeviceAttributeCooperativeMultiDeviceUnmatchedSharedMem ///< Supports cooperative launch on multiple
///devices with unmatched shared memories
} hipDeviceAttribute_t;
enum hipComputeMode {
tests/src/printf/hipPrintfAltForms.cpp
+76
View File
@@ -0,0 +1,76 @@
/*
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("%#o\n", 042);
printf("%#x\n", 0x42);
printf("%#X\n", 0x42);
printf("%#08x\n", 0x42);
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("%#a\n", 123.456);
printf("%#A\n", -123.456);
printf("%#.8x\n", 0x42);
printf("%#16.8x\n", 0x42);
printf("%-#16.8x\n", 0x42);
}
int main(int argc, char **argv) {
std::string reference(R"here(042
0x42
0X42
0x000042
-123.456000
123.456000
1.234560e+02
-1.234560E+02
-123.456
123.456
0x1.edd2f1a9fbe77p+6
-0X1.EDD2F1A9FBE77P+6
0x00000042
0x00000042
0x00000042
)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();
}
tests/src/printf/hipPrintfBasic.cpp
+238
View File
@@ -0,0 +1,238 @@
/*
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 <vector>
// 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<std::string, int> 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<std::string, int> 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<uint> 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<std::string, int> 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<std::string, int> 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<int *>(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();
}
tests/src/printf/hipPrintfFlags.cpp
+68
View File
@@ -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();
}
tests/src/printf/hipPrintfManyDevices.cpp
+77
View File
@@ -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<std::string, int> 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();
}
tests/src/printf/hipPrintfManyWaves.cpp
+301
View File
@@ -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 <vector>
// 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<std::string, int> 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<std::string, int> 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<std::string, int> 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<std::string, int> 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<uint> 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<int *>(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();
}
tests/src/printf/hipPrintfSpecifiers.cpp
+90
View File
@@ -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();
}
tests/src/printf/hipPrintfStar.cpp
+54
View File
@@ -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();
}
tests/src/printf/hipPrintfWidthPrecision.cpp
+74
View File
@@ -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();
}
tests/src/printf/printf_common.h
+94
View File
@@ -0,0 +1,94 @@
#ifndef COMMON_H
#define COMMON_H
#include <errno.h>
#include <error.h>
#include <fstream>
#include <iostream>
#include <map>
#include <stdlib.h>
#include <string>
#include <unistd.h>
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
tests/src/runtimeApi/module/hipLaunchCoopMultiKernel.cpp
+1 -3
View File
@@ -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";
tests/src/runtimeApi/module/hipLaunchCooperativeKernel.cpp
+1 -1
View File
@@ -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
*/
vdi/hip_conversions.hpp
+10 -10
View File
@@ -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;
vdi/hip_device.cpp
+5
View File
@@ -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_;
vdi/hip_device_runtime.cpp
+12
View File
@@ -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);
}
vdi/hip_hcc.def.in
+1
View File
@@ -238,6 +238,7 @@ hipFreeMipmappedArray
hipMipmappedArrayGetLevel
hipGetMipmappedArrayLevel
hipMallocHost
hipFreeHost
hipTexObjectCreate
hipTexObjectDestroy
hipTexObjectGetResourceDesc
vdi/hip_hcc.map.in
+1
View File
@@ -232,6 +232,7 @@ global:
hipMipmappedArrayGetLevel;
hipGetMipmappedArrayLevel;
hipMallocHost;
hipFreeHost;
hipTexObjectCreate;
hipTexObjectDestroy;
hipTexObjectGetResourceDesc;
vdi/hip_internal.hpp
+12 -3
View File
@@ -31,7 +31,11 @@
#include <stack>
#include <mutex>
#include <iterator>
#ifdef _WIN32
#include <process.h>
#else
#include <unistd.h>
#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<hip::Device*> g_devices;
vdi/hip_memory.cpp
+47 -53
View File
@@ -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<size_t*>(origin)[0] += xOffset;
static_cast<size_t*>(origin)[1] += yOffset;
static_cast<size_t*>(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<size_t*>(srcOrigin), static_cast<size_t*>(copyRegion), srcImage->getRowPitch(), srcImage->getSlicePitch())) {
@@ -930,6 +911,8 @@ hipError_t ihipMemcpyAtoD(hipArray* srcArray,
if (!dstRect.create(static_cast<size_t*>(dstOrigin), static_cast<size_t*>(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<size_t*>(srcOrigin), static_cast<size_t*>(copyRegion), srcRowPitch, srcSlicePitch)) {
return hipErrorInvalidValue;
}
srcRect.start_ += srcOffset;
srcRect.end_ += srcOffset;
amd::BufferRect dstRect;
if (!dstRect.create(static_cast<size_t*>(dstOrigin), static_cast<size_t*>(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<size_t*>(srcOrigin), static_cast<size_t*>(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<size_t*>(dstOrigin), static_cast<size_t*>(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<size_t*>(srcOrigin), static_cast<size_t*>(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<size_t*>(srcOrigin), static_cast<size_t*>(copyRegion), srcRowPitch, srcSlicePitch)) {
@@ -1162,6 +1148,8 @@ hipError_t ihipMemcpyHtoD(const void* srcHost,
if (!dstRect.create(static_cast<size_t*>(dstOrigin), static_cast<size_t*>(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));
}
vdi/hip_module.cpp
+61 -7
View File
@@ -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<const amd::Device*> 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<hip::Stream*>(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<hip::Stream*>(launchParamsList[i].stream)->asHostQueue();
queue->finish();
}
}
for (int i = 0; i < numDevices; ++i) {
const hipLaunchParams& launch = launchParamsList[i];
amd::HostQueue* queue = reinterpret_cast<hip::Stream*>(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<hip::Stream*>(launchParamsList[i].stream)->asHostQueue();
queue->finish();
}
}
return result;
}
vdi/hip_platform.cpp
+84 -52
View File
@@ -544,15 +544,15 @@ extern "C" void __hipRegisterFunction(
// executions.
extern "C" void __hipRegisterVar(
std::vector<std::pair<hipModule_t,bool> >* 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_t>(size), modules,
PlatformState::DeviceVar dvar{var, std::string{ hostVar }, size, modules,
std::vector<PlatformState::RegisteredVar>{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<int>(device->info().localMemSize_ / total_used_lds);
int lds_occupancy = static_cast<int>(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);
}
}
vdi/hip_platform.hpp
+29
View File
@@ -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);
}
vdi/hip_texture.cpp
+2 -2
View File
@@ -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;
}