fixed hipcc for new compiler flags
Change-Id: I49ec059be20ff26b7482c84d91ab7a43826c6a8d
This commit is contained in:
@@ -21,12 +21,51 @@ THE SOFTWARE.
|
||||
#include <hip/hip_fp16.h>
|
||||
#include "hip/hip_runtime_api.h"
|
||||
|
||||
__global__ void halfMath(hipLaunchParm lp, half *A, half *B, half *C) {
|
||||
#define LEN 64
|
||||
#define HALF_SIZE 64*sizeof(__half)
|
||||
#define HALF2_SIZE 64*sizeof(__half2)
|
||||
|
||||
__global__ void __halfMath(hipLaunchParm lp, __half *A, __half *B, __half *C) {
|
||||
int tx = hipThreadIdx_x;
|
||||
__half a = A[tx];
|
||||
|
||||
__half b = B[tx];
|
||||
__half c = C[tx];
|
||||
c = __hadd(a, c);
|
||||
c = __hadd_sat(b, c);
|
||||
c = __hfma(a, c, b);
|
||||
c = __hfma_sat(b, c, a);
|
||||
c = __hsub(a, c);
|
||||
c = __hsub_sat(b, c);
|
||||
c = __hmul(a, c);
|
||||
c = __hmul_sat(b, c);
|
||||
c = hdiv(a, c);
|
||||
}
|
||||
|
||||
__global__ void __half2Math(hipLaunchParm lp, __half2 *A, __half2 *B, __half2 *C) {
|
||||
int tx = hipThreadIdx_x;
|
||||
__half2 a = A[tx];
|
||||
__half2 b = B[tx];
|
||||
__half2 c = C[tx];
|
||||
c = __hadd2(a, c);
|
||||
c = __hadd2_sat(b, c);
|
||||
c = __hfma2(a, c, b);
|
||||
c = __hfma2_sat(b, c, a);
|
||||
c = __hsub2(a, c);
|
||||
c = __hsub2_sat(b, c);
|
||||
c = __hmul2(a, c);
|
||||
c = __hmul2_sat(b, c);
|
||||
}
|
||||
|
||||
int main(){
|
||||
__half *A, *B, *C;
|
||||
hipMalloc(&A, HALF_SIZE);
|
||||
hipMalloc(&B, HALF_SIZE);
|
||||
hipMalloc(&C, HALF_SIZE);
|
||||
hipLaunchKernel(__halfMath, dim3(1,1,1), dim3(LEN,1,1), 0, 0, A, B, C);
|
||||
__half2 *A2, *B2, *C2;
|
||||
hipMalloc(&A, HALF2_SIZE);
|
||||
hipMalloc(&B, HALF2_SIZE);
|
||||
hipMalloc(&C, HALF2_SIZE);
|
||||
hipLaunchKernel(__half2Math, dim3(1,1,1), dim3(LEN,1,1), 0, 0, A2, B2, C2);
|
||||
|
||||
}
|
||||
|
||||
Reference in New Issue
Block a user