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

fix merging conflicts

Browse Source 
[ROCm/hip commit: fb3b11774b]
This commit is contained in:
Peng Sun
2016-02-08 15:35:49 -06:00
parent 8347c2163d 473a40af80
commit f28e3eedc5
18 changed files with 441 additions and 109 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/include/hcc_detail/hip_runtime.h
+10 -1
View File
@@ -325,6 +325,16 @@ __device__ inline unsigned int __ffsll(unsigned long long int input)
return hc::__lastbit_u32_u64( input)+1;
}
__device__ inline unsigned int __ffs(int input)
{
return hc::__lastbit_u32_s32( input)+1;
}
__device__ inline unsigned int __ffsll(long long int input)
{
return hc::__lastbit_u32_s64( input)+1;
}
__device__ inline unsigned int __brev( unsigned int input)
{
return hc::__bitrev_b32( input);
@@ -336,7 +346,6 @@ __device__ inline unsigned long long int __brevll( unsigned long long int input)
}
// warp vote function __all __any __ballot
__device__ inline int __all( int input)
{
return hc::__all( input);
projects/hip/include/hcc_detail/hip_runtime_api.h
+13 -8
View File
@@ -230,6 +230,13 @@ hipError_t hipGetDevice(int *device);
*/
hipError_t hipGetDeviceCount(int *count);
/**
* @brief Query device attribute.
* @param [out] pi pointer to value to return
* @param [in] attr attribute to query
* @param [in] device which device to query for information
*/
hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device);
/**
* @brief Returns device properties.
@@ -687,11 +694,11 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind
/**
* @brief Copies count bytes from the memory area pointed to by src to the memory area pointed to by offset bytes from the start of symbol symbol.
* @brief Copies @p sizeBytes bytes from the memory area pointed to by @p src to the memory area pointed to by @p offset bytes from the start of symbol @p symbol.
*
* The memory areas may not overlap. Symbol can either be a variable that resides in global or constant memory space, or it can be a character string,
* naming a variable that resides in global or constant memory space. Kind can be either hipMemcpyHostToDevice or hipMemcpyDeviceToDevice
* TODO: cudaErrorInvalidSymbol and cudaErrorInvalidMemcpyDirection is not supported, use hipErrorUnknown for now
* TODO: cudaErrorInvalidSymbol and cudaErrorInvalidMemcpyDirection is not supported, use hipErrorUnknown for now.
*
* @param[in] symbolName - Symbol destination on device
* @param[in] src - Data being copy from
@@ -706,8 +713,6 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t siz
/**
* @brief Copy data from src to dst asynchronously.
*
* It supports memory from host to device,
* device to host, device to device and host to host.
* TODO: cudaErrorInvalidMemcpyDirection error code is not supported right now, use hipErrorUnknown for now
*
* @param[out] dst Data being copy to
@@ -735,7 +740,7 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes );
/**
* @brief Fills the first count bytes of the memory area pointed to by dev with the constant byte value value.
* @brief Fills the first sizeBytes bytes of the memory area pointed to by dev with the constant byte value value.
*
* hipMemsetAsync() is asynchronous with respect to the host, so the call may return before the memset is complete.
* The operation can optionally be associated to a stream by passing a non-zero stream argument.
@@ -743,7 +748,7 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes );
*
* @param[out] dst Pointer to device memory
* @param[in] value - Value to set for each byte of specified memory
* @param[in] count - Size in bytes to set
* @param[in] sizeBytes - Size in bytes to set
* @param[in] stream - Stream identifier
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree
*/
@@ -817,7 +822,7 @@ hipError_t hipDeviceEnablePeerAccess ( int peerDevice, unsigned int flags );
* @param [in] dstDevice - Destination device
* @param [in] src - Source device pointer
* @param [in] srcDevice - Source device
* @param [in] count - Size of memory copy in bytes
* @param [in] sizeBytes - Size of memory copy in bytes
*
* Returns #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidDevice
*/
@@ -830,7 +835,7 @@ hipError_t hipMemcpyPeer ( void* dst, int dstDevice, const void* src, int srcD
* @param [in] dstDevice - Destination device
* @param [in] src - Source device pointer
* @param [in] srcDevice - Source device
* @param [in] count - Size of memory copy in bytes
* @param [in] sizeBytes - Size of memory copy in bytes
* @param [in] stream - Stream identifier
*
* Returns #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidDevice
projects/hip/include/hip_runtime_api.h
+25 -1
View File
@@ -125,7 +125,31 @@ typedef enum hipError_t {
,hipErrorTbd ///< Marker that more error codes are needed.
} hipError_t;
/*
* @brief hipDeviceAttribute_t
* @enum
* @ingroup Enumerations
*/
typedef enum hipDeviceAttribute_t {
hipDeviceAttributeMaxThreadsPerBlock, ///< Maximum number of threads per block.
hipDeviceAttributeMaxBlockDimX, ///< Maximum x-dimension of a block.
hipDeviceAttributeMaxBlockDimY, ///< Maximum y-dimension of a block.
hipDeviceAttributeMaxBlockDimZ, ///< Maximum z-dimension of a block.
hipDeviceAttributeMaxGridDimX, ///< Maximum x-dimension of a grid.
hipDeviceAttributeMaxGridDimY, ///< Maximum y-dimension of a grid.
hipDeviceAttributeMaxGridDimZ, ///< Maximum z-dimension of a grid.
hipDeviceAttributeMaxSharedMemoryPerBlock, ///< Maximum shared memory available per block in bytes.
hipDeviceAttributeTotalConstantMemory, ///< Constant memory size in bytes.
hipDeviceAttributeWarpSize, ///< Warp size in threads.
hipDeviceAttributeMaxRegistersPerBlock, ///< Maximum number of 32-bit registers available to a thread block. This number is shared by all thread blocks simultaneously resident on a multiprocessor.
hipDeviceAttributeClockRate, ///< Peak clock frequency in kilohertz.
hipDeviceAttributeMultiprocessorCount, ///< Number of multiprocessors on the device.
hipDeviceAttributeComputeMode, ///< Compute mode that device is currently in.
hipDeviceAttributeL2CacheSize, ///< Size of L2 cache in bytes. 0 if the device doesn't have L2 cache.
hipDeviceAttributeMaxThreadsPerMultiProcessor, ///< Maximum resident threads per multiprocessor.
hipDeviceAttributeComputeCapabilityMajor, ///< Major compute capability version number.
hipDeviceAttributeComputeCapabilityMinor, ///< Minor compute capability version number.
} hipDeviceAttribute_t;
/**
* @}
projects/hip/include/nvcc_detail/hip_runtime_api.h
+51
View File
@@ -211,6 +211,57 @@ inline static hipError_t hipDeviceGetProperties(hipDeviceProp_t *p_prop, int dev
return hipCUDAErrorTohipError(cerror);
}
inline static hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device)
{
cudaDeviceAttr cdattr;
cudaError_t cerror;
switch (attr) {
case hipDeviceAttributeMaxThreadsPerBlock:
cdattr = cudaDevAttrMaxThreadsPerBlock; break;
case hipDeviceAttributeMaxBlockDimX:
cdattr = cudaDevAttrMaxBlockDimX; break;
case hipDeviceAttributeMaxBlockDimY:
cdattr = cudaDevAttrMaxBlockDimY; break;
case hipDeviceAttributeMaxBlockDimZ:
cdattr = cudaDevAttrMaxBlockDimZ; break;
case hipDeviceAttributeMaxGridDimX:
cdattr = cudaDevAttrMaxGridDimX; break;
case hipDeviceAttributeMaxGridDimY:
cdattr = cudaDevAttrMaxGridDimY; break;
case hipDeviceAttributeMaxGridDimZ:
cdattr = cudaDevAttrMaxGridDimZ; break;
case hipDeviceAttributeMaxSharedMemoryPerBlock:
cdattr = cudaDevAttrMaxSharedMemoryPerBlock; break;
case hipDeviceAttributeTotalConstantMemory:
cdattr = cudaDevAttrTotalConstantMemory; break;
case hipDeviceAttributeWarpSize:
cdattr = cudaDevAttrWarpSize; break;
case hipDeviceAttributeMaxRegistersPerBlock:
cdattr = cudaDevAttrMaxRegistersPerBlock; break;
case hipDeviceAttributeClockRate:
cdattr = cudaDevAttrClockRate; break;
case hipDeviceAttributeMultiprocessorCount:
cdattr = cudaDevAttrMultiProcessorCount; break;
case hipDeviceAttributeComputeMode:
cdattr = cudaDevAttrComputeMode; break;
case hipDeviceAttributeL2CacheSize:
cdattr = cudaDevAttrL2CacheSize; break;
case hipDeviceAttributeMaxThreadsPerMultiProcessor:
cdattr = cudaDevAttrMaxThreadsPerMultiProcessor; break;
case hipDeviceAttributeComputeCapabilityMajor:
cdattr = cudaDevAttrComputeCapabilityMajor; break;
case hipDeviceAttributeComputeCapabilityMinor:
cdattr = cudaDevAttrComputeCapabilityMinor; break;
default:
cerror = cudaErrorInvalidValue; break;
}
cerror = cudaDeviceGetAttribute(pi, cdattr, device);
return hipCUDAErrorTohipError(cerror);
}
inline static hipError_t hipMemGetInfo( size_t* free, size_t* total)
{
return hipCUDAErrorTohipError(cudaMemGetInfo(free,total));