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

Re-sync with upstream.

Browse Source
This commit is contained in:
Alex Voicu
2018-10-18 12:27:03 +01:00
parent 7839045d30 d133493669
commit fe959f7bd7
103 changed files with 1478 additions and 1308 deletions
Download Patch File Download Diff File Expand all files Collapse all files
include/hip/hcc_detail/device_functions.h
-74
View File
@@ -752,80 +752,6 @@ void *__amdgcn_get_dynamicgroupbaseptr() {
return __get_dynamicgroupbaseptr();
}
// hip.amdgcn.bc - sync threads
#define __CLK_LOCAL_MEM_FENCE 0x01
typedef unsigned __cl_mem_fence_flags;
typedef enum __memory_scope {
__memory_scope_work_item = __OPENCL_MEMORY_SCOPE_WORK_ITEM,
__memory_scope_work_group = __OPENCL_MEMORY_SCOPE_WORK_GROUP,
__memory_scope_device = __OPENCL_MEMORY_SCOPE_DEVICE,
__memory_scope_all_svm_devices = __OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES,
__memory_scope_sub_group = __OPENCL_MEMORY_SCOPE_SUB_GROUP
} __memory_scope;
// enum values aligned with what clang uses in EmitAtomicExpr()
typedef enum __memory_order
{
__memory_order_relaxed = __ATOMIC_RELAXED,
__memory_order_acquire = __ATOMIC_ACQUIRE,
__memory_order_release = __ATOMIC_RELEASE,
__memory_order_acq_rel = __ATOMIC_ACQ_REL,
__memory_order_seq_cst = __ATOMIC_SEQ_CST
} __memory_order;
__device__
inline
static void
__atomic_work_item_fence(__cl_mem_fence_flags flags, __memory_order order, __memory_scope scope)
{
// We're tying global-happens-before and local-happens-before together as does HSA
if (order != __memory_order_relaxed) {
switch (scope) {
case __memory_scope_work_item:
break;
case __memory_scope_sub_group:
switch (order) {
case __memory_order_relaxed: break;
case __memory_order_acquire: __llvm_fence_acq_sg(); break;
case __memory_order_release: __llvm_fence_rel_sg(); break;
case __memory_order_acq_rel: __llvm_fence_ar_sg(); break;
case __memory_order_seq_cst: __llvm_fence_sc_sg(); break;
}
break;
case __memory_scope_work_group:
switch (order) {
case __memory_order_relaxed: break;
case __memory_order_acquire: __llvm_fence_acq_wg(); break;
case __memory_order_release: __llvm_fence_rel_wg(); break;
case __memory_order_acq_rel: __llvm_fence_ar_wg(); break;
case __memory_order_seq_cst: __llvm_fence_sc_wg(); break;
}
break;
case __memory_scope_device:
switch (order) {
case __memory_order_relaxed: break;
case __memory_order_acquire: __llvm_fence_acq_dev(); break;
case __memory_order_release: __llvm_fence_rel_dev(); break;
case __memory_order_acq_rel: __llvm_fence_ar_dev(); break;
case __memory_order_seq_cst: __llvm_fence_sc_dev(); break;
}
break;
case __memory_scope_all_svm_devices:
switch (order) {
case __memory_order_relaxed: break;
case __memory_order_acquire: __llvm_fence_acq_sys(); break;
case __memory_order_release: __llvm_fence_rel_sys(); break;
case __memory_order_acq_rel: __llvm_fence_ar_sys(); break;
case __memory_order_seq_cst: __llvm_fence_sc_sys(); break;
}
break;
}
}
}
// Memory Fence Functions
__device__
inline
include/hip/hcc_detail/device_library_decls.h
+22 -18
View File
@@ -65,26 +65,30 @@ extern "C" __device__ __attribute__((const)) float __ocml_fmax_f32(float, float)
__device__ inline static __local void* __to_local(unsigned x) { return (__local void*)x; }
#endif //__HIP_DEVICE_COMPILE__
// __llvm_fence* functions from device-libs/irif/src/fence.ll
extern "C" __device__ void __llvm_fence_acq_sg(void);
extern "C" __device__ void __llvm_fence_acq_wg(void);
extern "C" __device__ void __llvm_fence_acq_dev(void);
extern "C" __device__ void __llvm_fence_acq_sys(void);
// Using hip.amdgcn.bc - sync threads
#define __CLK_LOCAL_MEM_FENCE 0x01
typedef unsigned __cl_mem_fence_flags;
extern "C" __device__ void __llvm_fence_rel_sg(void);
extern "C" __device__ void __llvm_fence_rel_wg(void);
extern "C" __device__ void __llvm_fence_rel_dev(void);
extern "C" __device__ void __llvm_fence_rel_sys(void);
typedef enum __memory_scope {
__memory_scope_work_item = __OPENCL_MEMORY_SCOPE_WORK_ITEM,
__memory_scope_work_group = __OPENCL_MEMORY_SCOPE_WORK_GROUP,
__memory_scope_device = __OPENCL_MEMORY_SCOPE_DEVICE,
__memory_scope_all_svm_devices = __OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES,
__memory_scope_sub_group = __OPENCL_MEMORY_SCOPE_SUB_GROUP
} __memory_scope;
extern "C" __device__ void __llvm_fence_ar_sg(void);
extern "C" __device__ void __llvm_fence_ar_wg(void);
extern "C" __device__ void __llvm_fence_ar_dev(void);
extern "C" __device__ void __llvm_fence_ar_sys(void);
// enum values aligned with what clang uses in EmitAtomicExpr()
typedef enum __memory_order
{
__memory_order_relaxed = __ATOMIC_RELAXED,
__memory_order_acquire = __ATOMIC_ACQUIRE,
__memory_order_release = __ATOMIC_RELEASE,
__memory_order_acq_rel = __ATOMIC_ACQ_REL,
__memory_order_seq_cst = __ATOMIC_SEQ_CST
} __memory_order;
extern "C" __device__ void __llvm_fence_sc_sg(void);
extern "C" __device__ void __llvm_fence_sc_wg(void);
extern "C" __device__ void __llvm_fence_sc_dev(void);
extern "C" __device__ void __llvm_fence_sc_sys(void);
// Linked from hip.amdgcn.bc
extern "C" __device__ void
__atomic_work_item_fence(__cl_mem_fence_flags, __memory_order, __memory_scope);
#endif
include/hip/hcc_detail/hip_runtime_api.h
+62 -9
View File
@@ -590,8 +590,6 @@ const char* hipGetErrorString(hipError_t hipError);
*
* The following Stream APIs are not (yet) supported in HIP:
* - cudaStreamAttachMemAsync
* - cudaStreamCreateWithPriority
* - cudaStreamGetPriority
*/
@@ -609,7 +607,7 @@ const char* hipGetErrorString(hipError_t hipError);
*
* @return #hipSuccess, #hipErrorInvalidValue
*
* @see hipStreamCreateWithFlags, hipStreamSynchronize, hipStreamWaitEvent, hipStreamDestroy
* @see hipStreamCreateWithFlags, hipStreamCreateWithPriority, hipStreamSynchronize, hipStreamWaitEvent, hipStreamDestroy
*/
hipError_t hipStreamCreate(hipStream_t* stream);
@@ -628,12 +626,50 @@ hipError_t hipStreamCreate(hipStream_t* stream);
* stream. See #hipStreamDefault, #hipStreamNonBlocking.
*
*
* @see hipStreamCreate, hipStreamSynchronize, hipStreamWaitEvent, hipStreamDestroy
* @see hipStreamCreate, hipStreamCreateWithPriority, hipStreamSynchronize, hipStreamWaitEvent, hipStreamDestroy
*/
hipError_t hipStreamCreateWithFlags(hipStream_t* stream, unsigned int flags);
/**
* @brief Create an asynchronous stream with the specified priority.
*
* @param[in, out] stream Pointer to new stream
* @param[in ] flags to control stream creation.
* @param[in ] priority of the stream. Lower numbers represent higher priorities.
* @return #hipSuccess, #hipErrorInvalidValue
*
* Create a new asynchronous stream with the specified priority. @p stream returns an opaque handle
* that can be used to reference the newly created stream in subsequent hipStream* commands. The
* stream is allocated on the heap and will remain allocated even if the handle goes out-of-scope.
* To release the memory used by the stream, applicaiton must call hipStreamDestroy. Flags controls
* behavior of the stream. See #hipStreamDefault, #hipStreamNonBlocking.
*
*
* @see hipStreamCreate, hipStreamSynchronize, hipStreamWaitEvent, hipStreamDestroy
*/
hipError_t hipStreamCreateWithPriority(hipStream_t* stream, unsigned int flags, int priority);
/**
* @brief Returns numerical values that correspond to the least and greatest stream priority.
*
* @param[in, out] leastPriority pointer in which value corresponding to least priority is returned.
* @param[in, out] greatestPriority pointer in which value corresponding to greatest priority is returned.
*
* Returns in *leastPriority and *greatestPriority the numerical values that correspond to the least
* and greatest stream priority respectively. Stream priorities follow a convention where lower numbers
* imply greater priorities. The range of meaningful stream priorities is given by
* [*greatestPriority, *leastPriority]. If the user attempts to create a stream with a priority value
* that is outside the the meaningful range as specified by this API, the priority is automatically
* clamped to within the valid range.
*/
hipError_t hipDeviceGetStreamPriorityRange(int* leastPriority, int* greatestPriority);
/**
* @brief Destroys the specified stream.
*
@@ -649,7 +685,7 @@ hipError_t hipStreamCreateWithFlags(hipStream_t* stream, unsigned int flags);
* The queue may be destroyed while some commands are still inflight, or may wait for all commands
* queued to the stream before destroying it.
*
* @see hipStreamCreate, hipStreamCreateWithFlags, hipStreamQuery, hipStreamWaitEvent,
* @see hipStreamCreate, hipStreamCreateWithFlags, hipStreamCreateWithPriority, hipStreamQuery, hipStreamWaitEvent,
* hipStreamSynchronize
*/
hipError_t hipStreamDestroy(hipStream_t stream);
@@ -667,7 +703,7 @@ hipError_t hipStreamDestroy(hipStream_t stream);
* host threads are sending work to the stream, the status may change immediately after the function
* is called. It is typically used for debug.
*
* @see hipStreamCreate, hipStreamCreateWithFlags, hipStreamWaitEvent, hipStreamSynchronize,
* @see hipStreamCreate, hipStreamCreateWithFlags, hipStreamCreateWithPriority, hipStreamWaitEvent, hipStreamSynchronize,
* hipStreamDestroy
*/
hipError_t hipStreamQuery(hipStream_t stream);
@@ -689,7 +725,7 @@ hipError_t hipStreamQuery(hipStream_t stream);
* This command honors the hipDeviceLaunchBlocking flag, which controls whether the wait is active
* or blocking.
*
* @see hipStreamCreate, hipStreamCreateWithFlags, hipStreamWaitEvent, hipStreamDestroy
* @see hipStreamCreate, hipStreamCreateWithFlags, hipStreamCreateWithPriority, hipStreamWaitEvent, hipStreamDestroy
*
*/
hipError_t hipStreamSynchronize(hipStream_t stream);
@@ -712,7 +748,7 @@ hipError_t hipStreamSynchronize(hipStream_t stream);
* does not impliciy wait for commands in the default stream to complete, even if the specified
* stream is created with hipStreamNonBlocking = 0.
*
* @see hipStreamCreate, hipStreamCreateWithFlags, hipStreamSynchronize, hipStreamDestroy
* @see hipStreamCreate, hipStreamCreateWithFlags, hipStreamCreateWithPriority, hipStreamSynchronize, hipStreamDestroy
*/
hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags);
@@ -732,6 +768,23 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int
*/
hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int* flags);
/**
* @brief Query the priority of a stream.
*
* @param[in] stream stream to be queried
* @param[in,out] priority Pointer to an unsigned integer in which the stream's priority is returned
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidResourceHandle
*
* @returns #hipSuccess #hipErrorInvalidValue #hipErrorInvalidResourceHandle
*
* Query the priority of a stream. The priority is returned in in priority.
*
* @see hipStreamCreateWithFlags
*/
hipError_t hipStreamGetPriority(hipStream_t stream, int* priority);
/**
* Stream CallBack struct
*/
@@ -749,7 +802,7 @@ typedef void (*hipStreamCallback_t)(hipStream_t stream, hipError_t status, void*
* @return #hipSuccess, #hipErrorInvalidResourceHandle, #hipErrorNotSupported
*
* @see hipStreamCreate, hipStreamCreateWithFlags, hipStreamQuery, hipStreamSynchronize,
* hipStreamWaitEvent, hipStreamDestroy
* hipStreamWaitEvent, hipStreamDestroy, hipStreamCreateWithPriority
*
*/
hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback, void* userData,
include/hip/hcc_detail/program_state.hpp
+5 -4
View File
@@ -93,11 +93,12 @@ public:
}
};
const std::unordered_map<hsa_agent_t, std::vector<hsa_executable_t>>& executables();
const std::unordered_map<hsa_agent_t, std::vector<hsa_executable_t>>& executables(
bool rebuild = false);
const std::unordered_map<std::uintptr_t, std::vector<std::pair<hsa_agent_t, Kernel_descriptor>>>&
functions();
const std::unordered_map<std::uintptr_t, std::string>& function_names();
std::unordered_map<std::string, void*>& globals();
functions(bool rebuild = false);
const std::unordered_map<std::uintptr_t, std::string>& function_names(bool rebuild = false);
std::unordered_map<std::string, void*>& globals(bool rebuild = false);
hsa_executable_t load_executable(const std::string& file, hsa_executable_t executable,
hsa_agent_t agent);
include/hip/nvcc_detail/hip_runtime_api.h
+16
View File
@@ -105,8 +105,10 @@ typedef enum hipChannelFormatKind {
#define hipHostMallocCoherent 0x0
#define hipHostMallocNonCoherent 0x0
#define hipHostRegisterDefault cudaHostRegisterDefault
#define hipHostRegisterPortable cudaHostRegisterPortable
#define hipHostRegisterMapped cudaHostRegisterMapped
#define hipHostRegisterIoMemory cudaHostRegisterIoMemory
#define HIP_LAUNCH_PARAM_BUFFER_POINTER CU_LAUNCH_PARAM_BUFFER_POINTER
#define HIP_LAUNCH_PARAM_BUFFER_SIZE CU_LAUNCH_PARAM_BUFFER_SIZE
@@ -890,6 +892,13 @@ inline static hipError_t hipStreamCreateWithFlags(hipStream_t* stream, unsigned
return hipCUDAErrorTohipError(cudaStreamCreateWithFlags(stream, flags));
}
inline static hipError_t hipStreamCreateWithPriority(hipStream_t* stream, unsigned int flags, int priority) {
return hipCUDAErrorTohipError(cudaStreamCreateWithPriority(stream, flags, priority));
}
inline static hipError_t hipDeviceGetStreamPriorityRange(int* leastPriority, int* greatestPriority) {
return hipCUDAErrorTohipError(cudaDeviceGetStreamPriorityRange(leastPriority, greatestPriority));
}
inline static hipError_t hipStreamCreate(hipStream_t* stream) {
return hipCUDAErrorTohipError(cudaStreamCreate(stream));
@@ -903,6 +912,13 @@ inline static hipError_t hipStreamDestroy(hipStream_t stream) {
return hipCUDAErrorTohipError(cudaStreamDestroy(stream));
}
inline static hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int *flags) {
return hipCUDAErrorTohipError(cudaStreamGetFlags(stream, flags));
}
inline static hipError_t hipStreamGetPriority(hipStream_t stream, int *priority) {
return hipCUDAErrorTohipError(cudaStreamGetPriority(stream, priority));
}
inline static hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event,
unsigned int flags) {