Merge branch 'amd-develop' into amd-master
Change-Id: I53d5a8916d769c4f0fe60d2ee3b240551da80b4f
(cherry picked from commit 01c523f6c9)
This commit is contained in:
Maneesh Gupta
zatwierdzone przez
Maneesh Gupta
Maneesh Gupta
rodzic
5366c23fea
commit
bb8c51a129
@@ -21,6 +21,7 @@ THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
#if GENERIC_GRID_LAUNCH == 1
|
||||
|
||||
#include "concepts.hpp"
|
||||
#include "helpers.hpp"
|
||||
@@ -840,14 +841,16 @@ namespace hip_impl
|
||||
group_mem_bytes,\
|
||||
stream,\
|
||||
...)\
|
||||
{\
|
||||
hipLaunchKernelGGL(\
|
||||
kernel_name,\
|
||||
num_blocks,\
|
||||
dim_blocks,\
|
||||
group_mem_bytes,\
|
||||
stream,\
|
||||
hipLaunchParm{},\
|
||||
##__VA_ARGS__);\
|
||||
}
|
||||
do {\
|
||||
hipLaunchKernelGGL(\
|
||||
kernel_name,\
|
||||
num_blocks,\
|
||||
dim_blocks,\
|
||||
group_mem_bytes,\
|
||||
stream,\
|
||||
hipLaunchParm{},\
|
||||
##__VA_ARGS__);\
|
||||
} while(0)
|
||||
|
||||
}
|
||||
#endif //GENERIC_GRID_LAUNCH
|
||||
|
||||
@@ -23,8 +23,7 @@ THE SOFTWARE.
|
||||
#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COMPLEX_H
|
||||
#define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_COMPLEX_H
|
||||
|
||||
#include "./hip_fp16.h"
|
||||
#include "./hip_vector_types.h"
|
||||
#include "hip/hcc_detail/hip_vector_types.h"
|
||||
|
||||
#if __cplusplus
|
||||
#define COMPLEX_ADD_OP_OVERLOAD(type) \
|
||||
@@ -177,45 +176,45 @@ COMPLEX_SCALAR_PRODUCT(hipDoubleComplex, unsigned long long)
|
||||
|
||||
#endif
|
||||
|
||||
__device__ static inline float hipCrealf(hipFloatComplex z){
|
||||
__device__ __host__ static inline float hipCrealf(hipFloatComplex z){
|
||||
return z.x;
|
||||
}
|
||||
|
||||
__device__ static inline float hipCimagf(hipFloatComplex z){
|
||||
__device__ __host__ static inline float hipCimagf(hipFloatComplex z){
|
||||
return z.y;
|
||||
}
|
||||
|
||||
__device__ static inline hipFloatComplex make_hipFloatComplex(float a, float b){
|
||||
__device__ __host__ static inline hipFloatComplex make_hipFloatComplex(float a, float b){
|
||||
hipFloatComplex z;
|
||||
z.x = a;
|
||||
z.y = b;
|
||||
return z;
|
||||
}
|
||||
|
||||
__device__ static inline hipFloatComplex hipConjf(hipFloatComplex z){
|
||||
__device__ __host__ static inline hipFloatComplex hipConjf(hipFloatComplex z){
|
||||
hipFloatComplex ret;
|
||||
ret.x = z.x;
|
||||
ret.y = -z.y;
|
||||
return ret;
|
||||
}
|
||||
|
||||
__device__ static inline float hipCsqabsf(hipFloatComplex z){
|
||||
__device__ __host__ static inline float hipCsqabsf(hipFloatComplex z){
|
||||
return z.x * z.x + z.y * z.y;
|
||||
}
|
||||
|
||||
__device__ static inline hipFloatComplex hipCaddf(hipFloatComplex p, hipFloatComplex q){
|
||||
__device__ __host__ static inline hipFloatComplex hipCaddf(hipFloatComplex p, hipFloatComplex q){
|
||||
return make_hipFloatComplex(p.x + q.x, p.y + q.y);
|
||||
}
|
||||
|
||||
__device__ static inline hipFloatComplex hipCsubf(hipFloatComplex p, hipFloatComplex q){
|
||||
__device__ __host__ static inline hipFloatComplex hipCsubf(hipFloatComplex p, hipFloatComplex q){
|
||||
return make_hipFloatComplex(p.x - q.x, p.y - q.y);
|
||||
}
|
||||
|
||||
__device__ static inline hipFloatComplex hipCmulf(hipFloatComplex p, hipFloatComplex q){
|
||||
__device__ __host__ static inline hipFloatComplex hipCmulf(hipFloatComplex p, hipFloatComplex q){
|
||||
return make_hipFloatComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y);
|
||||
}
|
||||
|
||||
__device__ static inline hipFloatComplex hipCdivf(hipFloatComplex p, hipFloatComplex q){
|
||||
__device__ __host__ static inline hipFloatComplex hipCdivf(hipFloatComplex p, hipFloatComplex q){
|
||||
float sqabs = hipCsqabsf(q);
|
||||
hipFloatComplex ret;
|
||||
ret.x = (p.x * q.x + p.y * q.y)/sqabs;
|
||||
@@ -223,51 +222,51 @@ __device__ static inline hipFloatComplex hipCdivf(hipFloatComplex p, hipFloatCom
|
||||
return ret;
|
||||
}
|
||||
|
||||
__device__ static inline float hipCabsf(hipFloatComplex z){
|
||||
__device__ __host__ static inline float hipCabsf(hipFloatComplex z){
|
||||
return sqrtf(hipCsqabsf(z));
|
||||
}
|
||||
|
||||
|
||||
|
||||
__device__ static inline double hipCreal(hipDoubleComplex z){
|
||||
__device__ __host__ static inline double hipCreal(hipDoubleComplex z){
|
||||
return z.x;
|
||||
}
|
||||
|
||||
__device__ static inline double hipCimag(hipDoubleComplex z){
|
||||
__device__ __host__ static inline double hipCimag(hipDoubleComplex z){
|
||||
return z.y;
|
||||
}
|
||||
|
||||
__device__ static inline hipDoubleComplex make_hipDoubleComplex(double a, double b){
|
||||
__device__ __host__ static inline hipDoubleComplex make_hipDoubleComplex(double a, double b){
|
||||
hipDoubleComplex z;
|
||||
z.x = a;
|
||||
z.y = b;
|
||||
return z;
|
||||
}
|
||||
|
||||
__device__ static inline hipDoubleComplex hipConj(hipDoubleComplex z){
|
||||
__device__ __host__ static inline hipDoubleComplex hipConj(hipDoubleComplex z){
|
||||
hipDoubleComplex ret;
|
||||
ret.x = z.x;
|
||||
ret.y = z.y;
|
||||
return ret;
|
||||
}
|
||||
|
||||
__device__ static inline double hipCsqabs(hipDoubleComplex z){
|
||||
__device__ __host__ static inline double hipCsqabs(hipDoubleComplex z){
|
||||
return z.x * z.x + z.y * z.y;
|
||||
}
|
||||
|
||||
__device__ static inline hipDoubleComplex hipCadd(hipDoubleComplex p, hipDoubleComplex q){
|
||||
__device__ __host__ static inline hipDoubleComplex hipCadd(hipDoubleComplex p, hipDoubleComplex q){
|
||||
return make_hipDoubleComplex(p.x + q.x, p.y + q.y);
|
||||
}
|
||||
|
||||
__device__ static inline hipDoubleComplex hipCsub(hipDoubleComplex p, hipDoubleComplex q){
|
||||
__device__ __host__ static inline hipDoubleComplex hipCsub(hipDoubleComplex p, hipDoubleComplex q){
|
||||
return make_hipDoubleComplex(p.x - q.x, p.y - q.y);
|
||||
}
|
||||
|
||||
__device__ static inline hipDoubleComplex hipCmul(hipDoubleComplex p, hipDoubleComplex q){
|
||||
__device__ __host__ static inline hipDoubleComplex hipCmul(hipDoubleComplex p, hipDoubleComplex q){
|
||||
return make_hipDoubleComplex(p.x * q.x - p.y * q.y, p.y * q.x + p.x * q.y);
|
||||
}
|
||||
|
||||
__device__ static inline hipDoubleComplex hipCdiv(hipDoubleComplex p, hipDoubleComplex q){
|
||||
__device__ __host__ static inline hipDoubleComplex hipCdiv(hipDoubleComplex p, hipDoubleComplex q){
|
||||
double sqabs = hipCsqabs(q);
|
||||
hipDoubleComplex ret;
|
||||
ret.x = (p.x * q.x + p.y * q.y)/sqabs;
|
||||
@@ -275,28 +274,28 @@ __device__ static inline hipDoubleComplex hipCdiv(hipDoubleComplex p, hipDoubleC
|
||||
return ret;
|
||||
}
|
||||
|
||||
__device__ static inline double hipCabs(hipDoubleComplex z){
|
||||
__device__ __host__ static inline double hipCabs(hipDoubleComplex z){
|
||||
return sqrtf(hipCsqabs(z));
|
||||
}
|
||||
|
||||
typedef hipFloatComplex hipComplex;
|
||||
|
||||
__device__ static inline hipComplex make_hipComplex(float x,
|
||||
__device__ __host__ static inline hipComplex make_hipComplex(float x,
|
||||
float y){
|
||||
return make_hipFloatComplex(x, y);
|
||||
}
|
||||
|
||||
__device__ static inline hipFloatComplex hipComplexDoubleToFloat
|
||||
__device__ __host__ static inline hipFloatComplex hipComplexDoubleToFloat
|
||||
(hipDoubleComplex z){
|
||||
return make_hipFloatComplex((float)z.x, (float)z.y);
|
||||
}
|
||||
|
||||
__device__ static inline hipDoubleComplex hipComplexFloatToDouble
|
||||
__device__ __host__ static inline hipDoubleComplex hipComplexFloatToDouble
|
||||
(hipFloatComplex z){
|
||||
return make_hipDoubleComplex((double)z.x, (double)z.y);
|
||||
}
|
||||
|
||||
__device__ static inline hipComplex hipCfmaf(hipComplex p, hipComplex q, hipComplex r){
|
||||
__device__ __host__ static inline hipComplex hipCfmaf(hipComplex p, hipComplex q, hipComplex r){
|
||||
float real = (p.x * q.x) + r.x;
|
||||
float imag = (q.x * p.y) + r.y;
|
||||
|
||||
@@ -306,7 +305,7 @@ __device__ static inline hipComplex hipCfmaf(hipComplex p, hipComplex q, hipComp
|
||||
return make_hipComplex(real, imag);
|
||||
}
|
||||
|
||||
__device__ static inline hipDoubleComplex hipCfma(hipDoubleComplex p, hipDoubleComplex q, hipDoubleComplex r){
|
||||
__device__ __host__ static inline hipDoubleComplex hipCfma(hipDoubleComplex p, hipDoubleComplex q, hipDoubleComplex r){
|
||||
float real = (p.x * q.x) + r.x;
|
||||
float imag = (q.x * p.y) + r.y;
|
||||
|
||||
|
||||
@@ -23,7 +23,7 @@ THE SOFTWARE.
|
||||
#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_FP16_H
|
||||
#define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_FP16_H
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "hip/hcc_detail/hip_vector_types.h"
|
||||
|
||||
#if __clang_major__ > 3
|
||||
|
||||
|
||||
@@ -28,6 +28,17 @@ THE SOFTWARE.
|
||||
#if __cplusplus
|
||||
#ifdef __HCC__
|
||||
#include <hc.hpp>
|
||||
|
||||
|
||||
/**
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
* @defgroup HCC-specific features
|
||||
* @warning These APIs provide access to special features of HCC compiler and are not available through the CUDA path.
|
||||
* @{
|
||||
*/
|
||||
|
||||
|
||||
/**
|
||||
* @brief Return hc::accelerator associated with the specified deviceId
|
||||
* @return #hipSuccess, #hipErrorInvalidDevice
|
||||
@@ -45,6 +56,29 @@ hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view **a
|
||||
|
||||
#endif // #ifdef __HCC__
|
||||
|
||||
/**
|
||||
* @brief launches kernel f with launch parameters and shared memory on stream with arguments passed to kernelparams or extra
|
||||
*
|
||||
* @param [in[ f Kernel to launch.
|
||||
* @param [in] gridDimX X grid dimension specified in work-items
|
||||
* @param [in] gridDimY Y grid dimension specified in work-items
|
||||
* @param [in] gridDimZ Z grid dimension specified in work-items
|
||||
* @param [in] blockDimX X block dimensions specified in work-items
|
||||
* @param [in] blockDimY Y grid dimension specified in work-items
|
||||
* @param [in] blockDimZ Z grid dimension specified in work-items
|
||||
* @param [in] sharedMemBytes Amount of dynamic shared memory to allocate for this kernel. The kernel can access this with HIP_DYNAMIC_SHARED.
|
||||
* @param [in] stream Stream where the kernel should be dispatched. May be 0, in which case th default stream is used with associated synchronization rules.
|
||||
* @param [in] kernelParams
|
||||
* @param [in] extra Pointer to kernel arguments. These are passed directly to the kernel and must be in the memory layout and alignment expected by the kernel.
|
||||
* @param [in] startEvent If non-null, specified event will be updated to track the start time of the kernel launch. The event must be created before calling this API.
|
||||
* @param [in] stopEvent If non-null, specified event will be updated to track the stop time of the kernel launch. The event must be created before calling this API.
|
||||
*
|
||||
* @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue
|
||||
*
|
||||
* @warning kernellParams argument is not yet implemented in HIP. Please use extra instead. Please refer to hip_porting_driver_api.md for sample usage.
|
||||
|
||||
* HIP/ROCm actually updates the start event when the associated kernel completes.
|
||||
*/
|
||||
hipError_t hipHccModuleLaunchKernel(hipFunction_t f,
|
||||
uint32_t globalWorkSizeX,
|
||||
uint32_t globalWorkSizeY,
|
||||
@@ -55,8 +89,15 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f,
|
||||
size_t sharedMemBytes,
|
||||
hipStream_t hStream,
|
||||
void **kernelParams,
|
||||
void **extra);
|
||||
void **extra,
|
||||
hipEvent_t startEvent=nullptr,
|
||||
hipEvent_t stopEvent=nullptr
|
||||
);
|
||||
|
||||
// doxygen end HCC-specific features
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
#endif // #if __cplusplus
|
||||
|
||||
#endif //
|
||||
|
||||
@@ -1913,19 +1913,18 @@ hipError_t hipModuleLoadData(hipModule_t *module, const void *image);
|
||||
/**
|
||||
* @brief launches kernel f with launch parameters and shared memory on stream with arguments passed to kernelparams or extra
|
||||
*
|
||||
* @param [in[ f
|
||||
* @param [in] gridDimX
|
||||
* @param [in] gridDimY
|
||||
* @param [in] gridDimZ
|
||||
* @param [in] blockDimX
|
||||
* @param [in] blockDimY
|
||||
* @param [in] blockDimZ
|
||||
* @param [in] sharedMemBytes
|
||||
* @param [in] stream
|
||||
* @param [in] kernelParams
|
||||
* @param [in] extraa
|
||||
* @param [in[ f Kernel to launch.
|
||||
* @param [in] gridDimX X grid dimension specified as multiple of blockDimX.
|
||||
* @param [in] gridDimY Y grid dimension specified as multiple of blockDimY.
|
||||
* @param [in] gridDimZ Z grid dimension specified as multiple of blockDimZ.
|
||||
* @param [in] blockDimX X block dimensions specified in work-items
|
||||
* @param [in] blockDimY Y grid dimension specified in work-items
|
||||
* @param [in] blockDimZ Z grid dimension specified in work-items
|
||||
* @param [in] sharedMemBytes Amount of dynamic shared memory to allocate for this kernel. The kernel can access this with HIP_DYNAMIC_SHARED.
|
||||
* @param [in] stream Stream where the kernel should be dispatched. May be 0, in which case th default stream is used with associated synchronization rules.
|
||||
* @param [in] kernelParams
|
||||
* @param [in] extra Pointer to kernel arguments. These are passed directly to the kernel and must be in the memory layout and alignment expected by the kernel.
|
||||
*
|
||||
* The function takes the above arguments and run the kernel in hipFunction_t f. with launch parameters specified in gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY and blockDimmZ. The amount of shared memory is specificed and can be used with HIP_DYNAMIC_SHARED. The arguemt extra is used to pass in the arguments for the kernel.
|
||||
* @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue
|
||||
*
|
||||
* @warning kernellParams argument is not yet implemented in HIP. Please use extra instead. Please refer to hip_porting_driver_api.md for sample usage.
|
||||
|
||||
@@ -32,7 +32,7 @@ THE SOFTWARE.
|
||||
#error("This version of HIP requires a newer version of HCC.");
|
||||
#endif
|
||||
|
||||
#include "host_defines.h"
|
||||
#include "hip/hcc_detail/host_defines.h"
|
||||
|
||||
#define MAKE_DEFAULT_CONSTRUCTOR_ONE_COMPONENT(type) \
|
||||
__device__ __host__ type() {} \
|
||||
|
||||
Reference in New Issue
Block a user