Added proper device data types
Change-Id: I42029635ff68c3c13a764a3eda6447e6c77878c6
[ROCm/hip commit: 39910029a6]
This commit is contained in:
Aditya Atluri
@@ -20,13 +20,18 @@ THE SOFTWARE.
|
||||
#ifndef HIP_HCC_DETAIL_DEVICE_FUNCTIONS_H
|
||||
#define HIP_HCC_DETAIL_DEVICE_FUNCTIONS_H
|
||||
|
||||
#include "hip_runtime.h"
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <hip/hip_vector_types.h>
|
||||
|
||||
__device__ float __int_as_float (int x);
|
||||
|
||||
__device__ double __hiloint2double (int hi, int lo);
|
||||
|
||||
extern __HIP_DEVICE__ double __longlong_as_double(long long int x);
|
||||
extern __HIP_DEVICE__ long long int __double_as_longlong(double x);
|
||||
__device__ char4 __hip_hc_add8pk(char4, char4);
|
||||
__device__ char4 __hip_hc_sub8pk(char4, char4);
|
||||
__device__ char4 __hip_hc_mul8pk(char4, char4);
|
||||
|
||||
extern __device__ double __longlong_as_double(long long int x);
|
||||
extern __device__ long long int __double_as_longlong(double x);
|
||||
|
||||
#endif
|
||||
|
||||
@@ -25,8 +25,8 @@ THE SOFTWARE.
|
||||
|
||||
#if __HCC__
|
||||
#if __hcc_workweek__ >= 16164
|
||||
#include "hip/hip_vector_types.h"
|
||||
#include "hip/hcc_detail/host_defines.h"
|
||||
#include "hip_vector_types.h"
|
||||
#include "host_defines.h"
|
||||
|
||||
|
||||
__device__ char __ldg(const char* );
|
||||
@@ -75,4 +75,3 @@ __device__ double2 __ldg(const double2* );
|
||||
#endif // __HCC__
|
||||
|
||||
#endif // HIP_LDG_H
|
||||
|
||||
|
||||
@@ -46,6 +46,7 @@ THE SOFTWARE.
|
||||
#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
|
||||
@@ -815,9 +816,6 @@ extern "C" __device__ void* __hip_hc_free(void *ptr);
|
||||
//extern "C" __device__ void* malloc(size_t size);
|
||||
//extern "C" __device__ void* free(void *ptr);
|
||||
|
||||
extern "C" __device__ char4 __hip_hc_add8pk(char4, char4);
|
||||
extern "C" __device__ char4 __hip_hc_sub8pk(char4, char4);
|
||||
extern "C" __device__ char4 __hip_hc_mul8pk(char4, char4);
|
||||
|
||||
#define __syncthreads() hc_barrier(CLK_LOCAL_MEM_FENCE)
|
||||
|
||||
|
||||
File diff suppressed because it is too large
Load Diff
@@ -14,7 +14,7 @@ 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
|
||||
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.
|
||||
@@ -129,34 +129,34 @@ __device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask
|
||||
|
||||
__device__ char4 __hip_hc_add8pk(char4 in1, char4 in2) {
|
||||
char4 out;
|
||||
unsigned one1 = in1.val & MASK1;
|
||||
unsigned one2 = in2.val & MASK1;
|
||||
out.val = (one1 + one2) & MASK1;
|
||||
one1 = in1.val & MASK2;
|
||||
one2 = in2.val & MASK2;
|
||||
out.val = out.val | ((one1 + one2) & MASK2);
|
||||
unsigned one1 = in1.a & MASK1;
|
||||
unsigned one2 = in2.a & MASK1;
|
||||
out.a = (one1 + one2) & MASK1;
|
||||
one1 = in1.a & MASK2;
|
||||
one2 = in2.a & MASK2;
|
||||
out.a = out.a | ((one1 + one2) & MASK2);
|
||||
return out;
|
||||
}
|
||||
|
||||
__device__ char4 __hip_hc_sub8pk(char4 in1, char4 in2) {
|
||||
char4 out;
|
||||
unsigned one1 = in1.val & MASK1;
|
||||
unsigned one2 = in2.val & MASK1;
|
||||
out.val = (one1 - one2) & MASK1;
|
||||
one1 = in1.val & MASK2;
|
||||
one2 = in2.val & MASK2;
|
||||
out.val = out.val | ((one1 - one2) & MASK2);
|
||||
unsigned one1 = in1.a & MASK1;
|
||||
unsigned one2 = in2.a & MASK1;
|
||||
out.a = (one1 - one2) & MASK1;
|
||||
one1 = in1.a & MASK2;
|
||||
one2 = in2.a & MASK2;
|
||||
out.a = out.a | ((one1 - one2) & MASK2);
|
||||
return out;
|
||||
}
|
||||
|
||||
__device__ char4 __hip_hc_mul8pk(char4 in1, char4 in2) {
|
||||
char4 out;
|
||||
unsigned one1 = in1.val & MASK1;
|
||||
unsigned one2 = in2.val & MASK1;
|
||||
out.val = (one1 * one2) & MASK1;
|
||||
one1 = in1.val & MASK2;
|
||||
one2 = in2.val & MASK2;
|
||||
out.val = out.val | ((one1 * one2) & MASK2);
|
||||
unsigned one1 = in1.a & MASK1;
|
||||
unsigned one2 = in2.a & MASK1;
|
||||
out.a = (one1 * one2) & MASK1;
|
||||
one1 = in1.a & MASK2;
|
||||
one2 = in2.a & MASK2;
|
||||
out.a = out.a | ((one1 * one2) & MASK2);
|
||||
return out;
|
||||
}
|
||||
|
||||
@@ -2179,426 +2179,17 @@ __device__ double __hip_fast_dsqrt_rz(double x) {
|
||||
return hc::fast_math::sqrt(x);
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ char1 make_char1(signed char x)
|
||||
{
|
||||
char1 c1;
|
||||
c1.x = x;
|
||||
return c1;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ char2 make_char2(signed char x, signed char y)
|
||||
{
|
||||
char2 c2;
|
||||
c2.x = x;
|
||||
c2.y = y;
|
||||
return c2;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ char3 make_char3(signed char x, signed char y, signed char z)
|
||||
{
|
||||
char3 c3;
|
||||
c3.x = x;
|
||||
c3.y = y;
|
||||
c3.z = z;
|
||||
return c3;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ char4 make_char4(signed char x, signed char y, signed char z, signed char w)
|
||||
{
|
||||
char4 c4;
|
||||
c4.x = x;
|
||||
c4.y = y;
|
||||
c4.z = z;
|
||||
c4.w = w;
|
||||
return c4;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ short1 make_short1(short x)
|
||||
{
|
||||
short1 s1;
|
||||
s1.x = x;
|
||||
return s1;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ short2 make_short2(short x, short y)
|
||||
{
|
||||
short2 s2;
|
||||
s2.x = x;
|
||||
s2.y = y;
|
||||
return s2;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ short3 make_short3(short x, short y, short z)
|
||||
{
|
||||
short3 s3;
|
||||
s3.x = x;
|
||||
s3.y = y;
|
||||
s3.z = z;
|
||||
return s3;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ short4 make_short4(short x, short y, short z, short w)
|
||||
{
|
||||
short4 s4;
|
||||
s4.x = x;
|
||||
s4.y = y;
|
||||
s4.z = z;
|
||||
s4.w = w;
|
||||
return s4;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ int1 make_int1(int x)
|
||||
{
|
||||
int1 i1;
|
||||
i1.x = x;
|
||||
return i1;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ int2 make_int2(int x, int y)
|
||||
{
|
||||
int2 i2;
|
||||
i2.x = x;
|
||||
i2.y = y;
|
||||
return i2;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ int3 make_int3(int x, int y, int z)
|
||||
{
|
||||
int3 i3;
|
||||
i3.x = x;
|
||||
i3.y = y;
|
||||
i3.z = z;
|
||||
return i3;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ int4 make_int4(int x, int y, int z, int w)
|
||||
{
|
||||
int4 i4;
|
||||
i4.x = x;
|
||||
i4.y = y;
|
||||
i4.z = z;
|
||||
i4.w = w;
|
||||
return i4;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ long1 make_long1(long x)
|
||||
{
|
||||
long1 l1;
|
||||
l1.x = x;
|
||||
return l1;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ long2 make_long2(long x, long y)
|
||||
{
|
||||
long2 l2;
|
||||
l2.x = x;
|
||||
l2.y = y;
|
||||
return l2;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ long3 make_long3(long x, long y, long z)
|
||||
{
|
||||
long3 l3;
|
||||
l3.x = x;
|
||||
l3.y = y;
|
||||
l3.z = z;
|
||||
return l3;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ long4 make_long4(long x, long y, long z, long w)
|
||||
{
|
||||
long4 l4;
|
||||
l4.x = x;
|
||||
l4.y = y;
|
||||
l4.z = z;
|
||||
l4.w = w;
|
||||
return l4;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ longlong1 make_longlong1(long long x)
|
||||
{
|
||||
longlong1 l1;
|
||||
l1.x = x;
|
||||
return l1;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ longlong2 make_longlong2(long long x, long long y)
|
||||
{
|
||||
longlong2 l2;
|
||||
l2.x = x;
|
||||
l2.y = y;
|
||||
return l2;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ longlong3 make_longlong3(long long x, long long y, long long z)
|
||||
{
|
||||
longlong3 l3;
|
||||
l3.x = x;
|
||||
l3.y = y;
|
||||
l3.z = z;
|
||||
return l3;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ longlong4 make_longlong4(long long x, long long y, long long z, long long w)
|
||||
{
|
||||
longlong4 l4;
|
||||
l4.x = x;
|
||||
l4.y = y;
|
||||
l4.z = z;
|
||||
l4.w = w;
|
||||
return l4;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ uchar1 make_uchar1(unsigned char x)
|
||||
{
|
||||
uchar1 c1;
|
||||
c1.x = x;
|
||||
return c1;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ uchar2 make_uchar2(unsigned char x, unsigned char y)
|
||||
{
|
||||
uchar2 c2;
|
||||
c2.x = x;
|
||||
c2.y = y;
|
||||
return c2;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ uchar3 make_uchar3(unsigned char x, unsigned char y, unsigned char z)
|
||||
{
|
||||
uchar3 c3;
|
||||
c3.x = x;
|
||||
c3.y = y;
|
||||
c3.z = z;
|
||||
return c3;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ uchar4 make_uchar4(unsigned char x, unsigned char y, unsigned char z, unsigned char w)
|
||||
{
|
||||
uchar4 c4;
|
||||
c4.x = x;
|
||||
c4.y = y;
|
||||
c4.z = z;
|
||||
c4.w = w;
|
||||
return c4;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ ushort1 make_ushort1(unsigned short x)
|
||||
{
|
||||
ushort1 s1;
|
||||
s1.x = x;
|
||||
return s1;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ ushort2 make_ushort2(unsigned short x, unsigned short y)
|
||||
{
|
||||
ushort2 s2;
|
||||
s2.x = x;
|
||||
s2.y = y;
|
||||
return s2;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ ushort3 make_ushort3(unsigned short x, unsigned short y, unsigned short z)
|
||||
{
|
||||
ushort3 s3;
|
||||
s3.x = x;
|
||||
s3.y = y;
|
||||
s3.z = z;
|
||||
return s3;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ ushort4 make_ushort4(unsigned short x, unsigned short y, unsigned short z, unsigned short w)
|
||||
{
|
||||
ushort4 s4;
|
||||
s4.x = x;
|
||||
s4.y = y;
|
||||
s4.z = z;
|
||||
s4.w = w;
|
||||
return s4;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ uint1 make_uint1(unsigned int x)
|
||||
{
|
||||
uint1 i1;
|
||||
i1.x = x;
|
||||
return i1;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ uint2 make_uint2(unsigned int x, unsigned int y)
|
||||
{
|
||||
uint2 i2;
|
||||
i2.x = x;
|
||||
i2.y = y;
|
||||
return i2;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ uint3 make_uint3(unsigned int x, unsigned int y, unsigned int z)
|
||||
{
|
||||
uint3 i3;
|
||||
i3.x = x;
|
||||
i3.y = y;
|
||||
i3.z = z;
|
||||
return i3;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ uint4 make_uint4(unsigned int x, unsigned int y, unsigned int z, unsigned int w)
|
||||
{
|
||||
uint4 i4;
|
||||
i4.x = x;
|
||||
i4.y = y;
|
||||
i4.z = z;
|
||||
i4.w = w;
|
||||
return i4;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ ulong1 make_ulong1(unsigned long x)
|
||||
{
|
||||
ulong1 l1;
|
||||
l1.x = x;
|
||||
return l1;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ ulong2 make_ulong2(unsigned long x, unsigned long y)
|
||||
{
|
||||
ulong2 l2;
|
||||
l2.x = x;
|
||||
l2.y = y;
|
||||
return l2;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ ulong3 make_ulong3(unsigned long x, unsigned long y, unsigned long z)
|
||||
{
|
||||
ulong3 l3;
|
||||
l3.x = x;
|
||||
l3.y = y;
|
||||
l3.z = z;
|
||||
return l3;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ ulong4 make_ulong4(unsigned long x, unsigned long y, unsigned long z, unsigned long w)
|
||||
{
|
||||
ulong4 l4;
|
||||
l4.x = x;
|
||||
l4.y = y;
|
||||
l4.z = z;
|
||||
l4.w = w;
|
||||
return l4;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ ulonglong1 make_ulonglong1(unsigned long long x)
|
||||
{
|
||||
ulonglong1 l1;
|
||||
l1.x = x;
|
||||
return l1;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ ulonglong2 make_ulonglong2(unsigned long long x, unsigned long long y)
|
||||
{
|
||||
ulonglong2 l2;
|
||||
l2.x = x;
|
||||
l2.y = y;
|
||||
return l2;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ ulonglong3 make_ulonglong3(unsigned long long x, unsigned long long y, unsigned long long z)
|
||||
{
|
||||
ulonglong3 l3;
|
||||
l3.x = x;
|
||||
l3.y = y;
|
||||
l3.z = z;
|
||||
return l3;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ ulonglong4 make_ulonglong4(unsigned long long x, unsigned long long y, unsigned long long z, unsigned long long w)
|
||||
{
|
||||
ulonglong4 l4;
|
||||
l4.x = x;
|
||||
l4.y = y;
|
||||
l4.z = z;
|
||||
l4.w = w;
|
||||
return l4;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ float1 make_float1(float x)
|
||||
{
|
||||
float1 f1;
|
||||
f1.x = x;
|
||||
return f1;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ float2 make_float2(float x, float y)
|
||||
{
|
||||
float2 f2;
|
||||
f2.x = x;
|
||||
f2.y = y;
|
||||
return f2;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ float3 make_float3(float x, float y, float z)
|
||||
{
|
||||
float3 f3;
|
||||
f3.x = x;
|
||||
f3.y = y;
|
||||
f3.z = z;
|
||||
return f3;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ float4 make_float4(float x, float y, float z, float w)
|
||||
{
|
||||
float4 f4;
|
||||
f4.x = x;
|
||||
f4.y = y;
|
||||
f4.z = z;
|
||||
f4.w = w;
|
||||
return f4;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ double1 make_double1(double x)
|
||||
{
|
||||
double1 d1;
|
||||
d1.x = x;
|
||||
return d1;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ double2 make_double2(double x, double y)
|
||||
{
|
||||
double2 d2;
|
||||
d2.x = x;
|
||||
d2.y = y;
|
||||
return d2;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ double3 make_double3(double x, double y, double z)
|
||||
{
|
||||
double3 d3;
|
||||
d3.x = x;
|
||||
d3.y = y;
|
||||
d3.z = z;
|
||||
return d3;
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ double4 make_double4(double x, double y, double z, double w)
|
||||
{
|
||||
double4 d4;
|
||||
d4.x = x;
|
||||
d4.y = y;
|
||||
d4.z = z;
|
||||
d4.w = w;
|
||||
return d4;
|
||||
}
|
||||
|
||||
|
||||
__HIP_DEVICE__ double __longlong_as_double(long long int x)
|
||||
__device__ double __longlong_as_double(long long int x)
|
||||
{
|
||||
return static_cast<double>(x);
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ long long __double_as_longlong(double x)
|
||||
__device__ long long __double_as_longlong(double x)
|
||||
{
|
||||
return static_cast<long long>(x);
|
||||
}
|
||||
|
||||
__HIP_DEVICE__ void __threadfence_system(void){
|
||||
__device__ void __threadfence_system(void){
|
||||
// no-op
|
||||
}
|
||||
|
||||
@@ -3391,5 +2982,3 @@ __host__ double norm4d(double a, double b, double c, double d)
|
||||
{
|
||||
return std::sqrt(a*a + b*b + c*c + d*d);
|
||||
}
|
||||
|
||||
|
||||
|
||||
@@ -20,23 +20,22 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include <hc.hpp>
|
||||
|
||||
#include "hip/hcc_detail/hip_ldg.h"
|
||||
#include "hip/hcc_detail/hip_vector_types.h"
|
||||
|
||||
__device__ char __ldg(const char* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
return *ptr;
|
||||
}
|
||||
|
||||
__device__ char2 __ldg(const char2* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
return *ptr;
|
||||
}
|
||||
|
||||
__device__ char4 __ldg(const char4* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
return *ptr;
|
||||
}
|
||||
|
||||
__device__ signed char __ldg(const signed char* ptr)
|
||||
@@ -169,6 +168,3 @@ __device__ double2 __ldg(const double2* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
}
|
||||
|
||||
|
||||
|
||||
|
||||
@@ -32,6 +32,7 @@ THE SOFTWARE.
|
||||
#include <stdlib.h>
|
||||
#include<iostream>
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "hip/hip_vector_types.h"
|
||||
#include "test_common.h"
|
||||
|
||||
#if (__hcc_workweek__ >= 16164) || defined (__HIP_PLATFORM_NVCC__)
|
||||
@@ -389,4 +390,3 @@ int main() {
|
||||
}
|
||||
|
||||
#endif
|
||||
|
||||
|
||||
Reference in New Issue
Block a user