diff --git a/include/hcc_detail/hip_vector_types.h b/include/hcc_detail/hip_vector_types.h index f18bad8b89..c2f7a4796a 100644 --- a/include/hcc_detail/hip_vector_types.h +++ b/include/hcc_detail/hip_vector_types.h @@ -115,7 +115,74 @@ typedef hc::short_vector::double2 double2; typedef hc::short_vector::double3 double3; typedef hc::short_vector::double4 double4; +#if __HCC__ +#include"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 (T x) { VT t; t.x = x; return t; }; @@ -198,7 +265,7 @@ ONE_COMPONENT_ACCESS (double, double1); TWO_COMPONENT_ACCESS (double, double2); THREE_COMPONENT_ACCESS(double, double3); FOUR_COMPONENT_ACCESS (double, double4); - +*/ #endif diff --git a/src/device_util.cpp b/src/device_util.cpp index 8a99282cc1..d577cb291b 100644 --- a/src/device_util.cpp +++ b/src/device_util.cpp @@ -799,4 +799,36 @@ __device__ float __dsqrt_rn(double x) {return hc::fast_math::sqrt(x); }; __device__ float __dsqrt_ru(double x) {return hc::fast_math::sqrt(x); }; __device__ float __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; +} diff --git a/tests/src/hip_test_make_type.cpp b/tests/src/hip_test_make_type.cpp new file mode 100644 index 0000000000..dec87ec712 --- /dev/null +++ b/tests/src/hip_test_make_type.cpp @@ -0,0 +1,393 @@ +/* +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. +*/ +#include +#include +#include +#include +#include +#include "hip_runtime.h" +#include "test_common.h" + +#define HIP_ASSERT(x) (assert((x)==hipSuccess)) + + +#define WIDTH 8 +#define HEIGHT 8 + +#define NUM (WIDTH*HEIGHT) + +#define THREADS_PER_BLOCK_X 8 +#define THREADS_PER_BLOCK_Y 8 +#define THREADS_PER_BLOCK_Z 1 + + +__global__ void +vectoradd_char1(hipLaunchParm lp, + char1* a, const char1* bm, const char1* cm, int width, int height) + + { + int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; + int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; + + int i = y * width + x; + if ( i < (width * height)) { + a[i] = make_char1(bm[i].x) + make_char1(cm[i].x); + } + } + +__global__ void +vectoradd_char2(hipLaunchParm lp, + char2* a, const char2* bm, const char2* cm, int width, int height) + + { + int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; + int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; + + int i = y * width + x; + if ( i < (width * height)) { + a[i] = make_char2(bm[i].x, bm[i].y) + make_char2(cm[i].x, cm[i].y); + } +} + +__global__ void +vectoradd_char3(hipLaunchParm lp, + char3* a, const char3* bm, const char3* cm, int width, int height) + + { + int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; + int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; + + int i = y * width + x; + if ( i < (width * height)) { + a[i] = make_char3(bm[i].x, bm[i].y, bm[i].z) + make_char3(cm[i].x, cm[i].y, cm[i].z); + } +} +__global__ void +vectoradd_char4(hipLaunchParm lp, + char4* a, const char4* bm, const char4* cm, int width, int height) + + { + int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; + int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; + + int i = y * width + x; + if ( i < (width * height)) { + a[i] = make_char4(bm[i].x, bm[i].y, bm[i].z, bm[i].w) + make_char4(cm[i].x, cm[i].y, cm[i].z, cm[i].w); + } +} + + +#if 0 +__kernel__ void vectoradd_float(float* a, const float* b, const float* c, int width, int height) { + + + int x = blockDimX * blockIdx.x + threadIdx.x; + int y = blockDimY * blockIdy.y + threadIdx.y; + + int i = y * width + x; + if ( i < (width * height)) { + a[i] = b[i] + c[i]; + } +} +#endif + +using namespace std; + +template +bool dataTypesRun(){ + T* hostA; + T* hostB; + T* hostC; + + T* deviceA; + T* deviceB; + T* deviceC; + + int i; + int errors; + + hostA = (T*)malloc(NUM * sizeof(T)); + hostB = (T*)malloc(NUM * sizeof(T)); + hostC = (T*)malloc(NUM * sizeof(T)); + + // initialize the input data + for (i = 0; i < NUM; i++) { + hostB[i] = (T)i; + hostC[i] = (T)i; + } + + HIP_ASSERT(hipMalloc((void**)&deviceA, NUM * sizeof(T))); + HIP_ASSERT(hipMalloc((void**)&deviceB, NUM * sizeof(T))); + HIP_ASSERT(hipMalloc((void**)&deviceC, NUM * sizeof(T))); + + HIP_ASSERT(hipMemcpy(deviceB, hostB, NUM*sizeof(T), hipMemcpyHostToDevice)); + HIP_ASSERT(hipMemcpy(deviceC, hostC, NUM*sizeof(T), hipMemcpyHostToDevice)); + + hipLaunchKernel(HIP_KERNEL_NAME(vectoradd_char1), + dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y), + dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), + 0, 0, + deviceA ,deviceB ,deviceC ,WIDTH ,HEIGHT); + + HIP_ASSERT(hipMemcpy(hostA, deviceA, NUM*sizeof(T), hipMemcpyDeviceToHost)); + + bool ret = false; + // verify the results + errors = 0; + for (i = 0; i < NUM; i++) { + if (hostA[i] != (hostB[i] + hostC[i])) { + errors++; + } + } + if (errors!=0) { + printf("FAILED: %d errors\n",errors); + ret = false; + } else { + ret = true; + } + + HIP_ASSERT(hipFree(deviceA)); + HIP_ASSERT(hipFree(deviceB)); + HIP_ASSERT(hipFree(deviceC)); + + free(hostA); + free(hostB); + free(hostC); + + return ret; +} + +template +bool dataTypesRun(){ + T* hostA; + T* hostB; + T* hostC; + + T* deviceA; + T* deviceB; + T* deviceC; + + int i; + int errors; + + hostA = (T*)malloc(NUM * sizeof(T)); + hostB = (T*)malloc(NUM * sizeof(T)); + hostC = (T*)malloc(NUM * sizeof(T)); + + // initialize the input data + for (i = 0; i < NUM; i++) { + hostB[i] = (T)i; + hostC[i] = (T)i; + } + + HIP_ASSERT(hipMalloc((void**)&deviceA, NUM * sizeof(T))); + HIP_ASSERT(hipMalloc((void**)&deviceB, NUM * sizeof(T))); + HIP_ASSERT(hipMalloc((void**)&deviceC, NUM * sizeof(T))); + + HIP_ASSERT(hipMemcpy(deviceB, hostB, NUM*sizeof(T), hipMemcpyHostToDevice)); + HIP_ASSERT(hipMemcpy(deviceC, hostC, NUM*sizeof(T), hipMemcpyHostToDevice)); + + hipLaunchKernel(HIP_KERNEL_NAME(vectoradd_char1), + dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y), + dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), + 0, 0, + deviceA ,deviceB ,deviceC ,WIDTH ,HEIGHT); + + HIP_ASSERT(hipMemcpy(hostA, deviceA, NUM*sizeof(T), hipMemcpyDeviceToHost)); + + bool ret = false; + // verify the results + errors = 0; + for (i = 0; i < NUM; i++) { + if (hostA[i] != (hostB[i] + hostC[i])) { + errors++; + } + } + if (errors!=0) { + printf("FAILED: %d errors\n",errors); + ret = false; + } else { + ret = true; + } + + HIP_ASSERT(hipFree(deviceA)); + HIP_ASSERT(hipFree(deviceB)); + HIP_ASSERT(hipFree(deviceC)); + + free(hostA); + free(hostB); + free(hostC); + + return ret; +} + +template +bool dataTypesRun(){ + T* hostA; + T* hostB; + T* hostC; + + T* deviceA; + T* deviceB; + T* deviceC; + + int i; + int errors; + + hostA = (T*)malloc(NUM * sizeof(T)); + hostB = (T*)malloc(NUM * sizeof(T)); + hostC = (T*)malloc(NUM * sizeof(T)); + + // initialize the input data + for (i = 0; i < NUM; i++) { + hostB[i] = (T)i; + hostC[i] = (T)i; + } + + HIP_ASSERT(hipMalloc((void**)&deviceA, NUM * sizeof(T))); + HIP_ASSERT(hipMalloc((void**)&deviceB, NUM * sizeof(T))); + HIP_ASSERT(hipMalloc((void**)&deviceC, NUM * sizeof(T))); + + HIP_ASSERT(hipMemcpy(deviceB, hostB, NUM*sizeof(T), hipMemcpyHostToDevice)); + HIP_ASSERT(hipMemcpy(deviceC, hostC, NUM*sizeof(T), hipMemcpyHostToDevice)); + + hipLaunchKernel(HIP_KERNEL_NAME(vectoradd_char1), + dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y), + dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), + 0, 0, + deviceA ,deviceB ,deviceC ,WIDTH ,HEIGHT); + + HIP_ASSERT(hipMemcpy(hostA, deviceA, NUM*sizeof(T), hipMemcpyDeviceToHost)); + + bool ret = false; + // verify the results + errors = 0; + for (i = 0; i < NUM; i++) { + if (hostA[i] != (hostB[i] + hostC[i])) { + errors++; + } + } + if (errors!=0) { + printf("FAILED: %d errors\n",errors); + ret = false; + } else { + ret = true; + } + + HIP_ASSERT(hipFree(deviceA)); + HIP_ASSERT(hipFree(deviceB)); + HIP_ASSERT(hipFree(deviceC)); + + free(hostA); + free(hostB); + free(hostC); + + return ret; +} + +bool dataTypesRunChar4(){ + char4* hostA; + char4* hostB; + char4* hostC; + + char4* deviceA; + char4* deviceB; + char4* deviceC; + + int i; + int errors; + + hostA = (T*)malloc(NUM * sizeof(T)); + hostB = (T*)malloc(NUM * sizeof(T)); + hostC = (T*)malloc(NUM * sizeof(T)); + + // initialize the input data + for (i = 0; i < NUM; i++) { + hostB[i] = (T)i; + hostC[i] = (T)i; + } + + HIP_ASSERT(hipMalloc((void**)&deviceA, NUM * sizeof(T))); + HIP_ASSERT(hipMalloc((void**)&deviceB, NUM * sizeof(T))); + HIP_ASSERT(hipMalloc((void**)&deviceC, NUM * sizeof(T))); + + HIP_ASSERT(hipMemcpy(deviceB, hostB, NUM*sizeof(T), hipMemcpyHostToDevice)); + HIP_ASSERT(hipMemcpy(deviceC, hostC, NUM*sizeof(T), hipMemcpyHostToDevice)); + + hipLaunchKernel(HIP_KERNEL_NAME(vectoradd_char1), + dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y), + dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), + 0, 0, + deviceA ,deviceB ,deviceC ,WIDTH ,HEIGHT); + + HIP_ASSERT(hipMemcpy(hostA, deviceA, NUM*sizeof(T), hipMemcpyDeviceToHost)); + + bool ret = false; + // verify the results + errors = 0; + for (i = 0; i < NUM; i++) { + if (hostA[i] != (hostB[i] + hostC[i])) { + errors++; + } + } + if (errors!=0) { + printf("FAILED: %d errors\n",errors); + ret = false; + } else { + ret = true; + } + + HIP_ASSERT(hipFree(deviceA)); + HIP_ASSERT(hipFree(deviceB)); + HIP_ASSERT(hipFree(deviceC)); + + free(hostA); + free(hostB); + free(hostC); + + return ret; +} + +int main() { + + hipDeviceProp_t devProp; + hipGetDeviceProperties(&devProp, 0); + cout << " System minor " << devProp.minor << endl; + cout << " System major " << devProp.major << endl; + cout << " agent prop name " << devProp.name << endl; + + int errors; + + errors = dataTypesRun() & + dataTypesRun() & + dataTypesRun() & + dataTypesRun(); + + + //hipResetDefaultAccelerator(); + if(errors == 1){ + passed(); + }else{ + std::cout<<"Failed Float"<