reorg make_datatype feature
Этот коммит содержится в:
Aditya Atluri
@@ -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
|
||||
|
||||
|
||||
@@ -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;
|
||||
}
|
||||
|
||||
@@ -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 <assert.h>
|
||||
#include <stdio.h>
|
||||
#include <algorithm>
|
||||
#include <stdlib.h>
|
||||
#include<iostream>
|
||||
#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<typename T>
|
||||
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<typename T>
|
||||
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<typename T>
|
||||
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<char1>() &
|
||||
dataTypesRun<char2>() &
|
||||
dataTypesRun<char3>() &
|
||||
dataTypesRun<char4>();
|
||||
|
||||
|
||||
//hipResetDefaultAccelerator();
|
||||
if(errors == 1){
|
||||
passed();
|
||||
}else{
|
||||
std::cout<<"Failed Float"<<std::endl;
|
||||
return -1;
|
||||
}
|
||||
}
|
||||
Ссылка в новой задаче
Block a user