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: I0230d43dca968ed6a8f7695064b98a3de703143d


[ROCm/hip commit: 1c20b5ba1e]
This commit is contained in:
Jenkins
2019-07-03 19:12:34 -04:00
parent 2375f318c9 01ee58e0c7
commit f2b8f662b9
11 changed files with 174 additions and 153 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/include/hip/hcc_detail/host_defines.h
+1
View File
@@ -72,6 +72,7 @@ THE SOFTWARE.
#define __noinline__ __attribute__((noinline))
#define __forceinline__ inline __attribute__((always_inline))
#define __hip_pinned_shadow__ __attribute__((hip_pinned_shadow))
#else
projects/hip/samples/0_Intro/module_api_global/runKernel.cpp
+5 -3
View File
@@ -31,7 +31,6 @@ THE SOFTWARE.
#define SIZE LEN * sizeof(float)
#define fileName "vcpy_kernel.code"
float myDeviceGlobalArray[16];
#define HIP_CHECK(cmd) \
{ \
hipError_t status = cmd; \
@@ -71,14 +70,17 @@ int main() {
float* deviceGlobal;
size_t deviceGlobalSize;
HIP_CHECK(hipModuleGetGlobal((void**)&deviceGlobal, &deviceGlobalSize, Module, "myDeviceGlobal"));
*deviceGlobal = 42.0;
HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(deviceGlobal), &myDeviceGlobal_h, deviceGlobalSize));
#define ARRAY_SIZE 16
float myDeviceGlobalArray_h[ARRAY_SIZE];
float *myDeviceGlobalArray;
size_t myDeviceGlobalArraySize;
HIP_CHECK(hipModuleGetGlobal((void**)&myDeviceGlobalArray, &myDeviceGlobalArraySize, Module, "myDeviceGlobalArray"));
for (int i = 0; i < ARRAY_SIZE; i++) {
myDeviceGlobalArray_h[i] = i * 1000.0f;
myDeviceGlobalArray[i] = i * 1000.0f;
HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(myDeviceGlobalArray), &myDeviceGlobalArray_h, myDeviceGlobalArraySize));
}
struct {
projects/hip/samples/0_Intro/module_api_global/vcpy_kernel.cpp
+1 -2
View File
@@ -25,8 +25,7 @@ THE SOFTWARE.
#define ARRAY_SIZE (16)
__device__ float myDeviceGlobal;
extern float myDeviceGlobalArray[16];
;
__device__ float myDeviceGlobalArray[16];
extern "C" __global__ void hello_world(const float* a, float* b) {
int tx = hipThreadIdx_x;
projects/hip/samples/2_Cookbook/9_unroll/Readme.md
+5 -7
View File
@@ -16,17 +16,15 @@ Programmers familiar with CUDA, OpenCL will be able to quickly learn and start c
## Simple Matrix Transpose
For this tutorial we will be using MatrixTranspose with shfl operation i.e., our 4_shfl tutorial since it is the only examples where we used loops inside the kernel.
For this tutorial we will be using an example which sums up the row of a 2D matrix and writes it in a 1D array.
In this tutorial, we'll use `#pragma unroll`. In the same sourcecode, we used for MatrixTranspose. We'll add it just before the for loop as following:
In this tutorial, we'll use `#pragma unroll`. In the same sourcecode, we used for gpuMatrixRowSum. We'll add it just before the for loop as following:
```
#pragma unroll
for(int i=0;i<width;i++)
{
for(int j=0;j<width;j++)
out[i*width + j] = __shfl(val,j*width + i);
}
for (int i = 0; i < width; i++) {
output[index] += input[index * width + i]
}
```
Specifying the optional parameter, #pragma unroll value, directs the unroller to unroll the loop value times. Be careful while using it.
projects/hip/samples/2_Cookbook/9_unroll/unroll.cpp
+59 -61
View File
@@ -25,100 +25,98 @@ THE SOFTWARE.
// hip header file
#include "hip/hip_runtime.h"
#define LENGTH 4
#define WIDTH 4
#define SIZE (LENGTH * LENGTH)
#define NUM (WIDTH * WIDTH)
#define THREADS_PER_BLOCK 1
#define BLOCKS_PER_GRID LENGTH
#define THREADS_PER_BLOCK_X 4
#define THREADS_PER_BLOCK_Y 4
#define THREADS_PER_BLOCK_Z 1
// Device (Kernel) function, it must be void
__global__ void matrixTranspose(float* out, float* in, const int width) {
int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
float val = in[x];
#pragma unroll
// CPU function - basically scan each row and save the output in array
void matrixRowSum(int* input, int* output, int width) {
for (int i = 0; i < width; i++) {
for (int j = 0; j < width; j++) out[i * width + j] = __shfl(val, j * width + i);
}
}
// CPU implementation of matrix transpose
void matrixTransposeCPUReference(float* output, float* input, const unsigned int width) {
for (unsigned int j = 0; j < width; j++) {
for (unsigned int i = 0; i < width; i++) {
output[i * width + j] = input[j * width + i];
for (int j = 0; j < width; j++) {
output[i] += input[i * width + j];
}
}
}
int main() {
float* Matrix;
float* TransposeMatrix;
float* cpuTransposeMatrix;
// Device (kernel) function
__global__ void gpuMatrixRowSum(int* input, int* output, int width) {
int index = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
#pragma unroll
for (int i = 0; i < width; i++) {
output[index] += input[index * width + i];
}
}
float* gpuMatrix;
float* gpuTransposeMatrix;
int main() {
int* Matrix;
int* sumMatrix;
int* cpuSumMatrix;
int* gpuMatrix;
int* gpuSumMatrix;
hipDeviceProp_t devProp;
hipGetDeviceProperties(&devProp, 0);
std::cout << "Device name " << devProp.name << std::endl;
int i;
int errors;
Matrix = (int*)malloc(sizeof(int) * SIZE);
sumMatrix = (int*)malloc(sizeof(int) * LENGTH);
cpuSumMatrix = (int*)malloc(sizeof(int) * LENGTH);
Matrix = (float*)malloc(NUM * sizeof(float));
TransposeMatrix = (float*)malloc(NUM * sizeof(float));
cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float));
// initialize the input data
for (i = 0; i < NUM; i++) {
Matrix[i] = (float)i * 10.0f;
for (int i = 0; i < SIZE; i++) {
Matrix[i] = i * 2;
}
// allocate the memory on the device side
hipMalloc((void**)&gpuMatrix, NUM * sizeof(float));
hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float));
for (int i = 0; i < LENGTH; i++) {
cpuSumMatrix[i] = 0;
}
// Memory transfer from host to device
hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice);
// Allocated Device Memory
hipMalloc((void**)&gpuMatrix, SIZE * sizeof(int));
hipMalloc((void**)&gpuSumMatrix, LENGTH * sizeof(int));
// Lauching kernel from host
hipLaunchKernelGGL(matrixTranspose, dim3(1), dim3(THREADS_PER_BLOCK_X * THREADS_PER_BLOCK_Y), 0, 0,
gpuTransposeMatrix, gpuMatrix, WIDTH);
// Memory Copy to Device
hipMemcpy(gpuMatrix, Matrix, SIZE * sizeof(int), hipMemcpyHostToDevice);
hipMemcpy(gpuSumMatrix, cpuSumMatrix, LENGTH * sizeof(float), hipMemcpyHostToDevice);
// Memory transfer from device to host
hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost);
// Launch device kernels
hipLaunchKernelGGL(gpuMatrixRowSum, dim3(BLOCKS_PER_GRID), dim3(THREADS_PER_BLOCK), 0, 0,
gpuMatrix, gpuSumMatrix, LENGTH);
// Memory copy back to device
hipMemcpy(sumMatrix, gpuSumMatrix, LENGTH * sizeof(int), hipMemcpyDeviceToHost);
// Cpu implementation
matrixRowSum(Matrix, cpuSumMatrix, LENGTH);
// CPU MatrixTranspose computation
matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH);
// verify the results
errors = 0;
double eps = 1.0E-6;
for (i = 0; i < NUM; i++) {
if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps) {
printf("%d cpu: %f gpu %f\n", i, cpuTransposeMatrix[i], TransposeMatrix[i]);
int errors = 0;
for (int i = 0; i < LENGTH; i++) {
if (sumMatrix[i] != cpuSumMatrix[i]) {
printf("%d - cpu: %d gpu: %d\n", i, sumMatrix[i], cpuSumMatrix[i]);
errors++;
}
}
if (errors != 0) {
printf("FAILED: %d errors\n", errors);
if (errors == 0) {
printf("PASSED\n");
} else {
printf("PASSED!\n");
printf("FAILED with %d errors\n", errors);
}
// free the resources on device side
// GPU Free
hipFree(gpuMatrix);
hipFree(gpuTransposeMatrix);
hipFree(gpuSumMatrix);
// free the resources on host side
// CPU Free
free(Matrix);
free(TransposeMatrix);
free(cpuTransposeMatrix);
free(sumMatrix);
free(cpuSumMatrix);
return errors;
}
projects/hip/src/hip_event.cpp
+61 -46
View File
@@ -85,7 +85,7 @@ hipError_t ihipEventCreate(hipEvent_t* event, unsigned flags) {
(flags & ~supportedFlags) || // can't set any unsupported flags.
(flags & releaseFlags) == releaseFlags; // can't set both release flags
if (!illegalFlags) {
if (event && !illegalFlags) {
*event = new ihipEvent_t(flags);
} else {
e = hipErrorInvalidValue;
@@ -109,46 +109,51 @@ hipError_t hipEventCreate(hipEvent_t* event) {
hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) {
HIP_INIT_SPECIAL_API(hipEventRecord, TRACE_SYNC, event, stream);
hipError_t status;
if (event){
auto ecd = event->locked_copyCrit();
if( ecd._state != hipEventStatusUnitialized) {
stream = ihipSyncAndResolveStream(stream);
auto ecd = event->locked_copyCrit();
if (HIP_SYNC_NULL_STREAM && stream->isDefaultStream()) {
// TODO-HIP_SYNC_NULL_STREAM : can remove this code when HIP_SYNC_NULL_STREAM = 0
//
// If default stream , then wait on all queues.
ihipCtx_t* ctx = ihipGetTlsDefaultCtx();
ctx->locked_syncDefaultStream(true, true);
if (event && ecd._state != hipEventStatusUnitialized) {
stream = ihipSyncAndResolveStream(stream);
{
LockedAccessor_EventCrit_t eCrit(event->criticalData());
eCrit->_eventData.marker(hc::completion_future()); // reset event
eCrit->_eventData._stream = stream;
eCrit->_eventData._timestamp = hc::get_system_ticks();
eCrit->_eventData._state = hipEventStatusComplete;
}
status = hipSuccess;
} else {
// Record the event in the stream:
// Keep a copy outside the critical section so we lock stream first, then event - to
// avoid deadlock
hc::completion_future cf = stream->locked_recordEvent(event);
if (HIP_SYNC_NULL_STREAM && stream->isDefaultStream()) {
// TODO-HIP_SYNC_NULL_STREAM : can remove this code when HIP_SYNC_NULL_STREAM = 0
//
// If default stream , then wait on all queues.
ihipCtx_t* ctx = ihipGetTlsDefaultCtx();
ctx->locked_syncDefaultStream(true, true);
{
LockedAccessor_EventCrit_t eCrit(event->criticalData());
eCrit->_eventData.marker(cf);
eCrit->_eventData._stream = stream;
eCrit->_eventData._timestamp = 0;
eCrit->_eventData._state = hipEventStatusRecording;
}
{
LockedAccessor_EventCrit_t eCrit(event->criticalData());
eCrit->_eventData.marker(hc::completion_future()); // reset event
eCrit->_eventData._stream = stream;
eCrit->_eventData._timestamp = hc::get_system_ticks();
eCrit->_eventData._state = hipEventStatusComplete;
status = hipSuccess;
}
return ihipLogStatus(hipSuccess);
} else {
// Record the event in the stream:
// Keep a copy outside the critical section so we lock stream first, then event - to
// avoid deadlock
hc::completion_future cf = stream->locked_recordEvent(event);
{
LockedAccessor_EventCrit_t eCrit(event->criticalData());
eCrit->_eventData.marker(cf);
eCrit->_eventData._stream = stream;
eCrit->_eventData._timestamp = 0;
eCrit->_eventData._state = hipEventStatusRecording;
}
return ihipLogStatus(hipSuccess);
status = hipErrorInvalidResourceHandle;
}
} else {
return ihipLogStatus(hipErrorInvalidResourceHandle);
status = hipErrorInvalidResourceHandle;
}
return ihipLogStatus(status);
}
@@ -202,11 +207,13 @@ hipError_t hipEventElapsedTime(float* ms, hipEvent_t start, hipEvent_t stop) {
hipError_t status = hipSuccess;
*ms = 0.0f;
if ((start == nullptr) || (stop == nullptr)) {
if (ms == nullptr) {
status = hipErrorInvalidValue;
}
else if ((start == nullptr) || (stop == nullptr)) {
status = hipErrorInvalidResourceHandle;
} else {
*ms = 0.0f;
auto startEcd = start->locked_copyCrit();
auto stopEcd = stop->locked_copyCrit();
@@ -256,18 +263,26 @@ hipError_t hipEventElapsedTime(float* ms, hipEvent_t start, hipEvent_t stop) {
hipError_t hipEventQuery(hipEvent_t event) {
HIP_INIT_SPECIAL_API(hipEventQuery, TRACE_QUERY, event);
hipError_t status = hipSuccess;
if (!(event->_flags & hipEventReleaseToSystem)) {
tprintf(DB_WARN,
"hipEventQuery on event without system-scope fence ; consider creating with "
"hipEventReleaseToSystem\n");
}
auto ecd = event->locked_copyCrit();
if ((ecd._state == hipEventStatusRecording) && !ecd._stream->locked_eventIsReady(event)) {
return ihipLogStatus(hipErrorNotReady);
if ( NULL == event)
{
status = hipErrorInvalidResourceHandle;
} else {
return ihipLogStatus(hipSuccess);
if (!(event->_flags & hipEventReleaseToSystem)) {
tprintf(DB_WARN,
"hipEventQuery on event without system-scope fence ; consider creating with "
"hipEventReleaseToSystem\n");
}
auto ecd = event->locked_copyCrit();
if ((ecd._state == hipEventStatusRecording) && !ecd._stream->locked_eventIsReady(event)) {
status = hipErrorNotReady;
} else {
status = hipSuccess;
}
}
return ihipLogStatus(status);
}
projects/hip/src/hip_memory.cpp
+11 -15
View File
@@ -243,13 +243,12 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) {
auto ctx = ihipGetTlsDefaultCtx();
// return NULL pointer when malloc size is 0
if (sizeBytes == 0) {
if ( nullptr == ctx || nullptr == ptr) {
hip_status = hipErrorInvalidValue;
}
else if (sizeBytes == 0) {
*ptr = NULL;
hip_status = hipSuccess;
} else if ((ctx == nullptr) || (ptr == nullptr)) {
hip_status = hipErrorInvalidValue;
} else {
auto device = ctx->getWriteableDevice();
*ptr = hip_internal::allocAndSharePtr("device_mem", sizeBytes, ctx, false /*shareWithAll*/,
@@ -309,12 +308,12 @@ hipError_t ihipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) {
}
auto ctx = ihipGetTlsDefaultCtx();
if (sizeBytes == 0) {
if ((ctx == nullptr) || (ptr == nullptr)) {
hip_status = hipErrorInvalidValue;
}
else if (sizeBytes == 0) {
hip_status = hipSuccess;
// TODO - should size of 0 return err or be siliently ignored?
} else if ((ctx == nullptr) || (ptr == nullptr)) {
hip_status = hipErrorInvalidValue;
} else {
unsigned trueFlags = flags;
if (flags == hipHostMallocDefault) {
@@ -400,7 +399,7 @@ hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags) {
// width in bytes
hipError_t ihipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height, size_t depth) {
hipError_t hip_status = hipSuccess;
if(ptr==NULL)
if(ptr==NULL || pitch == NULL)
{
hip_status=hipErrorInvalidValue;
return hip_status;
@@ -916,11 +915,8 @@ hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) {
am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, hostPtr);
if (status == AM_SUCCESS) {
*flagsPtr = amPointerInfo._appAllocationFlags;
if (*flagsPtr == 0) {
hip_status = hipErrorInvalidValue;
} else {
hip_status = hipSuccess;
}
//0 is valid flag hipHostMallocDefault, and during hipHostMalloc if unsupported flags are passed as parameter it throws error
hip_status = hipSuccess;
tprintf(DB_MEM, " %s: host ptr=%p\n", __func__, hostPtr);
} else {
hip_status = hipErrorInvalidValue;
projects/hip/src/hip_peer.cpp
+4 -2
View File
@@ -41,8 +41,10 @@ THE SOFTWARE.
hipError_t ihipDeviceCanAccessPeer(int* canAccessPeer, hipCtx_t thisCtx, hipCtx_t peerCtx) {
hipError_t err = hipSuccess;
if ((thisCtx != NULL) && (peerCtx != NULL)) {
if(canAccessPeer == NULL) {
err = hipErrorInvalidValue;
}
else if ((thisCtx != NULL) && (peerCtx != NULL)) {
if (thisCtx == peerCtx) {
*canAccessPeer = 0;
tprintf(DB_MEM, "Can't be peer to self. (this=%s, peer=%s)\n",
projects/hip/src/hip_stream.cpp
+21 -17
View File
@@ -56,6 +56,8 @@ hipError_t ihipStreamCreate(hipStream_t* stream, unsigned int flags, int priorit
if (ctx) {
if (HIP_FORCE_NULL_STREAM) {
*stream = 0;
} else if( NULL == stream ){
e = hipErrorInvalidValue;
} else {
hc::accelerator acc = ctx->getWriteableDevice()->_acc;
@@ -65,7 +67,7 @@ hipError_t ihipStreamCreate(hipStream_t* stream, unsigned int flags, int priorit
// CUDA stream behavior is that all kernels submitted will automatically
// wait for prev to complete, this behaviour will be mainatined by
// hipModuleLaunchKernel. execute_any_order will help
// hipExtModuleLaunchKernel , which uses a special flag
// hipExtModuleLaunchKernel , which uses a special flag
{
// Obtain mutex access to the device critical data, release by destructor
@@ -80,9 +82,9 @@ hipError_t ihipStreamCreate(hipStream_t* stream, unsigned int flags, int priorit
ctxCrit->addStream(istream);
*stream = istream;
}
tprintf(DB_SYNC, "hipStreamCreate, %s\n", ToString(*stream).c_str());
}
tprintf(DB_SYNC, "hipStreamCreate, %s\n", ToString(*stream).c_str());
} else {
e = hipErrorInvalidDevice;
}
@@ -94,8 +96,10 @@ hipError_t ihipStreamCreate(hipStream_t* stream, unsigned int flags, int priorit
//---
hipError_t hipStreamCreateWithFlags(hipStream_t* stream, unsigned int flags) {
HIP_INIT_API(hipStreamCreateWithFlags, stream, flags);
return ihipLogStatus(ihipStreamCreate(stream, flags, priority_normal));
if(flags == hipStreamDefault || flags == hipStreamNonBlocking)
return ihipLogStatus(ihipStreamCreate(stream, flags, priority_normal));
else
return ihipLogStatus(hipErrorInvalidValue);
}
//---
@@ -128,25 +132,25 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int
hipError_t e = hipSuccess;
auto ecd = event->locked_copyCrit();
if (event == nullptr) {
e = hipErrorInvalidResourceHandle;
} else if ((ecd._state != hipEventStatusUnitialized) && (ecd._state != hipEventStatusCreated)) {
if (HIP_SYNC_STREAM_WAIT || (HIP_SYNC_NULL_STREAM && (stream == 0))) {
// conservative wait on host for the specified event to complete:
// return _stream->locked_eventWaitComplete(this, waitMode);
//
ecd._stream->locked_eventWaitComplete(
} else {
auto ecd = event->locked_copyCrit();
if ((ecd._state != hipEventStatusUnitialized) && (ecd._state != hipEventStatusCreated)) {
if (HIP_SYNC_STREAM_WAIT || (HIP_SYNC_NULL_STREAM && (stream == 0))) {
// conservative wait on host for the specified event to complete:
// return _stream->locked_eventWaitComplete(this, waitMode);
//
ecd._stream->locked_eventWaitComplete(
ecd.marker(), (event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked
: hc::hcWaitModeActive);
} else {
stream = ihipSyncAndResolveStream(stream);
// This will use create_blocking_marker to wait on the specified queue.
stream->locked_streamWaitEvent(ecd);
} else {
stream = ihipSyncAndResolveStream(stream);
// This will use create_blocking_marker to wait on the specified queue.
stream->locked_streamWaitEvent(ecd);
}
}
} // else event not recorded, return immediately and don't create marker.
return ihipLogStatus(e);
projects/hip/tests/src/texture/hipBindTexRef1DFetch.cpp
+3
View File
@@ -32,6 +32,9 @@ THE SOFTWARE.
#define N 512
#if __HIP__
__hip_pinned_shadow__
#endif
texture<float, 1, hipReadModeElementType> tex;
__global__ void kernel(float *out) {
projects/hip/tests/src/texture/hipTextureRef2D.cpp
+3
View File
@@ -9,6 +9,9 @@
#include <hip/hip_runtime.h>
#include "test_common.h"
#if __HIP__
__hip_pinned_shadow__
#endif
texture<float, 2, hipReadModeElementType> tex;
__global__ void tex2DKernel(float* outputData,