AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

Merge pull request #482 from ROCm-Developer-Tools/feature_clean_up_hip_math

Просмотр исходного кода 
Switch to using ROCDL directly, as opposed to via HC. Add missing bits.
Этот коммит содержится в:
Maneesh Gupta
2018-06-06 16:07:22 +05:30
коммит произвёл GitHub
родитель 02ea7f13b3 14e449b5bb
Коммит 203dd6cb70
10 изменённых файлов: 1811 добавлений и 2148 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
tests/src/deviceLib/hipDoublePrecisionMathDevice.cpp
+16 -10
Просмотреть файл
@@ -50,8 +50,8 @@ __device__ void double_precision_math_functions() {
cos(0.0);
cosh(0.0);
cospi(0.0);
// cyl_bessel_i0(0.0);
// cyl_bessel_i1(0.0);
cyl_bessel_i0(0.0);
cyl_bessel_i1(0.0);
erf(0.0);
erfc(0.0);
erfcinv(2.0);
@@ -78,7 +78,7 @@ __device__ void double_precision_math_functions() {
j1(0.0);
jn(-1.0, 1.0);
ldexp(0.0, 0);
// lgamma(1.0);
lgamma(1.0);
llrint(0.0);
llround(0.0);
log(1.0);
@@ -88,7 +88,7 @@ __device__ void double_precision_math_functions() {
logb(1.0);
lrint(0.0);
lround(0.0);
// modf(0.0, &fX);
modf(0.0, &fX);
nan("1");
nearbyint(0.0);
nextafter(0.0, 0.0);
@@ -99,9 +99,9 @@ __device__ void double_precision_math_functions() {
normcdf(0.0);
normcdfinv(1.0);
pow(1.0, 0.0);
// rcbrt(1.0);
rcbrt(1.0);
remainder(2.0, 1.0);
// remquo(1.0, 2.0, &iX);
remquo(1.0, 2.0, &iX);
rhypot(0.0, 1.0);
rint(1.0);
fX = 1.0;
@@ -110,8 +110,8 @@ __device__ void double_precision_math_functions() {
rnorm4d(0.0, 0.0, 0.0, 1.0);
round(0.0);
rsqrt(1.0);
// scalbln(0.0, 1);
// scalbn(0.0, 1);
scalbln(0.0, 1);
scalbn(0.0, 1);
signbit(1.0);
sin(0.0);
sincos(0.0, &fX, &fY);
@@ -128,11 +128,17 @@ __device__ void double_precision_math_functions() {
yn(1, 1.0);
}
__global__ void compileDoublePrecisionMathOnDevice(hipLaunchParm lp, int ignored) {
__global__ void compileDoublePrecisionMathOnDevice(int) {
double_precision_math_functions();
}
int main() {
hipLaunchKernel(compileDoublePrecisionMathOnDevice, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, 1);
hipLaunchKernelGGL(
compileDoublePrecisionMathOnDevice,
dim3(1, 1, 1),
dim3(1, 1, 1),
0,
0,
1);
passed();
}
tests/src/deviceLib/hipFloatMathPrecise.cpp
+10 -10
Просмотреть файл
@@ -30,8 +30,8 @@ THE SOFTWARE.
#include <hip/math_functions.h>
#include "test_common.h"
__global__ void FloatMathPrecise(hipLaunchParm lp) {
// int iX; //uncomment this when remqouf() is enabled again
__global__ void FloatMathPrecise() {
int iX;
float fX, fY;
acosf(1.0f);
@@ -47,8 +47,8 @@ __global__ void FloatMathPrecise(hipLaunchParm lp) {
cosf(0.0f);
coshf(0.0f);
cospif(0.0f);
// cyl_bessel_i0f(0.0f);
// cyl_bessel_i1f(0.0f);
cyl_bessel_i0f(0.0f);
cyl_bessel_i1f(0.0f);
erfcf(0.0f);
erfcinvf(2.0f);
erfcxf(0.0f);
@@ -66,7 +66,7 @@ __global__ void FloatMathPrecise(hipLaunchParm lp) {
fX = fmaxf(0.0f, 0.0f);
fX = fminf(0.0f, 0.0f);
fmodf(0.0f, 1.0f);
// frexpf(0.0f, &iX);
frexpf(0.0f, &iX);
hypotf(1.0f, 0.0f);
ilogbf(1.0f);
isfinite(0.0f);
@@ -76,7 +76,7 @@ __global__ void FloatMathPrecise(hipLaunchParm lp) {
j1f(0.0f);
jnf(-1.0f, 1.0f);
ldexpf(0.0f, 0);
// lgammaf(1.0f);
lgammaf(1.0f);
llrintf(0.0f);
llroundf(0.0f);
log10f(1.0f);
@@ -86,10 +86,10 @@ __global__ void FloatMathPrecise(hipLaunchParm lp) {
logf(1.0f);
lrintf(0.0f);
lroundf(0.0f);
// modff(0.0f, &fX);
modff(0.0f, &fX);
fX = nanf("1");
fX = nearbyintf(0.0f);
// nextafterf(0.0f);
nextafterf(0.0f, 0.0f);
norm3df(1.0f, 0.0f, 0.0f);
norm4df(1.0f, 0.0f, 0.0f, 0.0f);
normcdff(0.0f);
@@ -99,7 +99,7 @@ __global__ void FloatMathPrecise(hipLaunchParm lp) {
powf(1.0f, 0.0f);
rcbrtf(1.0f);
remainderf(2.0f, 1.0f);
// remquof(1.0f, 2.0f, &iX);
remquof(1.0f, 2.0f, &iX);
rhypotf(0.0f, 1.0f);
fY = rintf(1.0f);
rnorm3df(0.0f, 0.0f, 1.0f);
@@ -127,6 +127,6 @@ __global__ void FloatMathPrecise(hipLaunchParm lp) {
}
int main() {
hipLaunchKernel(FloatMathPrecise, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0);
hipLaunchKernelGGL(FloatMathPrecise, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0);
passed();
}
tests/src/deviceLib/hipTestDevice.cpp
+32 -32
Просмотреть файл
@@ -31,82 +31,82 @@ THE SOFTWARE.
#define N 512
#define SIZE N * sizeof(float)
__global__ void test_sincosf(hipLaunchParm lp, float* a, float* b, float* c) {
__global__ void test_sincosf(float* a, float* b, float* c) {
int tid = threadIdx.x;
sincosf(a[tid], b + tid, c + tid);
}
__global__ void test_sincospif(hipLaunchParm lp, float* a, float* b, float* c) {
__global__ void test_sincospif(float* a, float* b, float* c) {
int tid = threadIdx.x;
sincospif(a[tid], b + tid, c + tid);
}
__global__ void test_fdividef(hipLaunchParm lp, float* a, float* b, float* c) {
__global__ void test_fdividef(float* a, float* b, float* c) {
int tid = threadIdx.x;
c[tid] = fdividef(a[tid], b[tid]);
}
__global__ void test_llrintf(hipLaunchParm lp, float* a, long long int* b) {
__global__ void test_llrintf(float* a, long long int* b) {
int tid = threadIdx.x;
b[tid] = llrintf(a[tid]);
}
__global__ void test_lrintf(hipLaunchParm lp, float* a, long int* b) {
__global__ void test_lrintf(float* a, long int* b) {
int tid = threadIdx.x;
b[tid] = lrintf(a[tid]);
}
__global__ void test_rintf(hipLaunchParm lp, float* a, float* b) {
__global__ void test_rintf(float* a, float* b) {
int tid = threadIdx.x;
b[tid] = rintf(a[tid]);
}
__global__ void test_llroundf(hipLaunchParm lp, float* a, long long int* b) {
__global__ void test_llroundf(float* a, long long int* b) {
int tid = threadIdx.x;
b[tid] = llroundf(a[tid]);
}
__global__ void test_lroundf(hipLaunchParm lp, float* a, long int* b) {
__global__ void test_lroundf(float* a, long int* b) {
int tid = threadIdx.x;
b[tid] = lroundf(a[tid]);
}
__global__ void test_rhypotf(hipLaunchParm lp, float* a, float* b, float* c) {
__global__ void test_rhypotf(float* a, float* b, float* c) {
int tid = threadIdx.x;
c[tid] = rhypotf(a[tid], b[tid]);
}
__global__ void test_norm3df(hipLaunchParm lp, float* a, float* b, float* c, float* d) {
__global__ void test_norm3df(float* a, float* b, float* c, float* d) {
int tid = threadIdx.x;
d[tid] = norm3df(a[tid], b[tid], c[tid]);
}
__global__ void test_norm4df(hipLaunchParm lp, float* a, float* b, float* c, float* d, float* e) {
__global__ void test_norm4df(float* a, float* b, float* c, float* d, float* e) {
int tid = threadIdx.x;
e[tid] = norm4df(a[tid], b[tid], c[tid], d[tid]);
}
__global__ void test_normf(hipLaunchParm lp, float* a, float* b) {
__global__ void test_normf(float* a, float* b) {
int tid = threadIdx.x;
b[tid] = normf(N, a);
}
__global__ void test_rnorm3df(hipLaunchParm lp, float* a, float* b, float* c, float* d) {
__global__ void test_rnorm3df(float* a, float* b, float* c, float* d) {
int tid = threadIdx.x;
d[tid] = rnorm3df(a[tid], b[tid], c[tid]);
}
__global__ void test_rnorm4df(hipLaunchParm lp, float* a, float* b, float* c, float* d, float* e) {
__global__ void test_rnorm4df(float* a, float* b, float* c, float* d, float* e) {
int tid = threadIdx.x;
e[tid] = rnorm4df(a[tid], b[tid], c[tid], d[tid]);
}
__global__ void test_rnormf(hipLaunchParm lp, float* a, float* b) {
__global__ void test_rnormf(float* a, float* b) {
int tid = threadIdx.x;
b[tid] = rnormf(N, a);
}
__global__ void test_erfinvf(hipLaunchParm lp, float* a, float* b) {
__global__ void test_erfinvf(float* a, float* b) {
int tid = threadIdx.x;
b[tid] = erff(erfinvf(a[tid]));
}
@@ -124,7 +124,7 @@ bool run_sincosf() {
hipMalloc((void**)&Bd, SIZE);
hipMalloc((void**)&Cd, SIZE);
hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
hipLaunchKernel(test_sincosf, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd);
hipLaunchKernelGGL(test_sincosf, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd);
hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost);
hipMemcpy(C, Cd, SIZE, hipMemcpyDeviceToHost);
int passed = 0;
@@ -166,7 +166,7 @@ bool run_sincospif() {
hipMalloc((void**)&Bd, SIZE);
hipMalloc((void**)&Cd, SIZE);
hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
hipLaunchKernel(test_sincospif, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd);
hipLaunchKernelGGL(test_sincospif, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd);
hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost);
hipMemcpy(C, Cd, SIZE, hipMemcpyDeviceToHost);
int passed = 0;
@@ -210,7 +210,7 @@ bool run_fdividef() {
hipMalloc((void**)&Cd, SIZE);
hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice);
hipLaunchKernel(test_fdividef, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd);
hipLaunchKernelGGL(test_fdividef, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd);
hipMemcpy(C, Cd, SIZE, hipMemcpyDeviceToHost);
int passed = 0;
for (int i = 0; i < 512; i++) {
@@ -244,7 +244,7 @@ bool run_llrintf() {
hipMalloc((void**)&Ad, SIZE);
hipMalloc((void**)&Bd, N * sizeof(long long int));
hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
hipLaunchKernel(test_llrintf, dim3(1), dim3(N), 0, 0, Ad, Bd);
hipLaunchKernelGGL(test_llrintf, dim3(1), dim3(N), 0, 0, Ad, Bd);
hipMemcpy(B, Bd, N * sizeof(long long int), hipMemcpyDeviceToHost);
int passed = 0;
for (int i = 0; i < 512; i++) {
@@ -277,7 +277,7 @@ bool run_lrintf() {
hipMalloc((void**)&Ad, SIZE);
hipMalloc((void**)&Bd, N * sizeof(long int));
hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
hipLaunchKernel(test_lrintf, dim3(1), dim3(N), 0, 0, Ad, Bd);
hipLaunchKernelGGL(test_lrintf, dim3(1), dim3(N), 0, 0, Ad, Bd);
hipMemcpy(B, Bd, N * sizeof(long int), hipMemcpyDeviceToHost);
int passed = 0;
for (int i = 0; i < 512; i++) {
@@ -310,7 +310,7 @@ bool run_rintf() {
hipMalloc((void**)&Ad, SIZE);
hipMalloc((void**)&Bd, SIZE);
hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
hipLaunchKernel(test_rintf, dim3(1), dim3(N), 0, 0, Ad, Bd);
hipLaunchKernelGGL(test_rintf, dim3(1), dim3(N), 0, 0, Ad, Bd);
hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost);
int passed = 0;
for (int i = 0; i < 512; i++) {
@@ -344,7 +344,7 @@ bool run_llroundf() {
hipMalloc((void**)&Ad, SIZE);
hipMalloc((void**)&Bd, N * sizeof(long long int));
hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
hipLaunchKernel(test_llroundf, dim3(1), dim3(N), 0, 0, Ad, Bd);
hipLaunchKernelGGL(test_llroundf, dim3(1), dim3(N), 0, 0, Ad, Bd);
hipMemcpy(B, Bd, N * sizeof(long long int), hipMemcpyDeviceToHost);
int passed = 0;
for (int i = 0; i < 512; i++) {
@@ -377,7 +377,7 @@ bool run_lroundf() {
hipMalloc((void**)&Ad, SIZE);
hipMalloc((void**)&Bd, N * sizeof(long int));
hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
hipLaunchKernel(test_lroundf, dim3(1), dim3(N), 0, 0, Ad, Bd);
hipLaunchKernelGGL(test_lroundf, dim3(1), dim3(N), 0, 0, Ad, Bd);
hipMemcpy(B, Bd, N * sizeof(long int), hipMemcpyDeviceToHost);
int passed = 0;
for (int i = 0; i < 512; i++) {
@@ -420,7 +420,7 @@ bool run_norm3df() {
hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice);
hipMemcpy(Cd, C, SIZE, hipMemcpyHostToDevice);
hipLaunchKernel(test_norm3df, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd, Dd);
hipLaunchKernelGGL(test_norm3df, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd, Dd);
hipMemcpy(D, Dd, SIZE, hipMemcpyDeviceToHost);
int passed = 0;
for (int i = 0; i < 512; i++) {
@@ -469,7 +469,7 @@ bool run_norm4df() {
hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice);
hipMemcpy(Cd, C, SIZE, hipMemcpyHostToDevice);
hipMemcpy(Dd, D, SIZE, hipMemcpyHostToDevice);
hipLaunchKernel(test_norm4df, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd, Dd, Ed);
hipLaunchKernelGGL(test_norm4df, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd, Dd, Ed);
hipMemcpy(E, Ed, SIZE, hipMemcpyDeviceToHost);
int passed = 0;
for (int i = 0; i < 512; i++) {
@@ -510,7 +510,7 @@ bool run_normf() {
hipMalloc((void**)&Ad, SIZE);
hipMalloc((void**)&Bd, SIZE);
hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
hipLaunchKernel(test_normf, dim3(1), dim3(N), 0, 0, Ad, Bd);
hipLaunchKernelGGL(test_normf, dim3(1), dim3(N), 0, 0, Ad, Bd);
hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost);
int passed = 0;
for (int i = 0; i < 512; i++) {
@@ -547,7 +547,7 @@ bool run_rhypotf() {
hipMalloc((void**)&Cd, SIZE);
hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice);
hipLaunchKernel(test_rhypotf, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd);
hipLaunchKernelGGL(test_rhypotf, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd);
hipMemcpy(C, Cd, SIZE, hipMemcpyDeviceToHost);
int passed = 0;
for (int i = 0; i < 512; i++) {
@@ -590,7 +590,7 @@ bool run_rnorm3df() {
hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice);
hipMemcpy(Cd, C, SIZE, hipMemcpyHostToDevice);
hipLaunchKernel(test_rnorm3df, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd, Dd);
hipLaunchKernelGGL(test_rnorm3df, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd, Dd);
hipMemcpy(D, Dd, SIZE, hipMemcpyDeviceToHost);
int passed = 0;
for (int i = 0; i < 512; i++) {
@@ -639,7 +639,7 @@ bool run_rnorm4df() {
hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice);
hipMemcpy(Cd, C, SIZE, hipMemcpyHostToDevice);
hipMemcpy(Dd, D, SIZE, hipMemcpyHostToDevice);
hipLaunchKernel(test_rnorm4df, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd, Dd, Ed);
hipLaunchKernelGGL(test_rnorm4df, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd, Dd, Ed);
hipMemcpy(E, Ed, SIZE, hipMemcpyDeviceToHost);
int passed = 0;
for (int i = 0; i < 512; i++) {
@@ -680,7 +680,7 @@ bool run_rnormf() {
hipMalloc((void**)&Ad, SIZE);
hipMalloc((void**)&Bd, SIZE);
hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
hipLaunchKernel(test_rnormf, dim3(1), dim3(N), 0, 0, Ad, Bd);
hipLaunchKernelGGL(test_rnormf, dim3(1), dim3(N), 0, 0, Ad, Bd);
hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost);
int passed = 0;
for (int i = 0; i < 512; i++) {
@@ -712,7 +712,7 @@ bool run_erfinvf() {
hipMalloc((void**)&Ad, SIZE);
hipMalloc((void**)&Bd, SIZE);
hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
hipLaunchKernel(test_erfinvf, dim3(1), dim3(N), 0, 0, Ad, Bd);
hipLaunchKernelGGL(test_erfinvf, dim3(1), dim3(N), 0, 0, Ad, Bd);
hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost);
int passed = 0;
for (int i = 0; i < 512; i++) {