Move include/* to include/hip/*
Change-Id: I7a7b2839b4df59c7a4c503550f99fdc9e45c0f54
Αυτή η υποβολή περιλαμβάνεται σε:
Maneesh Gupta
@@ -1 +0,0 @@
|
||||
../include
|
||||
@@ -0,0 +1,8 @@
|
||||
#ifndef HCC_H
|
||||
#define HCC_H
|
||||
|
||||
#if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__)
|
||||
#include "hip/hcc_detail/hcc_acc.h"
|
||||
#endif
|
||||
|
||||
#endif
|
||||
@@ -0,0 +1 @@
|
||||
|
||||
@@ -0,0 +1 @@
|
||||
|
||||
@@ -0,0 +1,24 @@
|
||||
#ifndef HCC_ACC_H
|
||||
#define HCC_ACC_H
|
||||
#include "hip/hip_runtime_api.h"
|
||||
|
||||
#if __cplusplus
|
||||
#ifdef __HCC__
|
||||
#include <hc.hpp>
|
||||
/**
|
||||
* @brief Return hc::accelerator associated with the specified deviceId
|
||||
* @return #hipSuccess, #hipErrorInvalidDevice
|
||||
*/
|
||||
hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator *acc);
|
||||
|
||||
/**
|
||||
* @brief Return hc::accelerator_view associated with the specified stream
|
||||
*
|
||||
* If stream is 0, the accelerator_view for the default stream is returned.
|
||||
* @return #hipSuccess
|
||||
*/
|
||||
hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view **av);
|
||||
#endif
|
||||
#endif
|
||||
|
||||
#endif
|
||||
@@ -0,0 +1,172 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
|
||||
|
||||
#ifndef HIPCOMPLEX_H
|
||||
#define HIPCOMPLEX_H
|
||||
|
||||
typedef struct{
|
||||
float x;
|
||||
float y;
|
||||
}hipFloatComplex;
|
||||
|
||||
__device__ static inline float hipCrealf(hipFloatComplex z){
|
||||
return z.x;
|
||||
}
|
||||
|
||||
__device__ static inline float hipCimagf(hipFloatComplex z){
|
||||
return z.y;
|
||||
}
|
||||
|
||||
__device__ 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){
|
||||
hipFloatComplex ret;
|
||||
ret.x = z.x;
|
||||
ret.y = -z.y;
|
||||
return ret;
|
||||
}
|
||||
|
||||
__device__ static inline float hipCsqabsf(hipFloatComplex z){
|
||||
return z.x * z.x + z.y * z.y;
|
||||
}
|
||||
|
||||
__device__ 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){
|
||||
return make_hipFloatComplex(p.x - q.x, p.y - q.y);
|
||||
}
|
||||
|
||||
__device__ 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){
|
||||
float sqabs = hipCsqabsf(q);
|
||||
hipFloatComplex ret;
|
||||
ret.x = (p.x * q.x + p.y * q.y)/sqabs;
|
||||
ret.y = (p.y * q.x - p.x * q.y)/sqabs;
|
||||
return ret;
|
||||
}
|
||||
|
||||
__device__ static inline float hipCabsf(hipFloatComplex z){
|
||||
return sqrtf(hipCsqabsf(z));
|
||||
}
|
||||
|
||||
|
||||
typedef struct{
|
||||
double x;
|
||||
double y;
|
||||
}hipDoubleComplex;
|
||||
|
||||
__device__ static inline double hipCreal(hipDoubleComplex z){
|
||||
return z.x;
|
||||
}
|
||||
|
||||
__device__ static inline double hipCimag(hipDoubleComplex z){
|
||||
return z.y;
|
||||
}
|
||||
|
||||
__device__ 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){
|
||||
hipDoubleComplex ret;
|
||||
ret.x = z.x;
|
||||
ret.y = z.y;
|
||||
return ret;
|
||||
}
|
||||
|
||||
__device__ static inline double hipCsqabs(hipDoubleComplex z){
|
||||
return z.x * z.x + z.y * z.y;
|
||||
}
|
||||
|
||||
__device__ 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){
|
||||
return make_hipDoubleComplex(p.x - q.x, p.y - q.y);
|
||||
}
|
||||
|
||||
__device__ 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){
|
||||
double sqabs = hipCsqabs(q);
|
||||
hipDoubleComplex ret;
|
||||
ret.x = (p.x * q.x + p.y * q.y)/sqabs;
|
||||
ret.y = (p.y * q.x - p.x * q.y)/sqabs;
|
||||
return ret;
|
||||
}
|
||||
|
||||
__device__ static inline double hipCabs(hipDoubleComplex z){
|
||||
return sqrtf(hipCsqabs(z));
|
||||
}
|
||||
|
||||
typedef hipFloatComplex hipComplex;
|
||||
|
||||
__device__ static inline hipComplex make_hipComplex(float x,
|
||||
float y){
|
||||
return make_hipFloatComplex(x, y);
|
||||
}
|
||||
|
||||
__device__ static inline hipFloatComplex hipComplexDoubleToFloat
|
||||
(hipDoubleComplex z){
|
||||
return make_hipFloatComplex((float)z.x, (float)z.y);
|
||||
}
|
||||
|
||||
__device__ 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){
|
||||
float real = (p.x * q.x) + r.x;
|
||||
float imag = (q.x * p.y) + r.y;
|
||||
|
||||
real = -(p.y * q.y) + real;
|
||||
imag = (p.x * q.y) + imag;
|
||||
|
||||
return make_hipComplex(real, imag);
|
||||
}
|
||||
|
||||
__device__ 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;
|
||||
|
||||
real = -(p.y * q.y) + real;
|
||||
imag = (p.x * q.y) + imag;
|
||||
|
||||
return make_hipDoubleComplex(real, imag);
|
||||
}
|
||||
|
||||
#endif
|
||||
@@ -0,0 +1,258 @@
|
||||
/*
|
||||
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 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.
|
||||
*/
|
||||
#pragma once
|
||||
|
||||
#include <hip_runtime_api.h>
|
||||
#include <hcblas.h>
|
||||
|
||||
//HGSOS for Kalmar leave it as C++, only cublas needs C linkage.
|
||||
|
||||
#ifdef __cplusplus
|
||||
extern "C" {
|
||||
#endif
|
||||
|
||||
typedef hcblasHandle_t* hipblasHandle_t;
|
||||
typedef hcComplex hipComplex ;
|
||||
|
||||
static hipblasHandle_t dummyGlobal;
|
||||
|
||||
/* Unsupported types
|
||||
"cublasFillMode_t",
|
||||
"cublasDiagType_t",
|
||||
"cublasSideMode_t",
|
||||
"cublasPointerMode_t",
|
||||
"cublasAtomicsMode_t",
|
||||
"cublasDataType_t"
|
||||
*/
|
||||
|
||||
inline static hcblasOperation_t hipOperationToHCCOperation( hipblasOperation_t op)
|
||||
{
|
||||
switch (op)
|
||||
{
|
||||
case HIPBLAS_OP_N:
|
||||
return HCBLAS_OP_N;
|
||||
|
||||
case HIPBLAS_OP_T:
|
||||
return HCBLAS_OP_T;
|
||||
|
||||
case HIPBLAS_OP_C:
|
||||
return HCBLAS_OP_C;
|
||||
|
||||
default:
|
||||
throw "Non existent OP";
|
||||
}
|
||||
}
|
||||
|
||||
inline static hipblasOperation_t HCCOperationToHIPOperation( hcblasOperation_t op)
|
||||
{
|
||||
switch (op)
|
||||
{
|
||||
case HCBLAS_OP_N :
|
||||
return HIPBLAS_OP_N;
|
||||
|
||||
case HCBLAS_OP_T :
|
||||
return HIPBLAS_OP_T;
|
||||
|
||||
case HCBLAS_OP_C :
|
||||
return HIPBLAS_OP_C;
|
||||
|
||||
default:
|
||||
throw "Non existent OP";
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
inline static hipblasStatus_t hipHCBLASStatusToHIPStatus(hcblasStatus_t hcStatus)
|
||||
{
|
||||
switch(hcStatus)
|
||||
{
|
||||
case HCBLAS_STATUS_SUCCESS:
|
||||
return HIPBLAS_STATUS_SUCCESS;
|
||||
case HCBLAS_STATUS_NOT_INITIALIZED:
|
||||
return HIPBLAS_STATUS_NOT_INITIALIZED;
|
||||
case HCBLAS_STATUS_ALLOC_FAILED:
|
||||
return HIPBLAS_STATUS_ALLOC_FAILED;
|
||||
case HCBLAS_STATUS_INVALID_VALUE:
|
||||
return HIPBLAS_STATUS_INVALID_VALUE;
|
||||
case HCBLAS_STATUS_MAPPING_ERROR:
|
||||
return HIPBLAS_STATUS_MAPPING_ERROR;
|
||||
case HCBLAS_STATUS_EXECUTION_FAILED:
|
||||
return HIPBLAS_STATUS_EXECUTION_FAILED;
|
||||
case HCBLAS_STATUS_INTERNAL_ERROR:
|
||||
return HIPBLAS_STATUS_INTERNAL_ERROR;
|
||||
default:
|
||||
throw "Unimplemented status";
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
inline static hipblasStatus_t hipblasCreate(hipblasHandle_t* handle) {
|
||||
hipblasStatus_t retVal = hipHCBLASStatusToHIPStatus(hcblasCreate(*handle));
|
||||
dummyGlobal = *handle;
|
||||
return retVal;
|
||||
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasDestroy(hipblasHandle_t& handle) {
|
||||
return hipHCBLASStatusToHIPStatus(hcblasDestroy(handle));
|
||||
}
|
||||
|
||||
//note: no handle
|
||||
inline static hipblasStatus_t hipblasSetVector(int n, int elemSize, const void *x, int incx, void *y, int incy){
|
||||
return hipHCBLASStatusToHIPStatus(hcblasSetVector(dummyGlobal, n, elemSize, x, incx, y, incy)); //HGSOS no need for handle moving forward
|
||||
}
|
||||
|
||||
//note: no handle
|
||||
inline static hipblasStatus_t hipblasGetVector(int n, int elemSize, const void *x, int incx, void *y, int incy){
|
||||
return hipHCBLASStatusToHIPStatus(hcblasGetVector(dummyGlobal, n, elemSize, x, incx, y, incy)); //HGSOS no need for handle
|
||||
}
|
||||
|
||||
//note: no handle
|
||||
inline static hipblasStatus_t hipblasSetMatrix(int rows, int cols, int elemSize, const void *A, int lda, void *B, int ldb){
|
||||
return hipHCBLASStatusToHIPStatus(hcblasSetMatrix(dummyGlobal, rows, cols, elemSize, A, lda, B, ldb));
|
||||
}
|
||||
|
||||
//note: no handle
|
||||
inline static hipblasStatus_t hipblasGetMatrix(int rows, int cols, int elemSize, const void *A, int lda, void *B, int ldb){
|
||||
return hipHCBLASStatusToHIPStatus(hcblasGetMatrix(dummyGlobal, rows, cols, elemSize, A, lda, B, ldb));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasSasum(hipblasHandle_t handle, int n, float *x, int incx, float *result){
|
||||
return hipHCBLASStatusToHIPStatus(hcblasSasum(handle, n, x, incx, result));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasDasum(hipblasHandle_t handle, int n, double *x, int incx, double *result){
|
||||
return hipHCBLASStatusToHIPStatus(hcblasDasum(handle, n, x, incx, result));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasSasumBatched(hipblasHandle_t handle, int n, float *x, int incx, float *result, int batchCount){
|
||||
return hipHCBLASStatusToHIPStatus(hcblasSasumBatched( handle, n, x, incx, result, batchCount));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasDasumBatched(hipblasHandle_t handle, int n, double *x, int incx, double *result, int batchCount){
|
||||
return hipHCBLASStatusToHIPStatus(hcblasDasumBatched(handle, n, x, incx, result, batchCount));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasSaxpy(hipblasHandle_t handle, int n, const float *alpha, const float *x, int incx, float *y, int incy) {
|
||||
return hipHCBLASStatusToHIPStatus(hcblasSaxpy(handle, n, alpha, x, incx, y, incy));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasSaxpyBatched(hipblasHandle_t handle, int n, const float *alpha, const float *x, int incx, float *y, int incy, int batchCount){
|
||||
return hipHCBLASStatusToHIPStatus(hcblasSaxpyBatched(handle, n, alpha, x, incx, y, incy, batchCount));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasScopy(hipblasHandle_t handle, int n, const float *x, int incx, float *y, int incy){
|
||||
return hipHCBLASStatusToHIPStatus(hcblasScopy( handle, n, x, incx, y, incy));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasDcopy(hipblasHandle_t handle, int n, const double *x, int incx, double *y, int incy){
|
||||
return hipHCBLASStatusToHIPStatus(hcblasDcopy( handle, n, x, incx, y, incy));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasScopyBatched(hipblasHandle_t handle, int n, const float *x, int incx, float *y, int incy, int batchCount){
|
||||
return hipHCBLASStatusToHIPStatus(hcblasScopyBatched( handle, n, x, incx, y, incy, batchCount));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasDcopyBatched(hipblasHandle_t handle, int n, const double *x, int incx, double *y, int incy, int batchCount){
|
||||
return hipHCBLASStatusToHIPStatus(hcblasDcopyBatched( handle, n, x, incx, y, incy, batchCount));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasSdot (hipblasHandle_t handle, int n, const float *x, int incx, const float *y, int incy, float *result){
|
||||
return hipHCBLASStatusToHIPStatus(hcblasSdot(handle, n, x, incx, y, incy, result));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasDdot (hipblasHandle_t handle, int n, const double *x, int incx, const double *y, int incy, double *result){
|
||||
return hipHCBLASStatusToHIPStatus(hcblasDdot(handle, n, x, incx, y, incy, result));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasSdotBatched (hipblasHandle_t handle, int n, const float *x, int incx, const float *y, int incy, float *result, int batchCount){
|
||||
return hipHCBLASStatusToHIPStatus(hcblasSdotBatched(handle, n, x, incx, y, incy, result, batchCount));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasDdotBatched (hipblasHandle_t handle, int n, const double *x, int incx, const double *y, int incy, double *result, int batchCount){
|
||||
return hipHCBLASStatusToHIPStatus(hcblasDdotBatched ( handle, n, x, incx, y, incy, result, batchCount));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasSscal(hipblasHandle_t handle, int n, const float *alpha, float *x, int incx){
|
||||
return hipHCBLASStatusToHIPStatus(hcblasSscal(handle, n, alpha, x, incx));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasDscal(hipblasHandle_t handle, int n, const double *alpha, double *x, int incx){
|
||||
return hipHCBLASStatusToHIPStatus(hcblasDscal(handle, n, alpha, x, incx));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasSscalBatched(hipblasHandle_t handle, int n, const float *alpha, float *x, int incx, int batchCount){
|
||||
return hipHCBLASStatusToHIPStatus(hcblasSscalBatched(handle, n, alpha, x, incx, batchCount));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasDscalBatched(hipblasHandle_t handle, int n, const double *alpha, double *x, int incx, int batchCount){
|
||||
return hipHCBLASStatusToHIPStatus(hcblasDscalBatched(handle, n, alpha, x, incx, batchCount));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasSgemv(hipblasHandle_t handle, hipblasOperation_t trans, int m, int n, const float *alpha, float *A, int lda,
|
||||
float *x, int incx, const float *beta, float *y, int incy){
|
||||
return hipHCBLASStatusToHIPStatus(hcblasSgemv(handle, hipOperationToHCCOperation(trans), m, n, alpha, A, lda, x, incx, beta, y, incy));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasSgemvBatched(hipblasHandle_t handle, hipblasOperation_t trans, int m, int n, const float *alpha, float *A, int lda,
|
||||
float *x, int incx, const float *beta, float *y, int incy, int batchCount){
|
||||
return hipHCBLASStatusToHIPStatus(hcblasSgemvBatched(handle, hipOperationToHCCOperation(trans), m, n, alpha, A, lda, x, incx, beta, y, incy, batchCount));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasSger(hipblasHandle_t handle, int m, int n, const float *alpha, const float *x, int incx, const float *y, int incy, float *A, int lda){
|
||||
return hipHCBLASStatusToHIPStatus(hcblasSger(handle, m, n, alpha, x, incx, y, incy, A, lda));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasSgerBatched(hipblasHandle_t handle, int m, int n, const float *alpha, const float *x, int incx, const float *y, int incy, float *A, int lda, int batchCount){
|
||||
return hipHCBLASStatusToHIPStatus(hcblasSgerBatched(handle, m, n, alpha, x, incx, y, incy, A, lda, batchCount));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasSgemm(hipblasHandle_t handle, hipblasOperation_t transa, hipblasOperation_t transb,
|
||||
int m, int n, int k, const float *alpha, float *A, int lda, float *B, int ldb, const float *beta, float *C, int ldc){
|
||||
return hipHCBLASStatusToHIPStatus(hcblasSgemm( handle, hipOperationToHCCOperation(transa), hipOperationToHCCOperation(transb), m, n, k, alpha, A, lda, B, ldb, beta, C, ldc));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasCgemm(hipblasHandle_t handle, hipblasOperation_t transa, hipblasOperation_t transb,
|
||||
int m, int n, int k, const hipComplex *alpha, hipComplex *A, int lda, hipComplex *B, int ldb, const hipComplex *beta, hipComplex *C, int ldc){
|
||||
return hipHCBLASStatusToHIPStatus(hcblasCgemm( handle, hipOperationToHCCOperation(transa), hipOperationToHCCOperation(transb), m, n, k, alpha, A, lda, B, ldb, beta, C, ldc));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasSgemmBatched(hipblasHandle_t handle, hipblasOperation_t transa, hipblasOperation_t transb,
|
||||
int m, int n, int k, const float *alpha, float *A, int lda, float *B, int ldb, const float *beta, float *C, int ldc, int batchCount){
|
||||
return hipHCBLASStatusToHIPStatus(hcblasSgemmBatched( handle, hipOperationToHCCOperation(transa), hipOperationToHCCOperation(transb), m, n, k, alpha, A, lda, B, ldb, beta, C, ldc, batchCount));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasCgemmBatched(hipblasHandle_t handle, hipblasOperation_t transa, hipblasOperation_t transb,
|
||||
int m, int n, int k, const hipComplex *alpha, hipComplex *A, int lda, hipComplex *B, int ldb, const hipComplex *beta, hipComplex *C, int ldc, int batchCount){
|
||||
return HIPBLAS_STATUS_NOT_SUPPORTED;
|
||||
//return hipHCBLASStatusToHIPStatus(hcblasCgemmBatched( handle, transa, transb, m, n, k, alpha, A, lda, B, ldb, beta, C, ldc, batchCount));
|
||||
}
|
||||
|
||||
|
||||
|
||||
|
||||
#ifdef __cplusplus
|
||||
}
|
||||
#endif
|
||||
|
||||
|
||||
@@ -0,0 +1,181 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
|
||||
#ifndef HIP_FP16_H
|
||||
#define HIP_FP16_H
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
|
||||
typedef struct{
|
||||
unsigned x: 16;
|
||||
} __half;
|
||||
|
||||
|
||||
typedef struct __attribute__((aligned(4))){
|
||||
__half p,q;
|
||||
} __half2;
|
||||
|
||||
typedef __half half;
|
||||
typedef __half2 half2;
|
||||
|
||||
typedef struct{
|
||||
union{
|
||||
float f;
|
||||
unsigned u;
|
||||
};
|
||||
} struct_float;
|
||||
|
||||
/*
|
||||
Arithmetic functions
|
||||
*/
|
||||
|
||||
__device__ __half __hadd(const __half a, const __half b);
|
||||
|
||||
__device__ __half __hadd_sat(const __half a, const __half b);
|
||||
|
||||
__device__ __half __hfma(const __half a, const __half b, const __half c);
|
||||
|
||||
__device__ __half __hfma_sat(const __half a, const __half b, const __half c);
|
||||
|
||||
__device__ __half __hmul(const __half a, const __half b);
|
||||
|
||||
__device__ __half __hmul_sat(const __half a, const __half b);
|
||||
|
||||
__device__ __half __hneq(const __half a);
|
||||
|
||||
__device__ __half __hsub(const __half a, const __half b);
|
||||
|
||||
__device__ __half __hsub_sat(const __half a, const __half b);
|
||||
|
||||
|
||||
|
||||
/*
|
||||
Half2 Arithmetic Instructions
|
||||
*/
|
||||
|
||||
__device__ __half2 __hadd2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ __half2 __hadd2_sat(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ __half2 __hfma2(const __half2 a, const __half2 b, const __half2 c);
|
||||
|
||||
__device__ __half2 __hfma2_sat(const __half2 a, const __half2 b, const __half2 c);
|
||||
|
||||
__device__ __half2 __hmul2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ __half2 __hmul2_sat(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ __half2 __hneq2(const __half2 a);
|
||||
|
||||
__device__ __half2 __hsub2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ __half2 __hsub2_sat(const __half2 a, const __half2 b);
|
||||
|
||||
/*
|
||||
Half Cmps
|
||||
*/
|
||||
|
||||
__device__ bool __heq(const __half a, const __half b);
|
||||
|
||||
__device__ bool __hge(const __half a, const __half b);
|
||||
|
||||
__device__ bool __hgt(const __half a, const __half b);
|
||||
|
||||
__device__ bool __hisinf(const __half a);
|
||||
|
||||
__device__ bool __hisnan(const __half a);
|
||||
|
||||
__device__ bool __hle(const __half a, const __half b);
|
||||
|
||||
__device__ bool __hlt(const __half a, const __half b);
|
||||
|
||||
__device__ bool __hne(const __half a, const __half b);
|
||||
|
||||
/*
|
||||
Half2 Cmps
|
||||
*/
|
||||
|
||||
__device__ bool __hbeq2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ bool __hbge2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ bool __hbgt2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ bool __hble2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ bool __hblt2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ bool __hbne2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ __half2 __heq2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ __half2 __hge2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ __half2 __hgt2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ __half2 __hisnan2(const __half2 a);
|
||||
|
||||
__device__ __half2 __hle2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ __half2 __hlt2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ __half2 __hne2(const __half2 a, const __half2 b);
|
||||
|
||||
|
||||
/*
|
||||
Half Cnvs and Data Mvmnt
|
||||
*/
|
||||
|
||||
__device__ __half2 __float22half2_rn(const float2 a);
|
||||
|
||||
__device__ __half __float2half(const float a);
|
||||
|
||||
__device__ __half2 __float2half2_rn(const float a);
|
||||
|
||||
__device__ __half2 __floats2half2_rn(const float a, const float b);
|
||||
|
||||
__device__ float2 __half22float2(const __half2 a);
|
||||
|
||||
__device__ float __half2float(const __half a);
|
||||
|
||||
__device__ __half2 __half2half2(const __half a);
|
||||
|
||||
__device__ __half2 __halves2half2(const __half a, const __half b);
|
||||
|
||||
__device__ float __high2float(const __half2 a);
|
||||
|
||||
__device__ __half __high2half(const __half2 a);
|
||||
|
||||
__device__ __half2 __high2half2(const __half2 a);
|
||||
|
||||
__device__ __half2 __highs2half2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ float __low2float(const __half2 a);
|
||||
|
||||
__device__ __half __low2half(const __half2 a);
|
||||
|
||||
__device__ __half2 __low2half2(const __half2 a);
|
||||
|
||||
__device__ __half2 __lows2half2(const __half2 a, const __half2 b);
|
||||
|
||||
__device__ __half2 __lowhigh2highlow(const __half2 a);
|
||||
|
||||
__device__ __half2 __low2half2(const __half2 a, const __half2 b);
|
||||
|
||||
#endif
|
||||
@@ -0,0 +1,706 @@
|
||||
/*
|
||||
Link errors represented as this: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.
|
||||
*/
|
||||
|
||||
#ifndef HIP_HCC_H
|
||||
#define HIP_HCC_H
|
||||
|
||||
#include <hc.hpp>
|
||||
#include <hsa/hsa.h>
|
||||
#include "hip/hcc_detail/hip_util.h"
|
||||
|
||||
|
||||
#if defined(__HCC__) && (__hcc_workweek__ < 16354)
|
||||
#error("This version of HIP requires a newer version of HCC.");
|
||||
#endif
|
||||
|
||||
// #define USE_MEMCPYTOSYMBOL
|
||||
//
|
||||
|
||||
|
||||
//---
|
||||
// Environment variables:
|
||||
|
||||
// Intended to distinguish whether an environment variable should be visible only in debug mode, or in debug+release.
|
||||
//static const int debug = 0;
|
||||
extern const int release;
|
||||
|
||||
extern int HIP_LAUNCH_BLOCKING;
|
||||
|
||||
extern int HIP_PRINT_ENV;
|
||||
extern int HIP_ATP_MARKER;
|
||||
//extern int HIP_TRACE_API;
|
||||
extern int HIP_ATP;
|
||||
extern int HIP_DB;
|
||||
extern int HIP_STAGING_SIZE; /* size of staging buffers, in KB */
|
||||
extern int HIP_STREAM_SIGNALS; /* number of signals to allocate at stream creation */
|
||||
extern int HIP_VISIBLE_DEVICES; /* Contains a comma-separated sequence of GPU identifiers */
|
||||
|
||||
|
||||
//---
|
||||
// Chicken bits for disabling functionality to work around potential issues:
|
||||
extern int HIP_DISABLE_HW_KERNEL_DEP;
|
||||
|
||||
//---
|
||||
//Extern tls
|
||||
extern thread_local hipError_t tls_lastHipError;
|
||||
|
||||
|
||||
//---
|
||||
//Forward defs:
|
||||
class ihipStream_t;
|
||||
class ihipDevice_t;
|
||||
class ihipCtx_t;
|
||||
|
||||
// Color defs for debug messages:
|
||||
#define KNRM "\x1B[0m"
|
||||
#define KRED "\x1B[31m"
|
||||
#define KGRN "\x1B[32m"
|
||||
#define KYEL "\x1B[33m"
|
||||
#define KBLU "\x1B[34m"
|
||||
#define KMAG "\x1B[35m"
|
||||
#define KCYN "\x1B[36m"
|
||||
#define KWHT "\x1B[37m"
|
||||
|
||||
extern const char *API_COLOR;
|
||||
extern const char *API_COLOR_END;
|
||||
|
||||
|
||||
// If set, thread-safety is enforced on all stream functions.
|
||||
// Stream functions will acquire a mutex before entering critical sections.
|
||||
#define STREAM_THREAD_SAFE 1
|
||||
|
||||
|
||||
#define CTX_THREAD_SAFE 1
|
||||
|
||||
|
||||
// Compile debug trace mode - this prints debug messages to stderr when env var HIP_DB is set.
|
||||
// May be set to 0 to remove debug if checks - possible code size and performance difference?
|
||||
#define COMPILE_HIP_DB 1
|
||||
|
||||
|
||||
// Compile HIP tracing capability.
|
||||
// 0x1 = print a string at function entry with arguments.
|
||||
// 0x2 = prints a simple message with function name + return code when function exits.
|
||||
// 0x3 = print both.
|
||||
// Must be enabled at runtime with HIP_TRACE_API
|
||||
#define COMPILE_HIP_TRACE_API 0x3
|
||||
|
||||
|
||||
// Compile code that generates trace markers for CodeXL ATP at HIP function begin/end.
|
||||
// ATP is standard CodeXL format that includes timestamps for kernels, HSA RT APIs, and HIP APIs.
|
||||
#ifndef COMPILE_HIP_ATP_MARKER
|
||||
#define COMPILE_HIP_ATP_MARKER 0
|
||||
#endif
|
||||
|
||||
|
||||
#define DB_SHOW_TID 0
|
||||
|
||||
#if DB_SHOW_TID
|
||||
#define COMPUTE_TID_STR \
|
||||
std::stringstream tid_ss;\
|
||||
std::stringstream tid_ss_num;\
|
||||
tid_ss_num << std::this_thread::get_id();\
|
||||
tid_ss << " tid:" << std::hex << std::stoull(tid_ss_num.str());
|
||||
#else
|
||||
#define COMPUTE_TID_STR std::stringstream tid_ss;
|
||||
#endif
|
||||
|
||||
|
||||
// Compile support for trace markers that are displayed on CodeXL GUI at start/stop of each function boundary.
|
||||
// TODO - currently we print the trace message at the beginning. if we waited, we could also include return codes, and any values returned
|
||||
// through ptr-to-args (ie the pointers allocated by hipMalloc).
|
||||
#if COMPILE_HIP_ATP_MARKER
|
||||
#include "CXLActivityLogger.h"
|
||||
#define SCOPED_MARKER(markerName,group,userString) amdtScopedMarker(markerName, group, userString)
|
||||
#else
|
||||
// Swallow scoped markers:
|
||||
#define SCOPED_MARKER(markerName,group,userString)
|
||||
#endif
|
||||
|
||||
|
||||
#if COMPILE_HIP_ATP_MARKER || (COMPILE_HIP_TRACE_API & 0x1)
|
||||
#define API_TRACE(...)\
|
||||
{\
|
||||
if (HIP_ATP_MARKER || (COMPILE_HIP_DB && HIP_TRACE_API)) {\
|
||||
std::string s = std::string(__func__) + " (" + ToString(__VA_ARGS__) + ')';\
|
||||
if (COMPILE_HIP_DB && HIP_TRACE_API) {\
|
||||
COMPUTE_TID_STR\
|
||||
fprintf (stderr, "%s<<hip-api:%s %s\n%s" , API_COLOR, tid_ss.str().c_str(), s.c_str(), API_COLOR_END);\
|
||||
}\
|
||||
SCOPED_MARKER(s.c_str(), "HIP", NULL);\
|
||||
}\
|
||||
}
|
||||
#else
|
||||
// Swallow API_TRACE
|
||||
#define API_TRACE(...)
|
||||
#endif
|
||||
|
||||
|
||||
// Just initialize the HIP runtime, but don't log any trace information.
|
||||
#define HIP_INIT()\
|
||||
std::call_once(hip_initialized, ihipInit);\
|
||||
ihipCtxStackUpdate();
|
||||
|
||||
|
||||
// This macro should be called at the beginning of every HIP API.
|
||||
// It initialies the hip runtime (exactly once), and
|
||||
// generate trace string that can be output to stderr or to ATP file.
|
||||
#define HIP_INIT_API(...) \
|
||||
HIP_INIT()\
|
||||
API_TRACE(__VA_ARGS__);
|
||||
|
||||
#define ihipLogStatus(hipStatus) \
|
||||
({\
|
||||
hipError_t localHipStatus = hipStatus; /*local copy so hipStatus only evaluated once*/ \
|
||||
tls_lastHipError = localHipStatus;\
|
||||
\
|
||||
if ((COMPILE_HIP_TRACE_API & 0x2) && HIP_TRACE_API) {\
|
||||
fprintf(stderr, " %ship-api: %-30s ret=%2d (%s)>>%s\n", (localHipStatus == 0) ? API_COLOR:KRED, __func__, localHipStatus, ihipErrorString(localHipStatus), API_COLOR_END);\
|
||||
}\
|
||||
localHipStatus;\
|
||||
})
|
||||
|
||||
|
||||
|
||||
|
||||
//---
|
||||
//HIP_DB Debug flags:
|
||||
#define DB_API 0 /* 0x01 - shortcut to enable HIP_TRACE_API on single switch */
|
||||
#define DB_SYNC 1 /* 0x02 - trace synchronization pieces */
|
||||
#define DB_MEM 2 /* 0x04 - trace memory allocation / deallocation */
|
||||
#define DB_COPY1 3 /* 0x08 - trace memory copy commands. . */
|
||||
#define DB_SIGNAL 4 /* 0x10 - trace signal pool commands */
|
||||
#define DB_COPY2 5 /* 0x20 - trace memory copy commands. Detailed. */
|
||||
// When adding a new debug flag, also add to the char name table below.
|
||||
|
||||
static const char *dbName [] =
|
||||
{
|
||||
KNRM "hip-api", // not used,
|
||||
KYEL "hip-sync",
|
||||
KCYN "hip-mem",
|
||||
KMAG "hip-copy1",
|
||||
KRED "hip-signal",
|
||||
KNRM "hip-copy2",
|
||||
};
|
||||
|
||||
|
||||
|
||||
#if COMPILE_HIP_DB
|
||||
#define tprintf(trace_level, ...) {\
|
||||
if (HIP_DB & (1<<(trace_level))) {\
|
||||
char msgStr[1000];\
|
||||
snprintf(msgStr, 2000, __VA_ARGS__);\
|
||||
COMPUTE_TID_STR\
|
||||
fprintf (stderr, " %s%s:%s%s", dbName[trace_level], tid_ss.str().c_str(), msgStr, KNRM); \
|
||||
}\
|
||||
}
|
||||
#else
|
||||
/* Compile to empty code */
|
||||
#define tprintf(trace_level, ...)
|
||||
#endif
|
||||
|
||||
|
||||
|
||||
|
||||
|
||||
class ihipException : public std::exception
|
||||
{
|
||||
public:
|
||||
ihipException(hipError_t e) : _code(e) {};
|
||||
|
||||
hipError_t _code;
|
||||
};
|
||||
|
||||
|
||||
#ifdef __cplusplus
|
||||
extern "C" {
|
||||
#endif
|
||||
|
||||
|
||||
#ifdef __cplusplus
|
||||
}
|
||||
#endif
|
||||
|
||||
const hipStream_t hipStreamNull = 0x0;
|
||||
|
||||
|
||||
// Used to remove lock, for performance or stimulating bugs.
|
||||
class FakeMutex
|
||||
{
|
||||
public:
|
||||
void lock() { }
|
||||
bool try_lock() {return true; }
|
||||
void unlock() { }
|
||||
};
|
||||
|
||||
|
||||
#if STREAM_THREAD_SAFE
|
||||
typedef std::mutex StreamMutex;
|
||||
#else
|
||||
#warning "Stream thread-safe disabled"
|
||||
typedef FakeMutex StreamMutex;
|
||||
#endif
|
||||
|
||||
// Pair Device and Ctx together, these could also be toggled separately if desired.
|
||||
#if CTX_THREAD_SAFE
|
||||
typedef std::mutex CtxMutex;
|
||||
#else
|
||||
typedef FakeMutex CtxMutex;
|
||||
#warning "Device thread-safe disabled"
|
||||
#endif
|
||||
|
||||
//
|
||||
//---
|
||||
// Protects access to the member _data with a lock acquired on contruction/destruction.
|
||||
// T must contain a _mutex field which meets the BasicLockable requirements (lock/unlock)
|
||||
template<typename T>
|
||||
class LockedAccessor
|
||||
{
|
||||
public:
|
||||
LockedAccessor(T &criticalData, bool autoUnlock=true) :
|
||||
_criticalData(&criticalData),
|
||||
_autoUnlock(autoUnlock)
|
||||
|
||||
{
|
||||
tprintf(DB_SYNC, "lock critical data %s.%p\n", typeid(T).name(), _criticalData);
|
||||
_criticalData->_mutex.lock();
|
||||
};
|
||||
|
||||
~LockedAccessor()
|
||||
{
|
||||
if (_autoUnlock) {
|
||||
tprintf(DB_SYNC, "auto-unlock critical data %s.%p\n",typeid(T).name(), _criticalData);
|
||||
_criticalData->_mutex.unlock();
|
||||
}
|
||||
}
|
||||
|
||||
void unlock()
|
||||
{
|
||||
tprintf(DB_SYNC, "unlock critical data %s.%p\n", typeid(T).name(), _criticalData);
|
||||
_criticalData->_mutex.unlock();
|
||||
}
|
||||
|
||||
// Syntactic sugar so -> can be used to get the underlying type.
|
||||
T *operator->() { return _criticalData; };
|
||||
|
||||
private:
|
||||
T *_criticalData;
|
||||
bool _autoUnlock;
|
||||
};
|
||||
|
||||
|
||||
template <typename MUTEX_TYPE>
|
||||
struct LockedBase {
|
||||
|
||||
// Experts-only interface for explicit locking.
|
||||
// Most uses should use the lock-accessor.
|
||||
void lock() { _mutex.lock(); }
|
||||
void unlock() { _mutex.unlock(); }
|
||||
|
||||
MUTEX_TYPE _mutex;
|
||||
};
|
||||
|
||||
|
||||
class ihipModule_t{
|
||||
public:
|
||||
hsa_executable_t executable;
|
||||
hsa_code_object_t object;
|
||||
std::string fileName;
|
||||
void *ptr;
|
||||
size_t size;
|
||||
};
|
||||
|
||||
|
||||
class ihipFunction_t{
|
||||
public:
|
||||
ihipFunction_t(const char *name) {
|
||||
size_t nameSz = strlen(name);
|
||||
char *kernelName = (char*)malloc(nameSz);
|
||||
strncpy(kernelName, name, nameSz);
|
||||
_kernelName = kernelName;
|
||||
};
|
||||
|
||||
~ihipFunction_t() {
|
||||
if (_kernelName) {
|
||||
free((void*)_kernelName);
|
||||
_kernelName = NULL;
|
||||
};
|
||||
};
|
||||
public:
|
||||
const char *_kernelName;
|
||||
hsa_executable_symbol_t _kernelSymbol;
|
||||
uint64_t _kernel;
|
||||
};
|
||||
|
||||
|
||||
template <typename MUTEX_TYPE>
|
||||
class ihipStreamCriticalBase_t : public LockedBase<MUTEX_TYPE>
|
||||
{
|
||||
public:
|
||||
ihipStreamCriticalBase_t(hc::accelerator_view av) :
|
||||
_kernelCnt(0),
|
||||
_av(av)
|
||||
{
|
||||
};
|
||||
|
||||
~ihipStreamCriticalBase_t() {
|
||||
}
|
||||
|
||||
ihipStreamCriticalBase_t<StreamMutex> * mlock() { LockedBase<MUTEX_TYPE>::lock(); return this;};
|
||||
|
||||
public:
|
||||
// TODO - remove _kernelCnt mechanism:
|
||||
uint32_t _kernelCnt; // Count of inflight kernels in this stream. Reset at ::wait().
|
||||
hc::accelerator_view _av;
|
||||
};
|
||||
|
||||
|
||||
// if HIP code needs to acquire locks for both ihipCtx_t and ihipStream_t, it should first acquire the lock
|
||||
// for the ihipCtx_t and then for the individual streams. The locks should not be acquired in reverse order
|
||||
// or deadlock may occur. In some cases, it may be possible to reduce the range where the locks must be held.
|
||||
// HIP routines should avoid acquiring and releasing the same lock during the execution of a single HIP API.
|
||||
|
||||
|
||||
typedef ihipStreamCriticalBase_t<StreamMutex> ihipStreamCritical_t;
|
||||
typedef LockedAccessor<ihipStreamCritical_t> LockedAccessor_StreamCrit_t;
|
||||
|
||||
|
||||
//---
|
||||
// Internal stream structure.
|
||||
class ihipStream_t {
|
||||
public:
|
||||
typedef uint64_t SeqNum_t ;
|
||||
ihipStream_t(ihipCtx_t *ctx, hc::accelerator_view av, unsigned int flags);
|
||||
~ihipStream_t();
|
||||
|
||||
// kind is hipMemcpyKind
|
||||
void locked_copySync (void* dst, const void* src, size_t sizeBytes, unsigned kind, bool resolveOn = true);
|
||||
|
||||
|
||||
void locked_copyAsync(void* dst, const void* src, size_t sizeBytes, unsigned kind);
|
||||
|
||||
|
||||
//---
|
||||
// Member functions that begin with locked_ are thread-safe accessors - these acquire / release the critical mutex.
|
||||
LockedAccessor_StreamCrit_t lockopen_preKernelCommand();
|
||||
void lockclose_postKernelCommand(hc::accelerator_view *av);
|
||||
|
||||
|
||||
void locked_wait(bool assertQueueEmpty=false);
|
||||
|
||||
hc::accelerator_view* locked_getAv() { LockedAccessor_StreamCrit_t crit(_criticalData); return &(crit->_av); };
|
||||
|
||||
void locked_waitEvent(hipEvent_t event);
|
||||
void locked_recordEvent(hipEvent_t event);
|
||||
|
||||
|
||||
//---
|
||||
|
||||
// Use this if we already have the stream critical data mutex:
|
||||
void wait(LockedAccessor_StreamCrit_t &crit, bool assertQueueEmpty=false);
|
||||
|
||||
void launchModuleKernel(hc::accelerator_view av, hsa_signal_t signal,
|
||||
uint32_t blockDimX, uint32_t blockDimY, uint32_t blockDimZ,
|
||||
uint32_t gridDimX, uint32_t gridDimY, uint32_t gridDimZ,
|
||||
uint32_t groupSegmentSize, uint32_t sharedMemBytes,
|
||||
void *kernarg, size_t kernSize, uint64_t kernel);
|
||||
|
||||
|
||||
|
||||
//-- Non-racy accessors:
|
||||
// These functions access fields set at initialization time and are non-racy (so do not acquire mutex)
|
||||
const ihipDevice_t * getDevice() const;
|
||||
ihipCtx_t * getCtx() const;
|
||||
|
||||
|
||||
public:
|
||||
//---
|
||||
//Public member vars - these are set at initialization and never change:
|
||||
SeqNum_t _id; // monotonic sequence ID
|
||||
unsigned _flags;
|
||||
|
||||
|
||||
private:
|
||||
|
||||
|
||||
// The unsigned return is hipMemcpyKind
|
||||
unsigned resolveMemcpyDirection(bool srcTracked, bool dstTracked, bool srcInDeviceMem, bool dstInDeviceMem);
|
||||
|
||||
bool canSeePeerMemory(const ihipCtx_t *thisCtx, ihipCtx_t *dstCtx, ihipCtx_t *srcCtx);
|
||||
|
||||
|
||||
private: // Data
|
||||
// Critical Data - MUST be accessed through LockedAccessor_StreamCrit_t
|
||||
ihipStreamCritical_t _criticalData;
|
||||
|
||||
ihipCtx_t *_ctx; // parent context that owns this stream.
|
||||
|
||||
// Friends:
|
||||
friend std::ostream& operator<<(std::ostream& os, const ihipStream_t& s);
|
||||
friend hipError_t hipStreamQuery(hipStream_t);
|
||||
};
|
||||
|
||||
|
||||
|
||||
//----
|
||||
// Internal event structure:
|
||||
enum hipEventStatus_t {
|
||||
hipEventStatusUnitialized = 0, // event is unutilized, must be "Created" before use.
|
||||
hipEventStatusCreated = 1,
|
||||
hipEventStatusRecording = 2, // event has been enqueued to record something.
|
||||
hipEventStatusRecorded = 3, // event has been recorded - timestamps are valid.
|
||||
} ;
|
||||
|
||||
|
||||
// internal hip event structure.
|
||||
struct ihipEvent_t {
|
||||
hipEventStatus_t _state;
|
||||
|
||||
hipStream_t _stream; // Stream where the event is recorded, or NULL if all streams.
|
||||
unsigned _flags;
|
||||
|
||||
hc::completion_future _marker;
|
||||
uint64_t _timestamp; // store timestamp, may be set on host or by marker.
|
||||
} ;
|
||||
|
||||
|
||||
|
||||
|
||||
|
||||
//----
|
||||
// Properties of the HIP device.
|
||||
// Multiple contexts can point to same device.
|
||||
class ihipDevice_t
|
||||
{
|
||||
public:
|
||||
ihipDevice_t(unsigned deviceId, unsigned deviceCnt, hc::accelerator &acc);
|
||||
~ihipDevice_t();
|
||||
|
||||
// Accessors:
|
||||
ihipCtx_t *getPrimaryCtx() const { return _primaryCtx; };
|
||||
|
||||
public:
|
||||
unsigned _deviceId; // device ID
|
||||
|
||||
hc::accelerator _acc;
|
||||
hsa_agent_t _hsaAgent; // hsa agent handle
|
||||
|
||||
//! Number of compute units supported by the device:
|
||||
unsigned _computeUnits;
|
||||
hipDeviceProp_t _props; // saved device properties.
|
||||
|
||||
// TODO - report this through device properties, base on HCC API call.
|
||||
int _isLargeBar;
|
||||
|
||||
ihipCtx_t *_primaryCtx;
|
||||
|
||||
private:
|
||||
hipError_t initProperties(hipDeviceProp_t* prop);
|
||||
};
|
||||
//=============================================================================
|
||||
|
||||
|
||||
|
||||
//=============================================================================
|
||||
//class ihipCtxCriticalBase_t
|
||||
template <typename MUTEX_TYPE>
|
||||
class ihipCtxCriticalBase_t : LockedBase<MUTEX_TYPE>
|
||||
{
|
||||
public:
|
||||
ihipCtxCriticalBase_t(unsigned deviceCnt) :
|
||||
_peerCnt(0)
|
||||
{
|
||||
_peerAgents = new hsa_agent_t[deviceCnt];
|
||||
};
|
||||
|
||||
~ihipCtxCriticalBase_t() {
|
||||
if (_peerAgents != nullptr) {
|
||||
delete _peerAgents;
|
||||
_peerAgents = nullptr;
|
||||
}
|
||||
_peerCnt = 0;
|
||||
}
|
||||
|
||||
// Streams:
|
||||
void addStream(ihipStream_t *stream);
|
||||
std::list<ihipStream_t*> &streams() { return _streams; };
|
||||
const std::list<ihipStream_t*> &const_streams() const { return _streams; };
|
||||
|
||||
|
||||
// Peer Accessor classes:
|
||||
bool isPeer(const ihipCtx_t *peer); // returns True if peer has access to memory physically located on this device.
|
||||
bool addPeer(ihipCtx_t *peer);
|
||||
bool removePeer(ihipCtx_t *peer);
|
||||
void resetPeers(ihipCtx_t *thisDevice);
|
||||
void printPeers(FILE *f) const;
|
||||
|
||||
uint32_t peerCnt() const { return _peerCnt; };
|
||||
hsa_agent_t *peerAgents() const { return _peerAgents; };
|
||||
|
||||
|
||||
|
||||
friend class LockedAccessor<ihipCtxCriticalBase_t>;
|
||||
private:
|
||||
//--- Stream Tracker:
|
||||
std::list< ihipStream_t* > _streams; // streams associated with this device.
|
||||
|
||||
|
||||
//--- Peer Tracker:
|
||||
// These reflect the currently Enabled set of peers for this GPU:
|
||||
// Enabled peers have permissions to access the memory physically allocated on this device.
|
||||
// Note the peers always contain the self agent for easy interfacing with HSA APIs.
|
||||
std::list<ihipCtx_t*> _peers; // list of enabled peer devices.
|
||||
uint32_t _peerCnt; // number of enabled peers
|
||||
hsa_agent_t *_peerAgents; // efficient packed array of enabled agents (to use for allocations.)
|
||||
private:
|
||||
void recomputePeerAgents();
|
||||
};
|
||||
// Note Mutex type Real/Fake selected based on CtxMutex
|
||||
typedef ihipCtxCriticalBase_t<CtxMutex> ihipCtxCritical_t;
|
||||
|
||||
// This type is used by functions that need access to the critical device structures.
|
||||
typedef LockedAccessor<ihipCtxCritical_t> LockedAccessor_CtxCrit_t;
|
||||
//=============================================================================
|
||||
|
||||
|
||||
//=============================================================================
|
||||
//class ihipCtx_t:
|
||||
// A HIP CTX (context) points at one of the existing devices and contains the streams,
|
||||
// peer-to-peer mappings, creation flags. Multiple contexts can point to the same
|
||||
// device.
|
||||
//
|
||||
class ihipCtx_t
|
||||
{
|
||||
public: // Functions:
|
||||
ihipCtx_t(ihipDevice_t *device, unsigned deviceCnt, unsigned flags); // note: calls constructor for _criticalData
|
||||
~ihipCtx_t();
|
||||
|
||||
// Functions which read or write the critical data are named locked_.
|
||||
// ihipCtx_t does not use recursive locks so the ihip implementation must avoid calling a locked_ function from within a locked_ function.
|
||||
// External functions which call several locked_ functions will acquire and release the lock for each function. if this occurs in
|
||||
// performance-sensitive code we may want to refactor by adding non-locked functions and creating a new locked_ member function to call them all.
|
||||
void locked_addStream(ihipStream_t *s);
|
||||
void locked_removeStream(ihipStream_t *s);
|
||||
void locked_reset();
|
||||
void locked_waitAllStreams();
|
||||
void locked_syncDefaultStream(bool waitOnSelf);
|
||||
|
||||
ihipCtxCritical_t &criticalData() { return _criticalData; }; // TODO, move private. Fix P2P.
|
||||
|
||||
const ihipDevice_t *getDevice() const { return _device; };
|
||||
|
||||
// TODO - review uses of getWriteableDevice(), can these be converted to getDevice()
|
||||
ihipDevice_t *getWriteableDevice() const { return _device; };
|
||||
|
||||
std::string toString() const;
|
||||
|
||||
public: // Data
|
||||
// The NULL stream is used if no other stream is specified.
|
||||
// Default stream has special synchronization properties with other streams.
|
||||
ihipStream_t *_defaultStream;
|
||||
|
||||
// Flags specified when the context is created:
|
||||
unsigned _ctxFlags;
|
||||
|
||||
private:
|
||||
ihipDevice_t *_device;
|
||||
|
||||
|
||||
private: // Critical data, protected with locked access:
|
||||
// Members of _protected data MUST be accessed through the LockedAccessor.
|
||||
// Search for LockedAccessor<ihipCtxCritical_t> for examples; do not access _criticalData directly.
|
||||
ihipCtxCritical_t _criticalData;
|
||||
|
||||
};
|
||||
|
||||
|
||||
|
||||
//=================================================================================================
|
||||
// Global variable definition:
|
||||
extern std::once_flag hip_initialized;
|
||||
extern unsigned g_deviceCnt;
|
||||
extern hsa_agent_t g_cpu_agent ; // the CPU agent.
|
||||
|
||||
//=================================================================================================
|
||||
// Extern functions:
|
||||
extern void ihipInit();
|
||||
extern const char *ihipErrorString(hipError_t);
|
||||
extern ihipCtx_t *ihipGetTlsDefaultCtx();
|
||||
extern void ihipSetTlsDefaultCtx(ihipCtx_t *ctx);
|
||||
extern hipError_t ihipSynchronize(void);
|
||||
extern void ihipCtxStackUpdate();
|
||||
|
||||
extern ihipDevice_t *ihipGetDevice(int);
|
||||
ihipCtx_t * ihipGetPrimaryCtx(unsigned deviceIndex);
|
||||
|
||||
extern void ihipSetTs(hipEvent_t e);
|
||||
|
||||
|
||||
hipStream_t ihipSyncAndResolveStream(hipStream_t);
|
||||
|
||||
// Stream printf functions:
|
||||
inline std::ostream& operator<<(std::ostream& os, const ihipStream_t& s)
|
||||
{
|
||||
os << "stream#";
|
||||
os << s.getDevice()->_deviceId;;
|
||||
os << '.';
|
||||
os << s._id;
|
||||
return os;
|
||||
}
|
||||
|
||||
inline std::ostream & operator<<(std::ostream& os, const dim3& s)
|
||||
{
|
||||
os << '{';
|
||||
os << s.x;
|
||||
os << ',';
|
||||
os << s.y;
|
||||
os << ',';
|
||||
os << s.z;
|
||||
os << '}';
|
||||
return os;
|
||||
}
|
||||
|
||||
inline std::ostream & operator<<(std::ostream& os, const gl_dim3& s)
|
||||
{
|
||||
os << '{';
|
||||
os << s.x;
|
||||
os << ',';
|
||||
os << s.y;
|
||||
os << ',';
|
||||
os << s.z;
|
||||
os << '}';
|
||||
return os;
|
||||
}
|
||||
|
||||
// Stream printf functions:
|
||||
inline std::ostream& operator<<(std::ostream& os, const hipEvent_t& e)
|
||||
{
|
||||
os << "event:" << std::hex << static_cast<void*> (e);
|
||||
return os;
|
||||
}
|
||||
|
||||
inline std::ostream& operator<<(std::ostream& os, const ihipCtx_t* c)
|
||||
{
|
||||
os << "ctx:" << static_cast<const void*> (c)
|
||||
<< " dev:" << c->getDevice()->_deviceId;
|
||||
return os;
|
||||
}
|
||||
|
||||
|
||||
#endif
|
||||
@@ -0,0 +1,75 @@
|
||||
/*
|
||||
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 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.
|
||||
*/
|
||||
|
||||
#ifndef HIP_LDG_H
|
||||
#define HIP_LDG_H
|
||||
|
||||
#if __HCC__
|
||||
#if __hcc_workweek__ >= 16164
|
||||
#include "hip/hip_vector_types.h"
|
||||
#include "hip/hcc_detail/host_defines.h"
|
||||
|
||||
|
||||
__device__ char __ldg(const char* );
|
||||
__device__ char2 __ldg(const char2* );
|
||||
__device__ char4 __ldg(const char4* );
|
||||
__device__ signed char __ldg(const signed char* );
|
||||
__device__ unsigned char __ldg(const unsigned char* );
|
||||
|
||||
__device__ short __ldg(const short* );
|
||||
__device__ short2 __ldg(const short2* );
|
||||
__device__ short4 __ldg(const short4* );
|
||||
__device__ unsigned short __ldg(const unsigned short* );
|
||||
|
||||
__device__ int __ldg(const int* );
|
||||
__device__ int2 __ldg(const int2* );
|
||||
__device__ int4 __ldg(const int4* );
|
||||
__device__ unsigned int __ldg(const unsigned int* );
|
||||
|
||||
|
||||
__device__ long __ldg(const long* );
|
||||
__device__ unsigned long __ldg(const unsigned long* );
|
||||
|
||||
__device__ long long __ldg(const long long* );
|
||||
__device__ longlong2 __ldg(const longlong2* );
|
||||
__device__ unsigned long long __ldg(const unsigned long long* );
|
||||
|
||||
__device__ uchar2 __ldg(const uchar2* );
|
||||
__device__ uchar4 __ldg(const uchar4* );
|
||||
|
||||
__device__ ushort2 __ldg(const ushort2* );
|
||||
|
||||
__device__ uint2 __ldg(const uint2* );
|
||||
__device__ uint4 __ldg(const uint4* );
|
||||
|
||||
__device__ ulonglong2 __ldg(const ulonglong2* );
|
||||
|
||||
__device__ float __ldg(const float* );
|
||||
__device__ float2 __ldg(const float2* );
|
||||
__device__ float4 __ldg(const float4* );
|
||||
|
||||
__device__ double __ldg(const double* );
|
||||
__device__ double2 __ldg(const double2* );
|
||||
|
||||
#endif // __hcc_workweek__
|
||||
|
||||
#endif // __HCC__
|
||||
|
||||
#endif // HIP_LDG_H
|
||||
|
||||
@@ -0,0 +1,694 @@
|
||||
/*
|
||||
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 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.
|
||||
*/
|
||||
/**
|
||||
* @file hcc_detail/hip_runtime.h
|
||||
* @brief Contains definitions of APIs for HIP runtime.
|
||||
*/
|
||||
|
||||
//#pragma once
|
||||
#ifndef HIP_RUNTIME_H
|
||||
#define HIP_RUNTIME_H
|
||||
|
||||
//---
|
||||
// Top part of file can be compiled with any compiler
|
||||
|
||||
|
||||
//#include <cstring>
|
||||
#if __cplusplus
|
||||
#include <cmath>
|
||||
#else
|
||||
#include <math.h>
|
||||
#include <string.h>
|
||||
#include <stddef.h>
|
||||
#endif
|
||||
// Define NVCC_COMPAT for CUDA compatibility
|
||||
#define NVCC_COMPAT
|
||||
#define CUDA_SUCCESS hipSuccess
|
||||
|
||||
#include <hip/hip_runtime_api.h>
|
||||
//#include "hip/hcc_detail/hip_hcc.h"
|
||||
//---
|
||||
// Remainder of this file only compiles with HCC
|
||||
#ifdef __HCC__
|
||||
#include <grid_launch.h>
|
||||
|
||||
#if defined (GRID_LAUNCH_VERSION) and (GRID_LAUNCH_VERSION >= 20)
|
||||
// Use field names for grid_launch 2.0 structure, if HCC supports GL 2.0.
|
||||
#else
|
||||
#error (HCC must support GRID_LAUNCH_20)
|
||||
#endif
|
||||
|
||||
#define HIP_LAUNCH_PARAM_BUFFER_POINTER ((void*) 0x01)
|
||||
#define HIP_LAUNCH_PARAM_BUFFER_SIZE ((void*) 0x02)
|
||||
#define HIP_LAUNCH_PARAM_END ((void*) 0x03)
|
||||
|
||||
extern int HIP_TRACE_API;
|
||||
|
||||
//TODO-HCC-GL - change this to typedef.
|
||||
//typedef grid_launch_parm hipLaunchParm ;
|
||||
#define hipLaunchParm grid_launch_parm
|
||||
#ifdef __cplusplus
|
||||
#include <hip/hcc_detail/hip_texture.h>
|
||||
#include <hip/hcc_detail/hip_ldg.h>
|
||||
#endif
|
||||
#include <hip/hcc_detail/host_defines.h>
|
||||
// TODO-HCC remove old definitions ; ~1602 hcc supports __HCC_ACCELERATOR__ define.
|
||||
#if defined (__KALMAR_ACCELERATOR__) && !defined (__HCC_ACCELERATOR__)
|
||||
#define __HCC_ACCELERATOR__ __KALMAR_ACCELERATOR__
|
||||
#endif
|
||||
|
||||
// Feature tests:
|
||||
#if defined(__HCC_ACCELERATOR__) && (__HCC_ACCELERATOR__ != 0)
|
||||
// Device compile and not host compile:
|
||||
|
||||
//TODO-HCC enable __HIP_ARCH_HAS_ATOMICS__ when HCC supports these.
|
||||
// 32-bit Atomics:
|
||||
#define __HIP_ARCH_HAS_GLOBAL_INT32_ATOMICS__ (1)
|
||||
#define __HIP_ARCH_HAS_GLOBAL_FLOAT_ATOMIC_EXCH__ (1)
|
||||
#define __HIP_ARCH_HAS_SHARED_INT32_ATOMICS__ (1)
|
||||
#define __HIP_ARCH_HAS_SHARED_FLOAT_ATOMIC_EXCH__ (1)
|
||||
#define __HIP_ARCH_HAS_FLOAT_ATOMIC_ADD__ (0)
|
||||
|
||||
// 64-bit Atomics:
|
||||
#define __HIP_ARCH_HAS_GLOBAL_INT64_ATOMICS__ (1)
|
||||
#define __HIP_ARCH_HAS_SHARED_INT64_ATOMICS__ (0)
|
||||
|
||||
// Doubles
|
||||
#define __HIP_ARCH_HAS_DOUBLES__ (1)
|
||||
|
||||
//warp cross-lane operations:
|
||||
#define __HIP_ARCH_HAS_WARP_VOTE__ (1)
|
||||
#define __HIP_ARCH_HAS_WARP_BALLOT__ (1)
|
||||
#define __HIP_ARCH_HAS_WARP_SHUFFLE__ (1)
|
||||
#define __HIP_ARCH_HAS_WARP_FUNNEL_SHIFT__ (0)
|
||||
|
||||
//sync
|
||||
#define __HIP_ARCH_HAS_THREAD_FENCE_SYSTEM__ (0)
|
||||
#define __HIP_ARCH_HAS_SYNC_THREAD_EXT__ (0)
|
||||
|
||||
// misc
|
||||
#define __HIP_ARCH_HAS_SURFACE_FUNCS__ (0)
|
||||
#define __HIP_ARCH_HAS_3DGRID__ (1)
|
||||
#define __HIP_ARCH_HAS_DYNAMIC_PARALLEL__ (0)
|
||||
|
||||
#endif /* Device feature flags */
|
||||
|
||||
|
||||
//TODO-HCC this is currently ignored by HCC target of HIP
|
||||
#define __launch_bounds__(requiredMaxThreadsPerBlock, minBlocksPerMultiprocessor)
|
||||
|
||||
// Detect if we are compiling C++ mode or C mode
|
||||
#if defined(__cplusplus)
|
||||
#define __HCC_CPP__
|
||||
#elif defined(__STDC_VERSION__)
|
||||
#define __HCC_C__
|
||||
#endif
|
||||
|
||||
__device__ float acosf(float x);
|
||||
__device__ float acoshf(float x);
|
||||
__device__ float asinf(float x);
|
||||
__device__ float asinhf(float x);
|
||||
__device__ float atan2f(float y, float x);
|
||||
__device__ float atanf(float x);
|
||||
__device__ float atanhf(float x);
|
||||
__device__ float cbrtf(float x);
|
||||
__device__ float ceilf(float x);
|
||||
__device__ float copysignf(float x, float y);
|
||||
__device__ float cosf(float x);
|
||||
__device__ float coshf(float x);
|
||||
__device__ float cyl_bessel_i0f(float x);
|
||||
__device__ float cyl_bessel_i1f(float x);
|
||||
__device__ float erfcf(float x);
|
||||
__device__ float erfcinvf(float y);
|
||||
__host__ float erfcinvf(float y);
|
||||
__device__ float erfcxf(float x);
|
||||
__host__ float erfcxf(float x);
|
||||
__device__ float erff(float x);
|
||||
__device__ float erfinvf(float y);
|
||||
__host__ float erfinvf(float y);
|
||||
__device__ float exp10f(float x);
|
||||
__device__ float exp2f(float x);
|
||||
__device__ float expf(float x);
|
||||
__device__ float expm1f(float x);
|
||||
__device__ float fabsf(float x);
|
||||
__device__ float fdimf(float x, float y);
|
||||
__device__ __host__ float fdividef(float x, float y);
|
||||
__device__ float floorf(float x);
|
||||
__device__ float fmaf(float x, float y, float z);
|
||||
__device__ float fmaxf(float x, float y);
|
||||
__device__ float fminf(float x, float y);
|
||||
__device__ float fmodf(float x, float y);
|
||||
__device__ float frexpf(float x, float y);
|
||||
__device__ float hypotf(float x, float y);
|
||||
__device__ float ilogbf(float x);
|
||||
__host__ __device__ unsigned isfinite(float a);
|
||||
__device__ unsigned isinf(float a);
|
||||
__device__ unsigned isnan(float a);
|
||||
__device__ float j0f(float x);
|
||||
__device__ float j1f(float x);
|
||||
__device__ float jnf(int n, float x);
|
||||
__device__ float ldexpf(float x, int exp);
|
||||
__device__ float lgammaf(float x);
|
||||
__device__ long long int llrintf(float x);
|
||||
__device__ long long int llroundf(float x);
|
||||
__device__ float log10f(float x);
|
||||
__device__ float log1pf(float x);
|
||||
__device__ float log2f(float x);
|
||||
__device__ float logbf(float x);
|
||||
__device__ float logf(float x);
|
||||
__device__ long int lrintf(float x);
|
||||
__device__ long int lroundf(float x);
|
||||
__device__ float modff(float x, float *iptr);
|
||||
__device__ float nanf(const char* tagp);
|
||||
__device__ float nearbyintf(float x);
|
||||
__device__ float nextafterf(float x, float y);
|
||||
__device__ float norm3df(float a, float b, float c);
|
||||
__host__ float norm3df(float a, float b, float c);
|
||||
__device__ float norm4df(float a, float b, float c, float d);
|
||||
__host__ float norm4df(float a, float b, float c, float d);
|
||||
__device__ float normcdff(float y);
|
||||
__host__ float normcdff(float y);
|
||||
__device__ float normcdfinvf(float y);
|
||||
__host__ float normcdfinvf(float y);
|
||||
__device__ float normf(int dim, const float *a);
|
||||
__device__ float powf(float x, float y);
|
||||
__device__ float rcbrtf(float x);
|
||||
__host__ float rcbrtf(float x);
|
||||
__device__ float remainderf(float x, float y);
|
||||
__device__ float remquof(float x, float y, int *quo);
|
||||
__device__ float rhypotf(float x, float y);
|
||||
__host__ float rhypotf(float x, float y);
|
||||
__device__ float rintf(float x);
|
||||
__device__ float rnorm3df(float a, float b, float c);
|
||||
__host__ float rnorm3df(float a, float b, float c);
|
||||
__device__ float rnorm4df(float a, float b, float c, float d);
|
||||
__host__ float rnorm4df(float a, float b, float c, float d);
|
||||
__device__ float rnormf(int dim, const float* a);
|
||||
__host__ float rnormf(int dim, const float* a);
|
||||
__device__ float roundf(float x);
|
||||
__device__ float rsqrtf(float x);
|
||||
__device__ float scalblnf(float x, long int n);
|
||||
__device__ float scalbnf(float x, int n);
|
||||
__host__ __device__ unsigned signbit(float a);
|
||||
__device__ void sincosf(float x, float *sptr, float *cptr);
|
||||
__device__ void sincospif(float x, float *sptr, float *cptr);
|
||||
__host__ void sincospif(float x, float *sptr, float *cptr);
|
||||
__device__ float sinf(float x);
|
||||
__device__ float sinhf(float x);
|
||||
__device__ float sinpif(float x);
|
||||
__device__ float sqrtf(float x);
|
||||
__device__ float tanf(float x);
|
||||
__device__ float tanhf(float x);
|
||||
__device__ float tgammaf(float x);
|
||||
__device__ float truncf(float x);
|
||||
__device__ float y0f(float x);
|
||||
__device__ float y1f(float x);
|
||||
__device__ float ynf(int n, float x);
|
||||
|
||||
__host__ __device__ float cospif(float x);
|
||||
__host__ __device__ float sinpif(float x);
|
||||
__device__ float sqrtf(float x);
|
||||
__host__ __device__ float rsqrtf(float x);
|
||||
|
||||
__device__ double acos(double x);
|
||||
__device__ double acosh(double x);
|
||||
__device__ double asin(double x);
|
||||
__device__ double asinh(double x);
|
||||
__device__ double atan(double x);
|
||||
__device__ double atan2(double y, double x);
|
||||
__device__ double atanh(double x);
|
||||
__device__ double cbrt(double x);
|
||||
__device__ double ceil(double x);
|
||||
__device__ double copysign(double x, double y);
|
||||
__device__ double cos(double x);
|
||||
__device__ double cosh(double x);
|
||||
__host__ __device__ double cospi(double x);
|
||||
__device__ double cyl_bessel_i0(double x);
|
||||
__device__ double cyl_bessel_i1(double x);
|
||||
__device__ double erf(double x);
|
||||
__device__ double erfc(double x);
|
||||
__device__ double erfcinv(double y);
|
||||
__device__ double erfcx(double x);
|
||||
__device__ double erfinv(double x);
|
||||
__device__ double exp(double x);
|
||||
__device__ double exp10(double x);
|
||||
__device__ double exp2(double x);
|
||||
__device__ double expm1(double x);
|
||||
__device__ double fabs(double x);
|
||||
__device__ double fdim(double x, double y);
|
||||
__device__ double fdivide(double x, double y);
|
||||
__device__ double floor(double x);
|
||||
__device__ double fma(double x, double y, double z);
|
||||
__device__ double fmax(double x, double y);
|
||||
__device__ double fmin(double x, double y);
|
||||
__device__ double fmod(double x, double y);
|
||||
__device__ double frexp(double x, int *nptr);
|
||||
__device__ double hypot(double x, double y);
|
||||
__device__ double ilogb(double x);
|
||||
__host__ __device__ unsigned isfinite(double x);
|
||||
__device__ unsigned isinf(double x);
|
||||
__device__ unsigned isnan(double x);
|
||||
__device__ double j0(double x);
|
||||
__device__ double j1(double x);
|
||||
__device__ double jn(int n, double x);
|
||||
__device__ double ldexp(double x, int exp);
|
||||
__device__ double lgamma(double x);
|
||||
__device__ long long llrint(double x);
|
||||
__device__ long long llround(double x);
|
||||
__device__ double log(double x);
|
||||
__device__ double log10(double x);
|
||||
__device__ double log1p(double x);
|
||||
__device__ double log2(double x);
|
||||
__device__ double logb(double x);
|
||||
__device__ long int lrint(double x);
|
||||
__device__ long int lround(double x);
|
||||
__device__ double modf(double x, double *iptr);
|
||||
__device__ double nan(const char* tagp);
|
||||
__device__ double nearbyint(double x);
|
||||
__device__ double nextafter(double x, double y);
|
||||
__device__ double norm(int dim, const double* t);
|
||||
__device__ double norm3d(double a, double b, double c);
|
||||
__host__ double norm3d(double a, double b, double c);
|
||||
__device__ double norm4d(double a, double b, double c, double d);
|
||||
__host__ double norm4d(double a, double b, double c, double d);
|
||||
__device__ double normcdf(double y);
|
||||
__host__ double normcdf(double y);
|
||||
__device__ double normcdfinv(double y);
|
||||
__host__ double normcdfinv(double y);
|
||||
__device__ double pow(double x, double y);
|
||||
__device__ double rcbrt(double x);
|
||||
__host__ double rcbrt(double x);
|
||||
__device__ double remainder(double x, double y);
|
||||
__device__ double remquo(double x, double y, int *quo);
|
||||
__device__ double rhypot(double x, double y);
|
||||
__host__ double rhypot(double x, double y);
|
||||
__device__ double rint(double x);
|
||||
__device__ double rnorm(int dim, const double* t);
|
||||
__host__ double rnorm(int dim, const double* t);
|
||||
__device__ double rnorm3d(double a, double b, double c);
|
||||
__host__ double rnorm3d(double a, double b, double c);
|
||||
__device__ double rnorm4d(double a, double b, double c, double d);
|
||||
__host__ double rnorm4d(double a, double b, double c, double d);
|
||||
__device__ double round(double x);
|
||||
__host__ __device__ double rsqrt(double x);
|
||||
__device__ double scalbln(double x, long int n);
|
||||
__device__ double scalbn(double x, int n);
|
||||
__host__ __device__ unsigned signbit(double a);
|
||||
__device__ double sin(double a);
|
||||
__device__ void sincos(double x, double *sptr, double *cptr);
|
||||
__device__ void sincospi(double x, double *sptr, double *cptr);
|
||||
__host__ void sincospi(double x, double *sptr, double *cptr);
|
||||
__device__ double sinh(double x);
|
||||
__host__ __device__ double sinpi(double x);
|
||||
__device__ double sqrt(double x);
|
||||
__device__ double tan(double x);
|
||||
__device__ double tanh(double x);
|
||||
__device__ double tgamma(double x);
|
||||
__device__ double trunc(double x);
|
||||
__device__ double y0(double x);
|
||||
__device__ double y1(double y);
|
||||
__device__ double yn(int n, double x);
|
||||
|
||||
__host__ double erfcinv(double y);
|
||||
__host__ double erfcx(double x);
|
||||
__host__ double erfinv(double y);
|
||||
__host__ double fdivide(double x, double y);
|
||||
|
||||
// TODO - hipify-clang - change to use the function call.
|
||||
//#define warpSize hc::__wavesize()
|
||||
extern const int warpSize;
|
||||
|
||||
|
||||
#define clock_t long long int
|
||||
__device__ long long int clock64();
|
||||
__device__ clock_t clock();
|
||||
|
||||
//atomicAdd()
|
||||
__device__ int atomicAdd(int* address, int val);
|
||||
__device__ unsigned int atomicAdd(unsigned int* address,
|
||||
unsigned int val);
|
||||
|
||||
__device__ unsigned long long int atomicAdd(unsigned long long int* address,
|
||||
unsigned long long int val);
|
||||
|
||||
__device__ float atomicAdd(float* address, float val);
|
||||
|
||||
|
||||
//atomicSub()
|
||||
__device__ int atomicSub(int* address, int val);
|
||||
|
||||
__device__ unsigned int atomicSub(unsigned int* address,
|
||||
unsigned int val);
|
||||
|
||||
|
||||
//atomicExch()
|
||||
__device__ int atomicExch(int* address, int val);
|
||||
|
||||
__device__ unsigned int atomicExch(unsigned int* address,
|
||||
unsigned int val);
|
||||
|
||||
__device__ unsigned long long int atomicExch(unsigned long long int* address,
|
||||
unsigned long long int val);
|
||||
|
||||
__device__ float atomicExch(float* address, float val);
|
||||
|
||||
|
||||
//atomicMin()
|
||||
__device__ int atomicMin(int* address, int val);
|
||||
__device__ unsigned int atomicMin(unsigned int* address,
|
||||
unsigned int val);
|
||||
__device__ unsigned long long int atomicMin(unsigned long long int* address,
|
||||
unsigned long long int val);
|
||||
|
||||
|
||||
//atomicMax()
|
||||
__device__ int atomicMax(int* address, int val);
|
||||
__device__ unsigned int atomicMax(unsigned int* address,
|
||||
unsigned int val);
|
||||
__device__ unsigned long long int atomicMax(unsigned long long int* address,
|
||||
unsigned long long int val);
|
||||
|
||||
|
||||
//atomicCAS()
|
||||
__device__ int atomicCAS(int* address, int compare, int val);
|
||||
__device__ unsigned int atomicCAS(unsigned int* address,
|
||||
unsigned int compare,
|
||||
unsigned int val);
|
||||
__device__ unsigned long long int atomicCAS(unsigned long long int* address,
|
||||
unsigned long long int compare,
|
||||
unsigned long long int val);
|
||||
|
||||
|
||||
//atomicAnd()
|
||||
__device__ int atomicAnd(int* address, int val);
|
||||
__device__ unsigned int atomicAnd(unsigned int* address,
|
||||
unsigned int val);
|
||||
__device__ unsigned long long int atomicAnd(unsigned long long int* address,
|
||||
unsigned long long int val);
|
||||
|
||||
|
||||
//atomicOr()
|
||||
__device__ int atomicOr(int* address, int val);
|
||||
__device__ unsigned int atomicOr(unsigned int* address,
|
||||
unsigned int val);
|
||||
__device__ unsigned long long int atomicOr(unsigned long long int* address,
|
||||
unsigned long long int val);
|
||||
|
||||
|
||||
//atomicXor()
|
||||
__device__ int atomicXor(int* address, int val);
|
||||
__device__ unsigned int atomicXor(unsigned int* address,
|
||||
unsigned int val);
|
||||
__device__ unsigned long long int atomicXor(unsigned long long int* address,
|
||||
unsigned long long int val);
|
||||
|
||||
//atomicInc()
|
||||
__device__ unsigned int atomicInc(unsigned int* address,
|
||||
unsigned int val);
|
||||
|
||||
|
||||
//atomicDec()
|
||||
__device__ unsigned int atomicDec(unsigned int* address,
|
||||
unsigned int val);
|
||||
|
||||
|
||||
// integer intrinsic function __poc __clz __ffs __brev
|
||||
__device__ unsigned int __popc( unsigned int input);
|
||||
__device__ unsigned int __popcll( unsigned long long int input);
|
||||
__device__ unsigned int __clz(unsigned int input);
|
||||
__device__ unsigned int __clzll(unsigned long long int input);
|
||||
__device__ unsigned int __clz(int input);
|
||||
__device__ unsigned int __clzll(long long int input);
|
||||
__device__ unsigned int __ffs(unsigned int input);
|
||||
__device__ unsigned int __ffsll(unsigned long long int input);
|
||||
__device__ unsigned int __ffs(int input);
|
||||
__device__ unsigned int __ffsll(long long int input);
|
||||
__device__ unsigned int __brev( unsigned int input);
|
||||
__device__ unsigned long long int __brevll( unsigned long long int input);
|
||||
|
||||
|
||||
// warp vote function __all __any __ballot
|
||||
__device__ int __all( int input);
|
||||
__device__ int __any( int input);
|
||||
__device__ unsigned long long int __ballot( int input);
|
||||
|
||||
// warp shuffle functions
|
||||
#ifdef __cplusplus
|
||||
__device__ int __shfl(int input, int lane, int width=warpSize);
|
||||
__device__ int __shfl_up(int input, unsigned int lane_delta, int width=warpSize);
|
||||
__device__ int __shfl_down(int input, unsigned int lane_delta, int width=warpSize);
|
||||
__device__ int __shfl_xor(int input, int lane_mask, int width=warpSize);
|
||||
__device__ float __shfl(float input, int lane, int width=warpSize);
|
||||
__device__ float __shfl_up(float input, unsigned int lane_delta, int width=warpSize);
|
||||
__device__ float __shfl_down(float input, unsigned int lane_delta, int width=warpSize);
|
||||
__device__ float __shfl_xor(float input, int lane_mask, int width=warpSize);
|
||||
#else
|
||||
__device__ int __shfl(int input, int lane, int width);
|
||||
__device__ int __shfl_up(int input, unsigned int lane_delta, int width);
|
||||
__device__ int __shfl_down(int input, unsigned int lane_delta, int width);
|
||||
__device__ int __shfl_xor(int input, int lane_mask, int width);
|
||||
__device__ float __shfl(float input, int lane, int width);
|
||||
__device__ float __shfl_up(float input, unsigned int lane_delta, int width);
|
||||
__device__ float __shfl_down(float input, unsigned int lane_delta, int width);
|
||||
__device__ float __shfl_xor(float input, int lane_mask, int width);
|
||||
#endif
|
||||
|
||||
__host__ __device__ int min(int arg1, int arg2);
|
||||
__host__ __device__ int max(int arg1, int arg2);
|
||||
|
||||
__device__ __attribute__((address_space(3))) void* __get_dynamicgroupbaseptr();
|
||||
|
||||
//TODO - add a couple fast math operations here, the set here will grow :
|
||||
__device__ float __cosf(float x);
|
||||
__device__ float __expf(float x);
|
||||
__device__ float __frsqrt_rn(float x);
|
||||
__device__ float __fsqrt_rd(float x);
|
||||
__device__ float __fsqrt_rn(float x);
|
||||
__device__ float __fsqrt_ru(float x);
|
||||
__device__ float __fsqrt_rz(float x);
|
||||
__device__ float __log10f(float x);
|
||||
__device__ float __log2f(float x);
|
||||
__device__ float __logf(float x);
|
||||
__device__ float __powf(float base, float exponent);
|
||||
__device__ void __sincosf(float x, float *s, float *c) ;
|
||||
__device__ float __sinf(float x);
|
||||
__device__ float __tanf(float x);
|
||||
__device__ float __dsqrt_rd(double x);
|
||||
__device__ float __dsqrt_rn(double x);
|
||||
__device__ float __dsqrt_ru(double x);
|
||||
__device__ float __dsqrt_rz(double x);
|
||||
|
||||
/**
|
||||
* CUDA 8 device function features
|
||||
|
||||
*/
|
||||
|
||||
|
||||
/**
|
||||
* Kernel launching
|
||||
*/
|
||||
|
||||
/**
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
* @defgroup Fence Fence Functions
|
||||
* @{
|
||||
*
|
||||
*
|
||||
* @warning The HIP memory fence functions are currently not supported yet.
|
||||
* If any of those threadfence stubs are reached by the application, you should set "export HSA_DISABLE_CACHE=1" to disable L1 and L2 caches.
|
||||
*
|
||||
*
|
||||
* On AMD platforms, the threadfence* routines are currently empty stubs.
|
||||
*/
|
||||
|
||||
/**
|
||||
* @brief threadfence_block makes writes visible to threads running in same block.
|
||||
*
|
||||
* @Returns void
|
||||
*
|
||||
* @param void
|
||||
*
|
||||
* @warning __threadfence_block is a stub and map to no-op.
|
||||
*/
|
||||
__device__ void __threadfence_block(void);
|
||||
|
||||
/**
|
||||
* @brief threadfence makes wirtes visible to other threads running on same GPU.
|
||||
*
|
||||
* @Returns void
|
||||
*
|
||||
* @param void
|
||||
*
|
||||
* @warning __threadfence is a stub and map to no-op, application should set "export HSA_DISABLE_CACHE=1" to disable both L1 and L2 caches.
|
||||
*/
|
||||
__device__ void __threadfence(void) __attribute__((deprecated("Provided for compile-time compatibility, not yet functional")));
|
||||
|
||||
/**
|
||||
* @brief threadfence_system makes writes to pinned system memory visible on host CPU.
|
||||
*
|
||||
* @Returns void
|
||||
*
|
||||
* @param void
|
||||
*
|
||||
* @warning __threadfence_system is a stub and map to no-op, application should set "export HSA_DISABLE_CACHE=1" to disable both L1 and L2 caches.
|
||||
*/
|
||||
__device__ void __threadfence_system(void) __attribute__((deprecated("Provided for compile-time compatibility, not yet functional")));
|
||||
|
||||
|
||||
// doxygen end Fence Fence
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
|
||||
|
||||
#define hipThreadIdx_x (hc_get_workitem_id(0))
|
||||
#define hipThreadIdx_y (hc_get_workitem_id(1))
|
||||
#define hipThreadIdx_z (hc_get_workitem_id(2))
|
||||
|
||||
#define hipBlockIdx_x (hc_get_group_id(0))
|
||||
#define hipBlockIdx_y (hc_get_group_id(1))
|
||||
#define hipBlockIdx_z (hc_get_group_id(2))
|
||||
|
||||
#define hipBlockDim_x (hc_get_group_size(0))
|
||||
#define hipBlockDim_y (hc_get_group_size(1))
|
||||
#define hipBlockDim_z (hc_get_group_size(2))
|
||||
|
||||
#define hipGridDim_x (hc_get_num_groups(0))
|
||||
#define hipGridDim_y (hc_get_num_groups(1))
|
||||
#define hipGridDim_z (hc_get_num_groups(2))
|
||||
|
||||
// loop unrolling
|
||||
__device__ static inline void* memcpy(void* dst, void* src, size_t size)
|
||||
{
|
||||
uint64_t i = 0;
|
||||
uint64_t totalLength = size/sizeof(uint32_t);
|
||||
for(i=hipThreadIdx_x+hipBlockIdx_x*hipBlockDim_x;
|
||||
i<(totalLength/4);
|
||||
i = i + hipBlockDim_x * hipGridDim_x)
|
||||
{
|
||||
((uint32_t*)dst)[4*i] = ((uint32_t*)src)[4*i];
|
||||
((uint32_t*)dst)[4*i+1] = ((uint32_t*)src)[4*i+1];
|
||||
((uint32_t*)dst)[4*i+2] = ((uint32_t*)src)[4*i+2];
|
||||
((uint32_t*)dst)[4*i+3] = ((uint32_t*)src)[4*i+3];
|
||||
}
|
||||
if(4*i < totalLength){
|
||||
((uint32_t*)dst)[4*i] = ((uint32_t*)src)[4*i];
|
||||
((uint32_t*)dst)[4*i+1] = ((uint32_t*)src)[4*i+1];
|
||||
((uint32_t*)dst)[4*i+2] = ((uint32_t*)src)[4*i+2];
|
||||
((uint32_t*)dst)[4*i+3] = ((uint32_t*)src)[4*i+3];
|
||||
|
||||
}
|
||||
return nullptr;
|
||||
}
|
||||
|
||||
__device__ static inline void* memset(void* ptr, uint8_t val, size_t size)
|
||||
{
|
||||
uint32_t _val = 0;
|
||||
_val = (val | val << 8 | val << 16 | val << 24);
|
||||
uint64_t totalLength = size/sizeof(uint32_t);
|
||||
uint64_t i = 0;
|
||||
for(i=hipThreadIdx_x+hipBlockIdx_x*hipBlockDim_x;
|
||||
i<(totalLength/4);
|
||||
i = i + hipBlockDim_x * hipGridDim_x)
|
||||
{
|
||||
((uint32_t*)ptr)[4*i] = _val;
|
||||
((uint32_t*)ptr)[4*i+1] = _val;
|
||||
((uint32_t*)ptr)[4*i+2] = _val;
|
||||
((uint32_t*)ptr)[4*i+3] = _val;
|
||||
}
|
||||
if(4*i < totalLength){
|
||||
((uint32_t*)ptr)[4*i] = _val;
|
||||
((uint32_t*)ptr)[4*i+1] = _val;
|
||||
((uint32_t*)ptr)[4*i+2] = _val;
|
||||
((uint32_t*)ptr)[4*i+3] = _val;
|
||||
|
||||
}
|
||||
return nullptr;
|
||||
}
|
||||
|
||||
#define __syncthreads() hc_barrier(CLK_LOCAL_MEM_FENCE)
|
||||
|
||||
#define HIP_KERNEL_NAME(...) __VA_ARGS__
|
||||
|
||||
#ifdef __HCC_CPP__
|
||||
extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_launch_parm *lp, const char *kernelNameStr);
|
||||
extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, grid_launch_parm *lp, const char *kernelNameStr);
|
||||
extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, grid_launch_parm *lp, const char *kernelNameStr);
|
||||
extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, grid_launch_parm *lp, const char *kernelNameStr);
|
||||
extern void ihipPostLaunchKernel(hipStream_t stream, grid_launch_parm &lp);
|
||||
|
||||
|
||||
// Due to multiple overloaded versions of ihipPreLaunchKernel, the numBlocks3D and blockDim3D can be either size_t or dim3 types
|
||||
#define hipLaunchKernel(_kernelName, _numBlocks3D, _blockDim3D, _groupMemBytes, _stream, ...) \
|
||||
do {\
|
||||
grid_launch_parm lp;\
|
||||
lp.dynamic_group_mem_bytes = _groupMemBytes; \
|
||||
hipStream_t trueStream = (ihipPreLaunchKernel(_stream, _numBlocks3D, _blockDim3D, &lp, #_kernelName)); \
|
||||
_kernelName (lp, ##__VA_ARGS__);\
|
||||
ihipPostLaunchKernel(trueStream, lp);\
|
||||
} while(0)
|
||||
|
||||
|
||||
#elif defined (__HCC_C__)
|
||||
|
||||
//TODO - develop C interface.
|
||||
|
||||
#endif
|
||||
|
||||
/**
|
||||
* extern __shared__
|
||||
*/
|
||||
|
||||
// Macro to replace extern __shared__ declarations
|
||||
// to local variable definitions
|
||||
#define HIP_DYNAMIC_SHARED(type, var) \
|
||||
__attribute__((address_space(3))) type* var = \
|
||||
(__attribute__((address_space(3))) type*)__get_dynamicgroupbaseptr(); \
|
||||
|
||||
#define HIP_DYNAMIC_SHARED_ATTRIBUTE __attribute__((address_space(3)))
|
||||
|
||||
#endif // __HCC__
|
||||
|
||||
|
||||
/**
|
||||
* @defgroup HIP-ENV HIP Environment Variables
|
||||
* @{
|
||||
*/
|
||||
//extern int HIP_PRINT_ENV ; ///< Print all HIP-related environment variables.
|
||||
//extern int HIP_TRACE_API; ///< Trace HIP APIs.
|
||||
//extern int HIP_LAUNCH_BLOCKING ; ///< Make all HIP APIs host-synchronous
|
||||
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
|
||||
|
||||
// End doxygen API:
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
|
||||
|
||||
|
||||
|
||||
#endif
|
||||
Το diff αρχείου καταστέλλεται επειδή είναι πολύ μεγάλο
Φόρτωση Διαφορών
@@ -0,0 +1,336 @@
|
||||
/*
|
||||
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 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.
|
||||
*/
|
||||
|
||||
//#pragma once
|
||||
|
||||
#ifndef HIP_TEXTURE_H
|
||||
#define HIP_TEXTURE_H
|
||||
|
||||
/**
|
||||
* @file hcc_detail/hip_texture.h
|
||||
* @brief HIP C++ Texture API for hcc compiler
|
||||
*/
|
||||
|
||||
#include <limits.h>
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
//----
|
||||
//Texture - TODO - likely need to move this to a separate file only included with kernel compilation.
|
||||
#define hipTextureType1D 1
|
||||
|
||||
typedef enum {
|
||||
hipChannelFormatKindSigned = 0,
|
||||
hipChannelFormatKindUnsigned,
|
||||
hipChannelFormatKindFloat,
|
||||
hipChannelFormatKindNone
|
||||
|
||||
} hipChannelFormatKind;
|
||||
|
||||
typedef struct hipChannelFormatDesc {
|
||||
int x;
|
||||
int y;
|
||||
int z;
|
||||
int w;
|
||||
hipChannelFormatKind f;
|
||||
} hipChannelFormatDesc;
|
||||
|
||||
typedef enum hipTextureReadMode
|
||||
{
|
||||
hipReadModeElementType, ///< Read texture as specified element type
|
||||
//! @warning cudaReadModeNormalizedFloat is not supported.
|
||||
} hipTextureReadMode;
|
||||
|
||||
typedef enum hipTextureFilterMode
|
||||
{
|
||||
hipFilterModePoint, ///< Point filter mode.
|
||||
//! @warning cudaFilterModeLinear is not supported.
|
||||
} hipTextureFilterMode;
|
||||
|
||||
struct textureReference {
|
||||
hipTextureFilterMode filterMode;
|
||||
bool normalized;
|
||||
hipChannelFormatDesc channelDesc;
|
||||
};
|
||||
#if __cplusplus
|
||||
template <class T, int texType=hipTextureType1D, enum hipTextureReadMode=hipReadModeElementType>
|
||||
struct texture : public textureReference {
|
||||
|
||||
const T * _dataPtr; // pointer to underlying data.
|
||||
|
||||
//texture() : filterMode(hipFilterModePoint), normalized(false), _dataPtr(NULL) {};
|
||||
unsigned int width;
|
||||
unsigned int height;
|
||||
|
||||
};
|
||||
#endif
|
||||
|
||||
typedef struct hipArray {
|
||||
unsigned int width;
|
||||
unsigned int height;
|
||||
hipChannelFormatKind f;
|
||||
void* data; //FIXME: generalize this
|
||||
} hipArray;
|
||||
|
||||
|
||||
#define tex1Dfetch(_tex, _addr) (_tex._dataPtr[_addr])
|
||||
|
||||
#define tex2D(_tex, _dx, _dy) \
|
||||
_tex._dataPtr[(unsigned int)_dx + (unsigned int)_dy*(_tex.width)]
|
||||
|
||||
/**
|
||||
* @brief Allocate an array on the device.
|
||||
*
|
||||
* @param[out] array Pointer to allocated array in device memory
|
||||
* @param[in] desc Requested channel format
|
||||
* @param[in] width Requested array allocation width
|
||||
* @param[in] height Requested array allocation height
|
||||
* @param[in] flags Requested properties of allocated array
|
||||
* @return #hipSuccess, #hipErrorMemoryAllocation
|
||||
*
|
||||
* @see hipMalloc, hipMallocPitch, hipFree, hipFreeArray, hipHostMalloc, hipHostFree
|
||||
*/
|
||||
hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc,
|
||||
size_t width, size_t height = 0, unsigned int flags = 0);
|
||||
|
||||
/**
|
||||
* @brief Frees an array on the device.
|
||||
*
|
||||
* @param[in] array Pointer to array to free
|
||||
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInitializationError
|
||||
*
|
||||
* @see hipMalloc, hipMallocPitch, hipFree, hipMallocArray, hipHostMalloc, hipHostFree
|
||||
*/
|
||||
hipError_t hipFreeArray(hipArray* array);
|
||||
|
||||
/**
|
||||
* @brief Copies data between host and device.
|
||||
*
|
||||
* @param[in] dst Destination memory address
|
||||
* @param[in] dpitch Pitch of destination memory
|
||||
* @param[in] src Source memory address
|
||||
* @param[in] spitch Pitch of source memory
|
||||
* @param[in] width Width of matrix transfer (columns in bytes)
|
||||
* @param[in] height Height of matrix transfer (rows)
|
||||
* @param[in] kind Type of transfer
|
||||
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue, #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection
|
||||
*
|
||||
* @see hipMemcpy, hipMemcpyToArray, hipMemcpy2DToArray, hipMemcpyFromArray, hipMemcpyToSymbol, hipMemcpyAsync
|
||||
*/
|
||||
hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind);
|
||||
|
||||
/**
|
||||
* @brief Copies data between host and device.
|
||||
*
|
||||
* @param[in] dst Destination memory address
|
||||
* @param[in] dpitch Pitch of destination memory
|
||||
* @param[in] src Source memory address
|
||||
* @param[in] spitch Pitch of source memory
|
||||
* @param[in] width Width of matrix transfer (columns in bytes)
|
||||
* @param[in] height Height of matrix transfer (rows)
|
||||
* @param[in] kind Type of transfer
|
||||
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue, #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection
|
||||
*
|
||||
* @see hipMemcpy, hipMemcpyToArray, hipMemcpy2D, hipMemcpyFromArray, hipMemcpyToSymbol, hipMemcpyAsync
|
||||
*/
|
||||
hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src,
|
||||
size_t spitch, size_t width, size_t height, hipMemcpyKind kind);
|
||||
|
||||
/**
|
||||
* @brief Copies data between host and device.
|
||||
*
|
||||
* @param[in] dst Destination memory address
|
||||
* @param[in] dpitch Pitch of destination memory
|
||||
* @param[in] src Source memory address
|
||||
* @param[in] spitch Pitch of source memory
|
||||
* @param[in] width Width of matrix transfer (columns in bytes)
|
||||
* @param[in] height Height of matrix transfer (rows)
|
||||
* @param[in] kind Type of transfer
|
||||
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidPitchValue, #hipErrorInvalidDevicePointer, #hipErrorInvalidMemcpyDirection
|
||||
*
|
||||
* @see hipMemcpy, hipMemcpy2DToArray, hipMemcpy2D, hipMemcpyFromArray, hipMemcpyToSymbol, hipMemcpyAsync
|
||||
*/
|
||||
hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset,
|
||||
const void* src, size_t count, hipMemcpyKind kind);
|
||||
|
||||
|
||||
/**
|
||||
* @addtogroup API HIP API
|
||||
* @{
|
||||
*
|
||||
* Defines the HIP API. See the individual sections for more information.
|
||||
*/
|
||||
|
||||
// These are C++ APIs - maybe belong in separate file.
|
||||
/**
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
*-------------------------------------------------------------------------------------------------
|
||||
* @defgroup Texture Texture Reference Management
|
||||
* @{
|
||||
*
|
||||
*
|
||||
* @warning The HIP texture API implements a small subset of full texture API. Known limitations include:
|
||||
* - Only point sampling is supported.
|
||||
* - Only C++ APIs are provided.
|
||||
* - Many APIs and modes are not implemented.
|
||||
*
|
||||
* The HIP texture support is intended to allow use of texture cache on hardware where this is beneficial.
|
||||
*
|
||||
* The following CUDA APIs are not currently supported:
|
||||
* - cudaBindTexture2D
|
||||
* - cudaBindTextureToArray
|
||||
* - cudaBindTextureToMipmappedArray
|
||||
* - cudaGetChannelDesc
|
||||
* - cudaGetTextureReference
|
||||
*
|
||||
*/
|
||||
|
||||
// C API:
|
||||
#if 0
|
||||
hipChannelFormatDesc hipBindTexture(size_t *offset, struct textureReference *tex, const void *devPtr, const struct hipChannelFormatDesc *desc, size_t size=UINT_MAX)
|
||||
{
|
||||
tex->_dataPtr = devPtr;
|
||||
}
|
||||
#endif
|
||||
|
||||
/**
|
||||
* @brief Returns a channel descriptor using the specified format.
|
||||
*
|
||||
* @param[in] x X component
|
||||
* @param[in] y Y component
|
||||
* @param[in] z Z component
|
||||
* @param[in] w W component
|
||||
* @param[in] f Channel format
|
||||
* @return Channel descriptor with format f
|
||||
*
|
||||
*/
|
||||
hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f);
|
||||
|
||||
// descriptors
|
||||
template <typename T> inline hipChannelFormatDesc hipCreateChannelDesc() {
|
||||
return hipCreateChannelDesc(0, 0, 0, 0, hipChannelFormatKindNone);
|
||||
}
|
||||
template <> inline hipChannelFormatDesc hipCreateChannelDesc<int>() {
|
||||
int e = (int)sizeof(int) * 8;
|
||||
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned);
|
||||
}
|
||||
template <> inline hipChannelFormatDesc hipCreateChannelDesc<unsigned int>() {
|
||||
int e = (int)sizeof(unsigned int) * 8;
|
||||
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned);
|
||||
}
|
||||
template <> inline hipChannelFormatDesc hipCreateChannelDesc<long>() {
|
||||
int e = (int)sizeof(long) * 8;
|
||||
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindSigned);
|
||||
}
|
||||
template <> inline hipChannelFormatDesc hipCreateChannelDesc<unsigned long>() {
|
||||
int e = (int)sizeof(unsigned long) * 8;
|
||||
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindUnsigned);
|
||||
}
|
||||
template <> inline hipChannelFormatDesc hipCreateChannelDesc<float>() {
|
||||
int e = (int)sizeof(float) * 8;
|
||||
return hipCreateChannelDesc(e, 0, 0, 0, hipChannelFormatKindFloat);
|
||||
}
|
||||
|
||||
/*
|
||||
* @brief hipBindTexture Binds size bytes of the memory area pointed to by @p devPtr to the texture reference tex.
|
||||
*
|
||||
* @p desc describes how the memory is interpreted when fetching values from the texture. The @p offset parameter is an optional byte offset as with the low-level
|
||||
* hipBindTexture() function. Any memory previously bound to tex is unbound.
|
||||
*
|
||||
* @param[in] offset - Offset in bytes
|
||||
* @param[out] tex - texture to bind
|
||||
* @param[in] devPtr - Memory area on device
|
||||
* @param[in] desc - Channel format
|
||||
* @param[in] size - Size of the memory area pointed to by devPtr
|
||||
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree, #hipErrorUnknown
|
||||
**/
|
||||
template <class T, int dim, enum hipTextureReadMode readMode>
|
||||
hipError_t hipBindTexture(size_t *offset,
|
||||
struct texture<T, dim, readMode> &tex,
|
||||
const void *devPtr,
|
||||
const struct hipChannelFormatDesc *desc,
|
||||
size_t size=UINT_MAX)
|
||||
{
|
||||
tex._dataPtr = static_cast<const T*>(devPtr);
|
||||
|
||||
return hipSuccess;
|
||||
}
|
||||
|
||||
/*
|
||||
* @brief hipBindTexture Binds size bytes of the memory area pointed to by @p devPtr to the texture reference tex.
|
||||
*
|
||||
* @p desc describes how the memory is interpreted when fetching values from the texture. The @p offset parameter is an optional byte offset as with the low-level
|
||||
* hipBindTexture() function. Any memory previously bound to tex is unbound.
|
||||
*
|
||||
* @param[in] offset - Offset in bytes
|
||||
* @param[in] tex - texture to bind
|
||||
* @param[in] devPtr - Memory area on device
|
||||
* @param[in] size - Size of the memory area pointed to by devPtr
|
||||
* @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree, #hipErrorUnknown
|
||||
**/
|
||||
template <class T, int dim, enum hipTextureReadMode readMode>
|
||||
hipError_t hipBindTexture(size_t *offset,
|
||||
struct texture<T, dim, readMode> &tex,
|
||||
const void *devPtr,
|
||||
size_t size=UINT_MAX)
|
||||
{
|
||||
return hipBindTexture(offset, tex, devPtr, &tex.channelDesc, size);
|
||||
}
|
||||
|
||||
template <class T, int dim, enum hipTextureReadMode readMode>
|
||||
hipError_t hipBindTextureToArray(struct texture<T, dim, readMode> &tex, hipArray* array) {
|
||||
tex.width = array->width;
|
||||
tex.height = array->height;
|
||||
tex._dataPtr = static_cast<const T*>(array->data);
|
||||
return hipSuccess;
|
||||
}
|
||||
|
||||
/*
|
||||
* @brief Unbinds the textuer bound to @p tex
|
||||
*
|
||||
* @param[in] tex - texture to unbind
|
||||
*
|
||||
* @return #hipSuccess
|
||||
**/
|
||||
template <class T, int dim, enum hipTextureReadMode readMode>
|
||||
hipError_t hipUnbindTexture(struct texture<T, dim, readMode> &tex)
|
||||
{
|
||||
tex._dataPtr = NULL;
|
||||
|
||||
return hipSuccess;
|
||||
}
|
||||
|
||||
|
||||
|
||||
// doxygen end Texture
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
|
||||
|
||||
// End doxygen API:
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
|
||||
#endif
|
||||
|
||||
@@ -0,0 +1,36 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
|
||||
|
||||
#ifndef HIP_UTIL_H
|
||||
#define HIP_UTIL_H
|
||||
|
||||
#include <assert.h>
|
||||
#include <stdint.h>
|
||||
#include <iostream>
|
||||
#include <sstream>
|
||||
#include <list>
|
||||
#include <sys/types.h>
|
||||
#include <unistd.h>
|
||||
#include <deque>
|
||||
#include <vector>
|
||||
#include <algorithm>
|
||||
|
||||
|
||||
#endif
|
||||
@@ -0,0 +1,416 @@
|
||||
/*
|
||||
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 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.
|
||||
*/
|
||||
|
||||
/**
|
||||
* @file hcc_detail/hip_vector_types.h
|
||||
* @brief Defines the different newt vector types for HIP runtime.
|
||||
*/
|
||||
|
||||
#ifndef HIP_VECTOR_TYPES_H
|
||||
#define HIP_VECTOR_TYPES_H
|
||||
|
||||
#if defined (__HCC__) && (__hcc_workweek__ < 16032)
|
||||
#error("This version of HIP requires a newer version of HCC.");
|
||||
#endif
|
||||
|
||||
#if __HCC__
|
||||
#include <hc_short_vector.hpp>
|
||||
|
||||
using namespace hc::short_vector;
|
||||
|
||||
|
||||
//-- Signed
|
||||
// Define char vector types
|
||||
typedef hc::short_vector::char1 char1;
|
||||
typedef hc::short_vector::char2 char2;
|
||||
typedef hc::short_vector::char3 char3;
|
||||
typedef hc::short_vector::char4 char4;
|
||||
|
||||
// Define short vector types
|
||||
typedef hc::short_vector::short1 short1;
|
||||
typedef hc::short_vector::short2 short2;
|
||||
typedef hc::short_vector::short3 short3;
|
||||
typedef hc::short_vector::short4 short4;
|
||||
|
||||
// Define int vector types
|
||||
typedef hc::short_vector::int1 int1;
|
||||
typedef hc::short_vector::int2 int2;
|
||||
typedef hc::short_vector::int3 int3;
|
||||
typedef hc::short_vector::int4 int4;
|
||||
|
||||
// Define long vector types
|
||||
typedef hc::short_vector::long1 long1;
|
||||
typedef hc::short_vector::long2 long2;
|
||||
typedef hc::short_vector::long3 long3;
|
||||
typedef hc::short_vector::long4 long4;
|
||||
|
||||
// Define longlong vector types
|
||||
typedef hc::short_vector::longlong1 longlong1;
|
||||
typedef hc::short_vector::longlong2 longlong2;
|
||||
typedef hc::short_vector::longlong3 longlong3;
|
||||
typedef hc::short_vector::longlong4 longlong4;
|
||||
|
||||
|
||||
//-- Unsigned
|
||||
// Define uchar vector types
|
||||
typedef hc::short_vector::uchar1 uchar1;
|
||||
typedef hc::short_vector::uchar2 uchar2;
|
||||
typedef hc::short_vector::uchar3 uchar3;
|
||||
typedef hc::short_vector::uchar4 uchar4;
|
||||
|
||||
// Define ushort vector types
|
||||
typedef hc::short_vector::ushort1 ushort1;
|
||||
typedef hc::short_vector::ushort2 ushort2;
|
||||
typedef hc::short_vector::ushort3 ushort3;
|
||||
typedef hc::short_vector::ushort4 ushort4;
|
||||
|
||||
// Define uint vector types
|
||||
typedef hc::short_vector::uint1 uint1;
|
||||
typedef hc::short_vector::uint2 uint2;
|
||||
typedef hc::short_vector::uint3 uint3;
|
||||
typedef hc::short_vector::uint4 uint4;
|
||||
|
||||
// Define ulong vector types
|
||||
typedef hc::short_vector::ulong1 ulong1;
|
||||
typedef hc::short_vector::ulong2 ulong2;
|
||||
typedef hc::short_vector::ulong3 ulong3;
|
||||
typedef hc::short_vector::ulong4 ulong4;
|
||||
|
||||
// Define ulonglong vector types
|
||||
typedef hc::short_vector::ulonglong1 ulonglong1;
|
||||
typedef hc::short_vector::ulonglong2 ulonglong2;
|
||||
typedef hc::short_vector::ulonglong3 ulonglong3;
|
||||
typedef hc::short_vector::ulonglong4 ulonglong4;
|
||||
|
||||
|
||||
//-- Floating point
|
||||
// Define float vector types
|
||||
typedef hc::short_vector::float1 float1;
|
||||
typedef hc::short_vector::float2 float2;
|
||||
typedef hc::short_vector::float3 float3;
|
||||
typedef hc::short_vector::float4 float4;
|
||||
|
||||
// Define double vector types
|
||||
typedef hc::short_vector::double1 double1;
|
||||
typedef hc::short_vector::double2 double2;
|
||||
typedef hc::short_vector::double3 double3;
|
||||
typedef hc::short_vector::double4 double4;
|
||||
|
||||
#else
|
||||
|
||||
#define __hip_align(name, val, data) \
|
||||
__attribute__((aligned(val))) name \
|
||||
{ data }
|
||||
|
||||
struct __hip_align(char1, 1, signed char x;);
|
||||
struct __hip_align(uchar1, 1, unsigned char x;);
|
||||
|
||||
struct __hip_align(char2, 2, signed char x; signed char y;);
|
||||
struct __hip_align(uchar2, 2, unsigned char x; unsigned char y;);
|
||||
|
||||
struct char3
|
||||
{
|
||||
signed char x, y, z;
|
||||
};
|
||||
|
||||
struct uchar3
|
||||
{
|
||||
unsigned char x, y, z;
|
||||
};
|
||||
|
||||
struct __hip_align(char4, 4, signed char x; signed char y; signed char z; signed char w;);
|
||||
struct __hip_align(uchar4, 4, unsigned char x; unsigned char y; unsigned char z; unsigned char w;);
|
||||
|
||||
struct __hip_align(short1, 2, signed short x;);
|
||||
struct __hip_align(ushort1, 2, unsigned short x;);
|
||||
|
||||
struct __hip_align(short2, 4, signed short x; signed short y;);
|
||||
struct __hip_align(ushort2, 4, unsigned short x; unsigned short y;);
|
||||
|
||||
struct short3
|
||||
{
|
||||
signed short x, y, z;
|
||||
};
|
||||
|
||||
struct ushort3
|
||||
{
|
||||
unsigned short x, y, z;
|
||||
};
|
||||
|
||||
struct __hip_align(short4, 8, signed short x; signed short y; signed short z; signed short w;);
|
||||
struct __hip_align(ushort4, 8, unsigned short x; unsigned short y; unsigned short z; unsigned short w;);
|
||||
|
||||
struct __hip_align(int1, 4, signed int x;);
|
||||
struct __hip_align(uint1, 4, unsigned int x;);
|
||||
|
||||
struct __hip_align(int2, 8, signed int x; signed int y;);
|
||||
struct __hip_align(uint2, 8, unsigned int x; unsigned int y;);
|
||||
|
||||
struct int3{
|
||||
signed int x, y, z;
|
||||
};
|
||||
struct uint3{
|
||||
unsigned int x, y, z;
|
||||
};
|
||||
|
||||
struct __hip_align(int4, 16, signed int x; signed int y; signed int z; signed int w;);
|
||||
struct __hip_align(uint4, 16, unsigned int x; unsigned int y; unsigned int z; unsigned int w;);
|
||||
|
||||
struct __hip_align(long1, 8, long int x;);
|
||||
struct __hip_align(ulong1, 8, unsigned long x;);
|
||||
|
||||
struct __hip_align(long2, 16, long int x; long int y;);
|
||||
struct __hip_align(ulong2, 16, unsigned long x; unsigned long y;);
|
||||
|
||||
struct long3{
|
||||
long int x, y, z;
|
||||
};
|
||||
struct ulong3{
|
||||
unsigned long x, y, z;
|
||||
};
|
||||
|
||||
struct __hip_align(long4, 32, long int x; long int y; long int z; long int w;);
|
||||
struct __hip_align(ulong4, 32, unsigned long x; unsigned long y; unsigned long z; unsigned long w;);
|
||||
|
||||
struct float1
|
||||
{
|
||||
float x;
|
||||
};
|
||||
|
||||
struct __hip_align(float2, 8, float x; float y;);
|
||||
|
||||
struct float3
|
||||
{
|
||||
float x, y, z;
|
||||
};
|
||||
|
||||
struct __hip_align(float4, 16, float x; float y; float z; float w;);
|
||||
|
||||
struct __hip_align(longlong1, 16, long long int x;);
|
||||
struct __hip_align(ulonglong1, 16, unsigned long long int x;);
|
||||
|
||||
struct __attribute__((aligned(32))) longlong2
|
||||
{
|
||||
long long int x, y;
|
||||
};
|
||||
|
||||
struct __attribute__((aligned(32))) ulonglong2
|
||||
{
|
||||
unsigned long long int x, y;
|
||||
};
|
||||
|
||||
struct longlong3
|
||||
{
|
||||
long long int x, y, z;
|
||||
};
|
||||
|
||||
struct ulonglong3
|
||||
{
|
||||
unsigned long long int x, y, z;
|
||||
};
|
||||
|
||||
struct __attribute__((aligned(64))) longlong4
|
||||
{
|
||||
long long int x, y, z, w;
|
||||
};
|
||||
|
||||
struct __attribute__((aligned(64))) ulonglong4
|
||||
{
|
||||
unsigned long long int x, y, z, w;
|
||||
};
|
||||
|
||||
struct double1
|
||||
{
|
||||
double x;
|
||||
};
|
||||
|
||||
struct __attribute__((aligned(16))) double2
|
||||
{
|
||||
double x, y;
|
||||
};
|
||||
|
||||
struct double3
|
||||
{
|
||||
double x, y, z;
|
||||
};
|
||||
|
||||
struct __attribute__((aligned(32))) double4
|
||||
{
|
||||
double x, y, z, w;
|
||||
};
|
||||
|
||||
#endif
|
||||
|
||||
#if __HCC__
|
||||
#include"hip/hcc_detail/host_defines.h"
|
||||
#define __HIP_DEVICE__ __device__ __host__
|
||||
#else
|
||||
#define __HIP_DEVICE__
|
||||
#endif
|
||||
|
||||
__HIP_DEVICE__ char1 make_char1(signed char );
|
||||
__HIP_DEVICE__ char2 make_char2(signed char, signed char );
|
||||
__HIP_DEVICE__ char3 make_char3(signed char, signed char, signed char );
|
||||
__HIP_DEVICE__ char4 make_char4(signed char, signed char, signed char, signed char );
|
||||
|
||||
__HIP_DEVICE__ short1 make_short1(short );
|
||||
__HIP_DEVICE__ short2 make_short2(short, short );
|
||||
__HIP_DEVICE__ short3 make_short3(short, short, short );
|
||||
__HIP_DEVICE__ short4 make_short4(short, short, short, short );
|
||||
|
||||
__HIP_DEVICE__ int1 make_int1(int );
|
||||
__HIP_DEVICE__ int2 make_int2(int, int );
|
||||
__HIP_DEVICE__ int3 make_int3(int, int, int );
|
||||
__HIP_DEVICE__ int4 make_int4(int, int, int, int );
|
||||
|
||||
__HIP_DEVICE__ long1 make_long1(long );
|
||||
__HIP_DEVICE__ long2 make_long2(long, long );
|
||||
__HIP_DEVICE__ long3 make_long3(long, long, long );
|
||||
__HIP_DEVICE__ long4 make_long4(long, long, long, long );
|
||||
|
||||
__HIP_DEVICE__ longlong1 make_longlong1(long long );
|
||||
__HIP_DEVICE__ longlong2 make_longlong2(long long, long long );
|
||||
__HIP_DEVICE__ longlong3 make_longlong3(long long, long long, long long );
|
||||
__HIP_DEVICE__ longlong4 make_longlong4(long long, long long, long long, long long );
|
||||
|
||||
__HIP_DEVICE__ uchar1 make_uchar1(unsigned char );
|
||||
__HIP_DEVICE__ uchar2 make_uchar2(unsigned char, unsigned char );
|
||||
__HIP_DEVICE__ uchar3 make_uchar3(unsigned char, unsigned char, unsigned char );
|
||||
__HIP_DEVICE__ uchar4 make_uchar4(unsigned char, unsigned char, unsigned char, unsigned char );
|
||||
|
||||
__HIP_DEVICE__ ushort1 make_ushort1(unsigned short );
|
||||
__HIP_DEVICE__ ushort2 make_ushort2(unsigned short, unsigned short );
|
||||
__HIP_DEVICE__ ushort3 make_ushort3(unsigned short, unsigned short, unsigned short );
|
||||
__HIP_DEVICE__ ushort4 make_ushort4(unsigned short, unsigned short, unsigned short, unsigned short );
|
||||
|
||||
__HIP_DEVICE__ uint1 make_uint1(unsigned int );
|
||||
__HIP_DEVICE__ uint2 make_uint2(unsigned int, unsigned int );
|
||||
__HIP_DEVICE__ uint3 make_uint3(unsigned int, unsigned int, unsigned int );
|
||||
__HIP_DEVICE__ uint4 make_uint4(unsigned int, unsigned int, unsigned int, unsigned int );
|
||||
|
||||
__HIP_DEVICE__ ulong1 make_ulong1(unsigned long );
|
||||
__HIP_DEVICE__ ulong2 make_ulong2(unsigned long, unsigned long );
|
||||
__HIP_DEVICE__ ulong3 make_ulong3(unsigned long, unsigned long, unsigned long );
|
||||
__HIP_DEVICE__ ulong4 make_ulong4(unsigned long, unsigned long, unsigned long, unsigned long );
|
||||
|
||||
__HIP_DEVICE__ ulonglong1 make_ulonglong1(unsigned long long );
|
||||
__HIP_DEVICE__ ulonglong2 make_ulonglong2(unsigned long long, unsigned long long);
|
||||
__HIP_DEVICE__ ulonglong3 make_ulonglong3(unsigned long long, unsigned long long, unsigned long long);
|
||||
__HIP_DEVICE__ ulonglong4 make_ulonglong4(unsigned long long, unsigned long long, unsigned long long, unsigned long long );
|
||||
|
||||
__HIP_DEVICE__ float1 make_float1(float );
|
||||
__HIP_DEVICE__ float2 make_float2(float, float );
|
||||
__HIP_DEVICE__ float3 make_float3(float, float, float );
|
||||
__HIP_DEVICE__ float4 make_float4(float, float, float, float );
|
||||
|
||||
__HIP_DEVICE__ double1 make_double1(double );
|
||||
__HIP_DEVICE__ double2 make_double2(double, double );
|
||||
__HIP_DEVICE__ double3 make_double3(double, double, double );
|
||||
__HIP_DEVICE__ double4 make_double4(double, double, double, double );
|
||||
|
||||
/*
|
||||
///---
|
||||
// Inline functions for creating vector types from basic types
|
||||
#define ONE_COMPONENT_ACCESS(T, VT) inline VT make_ ##VT [[hc]] [[cpu]] (T x) { VT t; t.x = x; return t; };
|
||||
#define TWO_COMPONENT_ACCESS(T, VT) inline VT make_ ##VT [[hc]] [[cpu]] (T x, T y) { VT t; t.x=x; t.y=y; return t; };
|
||||
#define THREE_COMPONENT_ACCESS(T, VT) inline VT make_ ##VT [[hc]] [[cpu]] (T x, T y, T z) { VT t; t.x=x; t.y=y; t.z=z; return t; };
|
||||
#define FOUR_COMPONENT_ACCESS(T, VT) inline VT make_ ##VT [[hc]] [[cpu]] (T x, T y, T z, T w) { VT t; t.x=x; t.y=y; t.z=z; t.w=w; return t; };
|
||||
|
||||
|
||||
//signed:
|
||||
ONE_COMPONENT_ACCESS (signed char, char1);
|
||||
TWO_COMPONENT_ACCESS (signed char, char2);
|
||||
THREE_COMPONENT_ACCESS(signed char, char3);
|
||||
FOUR_COMPONENT_ACCESS (signed char, char4);
|
||||
|
||||
ONE_COMPONENT_ACCESS (short, short1);
|
||||
TWO_COMPONENT_ACCESS (short, short2);
|
||||
THREE_COMPONENT_ACCESS(short, short3);
|
||||
FOUR_COMPONENT_ACCESS (short, short4);
|
||||
|
||||
ONE_COMPONENT_ACCESS (int, int1);
|
||||
TWO_COMPONENT_ACCESS (int, int2);
|
||||
THREE_COMPONENT_ACCESS(int, int3);
|
||||
FOUR_COMPONENT_ACCESS (int, int4);
|
||||
|
||||
ONE_COMPONENT_ACCESS (long int, long1);
|
||||
TWO_COMPONENT_ACCESS (long int, long2);
|
||||
THREE_COMPONENT_ACCESS(long int, long3);
|
||||
FOUR_COMPONENT_ACCESS (long int, long4);
|
||||
|
||||
ONE_COMPONENT_ACCESS (long long int, ulong1);
|
||||
TWO_COMPONENT_ACCESS (long long int, ulong2);
|
||||
THREE_COMPONENT_ACCESS(long long int, ulong3);
|
||||
FOUR_COMPONENT_ACCESS (long long int, ulong4);
|
||||
|
||||
ONE_COMPONENT_ACCESS (long long int, longlong1);
|
||||
TWO_COMPONENT_ACCESS (long long int, longlong2);
|
||||
THREE_COMPONENT_ACCESS(long long int, longlong3);
|
||||
FOUR_COMPONENT_ACCESS (long long int, longlong4);
|
||||
|
||||
|
||||
// unsigned:
|
||||
ONE_COMPONENT_ACCESS (unsigned char, uchar1);
|
||||
TWO_COMPONENT_ACCESS (unsigned char, uchar2);
|
||||
THREE_COMPONENT_ACCESS(unsigned char, uchar3);
|
||||
FOUR_COMPONENT_ACCESS (unsigned char, uchar4);
|
||||
|
||||
ONE_COMPONENT_ACCESS (unsigned short, ushort1);
|
||||
TWO_COMPONENT_ACCESS (unsigned short, ushort2);
|
||||
THREE_COMPONENT_ACCESS(unsigned short, ushort3);
|
||||
FOUR_COMPONENT_ACCESS (unsigned short, ushort4);
|
||||
|
||||
ONE_COMPONENT_ACCESS (unsigned int, uint1);
|
||||
TWO_COMPONENT_ACCESS (unsigned int, uint2);
|
||||
THREE_COMPONENT_ACCESS(unsigned int, uint3);
|
||||
FOUR_COMPONENT_ACCESS (unsigned int, uint4);
|
||||
|
||||
ONE_COMPONENT_ACCESS (unsigned long int, ulong1);
|
||||
TWO_COMPONENT_ACCESS (unsigned long int, ulong2);
|
||||
THREE_COMPONENT_ACCESS(unsigned long int, ulong3);
|
||||
FOUR_COMPONENT_ACCESS (unsigned long int, ulong4);
|
||||
|
||||
ONE_COMPONENT_ACCESS (unsigned long long int, ulong1);
|
||||
TWO_COMPONENT_ACCESS (unsigned long long int, ulong2);
|
||||
THREE_COMPONENT_ACCESS(unsigned long long int, ulong3);
|
||||
FOUR_COMPONENT_ACCESS (unsigned long long int, ulong4);
|
||||
|
||||
ONE_COMPONENT_ACCESS (unsigned long long int, ulonglong1);
|
||||
TWO_COMPONENT_ACCESS (unsigned long long int, ulonglong2);
|
||||
THREE_COMPONENT_ACCESS(unsigned long long int, ulonglong3);
|
||||
FOUR_COMPONENT_ACCESS (unsigned long long int, ulonglong4);
|
||||
|
||||
|
||||
//Floating point
|
||||
ONE_COMPONENT_ACCESS (float, float1);
|
||||
TWO_COMPONENT_ACCESS (float, float2);
|
||||
THREE_COMPONENT_ACCESS(float, float3);
|
||||
FOUR_COMPONENT_ACCESS (float, float4);
|
||||
|
||||
ONE_COMPONENT_ACCESS (double, double1);
|
||||
TWO_COMPONENT_ACCESS (double, double2);
|
||||
THREE_COMPONENT_ACCESS(double, double3);
|
||||
FOUR_COMPONENT_ACCESS (double, double4);
|
||||
*/
|
||||
|
||||
#endif
|
||||
|
||||
@@ -0,0 +1,70 @@
|
||||
/*
|
||||
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 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.
|
||||
*/
|
||||
|
||||
/**
|
||||
* @file hcc_detail/host_defines.h
|
||||
* @brief TODO-doc
|
||||
*/
|
||||
|
||||
#ifndef HOST_DEFINES_H
|
||||
#define HOST_DEFINES_H
|
||||
|
||||
#ifdef __HCC__
|
||||
/**
|
||||
* Function and kernel markers
|
||||
*/
|
||||
#define __host__ __attribute__((cpu))
|
||||
#define __device__ __attribute__((hc))
|
||||
|
||||
#define __global__ __attribute__((hc_grid_launch))
|
||||
|
||||
#define __noinline__ __attribute__((noinline))
|
||||
#define __forceinline__ __attribute__((always_inline))
|
||||
|
||||
|
||||
|
||||
/*
|
||||
* Variable Type Qualifiers:
|
||||
*/
|
||||
// _restrict is supported by the compiler
|
||||
#define __shared__ tile_static
|
||||
#define __constant__ __attribute__((address_space(2)))
|
||||
|
||||
#else
|
||||
// Non-HCC compiler
|
||||
/**
|
||||
* Function and kernel markers
|
||||
*/
|
||||
#define __host__
|
||||
#define __device__
|
||||
|
||||
#define __global__
|
||||
|
||||
#define __noinline__
|
||||
#define __forceinline__
|
||||
|
||||
#define __shared__
|
||||
#define __constant__
|
||||
|
||||
#endif
|
||||
|
||||
#endif
|
||||
@@ -0,0 +1,125 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
//#pragma once
|
||||
|
||||
#ifndef TRACE_HELPER_H
|
||||
#define TRACE_HELPER_H
|
||||
|
||||
#include <iostream>
|
||||
#include <iomanip>
|
||||
#include <string>
|
||||
|
||||
//---
|
||||
// Helper functions to convert HIP function arguments into strings.
|
||||
// Handles POD data types as well as enumerations (ie hipMemcpyKind).
|
||||
// The implementation uses C++11 variadic templates and template specialization.
|
||||
// The hipMemcpyKind example below is a good example that shows how to implement conversion for a new HSA type.
|
||||
|
||||
|
||||
// Handy macro to convert an enumeration to a stringified version of same:
|
||||
#define CASE_STR(x) case x: return #x;
|
||||
|
||||
|
||||
// Building block functions:
|
||||
template <typename T>
|
||||
inline std::string ToHexString(T v)
|
||||
{
|
||||
std::ostringstream ss;
|
||||
ss << "0x" << std::hex << v;
|
||||
return ss.str();
|
||||
};
|
||||
|
||||
|
||||
//---
|
||||
// Template overloads for ToString to handle specific types
|
||||
|
||||
// This is the default which works for most types:
|
||||
template <typename T>
|
||||
inline std::string ToString(T v)
|
||||
{
|
||||
std::ostringstream ss;
|
||||
ss << v;
|
||||
return ss.str();
|
||||
};
|
||||
|
||||
|
||||
// hipEvent_t specialization. TODO - maybe add an event ID for debug?
|
||||
template <>
|
||||
inline std::string ToString(hipEvent_t v)
|
||||
{
|
||||
std::ostringstream ss;
|
||||
ss << v;
|
||||
return ss.str();
|
||||
};
|
||||
|
||||
|
||||
|
||||
// hipStream_t
|
||||
template <>
|
||||
inline std::string ToString(hipStream_t v)
|
||||
{
|
||||
std::ostringstream ss;
|
||||
if (v == NULL) {
|
||||
ss << "stream:<null>";
|
||||
} else {
|
||||
ss << *v;
|
||||
}
|
||||
|
||||
return ss.str();
|
||||
};
|
||||
|
||||
// hipMemcpyKind specialization
|
||||
template <>
|
||||
inline std::string ToString(hipMemcpyKind v)
|
||||
{
|
||||
switch(v) {
|
||||
CASE_STR(hipMemcpyHostToHost);
|
||||
CASE_STR(hipMemcpyHostToDevice);
|
||||
CASE_STR(hipMemcpyDeviceToHost);
|
||||
CASE_STR(hipMemcpyDeviceToDevice);
|
||||
CASE_STR(hipMemcpyDefault);
|
||||
default : return ToHexString(v);
|
||||
};
|
||||
};
|
||||
|
||||
|
||||
template <>
|
||||
inline std::string ToString(hipError_t v)
|
||||
{
|
||||
return ihipErrorString(v);
|
||||
};
|
||||
|
||||
|
||||
// Catch empty arguments case
|
||||
inline std::string ToString()
|
||||
{
|
||||
return ("");
|
||||
}
|
||||
|
||||
|
||||
//---
|
||||
// C++11 variadic template - peels off first argument, converts to string, and calls itself again to peel the next arg.
|
||||
// Strings are automatically separated by comma+space.
|
||||
template <typename T, typename... Args>
|
||||
inline std::string ToString(T first, Args... args)
|
||||
{
|
||||
return ToString(first) + ", " + ToString(args...) ;
|
||||
}
|
||||
|
||||
#endif
|
||||
@@ -0,0 +1,31 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
|
||||
#include <hip/hip_common.h>
|
||||
|
||||
#if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__)
|
||||
#include <hip/hcc_detail/hipComplex.h>
|
||||
#elif defined(__HIP_PLATFORM_NVCC__) && !defined (__HIP_PLATFORM_HCC__)
|
||||
#include <hip/nvcc_detail/hipComplex.h>
|
||||
#else
|
||||
#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
|
||||
#endif
|
||||
|
||||
@@ -0,0 +1,86 @@
|
||||
/*
|
||||
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 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.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
|
||||
// Common code included at start of every hip file.
|
||||
// Auto enable __HIP_PLATFORM_HCC__ if compiling with HCC
|
||||
// Other compiler (GCC,ICC,etc) need to set one of these macros explicitly
|
||||
#if defined(__HCC__)
|
||||
#define __HIP_PLATFORM_HCC__
|
||||
#define __HIPCC__
|
||||
|
||||
#if defined(__HCC_ACCELERATOR__) && (__HCC_ACCELERATOR__ != 0)
|
||||
#define __HIP_DEVICE_COMPILE__ 1
|
||||
#else
|
||||
#define __HIP_DEVICE_COMPILE__ 0
|
||||
#endif
|
||||
#endif
|
||||
|
||||
// Auto enable __HIP_PLATFORM_NVCC__ if compiling with NVCC
|
||||
#if defined(__NVCC__)
|
||||
#define __HIP_PLATFORM_NVCC__
|
||||
# ifdef __CUDACC__
|
||||
# define __HIPCC__
|
||||
# endif
|
||||
|
||||
#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ != 0)
|
||||
#define __HIP_DEVICE_COMPILE__ 1
|
||||
#else
|
||||
#define __HIP_DEVICE_COMPILE__ 0
|
||||
#endif
|
||||
|
||||
#endif
|
||||
|
||||
|
||||
|
||||
|
||||
#if __HIP_DEVICE_COMPILE__ == 0
|
||||
// 32-bit Atomics
|
||||
#define __HIP_ARCH_HAS_GLOBAL_INT32_ATOMICS__ (0)
|
||||
#define __HIP_ARCH_HAS_GLOBAL_FLOAT_ATOMIC_EXCH__ (0)
|
||||
#define __HIP_ARCH_HAS_SHARED_INT32_ATOMICS__ (0)
|
||||
#define __HIP_ARCH_HAS_SHARED_FLOAT_ATOMIC_EXCH__ (0)
|
||||
#define __HIP_ARCH_HAS_FLOAT_ATOMIC_ADD__ (0)
|
||||
|
||||
// 64-bit Atomics
|
||||
#define __HIP_ARCH_HAS_GLOBAL_INT64_ATOMICS__ (0)
|
||||
#define __HIP_ARCH_HAS_SHARED_INT64_ATOMICS__ (0)
|
||||
|
||||
// Doubles
|
||||
#define __HIP_ARCH_HAS_DOUBLES__ (0)
|
||||
|
||||
// Warp cross-lane operations
|
||||
#define __HIP_ARCH_HAS_WARP_VOTE__ (0)
|
||||
#define __HIP_ARCH_HAS_WARP_BALLOT__ (0)
|
||||
#define __HIP_ARCH_HAS_WARP_SHUFFLE__ (0)
|
||||
#define __HIP_ARCH_HAS_WARP_FUNNEL_SHIFT__ (0)
|
||||
|
||||
// Sync
|
||||
#define __HIP_ARCH_HAS_THREAD_FENCE_SYSTEM__ (0)
|
||||
#define __HIP_ARCH_HAS_SYNC_THREAD_EXT__ (0)
|
||||
|
||||
// Misc
|
||||
#define __HIP_ARCH_HAS_SURFACE_FUNCS__ (0)
|
||||
#define __HIP_ARCH_HAS_3DGRID__ (0)
|
||||
#define __HIP_ARCH_HAS_DYNAMIC_PARALLEL__ (0)
|
||||
#endif
|
||||
@@ -0,0 +1,30 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
|
||||
#include <hip/hip_common.h>
|
||||
|
||||
#if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__)
|
||||
#include <hip/hcc_detail/hip_fp16.h>
|
||||
#elif defined(__HIP_PLATFORM_NVCC__) && !defined (__HIP_PLATFORM_HCC__)
|
||||
#include "cuda_fp16.h"
|
||||
#else
|
||||
#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
|
||||
#endif
|
||||
@@ -0,0 +1,64 @@
|
||||
/*
|
||||
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 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.
|
||||
*/
|
||||
|
||||
//! HIP = Heterogeneous-compute Interface for Portability
|
||||
//!
|
||||
//! Define a extremely thin runtime layer that allows source code to be compiled unmodified
|
||||
//! through either AMD HCC or NVCC. Key features tend to be in the spirit
|
||||
//! and terminology of CUDA, but with a portable path to other accelerators as well:
|
||||
//
|
||||
//! Both paths support rich C++ features including classes, templates, lambdas, etc.
|
||||
//! Runtime API is C
|
||||
//! Memory management is based on pure pointers and resembles malloc/free/copy.
|
||||
//
|
||||
//! hip_runtime.h : includes everything in hip_api.h, plus math builtins and kernel launch macros.
|
||||
//! hip_runtime_api.h : Defines HIP API. This is a C header file and does not use any C++ features.
|
||||
|
||||
#pragma once
|
||||
|
||||
// Some standard header files, these are included by hc.hpp and so want to make them avail on both
|
||||
// paths to provide a consistent include env and avoid "missing symbol" errors that only appears
|
||||
// on NVCC path:
|
||||
#include <stdint.h>
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
#include <assert.h>
|
||||
|
||||
#if __cplusplus > 199711L
|
||||
#include <thread>
|
||||
#endif
|
||||
|
||||
|
||||
#include <hip/hip_common.h>
|
||||
|
||||
#if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__)
|
||||
#include <hip/hcc_detail/hip_runtime.h>
|
||||
#elif defined(__HIP_PLATFORM_NVCC__) && !defined (__HIP_PLATFORM_HCC__)
|
||||
#include <hip/nvcc_detail/hip_runtime.h>
|
||||
#else
|
||||
#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
|
||||
#endif
|
||||
|
||||
|
||||
#include <hip/hip_runtime_api.h>
|
||||
#include <hip/hip_vector_types.h>
|
||||
|
||||
@@ -0,0 +1,281 @@
|
||||
/*
|
||||
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 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.
|
||||
*/
|
||||
/**
|
||||
* @file hip_runtime_api.h
|
||||
*
|
||||
* @brief Defines the API signatures for HIP runtime.
|
||||
* This file can be compiled with a standard compiler.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
|
||||
|
||||
#include <string.h> // for getDeviceProp
|
||||
#include <hip/hip_common.h>
|
||||
|
||||
enum {
|
||||
HIP_SUCCESS = 0,
|
||||
HIP_ERROR_INVALID_VALUE,
|
||||
HIP_ERROR_NOT_INITIALIZED,
|
||||
HIP_ERROR_LAUNCH_OUT_OF_RESOURCES
|
||||
};
|
||||
|
||||
typedef struct {
|
||||
// 32-bit Atomics
|
||||
unsigned hasGlobalInt32Atomics : 1; ///< 32-bit integer atomics for global memory.
|
||||
unsigned hasGlobalFloatAtomicExch : 1; ///< 32-bit float atomic exch for global memory.
|
||||
unsigned hasSharedInt32Atomics : 1; ///< 32-bit integer atomics for shared memory.
|
||||
unsigned hasSharedFloatAtomicExch : 1; ///< 32-bit float atomic exch for shared memory.
|
||||
unsigned hasFloatAtomicAdd : 1; ///< 32-bit float atomic add in global and shared memory.
|
||||
|
||||
// 64-bit Atomics
|
||||
unsigned hasGlobalInt64Atomics : 1; ///< 64-bit integer atomics for global memory.
|
||||
unsigned hasSharedInt64Atomics : 1; ///< 64-bit integer atomics for shared memory.
|
||||
|
||||
// Doubles
|
||||
unsigned hasDoubles : 1; ///< Double-precision floating point.
|
||||
|
||||
// Warp cross-lane operations
|
||||
unsigned hasWarpVote : 1; ///< Warp vote instructions (__any, __all).
|
||||
unsigned hasWarpBallot : 1; ///< Warp ballot instructions (__ballot).
|
||||
unsigned hasWarpShuffle : 1; ///< Warp shuffle operations. (__shfl_*).
|
||||
unsigned hasFunnelShift : 1; ///< Funnel two words into one with shift&mask caps.
|
||||
|
||||
// Sync
|
||||
unsigned hasThreadFenceSystem : 1; ///< __threadfence_system.
|
||||
unsigned hasSyncThreadsExt : 1; ///< __syncthreads_count, syncthreads_and, syncthreads_or.
|
||||
|
||||
// Misc
|
||||
unsigned hasSurfaceFuncs : 1; ///< Surface functions.
|
||||
unsigned has3dGrid : 1; ///< Grid and group dims are 3D (rather than 2D).
|
||||
unsigned hasDynamicParallelism : 1; ///< Dynamic parallelism.
|
||||
} hipDeviceArch_t;
|
||||
|
||||
|
||||
//---
|
||||
// Common headers for both NVCC and HCC paths:
|
||||
|
||||
/**
|
||||
* hipDeviceProp
|
||||
*
|
||||
*/
|
||||
typedef struct hipDeviceProp_t {
|
||||
char name[256]; ///< Device name.
|
||||
size_t totalGlobalMem; ///< Size of global memory region (in bytes).
|
||||
size_t sharedMemPerBlock; ///< Size of shared memory region (in bytes).
|
||||
int regsPerBlock; ///< Registers per block.
|
||||
int warpSize; ///< Warp size.
|
||||
int maxThreadsPerBlock; ///< Max work items per work group or workgroup max size.
|
||||
int maxThreadsDim[3]; ///< Max number of threads in each dimension (XYZ) of a block.
|
||||
int maxGridSize[3]; ///< Max grid dimensions (XYZ).
|
||||
int clockRate; ///< Max clock frequency of the multiProcessors in khz.
|
||||
int memoryClockRate; ///< Max global memory clock frequency in khz.
|
||||
int memoryBusWidth; ///< Global memory bus width in bits.
|
||||
size_t totalConstMem; ///< Size of shared memory region (in bytes).
|
||||
int major; ///< Major compute capability. On HCC, this is an approximation and features may differ from CUDA CC. See the arch feature flags for portable ways to query feature caps.
|
||||
int minor; ///< Minor compute capability. On HCC, this is an approximation and features may differ from CUDA CC. See the arch feature flags for portable ways to query feature caps.
|
||||
int multiProcessorCount; ///< Number of multi-processors (compute units).
|
||||
int l2CacheSize; ///< L2 cache size.
|
||||
int maxThreadsPerMultiProcessor; ///< Maximum resident threads per multi-processor.
|
||||
int computeMode; ///< Compute mode.
|
||||
int clockInstructionRate; ///< Frequency in khz of the timer used by the device-side "clock*" instructions. New for HIP.
|
||||
hipDeviceArch_t arch; ///< Architectural feature flags. New for HIP.
|
||||
int concurrentKernels; ///< Device can possibly execute multiple kernels concurrently.
|
||||
int pciBusID; ///< PCI Bus ID.
|
||||
int pciDeviceID; ///< PCI Device ID.
|
||||
size_t maxSharedMemoryPerMultiProcessor; ///< Maximum Shared Memory Per Multiprocessor.
|
||||
int isMultiGpuBoard; ///< 1 if device is on a multi-GPU board, 0 if not.
|
||||
int canMapHostMemory; ///< Check whether HIP can map host memory
|
||||
} hipDeviceProp_t;
|
||||
|
||||
|
||||
/**
|
||||
* Memory type (for pointer attributes)
|
||||
*/
|
||||
enum hipMemoryType {
|
||||
hipMemoryTypeHost, ///< Memory is physically located on host
|
||||
hipMemoryTypeDevice ///< Memory is physically located on device. (see deviceId for specific device)
|
||||
};
|
||||
|
||||
|
||||
|
||||
/**
|
||||
* Pointer attributes
|
||||
*/
|
||||
typedef struct hipPointerAttribute_t {
|
||||
enum hipMemoryType memoryType;
|
||||
int device;
|
||||
void *devicePointer;
|
||||
void *hostPointer;
|
||||
int isManaged;
|
||||
unsigned allocationFlags; /* flags specified when memory was allocated*/
|
||||
/* peers? */
|
||||
} hipPointerAttribute_t;
|
||||
|
||||
|
||||
// hack to get these to show up in Doxygen:
|
||||
/**
|
||||
* @defgroup GlobalDefs Global enum and defines
|
||||
* @{
|
||||
*
|
||||
*/
|
||||
|
||||
|
||||
/*
|
||||
* @brief hipError_t
|
||||
* @enum
|
||||
* @ingroup Enumerations
|
||||
*/
|
||||
// Developer note - when updating these, update the hipErrorName and hipErrorString functions in NVCC and HCC paths
|
||||
// Also update the hipCUDAErrorTohipError function in NVCC path.
|
||||
|
||||
typedef enum hipError_t {
|
||||
hipSuccess = 0, ///< Successful completion.
|
||||
hipErrorOutOfMemory = 2,
|
||||
hipErrorNotInitialized = 3,
|
||||
hipErrorDeinitialized = 4,
|
||||
hipErrorProfilerDisabled = 5,
|
||||
hipErrorProfilerNotInitialized = 6,
|
||||
hipErrorProfilerAlreadyStarted = 7,
|
||||
hipErrorProfilerAlreadyStopped = 8,
|
||||
hipErrorInvalidImage = 200,
|
||||
hipErrorInvalidContext = 201, ///< Produced when input context is invalid.
|
||||
hipErrorContextAlreadyCurrent = 202,
|
||||
hipErrorMapFailed = 205,
|
||||
hipErrorUnmapFailed = 206,
|
||||
hipErrorArrayIsMapped = 207,
|
||||
hipErrorAlreadyMapped = 208,
|
||||
hipErrorNoBinaryForGpu = 209,
|
||||
hipErrorAlreadyAcquired = 210,
|
||||
hipErrorNotMapped = 211,
|
||||
hipErrorNotMappedAsArray = 212,
|
||||
hipErrorNotMappedAsPointer = 213,
|
||||
hipErrorECCNotCorrectable = 214,
|
||||
hipErrorUnsupportedLimit = 215,
|
||||
hipErrorContextAlreadyInUse = 216,
|
||||
hipErrorPeerAccessUnsupported = 217,
|
||||
hipErrorInvalidKernelFile = 218, ///< In CUDA DRV, it is CUDA_ERROR_INVALID_PTX
|
||||
hipErrorInvalidGraphicsContext = 219,
|
||||
hipErrorInvalidSource = 300,
|
||||
hipErrorFileNotFound = 301,
|
||||
hipErrorSharedObjectSymbolNotFound = 302,
|
||||
hipErrorSharedObjectInitFailed = 303,
|
||||
hipErrorOperatingSystem = 304,
|
||||
hipErrorInvalidHandle = 400,
|
||||
hipErrorNotFound = 500,
|
||||
hipErrorIllegalAddress = 700,
|
||||
|
||||
// Runtime Error Codes start here.
|
||||
hipErrorMissingConfiguration = 1001,
|
||||
hipErrorMemoryAllocation = 1002, ///< Memory allocation error.
|
||||
hipErrorInitializationError = 1003, ///< TODO comment from hipErrorInitializationError
|
||||
hipErrorLaunchFailure = 1004, ///< An exception occurred on the device while executing a kernel.
|
||||
hipErrorPriorLaunchFailure = 1005,
|
||||
hipErrorLaunchTimeOut = 1006,
|
||||
hipErrorLaunchOutOfResources = 1007, ///< Out of resources error.
|
||||
hipErrorInvalidDeviceFunction = 1008,
|
||||
hipErrorInvalidConfiguration = 1009,
|
||||
hipErrorInvalidDevice = 1010, ///< DeviceID must be in range 0...#compute-devices.
|
||||
hipErrorInvalidValue = 1011, ///< One or more of the parameters passed to the API call is NULL or not in an acceptable range.
|
||||
hipErrorInvalidDevicePointer = 1017, ///< Invalid Device Pointer
|
||||
hipErrorInvalidMemcpyDirection = 1021, ///< Invalid memory copy direction
|
||||
hipErrorUnknown = 1030, ///< Unknown error.
|
||||
hipErrorInvalidResourceHandle = 1033, ///< Resource handle (hipEvent_t or hipStream_t) invalid.
|
||||
hipErrorNotReady = 1034, ///< Indicates that asynchronous operations enqueued earlier are not ready. This is not actually an error, but is used to distinguish from hipSuccess (which indicates completion). APIs that return this error include hipEventQuery and hipStreamQuery.
|
||||
hipErrorNoDevice = 1038, ///< Call to hipGetDeviceCount returned 0 devices
|
||||
hipErrorPeerAccessAlreadyEnabled = 1050, ///< Peer access was already enabled from the current device.
|
||||
|
||||
hipErrorPeerAccessNotEnabled = 1051, ///< Peer access was never enabled from the current device.
|
||||
hipErrorRuntimeMemory = 1052, ///< HSA runtime memory call returned error. Typically not seen in production systems.
|
||||
hipErrorRuntimeOther = 1053, ///< HSA runtime call other than memory returned error. Typically not seen in production systems.
|
||||
hipErrorHostMemoryAlreadyRegistered = 1061, ///< Produced when trying to lock a page-locked memory.
|
||||
hipErrorHostMemoryNotRegistered = 1062, ///< Produced when trying to unlock a non-page-locked memory.
|
||||
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.
|
||||
hipDeviceAttributeMemoryClockRate, ///< Peak memory clock frequency in kilohertz.
|
||||
hipDeviceAttributeMemoryBusWidth, ///< Global memory bus width in bits.
|
||||
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.
|
||||
hipDeviceAttributeConcurrentKernels, ///< Device can possibly execute multiple kernels concurrently.
|
||||
hipDeviceAttributePciBusId, ///< PCI Bus ID.
|
||||
hipDeviceAttributePciDeviceId, ///< PCI Device ID.
|
||||
hipDeviceAttributeMaxSharedMemoryPerMultiprocessor, ///< Maximum Shared Memory Per Multiprocessor.
|
||||
hipDeviceAttributeIsMultiGpuBoard, ///< Multiple GPU devices.
|
||||
} hipDeviceAttribute_t;
|
||||
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
|
||||
#if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__)
|
||||
#include "hip/hcc_detail/hip_runtime_api.h"
|
||||
#elif defined(__HIP_PLATFORM_NVCC__) && !defined (__HIP_PLATFORM_HCC__)
|
||||
#include "hip/nvcc_detail/hip_runtime_api.h"
|
||||
#else
|
||||
#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
|
||||
#endif
|
||||
|
||||
|
||||
/**
|
||||
* @brief: C++ wrapper for hipMalloc
|
||||
*
|
||||
* Perform automatic type conversion to eliminate need for excessive typecasting (ie void**)
|
||||
*
|
||||
* @see hipMalloc
|
||||
*/
|
||||
#ifdef __cplusplus
|
||||
template<class T>
|
||||
static inline hipError_t hipMalloc ( T** devPtr, size_t size)
|
||||
{
|
||||
return hipMalloc((void**)devPtr, size);
|
||||
}
|
||||
|
||||
// Provide an override to automatically typecast the pointer type from void**, and also provide a default for the flags.
|
||||
template<class T>
|
||||
static inline hipError_t hipHostMalloc( T** ptr, size_t size, unsigned int flags = hipHostMallocDefault)
|
||||
{
|
||||
return hipHostMalloc((void**)ptr, size, flags);
|
||||
}
|
||||
#endif
|
||||
@@ -0,0 +1,37 @@
|
||||
/*
|
||||
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 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.
|
||||
*/
|
||||
//! hip_vector_types.h : Defines the HIP vector types.
|
||||
|
||||
#pragma once
|
||||
|
||||
#include <hip/hip_common.h>
|
||||
|
||||
|
||||
#if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__)
|
||||
#if __cplusplus
|
||||
#include <hip/hcc_detail/hip_vector_types.h>
|
||||
#endif
|
||||
#elif defined(__HIP_PLATFORM_NVCC__) && !defined (__HIP_PLATFORM_HCC__)
|
||||
#include <vector_types.h>
|
||||
#else
|
||||
#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
|
||||
#endif
|
||||
@@ -0,0 +1,66 @@
|
||||
/*
|
||||
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 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.
|
||||
*/
|
||||
|
||||
//! HIP = Heterogeneous-compute Interface for Portability
|
||||
//!
|
||||
//! Define a extremely thin runtime layer that allows source code to be compiled unmodified
|
||||
//! through either AMD HCC or NVCC. Key features tend to be in the spirit
|
||||
//! and terminology of CUDA, but with a portable path to other accelerators as well.
|
||||
//!
|
||||
//! This is the master include file for hipblas, wrapping around hcblas and cublas "version 1"
|
||||
//
|
||||
|
||||
#pragma once
|
||||
|
||||
enum hipblasStatus_t {
|
||||
HIPBLAS_STATUS_SUCCESS, // Function succeeds
|
||||
HIPBLAS_STATUS_NOT_INITIALIZED, // HIPBLAS library not initialized
|
||||
HIPBLAS_STATUS_ALLOC_FAILED, // resource allocation failed
|
||||
HIPBLAS_STATUS_INVALID_VALUE, // unsupported numerical value was passed to function
|
||||
HIPBLAS_STATUS_MAPPING_ERROR, // access to GPU memory space failed
|
||||
HIPBLAS_STATUS_EXECUTION_FAILED, // GPU program failed to execute
|
||||
HIPBLAS_STATUS_INTERNAL_ERROR, // an internal HIPBLAS operation failed
|
||||
HIPBLAS_STATUS_NOT_SUPPORTED // cublas supports this, but not hcblas
|
||||
};
|
||||
|
||||
enum hipblasOperation_t {
|
||||
HIPBLAS_OP_N,
|
||||
HIPBLAS_OP_T,
|
||||
HIPBLAS_OP_C
|
||||
};
|
||||
|
||||
// Some standard header files, these are included by hc.hpp and so want to make them avail on both
|
||||
// paths to provide a consistent include env and avoid "missing symbol" errors that only appears
|
||||
// on NVCC path:
|
||||
|
||||
#if defined(__HIP_PLATFORM_HCC__) and not defined (__HIP_PLATFORM_NVCC__)
|
||||
#include <hcc_detail/hip_blas.h>
|
||||
#elif defined(__HIP_PLATFORM_NVCC__) and not defined (__HIP_PLATFORM_HCC__)
|
||||
#include <nvcc_detail/hip_blas.h>
|
||||
#else
|
||||
#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
|
||||
#endif
|
||||
|
||||
|
||||
|
||||
|
||||
|
||||
@@ -0,0 +1,108 @@
|
||||
#ifndef HIPCOMPLEX_H
|
||||
#define HIPCOMPLEX_H
|
||||
|
||||
#include"cuComplex.h"
|
||||
|
||||
typedef cuFloatComplex hipFloatComplex;
|
||||
|
||||
__device__ __host__ static inline float hipCrealf(hipFloatComplex z){
|
||||
return cuCrealf(z);
|
||||
}
|
||||
|
||||
__device__ __host__ static inline float hipCimagf(hipFloatComplex z){
|
||||
return cuCimagf(z);
|
||||
}
|
||||
|
||||
__device__ __host__ static inline hipFloatComplex make_hipFloatComplex(float a, float b){
|
||||
return make_cuFloatComplex(a, b);
|
||||
}
|
||||
|
||||
__device__ __host__ static inline hipFloatComplex hipConjf(hipFloatComplex z){
|
||||
return cuConjf(z);
|
||||
}
|
||||
|
||||
__device__ __host__ static inline float hipCsqabsf(hipFloatComplex z){
|
||||
return cuCabsf(z) * cuCabsf(z);
|
||||
}
|
||||
|
||||
__device__ __host__ static inline hipFloatComplex hipCaddf(hipFloatComplex p, hipFloatComplex q){
|
||||
return cuCaddf(p, q);
|
||||
}
|
||||
|
||||
__device__ __host__ static inline hipFloatComplex hipCsubf(hipFloatComplex p, hipFloatComplex q){
|
||||
return cuCsubf(p, q);
|
||||
}
|
||||
|
||||
__device__ __host__ static inline hipFloatComplex hipCmulf(hipFloatComplex p, hipFloatComplex q){
|
||||
return cuCmulf(p, q);
|
||||
}
|
||||
|
||||
__device__ __host__ static inline hipFloatComplex hipCdivf(hipFloatComplex p, hipFloatComplex q){
|
||||
return cuCdivf(p, q);
|
||||
}
|
||||
|
||||
__device__ __host__ static inline float hipCabsf(hipFloatComplex z){
|
||||
return cuCabsf(p, q);
|
||||
}
|
||||
|
||||
typedef cuDoubleComplex hipDoubleComplex;
|
||||
|
||||
__device__ __host__ static inline double hipCreal(hipDoubleComplex z){
|
||||
return cuCreal(z);
|
||||
}
|
||||
|
||||
__device__ __host__ static inline double hipCimag(hipDoubleComplex z){
|
||||
return cuCimag(z);
|
||||
}
|
||||
|
||||
__device__ __host__ static inline hipDoubleComplex make_hipDoubleComplex(double a, double b){
|
||||
return make_cuDoubleComplex(a, b);
|
||||
}
|
||||
|
||||
__device__ __host__ static inline hipDoubleComplex hipConj(hipDoubleComplex z){
|
||||
return cuConj(z);
|
||||
}
|
||||
|
||||
__device__ __host__ static inline hipDoubleComplex hipCsqabs(hipDoubleComplex z){
|
||||
return cuCabs(z) * cuCabs(z);
|
||||
}
|
||||
|
||||
__device__ __host__ static inline hipDoubleComplex hipCadd(hipDoubleComplex p, hipDoubleComplex q){
|
||||
return cuCadd(p, q);
|
||||
}
|
||||
|
||||
__device__ __host__ static inline hipDoubleComplex hipCsub(hipDoubleComplex p, hipDoubleComplex q){
|
||||
return cuCsub(p, q);
|
||||
}
|
||||
|
||||
__device__ __host__ static inline hipDoubleComplex hipCdiv(hipDoubleComplex p, hipDoubleComplex q){
|
||||
return cuCdiv(p, q);
|
||||
}
|
||||
|
||||
__device__ __host__ static inline double hipCabs(hipDoubleComplex z){
|
||||
return cuCabs(z);
|
||||
}
|
||||
|
||||
typedef cuFloatComplex hipComplex;
|
||||
|
||||
__device__ __host__ static inline hipComplex make_Complex(float x, float y){
|
||||
return make_cuComplex(x, y);
|
||||
}
|
||||
|
||||
__device__ __host__ static inline hipFloatComplex hipComplexDoubleToFloat(hipDoubleComplex z){
|
||||
return cuComplexDoubleToFloat(z);
|
||||
}
|
||||
|
||||
__device__ __host__ static inline hipDoubleComplex hipComplexFloatToDouble(hipFloatComplex z){
|
||||
return cuComplexFloatToDouble(z);
|
||||
}
|
||||
|
||||
__device__ __host__ static inline hipComplex hipCfmaf(hipComplex p, hipComplex q, hipComplex r){
|
||||
return cuCfmaf(p, q, r);
|
||||
}
|
||||
|
||||
__device__ __host__ static inline hipDoubleComplex hipCfma(hipComplex p, hipComplex q, hipComplex r){
|
||||
return cuCfma(p, q, r);
|
||||
}
|
||||
|
||||
#endif
|
||||
@@ -0,0 +1,268 @@
|
||||
/*
|
||||
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 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.
|
||||
*/
|
||||
#pragma once
|
||||
|
||||
#include <cuda_runtime_api.h>
|
||||
#include <cublas.h>
|
||||
#include <cublas_v2.h>
|
||||
|
||||
//HGSOS for Kalmar leave it as C++, only cublas needs C linkage.
|
||||
|
||||
#ifdef __cplusplus
|
||||
extern "C" {
|
||||
#endif
|
||||
|
||||
|
||||
typedef cublasHandle_t hipblasHandle_t ;
|
||||
typedef cuComplex hipComplex;
|
||||
|
||||
/* Unsupported types
|
||||
"cublasFillMode_t",
|
||||
"cublasDiagType_t",
|
||||
"cublasSideMode_t",
|
||||
"cublasPointerMode_t",
|
||||
"cublasAtomicsMode_t",
|
||||
"cublasDataType_t"
|
||||
*/
|
||||
|
||||
|
||||
inline static cublasOperation_t hipOperationToCudaOperation( hipblasOperation_t op)
|
||||
{
|
||||
switch (op)
|
||||
{
|
||||
case HIPBLAS_OP_N:
|
||||
return CUBLAS_OP_N;
|
||||
|
||||
case HIPBLAS_OP_T:
|
||||
return CUBLAS_OP_T;
|
||||
|
||||
case HIPBLAS_OP_C:
|
||||
return CUBLAS_OP_C;
|
||||
|
||||
default:
|
||||
throw "Non existent OP";
|
||||
}
|
||||
}
|
||||
|
||||
inline static hipblasOperation_t CudaOperationToHIPOperation( cublasOperation_t op)
|
||||
{
|
||||
switch (op)
|
||||
{
|
||||
case CUBLAS_OP_N :
|
||||
return HIPBLAS_OP_N;
|
||||
|
||||
case CUBLAS_OP_T :
|
||||
return HIPBLAS_OP_T;
|
||||
|
||||
case CUBLAS_OP_C :
|
||||
return HIPBLAS_OP_C;
|
||||
|
||||
default:
|
||||
throw "Non existent OP";
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
inline static hipblasStatus_t hipCUBLASStatusToHIPStatus(cublasStatus_t cuStatus)
|
||||
{
|
||||
switch(cuStatus)
|
||||
{
|
||||
case CUBLAS_STATUS_SUCCESS:
|
||||
return HIPBLAS_STATUS_SUCCESS;
|
||||
case CUBLAS_STATUS_NOT_INITIALIZED:
|
||||
return HIPBLAS_STATUS_NOT_INITIALIZED;
|
||||
case CUBLAS_STATUS_ALLOC_FAILED:
|
||||
return HIPBLAS_STATUS_ALLOC_FAILED;
|
||||
case CUBLAS_STATUS_INVALID_VALUE:
|
||||
return HIPBLAS_STATUS_INVALID_VALUE;
|
||||
case CUBLAS_STATUS_MAPPING_ERROR:
|
||||
return HIPBLAS_STATUS_MAPPING_ERROR;
|
||||
case CUBLAS_STATUS_EXECUTION_FAILED:
|
||||
return HIPBLAS_STATUS_EXECUTION_FAILED;
|
||||
case CUBLAS_STATUS_INTERNAL_ERROR:
|
||||
return HIPBLAS_STATUS_INTERNAL_ERROR;
|
||||
case CUBLAS_STATUS_NOT_SUPPORTED:
|
||||
return HIPBLAS_STATUS_NOT_SUPPORTED;
|
||||
default:
|
||||
throw "Unimplemented status";
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
inline static hipblasStatus_t hipblasCreate(hipblasHandle_t* handle) {
|
||||
return hipCUBLASStatusToHIPStatus(cublasCreate(&*handle));
|
||||
}
|
||||
|
||||
//TODO broke common API semantics, think about this again.
|
||||
inline static hipblasStatus_t hipblasDestroy(hipblasHandle_t handle) {
|
||||
return hipCUBLASStatusToHIPStatus(cublasDestroy(handle));
|
||||
}
|
||||
|
||||
//note: no handle
|
||||
inline static hipblasStatus_t hipblasSetVector(int n, int elemSize, const void *x, int incx, void *y, int incy){
|
||||
return hipCUBLASStatusToHIPStatus(cublasSetVector(n, elemSize, x, incx, y, incy)); //HGSOS no need for handle
|
||||
}
|
||||
|
||||
//note: no handle
|
||||
inline static hipblasStatus_t hipblasGetVector(int n, int elemSize, const void *x, int incx, void *y, int incy){
|
||||
return hipCUBLASStatusToHIPStatus(cublasGetVector(n, elemSize, x, incx, y, incy)); //HGSOS no need for handle
|
||||
}
|
||||
|
||||
//note: no handle
|
||||
inline static hipblasStatus_t hipblasSetMatrix(int rows, int cols, int elemSize, const void *A, int lda, void *B, int ldb){
|
||||
return hipCUBLASStatusToHIPStatus(cublasSetMatrix(rows, cols, elemSize, A, lda, B, ldb));
|
||||
}
|
||||
|
||||
//note: no handle
|
||||
inline static hipblasStatus_t hipblasGetMatrix(int rows, int cols, int elemSize, const void *A, int lda, void *B, int ldb){
|
||||
return hipCUBLASStatusToHIPStatus(cublasGetMatrix(rows, cols, elemSize, A, lda, B, ldb));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasSasum(hipblasHandle_t handle, int n, float *x, int incx, float *result){
|
||||
return hipCUBLASStatusToHIPStatus(cublasSasum(handle, n, x, incx, result));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasDasum(hipblasHandle_t handle, int n, double *x, int incx, double *result){
|
||||
return hipCUBLASStatusToHIPStatus(cublasDasum( handle, n, x, incx, result));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasSasumBatched(hipblasHandle_t handle, int n, float *x, int incx, float *result, int batchCount){
|
||||
//TODO warn user that function was demoted to ignore batch
|
||||
return hipCUBLASStatusToHIPStatus(cublasSasum( handle, n, x, incx, result));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasDasumBatched(hipblasHandle_t handle, int n, double *x, int incx, double *result, int batchCount){
|
||||
//TODO warn user that function was demoted to ignore batch
|
||||
return hipCUBLASStatusToHIPStatus(cublasDasum(handle, n, x, incx, result));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasSaxpy(hipblasHandle_t handle, int n, const float *alpha, const float *x, int incx, float *y, int incy) {
|
||||
return hipCUBLASStatusToHIPStatus(cublasSaxpy(handle, n, alpha, x, incx, y, incy));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasSaxpyBatched(hipblasHandle_t handle, int n, const float *alpha, const float *x, int incx, float *y, int incy, int batchCount){
|
||||
//TODO warn user that function was demoted to ignore batch
|
||||
return hipCUBLASStatusToHIPStatus(cublasSaxpy(handle, n, alpha, x, incx, y, incy));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasScopy(hipblasHandle_t handle, int n, const float *x, int incx, float *y, int incy){
|
||||
return hipCUBLASStatusToHIPStatus(cublasScopy( handle, n, x, incx, y, incy));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasDcopy(hipblasHandle_t handle, int n, const double *x, int incx, double *y, int incy){
|
||||
return hipCUBLASStatusToHIPStatus(cublasDcopy( handle, n, x, incx, y, incy));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasScopyBatched(hipblasHandle_t handle, int n, const float *x, int incx, float *y, int incy, int batchCount){
|
||||
//TODO warn user that function was demoted to ignore batch
|
||||
return hipCUBLASStatusToHIPStatus(cublasScopy( handle, n, x, incx, y, incy));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasDcopyBatched(hipblasHandle_t handle, int n, const double *x, int incx, double *y, int incy, int batchCount){
|
||||
//TODO warn user that function was demoted to ignore batch
|
||||
return hipCUBLASStatusToHIPStatus(cublasDcopy( handle, n, x, incx, y, incy));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasSdot (hipblasHandle_t handle, int n, const float *x, int incx, const float *y, int incy, float *result){
|
||||
return hipCUBLASStatusToHIPStatus(cublasSdot ( handle, n, x, incx, y, incy, result));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasDdot (hipblasHandle_t handle, int n, const double *x, int incx, const double *y, int incy, double *result){
|
||||
return hipCUBLASStatusToHIPStatus(cublasDdot ( handle, n, x, incx, y, incy, result));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasSdotBatched (hipblasHandle_t handle, int n, const float *x, int incx, const float *y, int incy, float *result, int batchCount){
|
||||
//TODO warn user that function was demoted to ignore batch
|
||||
return hipCUBLASStatusToHIPStatus(cublasSdot ( handle, n, x, incx, y, incy, result));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasDdotBatched (hipblasHandle_t handle, int n, const double *x, int incx, const double *y, int incy, double *result, int batchCount){
|
||||
//TODO warn user that function was demoted to ignore batch
|
||||
return hipCUBLASStatusToHIPStatus(cublasDdot ( handle, n, x, incx, y, incy, result));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasSscal(hipblasHandle_t handle, int n, const float *alpha, float *x, int incx){
|
||||
return hipCUBLASStatusToHIPStatus(cublasSscal(handle, n, alpha, x, incx));
|
||||
}
|
||||
inline static hipblasStatus_t hipblasDscal(hipblasHandle_t handle, int n, const double *alpha, double *x, int incx){
|
||||
return hipCUBLASStatusToHIPStatus(cublasDscal(handle, n, alpha, x, incx));
|
||||
}
|
||||
inline static hipblasStatus_t hipblasSscalBatched(hipblasHandle_t handle, int n, const float *alpha, float *x, int incx, int batchCount){
|
||||
//TODO warn user that function was demoted to ignore batch
|
||||
return hipCUBLASStatusToHIPStatus(cublasSscal(handle, n, alpha, x, incx));
|
||||
}
|
||||
inline static hipblasStatus_t hipblasDscalBatched(hipblasHandle_t handle, int n, const double *alpha, double *x, int incx, int batchCount){
|
||||
//TODO warn user that function was demoted to ignore batch
|
||||
return hipCUBLASStatusToHIPStatus(cublasDscal(handle, n, alpha, x, incx));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasSgemv(hipblasHandle_t handle, hipblasOperation_t trans, int m, int n, const float *alpha, float *A, int lda,
|
||||
float *x, int incx, const float *beta, float *y, int incy){
|
||||
return hipCUBLASStatusToHIPStatus(cublasSgemv(handle, hipOperationToCudaOperation(trans), m, n, alpha, A, lda, x, incx, beta, y, incy));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasSgemvBatched(hipblasHandle_t handle, hipblasOperation_t trans, int m, int n, const float *alpha, float *A, int lda,
|
||||
float *x, int incx, const float *beta, float *y, int incy, int batchCount){
|
||||
//TODO warn user that function was demoted to ignore batch
|
||||
return hipCUBLASStatusToHIPStatus(cublasSgemv(handle, hipOperationToCudaOperation(trans), m, n, alpha, A, lda, x, incx, beta, y, incy));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasSger(hipblasHandle_t handle, int m, int n, const float *alpha, const float *x, int incx, const float *y, int incy, float *A, int lda){
|
||||
return hipCUBLASStatusToHIPStatus(cublasSger(handle, m, n, alpha, x, incx, y, incy, A, lda));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasSgerBatched(hipblasHandle_t handle, int m, int n, const float *alpha, const float *x, int incx, const float *y, int incy, float *A, int lda, int batchCount){
|
||||
//TODO warn user that function was demoted to ignore batch
|
||||
return hipCUBLASStatusToHIPStatus(cublasSger(handle, m, n, alpha, x, incx, y, incy, A, lda));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasSgemm(hipblasHandle_t handle, hipblasOperation_t transa, hipblasOperation_t transb,
|
||||
int m, int n, int k, const float *alpha, float *A, int lda, float *B, int ldb, const float *beta, float *C, int ldc){
|
||||
return hipCUBLASStatusToHIPStatus(cublasSgemm( handle, hipOperationToCudaOperation(transa), hipOperationToCudaOperation(transb), m, n, k, alpha, A, lda, B, ldb, beta, C, ldc));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasCgemm(hipblasHandle_t handle, hipblasOperation_t transa, hipblasOperation_t transb,
|
||||
int m, int n, int k, const hipComplex *alpha, hipComplex *A, int lda, hipComplex *B, int ldb, const hipComplex *beta, hipComplex *C, int ldc){
|
||||
return hipCUBLASStatusToHIPStatus(cublasCgemm( handle, hipOperationToCudaOperation(transa), hipOperationToCudaOperation(transb), m, n, k, alpha, A, lda, B, ldb, beta, C, ldc));
|
||||
}
|
||||
|
||||
inline static hipblasStatus_t hipblasSgemmBatched(hipblasHandle_t handle, hipblasOperation_t transa, hipblasOperation_t transb,
|
||||
int m, int n, int k, const float *alpha, float *A, int lda, float *B, int ldb, const float *beta, float *C, int ldc, int batchCount){
|
||||
//TODO incompatible API
|
||||
return HIPBLAS_STATUS_NOT_SUPPORTED;
|
||||
//return hipCUBLASStatusToHIPStatus(cublasSgemmBatched( handle, hipOperationToCudaOperation(transa), hipOperationToCudaOperation(transb), m, n, k, alpha, A, lda, B, ldb, beta, C, ldc, batchCount));
|
||||
}
|
||||
|
||||
|
||||
inline static hipblasStatus_t hipblasCgemmBatched(hipblasHandle_t handle, hipblasOperation_t transa, hipblasOperation_t transb,
|
||||
int m, int n, int k, const hipComplex *alpha, hipComplex *A, int lda, hipComplex *B, int ldb, const hipComplex *beta, hipComplex *C, int ldc, int batchCount){
|
||||
|
||||
//TODO incompatible API
|
||||
return HIPBLAS_STATUS_NOT_SUPPORTED;
|
||||
//return hipCUBLASStatusToHIPStatus(cublasCgemmBatched( handle, hipOperationToCudaOperation(transa), hipOperationToCudaOperation(transb), m, n, k, alpha, A, lda, B, ldb, beta, C, ldc, batchCount));
|
||||
}
|
||||
|
||||
#ifdef __cplusplus
|
||||
}
|
||||
#endif
|
||||
|
||||
|
||||
@@ -0,0 +1,106 @@
|
||||
/*
|
||||
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 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.
|
||||
*/
|
||||
#pragma once
|
||||
|
||||
#include <cuda_runtime.h>
|
||||
|
||||
#include <hip/hip_runtime_api.h>
|
||||
|
||||
#define HIP_KERNEL_NAME(...) __VA_ARGS__
|
||||
|
||||
typedef int hipLaunchParm ;
|
||||
|
||||
#define hipLaunchKernel(kernelName, numblocks, numthreads, memperblock, streamId, ...) \
|
||||
do {\
|
||||
kernelName<<<numblocks,numthreads,memperblock,streamId>>>(0, ##__VA_ARGS__);\
|
||||
} while(0)
|
||||
|
||||
|
||||
#define hipReadModeElementType cudaReadModeElementType
|
||||
|
||||
#ifdef __CUDA_ARCH__
|
||||
|
||||
|
||||
// 32-bit Atomics:
|
||||
#define __HIP_ARCH_HAS_GLOBAL_INT32_ATOMICS__ (__CUDA_ARCH__ >= 110)
|
||||
#define __HIP_ARCH_HAS_GLOBAL_FLOAT_ATOMIC_EXCH__ (__CUDA_ARCH__ >= 110)
|
||||
#define __HIP_ARCH_HAS_SHARED_INT32_ATOMICS__ (__CUDA_ARCH__ >= 120)
|
||||
#define __HIP_ARCH_HAS_SHARED_FLOAT_ATOMIC_EXCH__ (__CUDA_ARCH__ >= 120)
|
||||
#define __HIP_ARCH_HAS_FLOAT_ATOMIC_ADD__
|
||||
|
||||
// 64-bit Atomics:
|
||||
#define __HIP_ARCH_HAS_GLOBAL_INT64_ATOMICS__ (__CUDA_ARCH__ >= 200)
|
||||
#define __HIP_ARCH_HAS_SHARED_INT64_ATOMICS__ (__CUDA_ARCH__ >= 120)
|
||||
|
||||
// Doubles
|
||||
#define __HIP_ARCH_HAS_DOUBLES__ (__CUDA_ARCH__ >= 120)
|
||||
|
||||
//warp cross-lane operations:
|
||||
#define __HIP_ARCH_HAS_WARP_VOTE__ (__CUDA_ARCH__ >= 120)
|
||||
#define __HIP_ARCH_HAS_WARP_BALLOT__ (__CUDA_ARCH__ >= 200)
|
||||
#define __HIP_ARCH_HAS_WARP_SHUFFLE__ (__CUDA_ARCH__ >= 300)
|
||||
#define __HIP_ARCH_HAS_WARP_FUNNEL_SHIFT__ (__CUDA_ARCH__ >= 350)
|
||||
|
||||
//sync
|
||||
#define __HIP_ARCH_HAS_THREAD_FENCE_SYSTEM__ (__CUDA_ARCH__ >= 200)
|
||||
#define __HIP_ARCH_HAS_SYNC_THREAD_EXT__ (__CUDA_ARCH__ >= 200)
|
||||
|
||||
// misc
|
||||
#define __HIP_ARCH_HAS_SURFACE_FUNCS__ (__CUDA_ARCH__ >= 200)
|
||||
#define __HIP_ARCH_HAS_3DGRID__ (__CUDA_ARCH__ >= 200)
|
||||
#define __HIP_ARCH_HAS_DYNAMIC_PARALLEL__ (__CUDA_ARCH__ >= 350)
|
||||
|
||||
#endif
|
||||
|
||||
#ifdef __CUDACC__
|
||||
|
||||
|
||||
|
||||
|
||||
#define hipThreadIdx_x threadIdx.x
|
||||
#define hipThreadIdx_y threadIdx.y
|
||||
#define hipThreadIdx_z threadIdx.z
|
||||
|
||||
#define hipBlockIdx_x blockIdx.x
|
||||
#define hipBlockIdx_y blockIdx.y
|
||||
#define hipBlockIdx_z blockIdx.z
|
||||
|
||||
#define hipBlockDim_x blockDim.x
|
||||
#define hipBlockDim_y blockDim.y
|
||||
#define hipBlockDim_z blockDim.z
|
||||
|
||||
#define hipGridDim_x gridDim.x
|
||||
#define hipGridDim_y gridDim.y
|
||||
#define hipGridDim_z gridDim.z
|
||||
|
||||
/**
|
||||
* extern __shared__
|
||||
*/
|
||||
|
||||
#define HIP_DYNAMIC_SHARED(type, var) \
|
||||
extern __shared__ type var[]; \
|
||||
|
||||
#define HIP_DYNAMIC_SHARED_ATTRIBUTE
|
||||
|
||||
#endif
|
||||
|
||||
|
||||
@@ -0,0 +1,812 @@
|
||||
/*
|
||||
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 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.
|
||||
*/
|
||||
#pragma once
|
||||
|
||||
#include <cuda_runtime_api.h>
|
||||
#include <cuda.h>
|
||||
|
||||
#ifdef __cplusplus
|
||||
extern "C" {
|
||||
#endif
|
||||
|
||||
//TODO -move to include/hip_runtime_api.h as a common implementation.
|
||||
/**
|
||||
* Memory copy types
|
||||
*
|
||||
*/
|
||||
typedef enum hipMemcpyKind {
|
||||
hipMemcpyHostToHost
|
||||
,hipMemcpyHostToDevice
|
||||
,hipMemcpyDeviceToHost
|
||||
,hipMemcpyDeviceToDevice
|
||||
,hipMemcpyDefault
|
||||
} hipMemcpyKind ;
|
||||
|
||||
// hipErrorNoDevice.
|
||||
|
||||
/*typedef enum hipTextureFilterMode
|
||||
{
|
||||
hipFilterModePoint = cudaFilterModePoint, ///< Point filter mode.
|
||||
//! @warning cudaFilterModeLinear is not supported.
|
||||
} hipTextureFilterMode;*/
|
||||
#define hipFilterModePoint cudaFilterModePoint
|
||||
|
||||
//! Flags that can be used with hipEventCreateWithFlags:
|
||||
#define hipEventDefault cudaEventDefault
|
||||
#define hipEventBlockingSync cudaEventBlockingSync
|
||||
#define hipEventDisableTiming cudaEventDisableTiming
|
||||
#define hipEventInterprocess cudaEventInterprocess
|
||||
|
||||
#define hipHostMallocDefault cudaHostAllocDefault
|
||||
#define hipHostMallocPortable cudaHostAllocPortable
|
||||
#define hipHostMallocMapped cudaHostAllocMapped
|
||||
#define hipHostMallocWriteCombined cudaHostAllocWriteCombined
|
||||
|
||||
#define hipHostRegisterPortable cudaHostRegisterPortable
|
||||
#define hipHostRegisterMapped cudaHostRegisterMapped
|
||||
|
||||
#define HIP_LAUNCH_PARAM_BUFFER_POINTER CU_LAUNCH_PARAM_BUFFER_POINTER
|
||||
#define HIP_LAUNCH_PARAM_BUFFER_SIZE CU_LAUNCH_PARAM_BUFFER_SIZE
|
||||
#define HIP_LAUNCH_PARAM_END CU_LAUNCH_PARAM_END
|
||||
|
||||
typedef cudaEvent_t hipEvent_t;
|
||||
typedef cudaStream_t hipStream_t;
|
||||
typedef cudaIpcEventHandle_t hipIpcEventHandle_t;
|
||||
typedef cudaIpcMemHandle_t hipIpcMemHandle_t;
|
||||
typedef CUcontext hipCtx_t;
|
||||
typedef CUsharedconfig hipSharedMemConfig;
|
||||
typedef CUfunc_cache hipFuncCache;
|
||||
typedef CUdevice hipDevice_t;
|
||||
typedef CUmodule hipModule_t;
|
||||
typedef CUfunction hipFunction_t;
|
||||
typedef CUdeviceptr hipDeviceptr_t;
|
||||
|
||||
//typedef cudaChannelFormatDesc hipChannelFormatDesc;
|
||||
#define hipChannelFormatDesc cudaChannelFormatDesc
|
||||
|
||||
inline static hipError_t hipCUDAErrorTohipError(cudaError_t cuError) {
|
||||
switch(cuError) {
|
||||
case cudaSuccess : return hipSuccess;
|
||||
case cudaErrorMemoryAllocation : return hipErrorMemoryAllocation ;
|
||||
case cudaErrorLaunchOutOfResources : return hipErrorLaunchOutOfResources ;
|
||||
case cudaErrorInvalidValue : return hipErrorInvalidValue ;
|
||||
case cudaErrorInvalidResourceHandle : return hipErrorInvalidResourceHandle ;
|
||||
case cudaErrorInvalidDevice : return hipErrorInvalidDevice ;
|
||||
case cudaErrorInvalidMemcpyDirection : return hipErrorInvalidMemcpyDirection ;
|
||||
case cudaErrorInvalidDevicePointer : return hipErrorInvalidDevicePointer ;
|
||||
case cudaErrorInitializationError : return hipErrorInitializationError ;
|
||||
case cudaErrorNoDevice : return hipErrorNoDevice ;
|
||||
case cudaErrorNotReady : return hipErrorNotReady ;
|
||||
case cudaErrorUnknown : return hipErrorUnknown ;
|
||||
case cudaErrorPeerAccessNotEnabled : return hipErrorPeerAccessNotEnabled ;
|
||||
case cudaErrorPeerAccessAlreadyEnabled : return hipErrorPeerAccessAlreadyEnabled ;
|
||||
case cudaErrorHostMemoryAlreadyRegistered : return hipErrorHostMemoryAlreadyRegistered ;
|
||||
case cudaErrorHostMemoryNotRegistered : return hipErrorHostMemoryNotRegistered ;
|
||||
default : return hipErrorUnknown; // Note - translated error.
|
||||
};
|
||||
}
|
||||
|
||||
inline static hipError_t hipCUResultTohipError(CUresult cuError) { //TODO Populate further
|
||||
switch(cuError) {
|
||||
case CUDA_SUCCESS : return hipSuccess;
|
||||
case CUDA_ERROR_OUT_OF_MEMORY : return hipErrorMemoryAllocation ;
|
||||
case CUDA_ERROR_INVALID_VALUE : return hipErrorInvalidValue ;
|
||||
case CUDA_ERROR_INVALID_DEVICE : return hipErrorInvalidDevice ;
|
||||
case CUDA_ERROR_DEINITIALIZED : return hipErrorDeinitialized ;
|
||||
case CUDA_ERROR_NO_DEVICE : return hipErrorNoDevice ;
|
||||
case CUDA_ERROR_INVALID_CONTEXT : return hipErrorInvalidContext ;
|
||||
case CUDA_ERROR_NOT_INITIALIZED : return hipErrorNotInitialized ;
|
||||
default : return hipErrorUnknown; // Note - translated error.
|
||||
};
|
||||
}
|
||||
|
||||
// TODO match the error enum names of hip and cuda
|
||||
inline static cudaError_t hipErrorToCudaError(hipError_t hError) {
|
||||
switch(hError) {
|
||||
case hipSuccess : return cudaSuccess;
|
||||
case hipErrorMemoryAllocation : return cudaErrorMemoryAllocation ;
|
||||
case hipErrorLaunchOutOfResources : return cudaErrorLaunchOutOfResources ;
|
||||
case hipErrorInvalidValue : return cudaErrorInvalidValue ;
|
||||
case hipErrorInvalidResourceHandle : return cudaErrorInvalidResourceHandle ;
|
||||
case hipErrorInvalidDevice : return cudaErrorInvalidDevice ;
|
||||
case hipErrorInvalidMemcpyDirection : return cudaErrorInvalidMemcpyDirection ;
|
||||
case hipErrorInvalidDevicePointer : return cudaErrorInvalidDevicePointer ;
|
||||
case hipErrorInitializationError : return cudaErrorInitializationError ;
|
||||
case hipErrorNoDevice : return cudaErrorNoDevice ;
|
||||
case hipErrorNotReady : return cudaErrorNotReady ;
|
||||
case hipErrorUnknown : return cudaErrorUnknown ;
|
||||
case hipErrorPeerAccessNotEnabled : return cudaErrorPeerAccessNotEnabled ;
|
||||
case hipErrorPeerAccessAlreadyEnabled : return cudaErrorPeerAccessAlreadyEnabled ;
|
||||
case hipErrorRuntimeMemory : return cudaErrorUnknown ; // Does not exist in CUDA
|
||||
case hipErrorRuntimeOther : return cudaErrorUnknown ; // Does not exist in CUDA
|
||||
case hipErrorHostMemoryAlreadyRegistered : return cudaErrorHostMemoryAlreadyRegistered ;
|
||||
case hipErrorHostMemoryNotRegistered : return cudaErrorHostMemoryNotRegistered ;
|
||||
case hipErrorTbd : return cudaErrorUnknown; // Note - translated error.
|
||||
default : return cudaErrorUnknown; // Note - translated error.
|
||||
}
|
||||
}
|
||||
|
||||
inline static cudaMemcpyKind hipMemcpyKindToCudaMemcpyKind(hipMemcpyKind kind) {
|
||||
switch(kind) {
|
||||
case hipMemcpyHostToHost:
|
||||
return cudaMemcpyHostToHost;
|
||||
case hipMemcpyHostToDevice:
|
||||
return cudaMemcpyHostToDevice;
|
||||
case hipMemcpyDeviceToHost:
|
||||
return cudaMemcpyDeviceToHost;
|
||||
case hipMemcpyDeviceToDevice:
|
||||
return cudaMemcpyDeviceToDevice;
|
||||
default:
|
||||
return cudaMemcpyDefault;
|
||||
}
|
||||
}
|
||||
|
||||
inline static hipError_t hipInit(unsigned int flags)
|
||||
{
|
||||
return hipCUResultTohipError(cuInit(flags));
|
||||
}
|
||||
|
||||
inline static hipError_t hipDeviceReset() {
|
||||
return hipCUDAErrorTohipError(cudaDeviceReset());
|
||||
}
|
||||
|
||||
inline static hipError_t hipGetLastError() {
|
||||
return hipCUDAErrorTohipError(cudaGetLastError());
|
||||
}
|
||||
|
||||
inline static hipError_t hipPeekAtLastError() {
|
||||
return hipCUDAErrorTohipError(cudaPeekAtLastError());
|
||||
}
|
||||
|
||||
inline static hipError_t hipMalloc(void** ptr, size_t size) {
|
||||
return hipCUDAErrorTohipError(cudaMalloc(ptr, size));
|
||||
}
|
||||
|
||||
inline static hipError_t hipFree(void* ptr) {
|
||||
return hipCUDAErrorTohipError(cudaFree(ptr));
|
||||
}
|
||||
|
||||
inline static hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int flags){
|
||||
return hipCUDAErrorTohipError(cudaHostAlloc(ptr, size, flags));
|
||||
}
|
||||
|
||||
inline static hipError_t hipHostGetDevicePointer(void** devPtr, void* hostPtr, unsigned int flags){
|
||||
return hipCUDAErrorTohipError(cudaHostGetDevicePointer(devPtr, hostPtr, flags));
|
||||
}
|
||||
|
||||
inline static hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr){
|
||||
return hipCUDAErrorTohipError(cudaHostGetFlags(flagsPtr, hostPtr));
|
||||
}
|
||||
|
||||
inline static hipError_t hipHostRegister(void* ptr, size_t size, unsigned int flags){
|
||||
return hipCUDAErrorTohipError(cudaHostRegister(ptr, size, flags));
|
||||
}
|
||||
|
||||
inline static hipError_t hipHostUnregister(void* ptr){
|
||||
return hipCUDAErrorTohipError(cudaHostUnregister(ptr));
|
||||
}
|
||||
|
||||
inline static hipError_t hipHostFree(void* ptr) {
|
||||
return hipCUDAErrorTohipError(cudaFreeHost(ptr));
|
||||
}
|
||||
|
||||
inline static hipError_t hipSetDevice(int device) {
|
||||
return hipCUDAErrorTohipError(cudaSetDevice(device));
|
||||
}
|
||||
|
||||
inline static hipError_t hipChooseDevice( int* device, const hipDeviceProp_t* prop )
|
||||
{
|
||||
cudaDeviceProp cdprop;
|
||||
memset(&cdprop,0x0,sizeof(cudaDeviceProp));
|
||||
cdprop.major= prop->major;
|
||||
cdprop.minor = prop->minor;
|
||||
cdprop.totalGlobalMem = prop->totalGlobalMem ;
|
||||
cdprop.sharedMemPerBlock = prop->sharedMemPerBlock;
|
||||
cdprop.regsPerBlock = prop->regsPerBlock;
|
||||
cdprop.warpSize = prop->warpSize ;
|
||||
cdprop.maxThreadsPerBlock = prop->maxThreadsPerBlock ;
|
||||
cdprop.clockRate = prop->clockRate;
|
||||
cdprop.totalConstMem = prop->totalConstMem ;
|
||||
cdprop.multiProcessorCount = prop->multiProcessorCount ;
|
||||
cdprop.l2CacheSize = prop->l2CacheSize ;
|
||||
cdprop.maxThreadsPerMultiProcessor = prop->maxThreadsPerMultiProcessor ;
|
||||
cdprop.computeMode = prop->computeMode ;
|
||||
cdprop.canMapHostMemory = prop->canMapHostMemory;
|
||||
cdprop.memoryClockRate = prop->memoryClockRate;
|
||||
cdprop.memoryBusWidth = prop->memoryBusWidth;
|
||||
return hipCUDAErrorTohipError(cudaChooseDevice(device,&cdprop));
|
||||
}
|
||||
|
||||
inline static hipError_t hipMemcpyHtoD(hipDeviceptr_t dst,
|
||||
void* src, size_t size)
|
||||
{
|
||||
return hipCUResultTohipError(cuMemcpyHtoD(dst, src, size));
|
||||
}
|
||||
|
||||
inline static hipError_t hipMemcpyDtoH(void* dst,
|
||||
hipDeviceptr_t src, size_t size)
|
||||
{
|
||||
return hipCUResultTohipError(cuMemcpyDtoH(dst, src, size));
|
||||
}
|
||||
|
||||
inline static hipError_t hipMemcpyDtoD(hipDeviceptr_t dst,
|
||||
hipDeviceptr_t src, size_t size)
|
||||
{
|
||||
return hipCUResultTohipError(cuMemcpyDtoD(dst, src, size));
|
||||
}
|
||||
|
||||
inline static hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst,
|
||||
void* src, size_t size, hipStream_t stream)
|
||||
{
|
||||
return hipCUResultTohipError(cuMemcpyHtoDAsync(dst, src, size, stream));
|
||||
}
|
||||
|
||||
inline static hipError_t hipMemcpyDtoHAsync(void* dst,
|
||||
hipDeviceptr_t src, size_t size, hipStream_t stream)
|
||||
{
|
||||
return hipCUResultTohipError(cuMemcpyDtoH(dst, src, size));
|
||||
}
|
||||
|
||||
inline static hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst,
|
||||
hipDeviceptr_t src, size_t size, hipStream_t stream)
|
||||
{
|
||||
return hipCUResultTohipError(cuMemcpyDtoD(dst, src, size));
|
||||
}
|
||||
|
||||
inline static hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind copyKind) {
|
||||
return hipCUDAErrorTohipError(cudaMemcpy(dst, src, sizeBytes, hipMemcpyKindToCudaMemcpyKind(copyKind)));
|
||||
}
|
||||
|
||||
|
||||
inline static hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind copyKind, hipStream_t stream=0) {
|
||||
return hipCUDAErrorTohipError(cudaMemcpyAsync(dst, src, sizeBytes, hipMemcpyKindToCudaMemcpyKind(copyKind), stream));
|
||||
}
|
||||
|
||||
|
||||
inline static hipError_t hipMemcpyToSymbol(const char * symbolName, const void* src, size_t sizeBytes, size_t offset = 0, hipMemcpyKind copyType = hipMemcpyHostToDevice) {
|
||||
return hipCUDAErrorTohipError(cudaMemcpyToSymbol(symbolName, src, sizeBytes, offset, hipMemcpyKindToCudaMemcpyKind(copyType)));
|
||||
}
|
||||
inline static hipError_t hipDeviceSynchronize() {
|
||||
return hipCUDAErrorTohipError(cudaDeviceSynchronize());
|
||||
}
|
||||
|
||||
inline static const char* hipGetErrorString(hipError_t error){
|
||||
return cudaGetErrorString( hipErrorToCudaError(error) );
|
||||
}
|
||||
|
||||
inline static const char* hipGetErrorName(hipError_t error){
|
||||
return cudaGetErrorName( hipErrorToCudaError(error) );
|
||||
}
|
||||
|
||||
inline static hipError_t hipGetDeviceCount(int * count){
|
||||
return hipCUDAErrorTohipError(cudaGetDeviceCount(count));
|
||||
}
|
||||
|
||||
inline static hipError_t hipGetDevice(int * device){
|
||||
return hipCUDAErrorTohipError(cudaGetDevice(device));
|
||||
}
|
||||
|
||||
inline static hipError_t hipIpcCloseMemHandle(void *devPtr){
|
||||
return hipCUDAErrorTohipError(cudaIpcCloseMemHandle(devPtr));
|
||||
}
|
||||
|
||||
inline static hipError_t hipIpcGetEventHandle(hipIpcEventHandle_t* handle, hipEvent_t event){
|
||||
return hipCUDAErrorTohipError(cudaIpcGetEventHandle(handle, event));
|
||||
}
|
||||
|
||||
inline static hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr){
|
||||
return hipCUDAErrorTohipError(cudaIpcGetMemHandle(handle, devPtr));
|
||||
}
|
||||
|
||||
inline static hipError_t hipIpcOpenEventHandle(hipEvent_t* event, hipIpcEventHandle_t handle){
|
||||
return hipCUDAErrorTohipError(cudaIpcOpenEventHandle(event, handle));
|
||||
}
|
||||
|
||||
inline static hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned int flags){
|
||||
return hipCUDAErrorTohipError(cudaIpcOpenMemHandle(devPtr, handle, flags));
|
||||
}
|
||||
|
||||
inline static hipError_t hipMemset(void* devPtr,int value, size_t count) {
|
||||
return hipCUDAErrorTohipError(cudaMemset(devPtr, value, count));
|
||||
}
|
||||
|
||||
inline static hipError_t hipMemsetAsync(void* devPtr,int value, size_t count, hipStream_t stream = 0) {
|
||||
return hipCUDAErrorTohipError(cudaMemsetAsync(devPtr, value, count, stream));
|
||||
}
|
||||
|
||||
inline static hipError_t hipGetDeviceProperties(hipDeviceProp_t *p_prop, int device)
|
||||
{
|
||||
cudaDeviceProp cdprop;
|
||||
cudaError_t cerror;
|
||||
cerror = cudaGetDeviceProperties(&cdprop,device);
|
||||
strncpy(p_prop->name,cdprop.name, 256);
|
||||
p_prop->totalGlobalMem = cdprop.totalGlobalMem ;
|
||||
p_prop->sharedMemPerBlock = cdprop.sharedMemPerBlock;
|
||||
p_prop->regsPerBlock = cdprop.regsPerBlock;
|
||||
p_prop->warpSize = cdprop.warpSize ;
|
||||
for (int i=0 ; i<3; i++) {
|
||||
p_prop->maxThreadsDim[i] = cdprop.maxThreadsDim[i];
|
||||
p_prop->maxGridSize[i] = cdprop.maxGridSize[i];
|
||||
}
|
||||
p_prop->maxThreadsPerBlock = cdprop.maxThreadsPerBlock ;
|
||||
p_prop->clockRate = cdprop.clockRate;
|
||||
p_prop->totalConstMem = cdprop.totalConstMem ;
|
||||
p_prop->major = cdprop.major ;
|
||||
p_prop->minor = cdprop. minor ;
|
||||
p_prop->multiProcessorCount = cdprop.multiProcessorCount ;
|
||||
p_prop->l2CacheSize = cdprop.l2CacheSize ;
|
||||
p_prop->maxThreadsPerMultiProcessor = cdprop.maxThreadsPerMultiProcessor ;
|
||||
p_prop->computeMode = cdprop.computeMode ;
|
||||
p_prop->canMapHostMemory = cdprop.canMapHostMemory;
|
||||
p_prop->memoryClockRate = cdprop.memoryClockRate;
|
||||
p_prop->memoryBusWidth = cdprop.memoryBusWidth;
|
||||
|
||||
// Same as clock-rate:
|
||||
p_prop->clockInstructionRate = cdprop.clockRate;
|
||||
|
||||
int ccVers = p_prop->major*100 + p_prop->minor * 10;
|
||||
|
||||
p_prop->arch.hasGlobalInt32Atomics = (ccVers >= 110);
|
||||
p_prop->arch.hasGlobalFloatAtomicExch = (ccVers >= 110);
|
||||
p_prop->arch.hasSharedInt32Atomics = (ccVers >= 120);
|
||||
p_prop->arch.hasSharedFloatAtomicExch = (ccVers >= 120);
|
||||
|
||||
p_prop->arch.hasFloatAtomicAdd = (ccVers >= 200);
|
||||
|
||||
p_prop->arch.hasGlobalInt64Atomics = (ccVers >= 120);
|
||||
p_prop->arch.hasSharedInt64Atomics = (ccVers >= 110);
|
||||
|
||||
p_prop->arch.hasDoubles = (ccVers >= 130);
|
||||
|
||||
p_prop->arch.hasWarpVote = (ccVers >= 120);
|
||||
p_prop->arch.hasWarpBallot = (ccVers >= 200);
|
||||
p_prop->arch.hasWarpShuffle = (ccVers >= 300);
|
||||
p_prop->arch.hasFunnelShift = (ccVers >= 350);
|
||||
|
||||
p_prop->arch.hasThreadFenceSystem = (ccVers >= 200);
|
||||
p_prop->arch.hasSyncThreadsExt = (ccVers >= 200);
|
||||
|
||||
p_prop->arch.hasSurfaceFuncs = (ccVers >= 200);
|
||||
p_prop->arch.has3dGrid = (ccVers >= 200);
|
||||
p_prop->arch.hasDynamicParallelism = (ccVers >= 350);
|
||||
|
||||
p_prop->concurrentKernels = cdprop.concurrentKernels;
|
||||
|
||||
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 hipDeviceAttributeMemoryClockRate:
|
||||
cdattr = cudaDevAttrMemoryClockRate; break;
|
||||
case hipDeviceAttributeMemoryBusWidth:
|
||||
cdattr = cudaDevAttrGlobalMemoryBusWidth; 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 hipDeviceAttributeConcurrentKernels:
|
||||
cdattr = cudaDevAttrConcurrentKernels; break;
|
||||
case hipDeviceAttributePciBusId:
|
||||
cdattr = cudaDevAttrPciBusId; break;
|
||||
case hipDeviceAttributePciDeviceId:
|
||||
cdattr = cudaDevAttrPciDeviceId; break;
|
||||
case hipDeviceAttributeMaxSharedMemoryPerMultiprocessor:
|
||||
cdattr = cudaDevAttrMaxSharedMemoryPerMultiprocessor; break;
|
||||
case hipDeviceAttributeIsMultiGpuBoard:
|
||||
cdattr = cudaDevAttrIsMultiGpuBoard; break;
|
||||
default:
|
||||
cerror = cudaErrorInvalidValue; break;
|
||||
}
|
||||
|
||||
cerror = cudaDeviceGetAttribute(pi, cdattr, device);
|
||||
|
||||
return hipCUDAErrorTohipError(cerror);
|
||||
}
|
||||
|
||||
inline static hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor(
|
||||
int *numBlocks,
|
||||
const void* func,
|
||||
int blockSize,
|
||||
size_t dynamicSMemSize
|
||||
)
|
||||
{
|
||||
cudaError_t cerror;
|
||||
cerror = cudaOccupancyMaxActiveBlocksPerMultiprocessor(numBlocks, func, blockSize, dynamicSMemSize);
|
||||
return hipCUDAErrorTohipError(cerror);
|
||||
}
|
||||
|
||||
inline static hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, void* ptr){
|
||||
cudaPointerAttributes cPA;
|
||||
hipError_t err = hipCUDAErrorTohipError(cudaPointerGetAttributes(&cPA, ptr));
|
||||
if(err == hipSuccess){
|
||||
switch (cPA.memoryType){
|
||||
case cudaMemoryTypeDevice:
|
||||
attributes->memoryType = hipMemoryTypeDevice; break;
|
||||
case cudaMemoryTypeHost:
|
||||
attributes->memoryType = hipMemoryTypeHost; break;
|
||||
default:
|
||||
return hipErrorUnknown;
|
||||
}
|
||||
attributes->device = cPA.device;
|
||||
attributes->devicePointer = cPA.devicePointer;
|
||||
attributes->hostPointer = cPA.hostPointer;
|
||||
attributes->isManaged = 0;
|
||||
attributes->allocationFlags = 0;
|
||||
}
|
||||
return err;
|
||||
}
|
||||
|
||||
|
||||
inline static hipError_t hipMemGetInfo( size_t* free, size_t* total)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaMemGetInfo(free,total));
|
||||
}
|
||||
|
||||
inline static hipError_t hipEventCreate( hipEvent_t* event)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaEventCreate(event));
|
||||
}
|
||||
|
||||
inline static hipError_t hipEventRecord( hipEvent_t event, hipStream_t stream = NULL)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaEventRecord(event,stream));
|
||||
}
|
||||
|
||||
inline static hipError_t hipEventSynchronize( hipEvent_t event)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaEventSynchronize(event));
|
||||
}
|
||||
|
||||
inline static hipError_t hipEventElapsedTime( float *ms, hipEvent_t start, hipEvent_t stop)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaEventElapsedTime(ms,start,stop));
|
||||
}
|
||||
|
||||
inline static hipError_t hipEventDestroy( hipEvent_t event)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaEventDestroy(event));
|
||||
}
|
||||
|
||||
|
||||
inline static hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaStreamCreateWithFlags(stream, flags));
|
||||
}
|
||||
|
||||
|
||||
inline static hipError_t hipStreamCreate(hipStream_t *stream)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaStreamCreate(stream));
|
||||
}
|
||||
|
||||
inline static hipError_t hipStreamSynchronize(hipStream_t stream)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaStreamSynchronize(stream));
|
||||
}
|
||||
|
||||
inline static hipError_t hipStreamDestroy(hipStream_t stream)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaStreamDestroy(stream));
|
||||
}
|
||||
|
||||
|
||||
inline static hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaStreamWaitEvent(stream, event, flags));
|
||||
}
|
||||
|
||||
inline static hipError_t hipStreamQuery(hipStream_t stream)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaStreamQuery(stream));
|
||||
}
|
||||
|
||||
|
||||
inline static hipError_t hipDriverGetVersion(int *driverVersion)
|
||||
{
|
||||
cudaError_t err = cudaDriverGetVersion(driverVersion);
|
||||
|
||||
// Override driver version to match version reported on HCC side.
|
||||
*driverVersion = 4;
|
||||
|
||||
return hipCUDAErrorTohipError(err);
|
||||
}
|
||||
|
||||
inline static hipError_t hipRuntimeGetVersion(int *runtimeVersion)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaRuntimeGetVersion(runtimeVersion));
|
||||
}
|
||||
|
||||
inline static hipError_t hipDeviceCanAccessPeer ( int* canAccessPeer, int device, int peerDevice )
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaDeviceCanAccessPeer(canAccessPeer, device, peerDevice));
|
||||
}
|
||||
|
||||
inline static hipError_t hipDeviceDisablePeerAccess ( int peerDevice )
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaDeviceDisablePeerAccess ( peerDevice ));
|
||||
};
|
||||
|
||||
inline static hipError_t hipDeviceEnablePeerAccess ( int peerDevice, unsigned int flags )
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaDeviceEnablePeerAccess ( peerDevice, flags ));
|
||||
}
|
||||
|
||||
inline static hipError_t hipCtxDisablePeerAccess ( hipCtx_t peerCtx )
|
||||
{
|
||||
return hipCUResultTohipError(cuCtxDisablePeerAccess ( peerCtx ));
|
||||
}
|
||||
|
||||
inline static hipError_t hipCtxEnablePeerAccess ( hipCtx_t peerCtx, unsigned int flags )
|
||||
{
|
||||
return hipCUResultTohipError(cuCtxEnablePeerAccess ( peerCtx, flags ));
|
||||
}
|
||||
|
||||
inline static hipError_t hipMemcpyPeer ( void* dst, int dstDevice, const void* src, int srcDevice, size_t count )
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaMemcpyPeer ( dst, dstDevice, src, srcDevice, count ));
|
||||
};
|
||||
|
||||
inline static hipError_t hipMemcpyPeerAsync ( void* dst, int dstDevice, const void* src, int srcDevice, size_t count, hipStream_t stream=0 )
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaMemcpyPeerAsync ( dst, dstDevice, src, srcDevice, count, stream ));
|
||||
};
|
||||
|
||||
inline static hipError_t hipSetDeviceFlags (unsigned int flags)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaSetDeviceFlags( flags ));
|
||||
}
|
||||
|
||||
inline static hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned int flags)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaEventCreateWithFlags(event, flags));
|
||||
}
|
||||
|
||||
inline static hipError_t hipEventQuery(hipEvent_t event)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaEventQuery(event));
|
||||
}
|
||||
|
||||
inline static hipError_t hipCtxCreate(hipCtx_t *ctx, unsigned int flags, hipDevice_t device)
|
||||
{
|
||||
return hipCUResultTohipError(cuCtxCreate ( ctx,flags,device ));
|
||||
}
|
||||
|
||||
inline static hipError_t hipCtxDestroy(hipCtx_t ctx)
|
||||
{
|
||||
return hipCUResultTohipError(cuCtxDestroy ( ctx ));
|
||||
}
|
||||
|
||||
inline static hipError_t hipCtxPopCurrent(hipCtx_t* ctx)
|
||||
{
|
||||
return hipCUResultTohipError(cuCtxPopCurrent ( ctx ));
|
||||
}
|
||||
|
||||
inline static hipError_t hipCtxPushCurrent(hipCtx_t ctx)
|
||||
{
|
||||
return hipCUResultTohipError(cuCtxPushCurrent ( ctx ));
|
||||
}
|
||||
|
||||
inline static hipError_t hipCtxSetCurrent(hipCtx_t ctx)
|
||||
{
|
||||
return hipCUResultTohipError(cuCtxSetCurrent ( ctx ));
|
||||
}
|
||||
|
||||
inline static hipError_t hipCtxGetCurrent(hipCtx_t* ctx)
|
||||
{
|
||||
return hipCUResultTohipError(cuCtxGetCurrent ( ctx ));
|
||||
}
|
||||
|
||||
inline static hipError_t hipCtxGetDevice(hipDevice_t *device)
|
||||
{
|
||||
return hipCUResultTohipError(cuCtxGetDevice ( device ));
|
||||
}
|
||||
|
||||
inline static hipError_t hipCtxGetApiVersion (hipCtx_t ctx,int *apiVersion)
|
||||
{
|
||||
return hipCUResultTohipError(cuCtxGetApiVersion ( ctx,(unsigned int*)apiVersion ));
|
||||
}
|
||||
|
||||
inline static hipError_t hipCtxGetCacheConfig ( hipFuncCache *cacheConfig )
|
||||
{
|
||||
return hipCUResultTohipError(cuCtxGetCacheConfig ( cacheConfig ));
|
||||
}
|
||||
|
||||
inline static hipError_t hipCtxSetCacheConfig ( hipFuncCache cacheConfig )
|
||||
{
|
||||
return hipCUResultTohipError(cuCtxSetCacheConfig ( cacheConfig ));
|
||||
}
|
||||
|
||||
inline static hipError_t hipCtxSetSharedMemConfig ( hipSharedMemConfig config )
|
||||
{
|
||||
return hipCUResultTohipError(cuCtxSetSharedMemConfig ( config ));
|
||||
}
|
||||
|
||||
inline static hipError_t hipCtxGetSharedMemConfig ( hipSharedMemConfig * pConfig )
|
||||
{
|
||||
return hipCUResultTohipError(cuCtxGetSharedMemConfig ( pConfig ));
|
||||
}
|
||||
|
||||
inline static hipError_t hipCtxSynchronize ( void )
|
||||
{
|
||||
return hipCUResultTohipError(cuCtxSynchronize ( ));
|
||||
}
|
||||
|
||||
inline static hipError_t hipCtxGetFlags ( unsigned int* flags )
|
||||
{
|
||||
return hipCUResultTohipError(cuCtxGetFlags ( flags ));
|
||||
}
|
||||
|
||||
inline static hipError_t hipCtxDetach(hipCtx_t ctx)
|
||||
{
|
||||
return hipCUResultTohipError(cuCtxDetach(ctx));
|
||||
}
|
||||
|
||||
inline static hipError_t hipDeviceGet(hipDevice_t *device, int ordinal)
|
||||
{
|
||||
return hipCUResultTohipError(cuDeviceGet(device, ordinal));
|
||||
}
|
||||
|
||||
inline static hipError_t hipDeviceComputeCapability(int *major, int *minor, hipDevice_t device)
|
||||
{
|
||||
return hipCUResultTohipError(cuDeviceComputeCapability(major,minor,device));
|
||||
}
|
||||
|
||||
inline static hipError_t hipDeviceGetName(char *name,int len,hipDevice_t device)
|
||||
{
|
||||
return hipCUResultTohipError(cuDeviceGetName(name,len,device));
|
||||
}
|
||||
|
||||
inline static hipError_t hipDeviceGetPCIBusId (int *pciBusId,int len,hipDevice_t device)
|
||||
{
|
||||
return hipCUResultTohipError(cuDeviceGetPCIBusId((char*)pciBusId,len,device));
|
||||
}
|
||||
|
||||
inline static hipError_t hipDeviceTotalMem (size_t *bytes,hipDevice_t device)
|
||||
{
|
||||
return hipCUResultTohipError(cuDeviceTotalMem(bytes,device));
|
||||
}
|
||||
|
||||
inline static hipError_t hipModuleLoad(hipModule_t *module, const char* fname)
|
||||
{
|
||||
return hipCUResultTohipError(cuModuleLoad(module, fname));
|
||||
}
|
||||
|
||||
inline static hipError_t hipModuleUnload(hipModule_t hmod)
|
||||
{
|
||||
return hipCUResultTohipError(cuModuleUnload(hmod));
|
||||
}
|
||||
|
||||
inline static hipError_t hipModuleGetFunction(hipFunction_t *function,
|
||||
hipModule_t module, const char *kname)
|
||||
{
|
||||
return hipCUResultTohipError(cuModuleGetFunction(function, module, kname));
|
||||
}
|
||||
|
||||
inline static hipError_t hipModuleGetGlobal(hipDeviceptr_t *dptr, size_t *bytes,
|
||||
hipModule_t hmod, const char* name)
|
||||
{
|
||||
return hipCUResultTohipError(cuModuleGetGlobal(dptr, bytes, hmod, name));
|
||||
}
|
||||
|
||||
inline static hipError_t hipModuleLoadData(hipModule_t *module, const void *image)
|
||||
{
|
||||
return hipCUResultTohipError(cuModuleLoadData(module, image));
|
||||
}
|
||||
|
||||
inline static hipError_t hipModuleLaunchKernel(hipFunction_t f,
|
||||
unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ,
|
||||
unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ,
|
||||
unsigned int sharedMemBytes, hipStream_t stream,
|
||||
void **kernelParams, void **extra)
|
||||
{
|
||||
return hipCUResultTohipError(cuLaunchKernel(f,
|
||||
gridDimX, gridDimY, gridDimZ,
|
||||
blockDimX, blockDimY, blockDimZ,
|
||||
sharedMemBytes, stream, kernelParams, extra));
|
||||
}
|
||||
|
||||
#ifdef __cplusplus
|
||||
}
|
||||
#endif
|
||||
|
||||
#ifdef __CUDACC__
|
||||
|
||||
template<class T>
|
||||
inline static hipError_t hipOccupancyMaxPotentialBlockSize(
|
||||
int *minGridSize,
|
||||
int *blockSize,
|
||||
T func,
|
||||
size_t dynamicSMemSize = 0,
|
||||
int blockSizeLimit = 0,
|
||||
unsigned int flags = 0
|
||||
){
|
||||
cudaError_t cerror;
|
||||
cerror = cudaOccupancyMaxPotentialBlockSize(minGridSize, blockSize, func, dynamicSMemSize, blockSizeLimit, flags);
|
||||
return hipCUDAErrorTohipError(cerror);
|
||||
}
|
||||
|
||||
template <class T, int dim, enum cudaTextureReadMode readMode>
|
||||
inline static hipError_t hipBindTexture(size_t *offset,
|
||||
const struct texture<T, dim, readMode> &tex,
|
||||
const void *devPtr,
|
||||
size_t size=UINT_MAX)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, size));
|
||||
}
|
||||
|
||||
template <class T, int dim, enum cudaTextureReadMode readMode>
|
||||
inline static hipError_t hipBindTexture(size_t *offset,
|
||||
struct texture<T, dim, readMode> *tex,
|
||||
const void *devPtr,
|
||||
const struct hipChannelFormatDesc *desc,
|
||||
size_t size=UINT_MAX)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, desc, size));
|
||||
}
|
||||
|
||||
template <class T, int dim, enum cudaTextureReadMode readMode>
|
||||
inline static hipError_t hipUnbindTexture(struct texture<T, dim, readMode> *tex)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaUnbindTexture(tex));
|
||||
}
|
||||
|
||||
template <class T>
|
||||
inline static hipChannelFormatDesc hipCreateChannelDesc()
|
||||
{
|
||||
return cudaCreateChannelDesc<T>();
|
||||
}
|
||||
#endif
|
||||
Αναφορά σε νέο ζήτημα
Block a user