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

Merge 'master' into 'amd-master'

Browse Source 
Change-Id: I79fe77319a0627dc35a1a88fe12f0e7115cd679f
This commit is contained in:
Jenkins
2019-07-19 05:10:30 -04:00
parent 98ae8eca8f 54ad90f2fa
commit 527b73618d
10 changed files with 485 additions and 5 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/docs/markdown/hip_programming_guide.md
+1 -1
View File
@@ -3,7 +3,7 @@
## Host Memory
### Introduction
hipHostMemory allocates pinned host memory which is mapped into the address space of all GPUs in the system.
hipHostMalloc allocates pinned host memory which is mapped into the address space of all GPUs in the system.
There are two use cases for this host memory:
- Faster HostToDevice and DeviceToHost Data Transfers:
The runtime tracks the hipHostMalloc allocations and can avoid some of the setup required for regular unpinned memory. For exact measurements on a specific system, experiment with --unpinned and --pinned switches for the hipBusBandwidth tool.
hipamd/src/hip_module.cpp
+1 -1
View File
@@ -56,7 +56,7 @@ using namespace ELFIO;
using namespace std;
// For HIP implicit kernargs.
static const size_t HIP_IMPLICIT_KERNARG_SIZE = 48;
static const size_t HIP_IMPLICIT_KERNARG_SIZE = 56;
static const size_t HIP_IMPLICIT_KERNARG_ALIGNMENT = 8;
struct amd_kernel_code_v3_t {
hipamd/tests/README.md
+3 -1
View File
@@ -66,7 +66,9 @@ BUILD_CMD: <targetname> <build_command> EXCLUDE_HIP_PLATFORM <hcc|nvcc|all>
%s: refers to current source file name. Additional source files needed for the test can be specified by name (including relative path).
%t: refers to target executable named derived by removing the extension from the current source file. Alternatively a target executable name can be specified.
%hc: refers to hipcc pointed to by $CMAKE_INSTALL_PREFIX/bin/hipcc.
%cc: refers to system c compiler pointed to by $CC.
%hip-path: refers to hip installed location pointed to by $CMAKE_INSTALL_PREFIX
%cc: refers to system c compiler pointed to by /usr/bin/cc.
%cxx: refers to system c compiler pointed to by /usr/bin/c++.
%S: refers to path to current source file.
%T: refers to path to current build target.
hipamd/tests/hit/HIT.cmake
+3 -1
View File
@@ -201,7 +201,9 @@ macro(HIT_ADD_FILES _dir _label _parent)
OUTPUT_STRIP_TRAILING_WHITESPACE)
string(REGEX REPLACE "\n" ";" _contents "${_contents}")
string(REGEX REPLACE "%hc" "${HIP_HIPCC_EXECUTABLE}" _contents "${_contents}")
string(REGEX REPLACE "%cc" "${CC}" _contents "${_contents}")
string(REGEX REPLACE "%hip-path" "${HIP_ROOT_DIR}" _contents "${_contents}")
string(REGEX REPLACE "%cc" "/usr/bin/cc" _contents "${_contents}")
string(REGEX REPLACE "%cxx" "/usr/bin/c++" _contents "${_contents}")
string(REGEX REPLACE "%S" ${_dir} _contents "${_contents}")
string(REGEX REPLACE "%T" ${_label} _contents "${_contents}")
foreach(_cmd ${_contents})
hipamd/tests/hit/parser
+2 -1
View File
@@ -46,7 +46,8 @@ sub parse_file {
$line =~ s/^ \* BUILD_CMD: //g; # Remove " * BUILD_CMD: "
$line =~ s/%s/$file/g; # Substitute %s -> filename
$line =~ s/%t/$exe/g; # Substitute %t -> targetname
# Substitute %hc -> /path/to/hipcc and %cc -> /path/to/cc happens in cmake
# Substitute %hc -> /path/to/hipcc and %hip-path -> /path/to/hip happens in cmake
# Substitute %cc -> /usr/bin/cc and %cxx -> /usr/bin/c++ happens in cmake
# Substitute %S -> src dir and %T -> target build dir happens in cmake
$line =~ s/\R//g; # Remove line endings
push @customBuildCMDs, $line;
hipamd/tests/src/runtimeApi/module/global_kernel.cpp
+39
View File
@@ -0,0 +1,39 @@
/*
Copyright (c) 2017-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.
*/
#include "hip/hip_runtime.h"
#define ARRAY_SIZE (16)
__device__ float myDeviceGlobal;
__device__ float myDeviceGlobalArray[16];
extern "C" __global__ void hello_world(const float* a, float* b) {
int tx = hipThreadIdx_x;
b[tx] = a[tx];
}
extern "C" __global__ void test_globals(const float* a, float* b) {
int tx = hipThreadIdx_x;
b[tx] = a[tx] + myDeviceGlobal + myDeviceGlobalArray[tx % ARRAY_SIZE];
}
hipamd/tests/src/runtimeApi/module/hipModuleGetGlobal.cpp
+159
View File
@@ -0,0 +1,159 @@
/*
Copyright (c) 2017-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_CMD: global_kernel.code %hc --genco %S/global_kernel.cpp -o global_kernel.code
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
* TEST: %t
* HIT_END
*/
#include "hip/hip_runtime.h"
#include "hip/hip_runtime_api.h"
#include <iostream>
#include <fstream>
#include <vector>
#include <hip/hip_hcc.h>
#define LEN 64
#define SIZE LEN * sizeof(float)
#define fileName "global_kernel.code"
#define HIP_CHECK(cmd) \
{ \
hipError_t status = cmd; \
if (status != hipSuccess) { \
std::cout << "error: #" << status << " (" << hipGetErrorString(status) \
<< ") at line:" << __LINE__ << ": " << #cmd << std::endl; \
abort(); \
} \
}
int main() {
float *A, *B;
float *Ad, *Bd;
A = new float[LEN];
B = new float[LEN];
for (uint32_t i = 0; i < LEN; i++) {
A[i] = i * 1.0f;
B[i] = 0.0f;
}
hipInit(0);
hipDevice_t device;
hipCtx_t context;
hipDeviceGet(&device, 0);
hipCtxCreate(&context, 0, device);
hipMalloc((void**)&Ad, SIZE);
hipMalloc((void**)&Bd, SIZE);
hipMemcpyHtoD(hipDeviceptr_t(Ad), A, SIZE);
hipMemcpyHtoD((hipDeviceptr_t)(Bd), B, SIZE);
hipModule_t Module;
HIP_CHECK(hipModuleLoad(&Module, fileName));
float myDeviceGlobal_h = 42.0;
hipDeviceptr_t deviceGlobal;
size_t deviceGlobalSize;
HIP_CHECK(hipModuleGetGlobal(&deviceGlobal, &deviceGlobalSize, Module, "myDeviceGlobal"));
HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(deviceGlobal), &myDeviceGlobal_h, deviceGlobalSize));
#define ARRAY_SIZE 16
float myDeviceGlobalArray_h[ARRAY_SIZE];
hipDeviceptr_t myDeviceGlobalArray;
size_t myDeviceGlobalArraySize;
HIP_CHECK(hipModuleGetGlobal((hipDeviceptr_t*)&myDeviceGlobalArray, &myDeviceGlobalArraySize, Module, "myDeviceGlobalArray"));
for (int i = 0; i < ARRAY_SIZE; i++) {
myDeviceGlobalArray_h[i] = i * 1000.0f;
HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(myDeviceGlobalArray), &myDeviceGlobalArray_h, myDeviceGlobalArraySize));
}
struct {
void* _Ad;
void* _Bd;
} args;
args._Ad = (void*) Ad;
args._Bd = (void*) Bd;
size_t size = sizeof(args);
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
HIP_LAUNCH_PARAM_END};
{
hipFunction_t Function;
HIP_CHECK(hipModuleGetFunction(&Function, Module, "hello_world"));
HIP_CHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config));
hipMemcpyDtoH(B, hipDeviceptr_t(Bd), SIZE);
int mismatchCount = 0;
for (uint32_t i = 0; i < LEN; i++) {
if (A[i] != B[i]) {
mismatchCount++;
std::cout << "error: mismatch " << A[i] << " != " << B[i] << std::endl;
if (mismatchCount >= 10) {
break;
}
}
}
if (mismatchCount == 0) {
std::cout << "PASSED!\n";
} else {
std::cout << "FAILED!\n";
};
}
{
hipFunction_t Function;
HIP_CHECK(hipModuleGetFunction(&Function, Module, "test_globals"));
HIP_CHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config));
hipMemcpyDtoH(B, hipDeviceptr_t(Bd), SIZE);
int mismatchCount = 0;
for (uint32_t i = 0; i < LEN; i++) {
float expected = A[i] + myDeviceGlobal_h + + myDeviceGlobalArray_h[i % 16];
if (expected != B[i]) {
mismatchCount++;
std::cout << "error: mismatch " << expected << " != " << B[i] << std::endl;
if (mismatchCount >= 10) {
break;
}
}
}
if (mismatchCount == 0) {
std::cout << "PASSED!\n";
} else {
std::cout << "FAILED!\n";
};
}
hipCtxDestroy(context);
return 0;
}
hipamd/tests/src/runtimeApi/module/hipModuleLoadData.cpp
+98
View File
@@ -0,0 +1,98 @@
/*
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
* TEST: %t
* HIT_END
*/
#include "hip/hip_runtime.h"
#include "hip/hip_runtime_api.h"
#include <iostream>
#include <fstream>
#include <vector>
#include <thread>
#include <chrono>
#include "test_common.h"
#define LEN 64
#define SIZE LEN << 2
#define FILENAME "vcpy_kernel.code"
#define kernel_name "hello_world"
int main() {
float *A, *B, *Ad, *Bd;
A = new float[LEN];
B = new float[LEN];
for (uint32_t i = 0; i < LEN; i++) {
A[i] = i * 1.0f;
B[i] = 0.0f;
}
HIPCHECK(hipInit(0));
HIPCHECK(hipMalloc((void**)&Ad, SIZE));
HIPCHECK(hipMalloc((void**)&Bd, SIZE));
HIPCHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
HIPCHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice));
hipModule_t Module;
hipFunction_t Function;
std::ifstream file(FILENAME, std::ios::binary | std::ios::ate);
std::streamsize fsize = file.tellg();
file.seekg(0, std::ios::beg);
std::vector<char> buffer(fsize);
if (file.read(buffer.data(), fsize)) {
HIPCHECK(hipModuleLoadData(&Module, &buffer[0]));
HIPCHECK(hipModuleGetFunction(&Function, Module, kernel_name));
}
else {
failed("could not open code object '%s'\n", FILENAME);
}
hipStream_t stream;
HIPCHECK(hipStreamCreate(&stream));
struct {
void* _Ad;
void* _Bd;
} args;
args._Ad = (void*) Ad;
args._Bd = (void*) Bd;
size_t size = sizeof(args);
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
HIP_LAUNCH_PARAM_END};
HIPCHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, stream, NULL, (void**)&config));
HIPCHECK(hipStreamDestroy(stream));
HIPCHECK(hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost));
for (uint32_t i = 0; i < LEN; i++) {
assert(A[i] == B[i]);
}
passed();
}
hipamd/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp
+144
View File
@@ -0,0 +1,144 @@
/*
Copyright (c) 2015-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 ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
* TEST: %t
* HIT_END
*/
#include "hip/hip_runtime.h"
//#include "hip/hip_runtime_api.h"
#include <iostream>
#include <fstream>
#include <vector>
//#include <hip/hip_hcc.h>
#define fileName "tex2d_kernel.code"
texture<float, 2, hipReadModeElementType> tex;
bool testResult = false;
#define HIP_CHECK(cmd) \
{ \
hipError_t status = cmd; \
if (status != hipSuccess) { \
std::cout << "error: #" << status << " (" << hipGetErrorString(status) \
<< ") at line:" << __LINE__ << ": " << #cmd << std::endl; \
abort(); \
} \
}
bool runTest(int argc, char** argv) {
unsigned int width = 256;
unsigned int height = 256;
unsigned int size = width * height * sizeof(float);
float* hData = (float*)malloc(size);
memset(hData, 0, size);
for (int i = 0; i < height; i++) {
for (int j = 0; j < width; j++) {
hData[i * width + j] = i * width + j;
}
}
hipModule_t Module;
HIP_CHECK(hipModuleLoad(&Module, fileName));
hipArray* array;
HIP_ARRAY_DESCRIPTOR desc;
desc.Format = HIP_AD_FORMAT_FLOAT;
desc.NumChannels = 1;
desc.Width = width;
desc.Height = height;
hipArrayCreate(&array, &desc);
hip_Memcpy2D copyParam;
memset(&copyParam, 0, sizeof(copyParam));
copyParam.dstMemoryType = hipMemoryTypeArray;
copyParam.dstArray = array;
copyParam.srcMemoryType = hipMemoryTypeHost;
copyParam.srcHost = hData;
copyParam.srcPitch = width * sizeof(float);
copyParam.WidthInBytes = copyParam.srcPitch;
copyParam.Height = height;
hipMemcpyParam2D(&copyParam);
textureReference* texref;
hipModuleGetTexRef(&texref, Module, "tex");
hipTexRefSetAddressMode(texref, 0, hipAddressModeWrap);
hipTexRefSetAddressMode(texref, 1, hipAddressModeWrap);
hipTexRefSetFilterMode(texref, hipFilterModePoint);
hipTexRefSetFlags(texref, 0);
hipTexRefSetFormat(texref, HIP_AD_FORMAT_FLOAT, 1);
hipTexRefSetArray(texref, array, HIP_TRSA_OVERRIDE_FORMAT);
float* dData = NULL;
hipMalloc((void**)&dData, size);
struct {
void* _Ad;
unsigned int _Bd;
unsigned int _Cd;
} args;
args._Ad = (void*) dData;
args._Bd = width;
args._Cd = height;
size_t sizeTemp = sizeof(args);
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE,
&sizeTemp, HIP_LAUNCH_PARAM_END};
hipFunction_t Function;
HIP_CHECK(hipModuleGetFunction(&Function, Module, "tex2dKernel"));
int temp1 = width / 16;
int temp2 = height / 16;
HIP_CHECK(
hipModuleLaunchKernel(Function, 16, 16, 1, temp1, temp2, 1, 0, 0, NULL, (void**)&config));
hipDeviceSynchronize();
float* hOutputData = (float*)malloc(size);
memset(hOutputData, 0, size);
hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost);
for (int i = 0; i < height; i++) {
for (int j = 0; j < width; j++) {
if (hData[i * width + j] != hOutputData[i * width + j]) {
printf("Difference [ %d %d ]:%f ----%f\n", i, j, hData[i * width + j],
hOutputData[i * width + j]);
testResult = false;
break;
}
}
}
hipFree(dData);
hipFreeArray(array);
return true;
}
int main(int argc, char** argv) {
hipInit(0);
testResult = runTest(argc, argv);
printf("%s ...\n", testResult ? "PASSED" : "FAILED");
exit(testResult ? EXIT_SUCCESS : EXIT_FAILURE);
return 0;
}
hipamd/tests/src/runtimeApi/module/tex2d_kernel.cpp
+35
View File
@@ -0,0 +1,35 @@
/*
Copyright (c) 2015-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_CMD: tex2d_kernel.code %hc --genco %S/tex2d_kernel.cpp -o tex2d_kernel.code
* HIT_END
*/
#include "hip/hip_runtime.h"
extern texture<float, 2, hipReadModeElementType> tex;
extern "C" __global__ void tex2dKernel(float* outputData, int width, int height) {
int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
outputData[y * width + x] = tex2D(tex, x, y);
}