Merge 'develop' into 'amd-staging'
Change-Id: Ia94cdfc147fe4e14c161eb4a1a9ff246754dba25
Этот коммит содержится в:
Jenkins
@@ -10,9 +10,7 @@
|
||||
"Unit_hipGetDeviceFlags_Positive_Context",
|
||||
"Unit_hipIpcCloseMemHandle_Negative_Close_In_Originating_Process",
|
||||
"Unit_hipIpcOpenMemHandle_Negative_Open_In_Creating_Process",
|
||||
"Unit_hipDeviceGetPCIBusId_Negative_PartialFill",
|
||||
"Unit_hipInit_Negative",
|
||||
"Unit_hipMemset_Negative_OutOfBoundsPtr",
|
||||
"Unit_hipDeviceReset_Positive_Basic",
|
||||
"Unit_hipDeviceReset_Positive_Threaded",
|
||||
"Unit_hipGraphMemcpyNodeSetParamsToSymbol_Positive_Basic",
|
||||
@@ -40,15 +38,8 @@
|
||||
"Note: Following four tests disabled due to defect - EXSWHTEC-203",
|
||||
"Unit_hipGraphAddMemsetNode_Positive_Basic - uint16_t",
|
||||
"Unit_hipGraphAddMemsetNode_Positive_Basic - uint32_t",
|
||||
"Unit_hipGraphMemsetNodeSetParams_Positive_Basic - uint8_t",
|
||||
"Unit_hipGraphMemsetNodeSetParams_Positive_Basic - uint16_t",
|
||||
"Unit_hipGraphMemsetNodeSetParams_Positive_Basic - uint32_t",
|
||||
"Unit_hipGraphExecMemsetNodeSetParams_Positive_Basic - uint8_t",
|
||||
"Unit_hipGraphExecMemsetNodeSetParams_Positive_Basic - uint16_t",
|
||||
"Unit_hipGraphExecMemsetNodeSetParams_Positive_Basic - uint32_t",
|
||||
"Unit_hipStreamSetCaptureDependencies_Positive_Functional",
|
||||
"Note: Test disabled due to defect - EXSWHTEC-207",
|
||||
"Unit_hipGraphExecMemsetNodeSetParams_Negative_Updating_Non1D_Node",
|
||||
"Unit_hipIpcGetMemHandle_Positive_Unique_Handles_Separate_Allocations",
|
||||
"Unit_hipStreamCreateWithFlags_DefaultStreamInteraction",
|
||||
"Unit_hipStreamWaitEvent_UninitializedStream_Negative",
|
||||
|
||||
@@ -38,6 +38,7 @@ set(AMD_TEST_SRC
|
||||
floatTM.cc
|
||||
hipMathFunctions.cc
|
||||
hmax_hmin.cc
|
||||
bfloat16.cc
|
||||
)
|
||||
set(AMD_ARCH_SPEC_TEST_SRC
|
||||
AtomicAdd_Coherent_withunsafeflag.cc
|
||||
|
||||
@@ -0,0 +1,511 @@
|
||||
/*
|
||||
Copyright (c) 2023 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 <hip_test_common.hh>
|
||||
#include <hip/hip_bf16.h>
|
||||
|
||||
#include <cmath>
|
||||
#include <memory>
|
||||
#include <limits>
|
||||
#include <algorithm>
|
||||
|
||||
// Struct used to generate floats from combination of various componenets
|
||||
union float_holder {
|
||||
float fp32;
|
||||
struct parts_ {
|
||||
unsigned int fp32_mantisa : 16; // ignored for bf16
|
||||
unsigned int bf16_mantisa : 7; // bf16 mantisa
|
||||
unsigned int exponent : 8;
|
||||
unsigned int sign : 1;
|
||||
} parts;
|
||||
unsigned int u32;
|
||||
};
|
||||
|
||||
std::vector<float> getAllBF16() {
|
||||
constexpr unsigned char max_mantissa = std::numeric_limits<unsigned char>::max() >> 1; // 7 bits
|
||||
const size_t max_bf16_num =
|
||||
2 /*sign*/ * std::pow(2, 8) /*exponent*/ * std::pow(2, 7) /*mantissa*/;
|
||||
|
||||
std::vector<float> f_in;
|
||||
f_in.reserve(max_bf16_num);
|
||||
|
||||
|
||||
for (size_t s = 0; s <= 1; s++) { // sign
|
||||
for (size_t e = 0; e <= std::numeric_limits<unsigned char>::max(); e++) { // expo
|
||||
for (size_t m = 0; m <= max_mantissa; m++) { // man
|
||||
float_holder hold;
|
||||
hold.u32 = 0; // Init - clear all bits
|
||||
|
||||
hold.parts.sign = s;
|
||||
hold.parts.exponent = e;
|
||||
hold.parts.bf16_mantisa = m;
|
||||
|
||||
f_in.push_back(hold.fp32);
|
||||
}
|
||||
}
|
||||
}
|
||||
return f_in;
|
||||
}
|
||||
|
||||
enum MathOp { Add = 0, Sub, Mul, Div, LastOp = Div };
|
||||
|
||||
__device__ __hip_bfloat16 bf16_math(__hip_bfloat16 a, __hip_bfloat16 b, MathOp op) {
|
||||
switch (op) {
|
||||
case Add:
|
||||
return __hadd(a, b);
|
||||
case Sub:
|
||||
return __hsub(a, b);
|
||||
case Mul:
|
||||
return __hmul(a, b);
|
||||
case Div:
|
||||
return __hdiv(a, b);
|
||||
}
|
||||
}
|
||||
|
||||
__device__ float fp32_math(float a, float b, MathOp op) {
|
||||
switch (op) {
|
||||
case Add:
|
||||
return a + b;
|
||||
case Sub:
|
||||
return a - b;
|
||||
case Mul:
|
||||
return a * b;
|
||||
case Div:
|
||||
return a / b;
|
||||
}
|
||||
}
|
||||
|
||||
__device__ void bf16_math_op_kernel(float* a, float* b, float* c) {
|
||||
for (int i = Add; i < LastOp; i++) {
|
||||
auto op = static_cast<MathOp>(i);
|
||||
c[i] = __bfloat162float(bf16_math(__float2bfloat16(a[i]), __float2bfloat16(b[i]), op));
|
||||
}
|
||||
}
|
||||
|
||||
__device__ void fp32_math_op_kernel(float* a, float* b, float* c) {
|
||||
for (int i = Add; i < LastOp; i++) {
|
||||
auto op = static_cast<MathOp>(i);
|
||||
c[i] = fp32_math(a[i], b[i], op);
|
||||
}
|
||||
}
|
||||
|
||||
// c = a MathOp b in fp32
|
||||
// conv_res = a MathOp b in bf16 converted back to fp32
|
||||
__global__ void do_math(float* a, float* b, float* c, float* conv_res) {
|
||||
fp32_math_op_kernel(a, b, c);
|
||||
bf16_math_op_kernel(a, b, conv_res);
|
||||
}
|
||||
|
||||
// Convert float -> bfloat16 -> float
|
||||
__global__ void fp32_bf16_fp32(float* a, float* c, size_t size) {
|
||||
auto i = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (i < size) {
|
||||
auto b = __float2bfloat16(a[i]);
|
||||
c[i] = __bfloat162float(b);
|
||||
}
|
||||
}
|
||||
|
||||
__device__ unsigned bool_to_unsigned(bool in) { return in ? 1u : 0u; }
|
||||
|
||||
// Test equal compare
|
||||
__global__ void bf16_is_equal(float* val, unsigned* res, size_t size) {
|
||||
auto i = threadIdx.x;
|
||||
if (i < size) {
|
||||
auto v1 = __float2bfloat16(val[i]);
|
||||
auto v2 = __float2bfloat16(val[i]);
|
||||
res[i] =
|
||||
bool_to_unsigned((__heq(v1, v2) && __hge(v1, v2) &&
|
||||
__hle(v1, v2))); // Equal, Greater Equal, Less Equal should all have true
|
||||
}
|
||||
}
|
||||
|
||||
// Test other compare functions
|
||||
__global__ void bf16_compare(float* val, unsigned* res, size_t size) {
|
||||
auto i = threadIdx.x;
|
||||
if (i < size) {
|
||||
__hip_bfloat16 v1 = __float2bfloat16(val[i]);
|
||||
__hip_bfloat16 v2 = __float2bfloat16(val[i]);
|
||||
|
||||
v1 = __hadd(v1, v2); // v1 = v1 + v2
|
||||
bool r1 = (__hlt(v2, v1) && __hgt(v1, v2) && // v1 > v2
|
||||
__hne(v1, v2) && // v1 != v2
|
||||
__heq(__hmax(v1, v2), v1) && // max(v1,v2) == v1
|
||||
__heq(__hmin(v1, v2), v2)); // min(v1,v2) == v2
|
||||
|
||||
v1 = __hsub(v1, v2); // Back to v1's original value
|
||||
bool r2 = __heq(v1, v2); // v1 == v2
|
||||
|
||||
v1 = __hmul(v1, v2); // v1 = v1 * v2, both have same values so square it
|
||||
bool r3 = __heq(v1, __hmul(v2, v2)); // v1 == (v2 * v2)
|
||||
|
||||
v1 = hsqrt(v1); // Back to v1's original value
|
||||
bool r4 = __heq(v1, v2); // v1 == v2
|
||||
|
||||
v1 = __hdiv(v1, v2); // v1 = v1/v2, both have same values
|
||||
bool r5 = __heq(v1, __float2bfloat16(1.0f)); // v1 == 1.0f
|
||||
|
||||
// Uncomment to debug
|
||||
// printf("%u - %u - %u - %u - %u\n", bool_to_unsigned(r1), bool_to_unsigned(r2),
|
||||
// bool_to_unsigned(r3), bool_to_unsigned(r4), bool_to_unsigned(r5));
|
||||
|
||||
res[i] = bool_to_unsigned(r1 && r2 && r3 && r4 && r5);
|
||||
}
|
||||
}
|
||||
|
||||
// Convert to bits
|
||||
__global__ void bf16_conv_bits(float* val, unsigned short* res, size_t size) {
|
||||
auto i = threadIdx.x;
|
||||
if (i < size) {
|
||||
__hip_bfloat16 v1 = __float2bfloat16(val[i]);
|
||||
res[i] = *reinterpret_cast<unsigned short*>(&v1);
|
||||
}
|
||||
}
|
||||
|
||||
__global__ void bf16_neg(float* in, float* out, size_t size) {
|
||||
auto i = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (i < size) {
|
||||
out[i] = __bfloat162float(__hneg(__float2bfloat16(in[i])));
|
||||
}
|
||||
}
|
||||
|
||||
__global__ void bf16_to_short(float* in, short* s_res, unsigned short* u_res, size_t size) {
|
||||
auto i = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (i < size) {
|
||||
s_res[i] = __bfloat16_as_short(__float2bfloat16(in[i]));
|
||||
u_res[i] = __bfloat16_as_ushort(__float2bfloat16(in[i]));
|
||||
}
|
||||
}
|
||||
|
||||
__global__ void short_to_bf16(short* in, float* out, size_t size) {
|
||||
auto i = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (i < size) {
|
||||
out[i] = __bfloat162float(__short_as_bfloat16(in[i]));
|
||||
}
|
||||
}
|
||||
|
||||
__global__ void ushort_to_bf16(unsigned short* in, float* out, size_t size) {
|
||||
auto i = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (i < size) {
|
||||
out[i] = __bfloat162float(__ushort_as_bfloat16(in[i]));
|
||||
}
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_bf16_basic") {
|
||||
auto f_in = getAllBF16();
|
||||
auto max_bf16_num = f_in.size();
|
||||
|
||||
SECTION("Conversion float to bfloat16 to float") {
|
||||
constexpr size_t size = 256;
|
||||
|
||||
float *d_a, *d_c;
|
||||
HIP_CHECK(hipMalloc(&d_a, sizeof(float) * size));
|
||||
HIP_CHECK(hipMalloc(&d_c, sizeof(float) * size));
|
||||
|
||||
auto h_a = std::make_unique<float[]>(size);
|
||||
auto h_c = std::make_unique<float[]>(size);
|
||||
|
||||
for (size_t i = 0; i < size; i++) {
|
||||
h_a[i] = i + 1.25;
|
||||
}
|
||||
|
||||
HIP_CHECK(hipMemcpy(d_a, h_a.get(), sizeof(float) * size, hipMemcpyHostToDevice));
|
||||
|
||||
fp32_bf16_fp32<<<1, size>>>(d_a, d_c, size);
|
||||
|
||||
HIP_CHECK(hipMemcpy(h_c.get(), d_c, sizeof(float) * size, hipMemcpyDeviceToHost));
|
||||
|
||||
for (size_t i = 0; i < size; i++) {
|
||||
INFO("Initial: " << h_a[i] << " - After Conv: " << h_c[i]);
|
||||
// The relative error should be less than 1/(2^7) since bfloat16 has 7 bits mantissa.
|
||||
REQUIRE((std::fabs(h_c[i] - h_a[i]) / h_a[i]) < (1.0 / 128.0f));
|
||||
}
|
||||
|
||||
HIP_CHECK(hipFree(d_a));
|
||||
HIP_CHECK(hipFree(d_c));
|
||||
}
|
||||
|
||||
SECTION("Math Op Accuracy") {
|
||||
constexpr size_t size = static_cast<size_t>(LastOp);
|
||||
|
||||
float f_val1[size], f_val2[size], f_res[size], bf_res[size];
|
||||
for (size_t i = 0; i < size; i++) {
|
||||
f_val1[i] = i + 1.50;
|
||||
f_val2[i] = i + 1.25;
|
||||
}
|
||||
|
||||
float *df_val1, *df_val2, *float_res, *bf16_res;
|
||||
HIP_CHECK(hipMalloc(&df_val1, sizeof(float) * size));
|
||||
HIP_CHECK(hipMalloc(&df_val2, sizeof(float) * size));
|
||||
HIP_CHECK(hipMalloc(&float_res, sizeof(float) * size));
|
||||
HIP_CHECK(hipMalloc(&bf16_res, sizeof(float) * size));
|
||||
|
||||
HIP_CHECK(hipMemcpy(df_val1, f_val1, sizeof(float) * size, hipMemcpyHostToDevice));
|
||||
HIP_CHECK(hipMemcpy(df_val2, f_val2, sizeof(float) * size, hipMemcpyHostToDevice));
|
||||
|
||||
do_math<<<1, 1>>>(df_val1, df_val2, float_res, bf16_res);
|
||||
|
||||
HIP_CHECK(hipMemcpy(f_res, float_res, sizeof(float) * size, hipMemcpyDeviceToHost));
|
||||
HIP_CHECK(hipMemcpy(bf_res, bf16_res, sizeof(float) * size, hipMemcpyDeviceToHost));
|
||||
|
||||
HIP_CHECK(hipFree(df_val1));
|
||||
HIP_CHECK(hipFree(df_val2));
|
||||
HIP_CHECK(hipFree(float_res));
|
||||
HIP_CHECK(hipFree(bf16_res));
|
||||
|
||||
for (size_t i = 0; i < size; i++) {
|
||||
INFO("FP res: " << f_res[i] << " - BF16 res: " << bf_res[i]);
|
||||
// The relative error should be less than 1/(2^7) since bfloat16 has 7 bits mantissa.
|
||||
REQUIRE((std::fabs(bf_res[i] - f_res[i]) / f_res[i]) < (1.0 / 128.0f));
|
||||
}
|
||||
}
|
||||
|
||||
SECTION("Equal bfloat16") {
|
||||
constexpr size_t size = 5;
|
||||
float in[size] = {1.0f, 0.5f, -0.33333f, 0.0f, -0.0f}, *d_in;
|
||||
unsigned* d_res;
|
||||
HIP_CHECK(hipMalloc(&d_in, sizeof(float) * size));
|
||||
HIP_CHECK(hipMalloc(&d_res, sizeof(unsigned) * size));
|
||||
|
||||
HIP_CHECK(hipMemcpy(d_in, in, sizeof(float) * size, hipMemcpyHostToDevice));
|
||||
|
||||
bf16_is_equal<<<1, size>>>(d_in, d_res, size);
|
||||
|
||||
std::vector<unsigned> res(size, 0);
|
||||
HIP_CHECK(hipMemcpy(res.data(), d_res, sizeof(unsigned) * size, hipMemcpyDeviceToHost));
|
||||
|
||||
REQUIRE(std::all_of(res.begin(), res.end(), [](unsigned n) { return n == 1; }));
|
||||
}
|
||||
|
||||
SECTION("MathOp Compare") {
|
||||
constexpr size_t size = 7;
|
||||
float in[size] = {0.5f, 1.0f, 1.5f, 0.33333f, 2.5f, 3.0f, 3.5f}, *d_in;
|
||||
unsigned* d_res;
|
||||
HIP_CHECK(hipMalloc(&d_in, sizeof(float) * size));
|
||||
HIP_CHECK(hipMalloc(&d_res, sizeof(unsigned) * size));
|
||||
|
||||
HIP_CHECK(hipMemcpy(d_in, in, sizeof(float) * size, hipMemcpyHostToDevice));
|
||||
|
||||
bf16_compare<<<1, size>>>(d_in, d_res, size);
|
||||
|
||||
std::vector<unsigned> res(size, 0);
|
||||
HIP_CHECK(hipMemcpy(res.data(), d_res, sizeof(unsigned) * size, hipMemcpyDeviceToHost));
|
||||
|
||||
REQUIRE(std::all_of(res.begin(), res.end(), [](unsigned n) { return n == 1; }));
|
||||
|
||||
HIP_CHECK(hipFree(d_in));
|
||||
HIP_CHECK(hipFree(d_res));
|
||||
}
|
||||
|
||||
SECTION("Bits equal") {
|
||||
constexpr size_t size = 5;
|
||||
float in[size] = {1.0f, 0.5f, 0.33333f, 3.38e38f, 3.40e38f}, *d_in;
|
||||
unsigned short* d_res;
|
||||
HIP_CHECK(hipMalloc(&d_in, sizeof(float) * size));
|
||||
HIP_CHECK(hipMalloc(&d_res, sizeof(unsigned short) * size));
|
||||
|
||||
HIP_CHECK(hipMemcpy(d_in, in, sizeof(float) * size, hipMemcpyHostToDevice));
|
||||
|
||||
bf16_conv_bits<<<1, size>>>(d_in, d_res, size);
|
||||
|
||||
std::vector<unsigned short> res_cmp = {0x3f80, 0x3f00, 0x3eab, 0x7f7e, 0x7f80 /*Inf*/};
|
||||
std::vector<unsigned short> res(size, 0);
|
||||
HIP_CHECK(hipMemcpy(res.data(), d_res, sizeof(unsigned short) * size, hipMemcpyDeviceToHost));
|
||||
|
||||
HIP_CHECK(hipFree(d_in));
|
||||
HIP_CHECK(hipFree(d_res));
|
||||
REQUIRE(res == res_cmp);
|
||||
}
|
||||
|
||||
SECTION("Round trip equal") {
|
||||
constexpr size_t size = 7;
|
||||
float *d_in, *d_out;
|
||||
std::vector<float> in = {
|
||||
std::numeric_limits<float>::infinity(), -1.0f, -0.5f, -0.0f, 0.0f, 0.5f, 1.0f};
|
||||
HIP_CHECK(hipMalloc(&d_in, sizeof(float) * size));
|
||||
HIP_CHECK(hipMalloc(&d_out, sizeof(float) * size));
|
||||
|
||||
HIP_CHECK(hipMemcpy(d_in, in.data(), sizeof(float) * size, hipMemcpyHostToDevice));
|
||||
|
||||
fp32_bf16_fp32<<<1, size>>>(d_in, d_out, size);
|
||||
|
||||
std::vector<float> res(size, 0.0f);
|
||||
|
||||
HIP_CHECK(hipMemcpy(res.data(), d_out, sizeof(unsigned) * size, hipMemcpyDeviceToHost));
|
||||
|
||||
HIP_CHECK(hipFree(d_in));
|
||||
HIP_CHECK(hipFree(d_out));
|
||||
REQUIRE(in == res);
|
||||
}
|
||||
|
||||
SECTION("Round trip subsection") {
|
||||
float *in, *out;
|
||||
HIP_CHECK(hipMalloc(&in, sizeof(float) * max_bf16_num));
|
||||
HIP_CHECK(hipMalloc(&out, sizeof(float) * max_bf16_num));
|
||||
HIP_CHECK(hipMemcpy(in, f_in.data(), sizeof(float) * max_bf16_num, hipMemcpyHostToDevice));
|
||||
|
||||
fp32_bf16_fp32<<<(max_bf16_num / 256) + 1, 256>>>(in, out, max_bf16_num); // round-trip
|
||||
|
||||
std::vector<float> f_out(f_in.size(), 0.0f);
|
||||
HIP_CHECK(hipMemcpy(f_out.data(), out, sizeof(float) * max_bf16_num, hipMemcpyDeviceToHost));
|
||||
HIP_CHECK(hipFree(in));
|
||||
HIP_CHECK(hipFree(out));
|
||||
|
||||
REQUIRE(f_in.size() == f_out.size()); // Size should be equal
|
||||
for (size_t i = 0; i < f_in.size(); i++) {
|
||||
INFO("Initial: " << f_in[i] << " After Conversion: " << f_out[i]);
|
||||
if (std::isnan(f_in[i])) { // NaNs can't be compared
|
||||
REQUIRE(std::isnan(f_out[i]));
|
||||
} else {
|
||||
REQUIRE(f_in[i] == f_out[i]);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
SECTION("Conversion to short") {
|
||||
float* in;
|
||||
HIP_CHECK(hipMalloc(&in, sizeof(float) * max_bf16_num));
|
||||
short* s_res;
|
||||
HIP_CHECK(hipMalloc(&s_res, sizeof(short) * max_bf16_num));
|
||||
unsigned short* u_res;
|
||||
HIP_CHECK(hipMalloc(&u_res, sizeof(unsigned short) * max_bf16_num));
|
||||
|
||||
HIP_CHECK(hipMemcpy(in, f_in.data(), sizeof(float) * max_bf16_num, hipMemcpyHostToDevice));
|
||||
HIP_CHECK(hipMemset(s_res, 0, sizeof(short) * max_bf16_num));
|
||||
HIP_CHECK(hipMemset(u_res, 0, sizeof(unsigned short) * max_bf16_num));
|
||||
|
||||
bf16_to_short<<<(max_bf16_num / 256) + 1, 256>>>(in, s_res, u_res, max_bf16_num);
|
||||
float* s_out;
|
||||
HIP_CHECK(hipMalloc(&s_out, sizeof(float) * max_bf16_num));
|
||||
float* u_out;
|
||||
HIP_CHECK(hipMalloc(&u_out, sizeof(float) * max_bf16_num));
|
||||
|
||||
short_to_bf16<<<(max_bf16_num / 256) + 1, 256>>>(s_res, s_out, max_bf16_num);
|
||||
ushort_to_bf16<<<(max_bf16_num / 256) + 1, 256>>>(u_res, u_out, max_bf16_num);
|
||||
|
||||
std::vector<float> f_res_s(max_bf16_num, 0.0f);
|
||||
std::vector<float> f_res_u(max_bf16_num, 0.0f);
|
||||
|
||||
HIP_CHECK(
|
||||
hipMemcpy(f_res_s.data(), s_out, sizeof(float) * max_bf16_num, hipMemcpyDeviceToHost));
|
||||
HIP_CHECK(
|
||||
hipMemcpy(f_res_u.data(), u_out, sizeof(float) * max_bf16_num, hipMemcpyDeviceToHost));
|
||||
|
||||
for (size_t i = 0; i < f_in.size(); i++) {
|
||||
if (std::isnan(f_res_s[i])) { // NaNs can't be compared
|
||||
REQUIRE(std::isnan(f_res_u[i]));
|
||||
} else {
|
||||
REQUIRE(f_res_s[i] == f_res_u[i]);
|
||||
}
|
||||
}
|
||||
|
||||
HIP_CHECK(hipFree(in));
|
||||
HIP_CHECK(hipFree(s_res));
|
||||
HIP_CHECK(hipFree(u_res));
|
||||
HIP_CHECK(hipFree(s_out));
|
||||
HIP_CHECK(hipFree(u_out));
|
||||
}
|
||||
|
||||
SECTION("Neg Subsection") {
|
||||
float *in, *out;
|
||||
HIP_CHECK(hipMalloc(&in, sizeof(float) * max_bf16_num));
|
||||
HIP_CHECK(hipMalloc(&out, sizeof(float) * max_bf16_num));
|
||||
HIP_CHECK(hipMemcpy(in, f_in.data(), sizeof(float) * max_bf16_num, hipMemcpyHostToDevice));
|
||||
|
||||
bf16_neg<<<(max_bf16_num / 256) + 1, 256>>>(in, out, max_bf16_num); // round-trip
|
||||
|
||||
std::vector<float> f_out(f_in.size(), 0.0f);
|
||||
HIP_CHECK(hipMemcpy(f_out.data(), out, sizeof(float) * max_bf16_num, hipMemcpyDeviceToHost));
|
||||
HIP_CHECK(hipFree(in));
|
||||
HIP_CHECK(hipFree(out));
|
||||
|
||||
REQUIRE(f_in.size() == f_out.size()); // Size should be equal
|
||||
for (size_t i = 0; i < f_in.size(); i++) {
|
||||
INFO("Initial: " << f_in[i] << " After Conversion: " << f_out[i]);
|
||||
if (std::isnan(f_in[i])) { // NaNs can't be compared
|
||||
REQUIRE(std::isnan(f_out[i]));
|
||||
} else {
|
||||
REQUIRE(f_in[i] == -f_out[i]);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
__global__ void bf162_eq(float* in, char* out, size_t size) {
|
||||
auto i = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (i < size) {
|
||||
auto temp = __bfloat162bfloat162(__float2bfloat16(in[i]));
|
||||
out[i] = __hbequ2(temp, temp) ? 1 : 0;
|
||||
}
|
||||
}
|
||||
|
||||
__global__ void bf162_neq(float* in, char* out, size_t size) {
|
||||
auto i = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (i < size) {
|
||||
auto val = __float2bfloat16(in[i]);
|
||||
auto other_val =
|
||||
__heq(__float2bfloat16(1.0f), val) ? __float2bfloat16(2.0f) : __float2bfloat16(1.0f);
|
||||
auto temp1 = __halves2bfloat162(val, other_val);
|
||||
auto temp2 = __halves2bfloat162(other_val, val);
|
||||
out[i] = (__hbneu2(temp1, temp2)) ? 1 : 0;
|
||||
}
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_bf162_basic") {
|
||||
auto f_in = getAllBF16();
|
||||
auto max_bf16_num = f_in.size();
|
||||
|
||||
SECTION("Eq Operation") {
|
||||
float* in;
|
||||
HIP_CHECK(hipMalloc(&in, sizeof(float) * max_bf16_num));
|
||||
HIP_CHECK(hipMemcpy(in, f_in.data(), sizeof(float) * max_bf16_num, hipMemcpyHostToDevice));
|
||||
char* out;
|
||||
HIP_CHECK(hipMalloc(&out, sizeof(char) * max_bf16_num));
|
||||
bf162_eq<<<(max_bf16_num / 256) + 1, 256>>>(in, out, max_bf16_num);
|
||||
std::vector<char> result(max_bf16_num, 0);
|
||||
HIP_CHECK(hipMemcpy(result.data(), out, sizeof(char) * max_bf16_num, hipMemcpyDeviceToHost));
|
||||
// Cant use allof, incase of mismatch we need to show which value had a mismatch
|
||||
for (size_t i = 0; i < max_bf16_num; i++) {
|
||||
INFO("Comparing: " << f_in[i] << " for iter: " << i);
|
||||
REQUIRE(result[i] == 1);
|
||||
}
|
||||
HIP_CHECK(hipFree(in));
|
||||
HIP_CHECK(hipFree(out));
|
||||
}
|
||||
|
||||
SECTION("Neq Operation") {
|
||||
float* in;
|
||||
HIP_CHECK(hipMalloc(&in, sizeof(float) * max_bf16_num));
|
||||
HIP_CHECK(hipMemcpy(in, f_in.data(), sizeof(float) * max_bf16_num, hipMemcpyHostToDevice));
|
||||
char* out;
|
||||
HIP_CHECK(hipMalloc(&out, sizeof(char) * max_bf16_num));
|
||||
bf162_neq<<<(max_bf16_num / 256) + 1, 256>>>(in, out, max_bf16_num);
|
||||
std::vector<char> result(max_bf16_num, 0);
|
||||
HIP_CHECK(hipMemcpy(result.data(), out, sizeof(char) * max_bf16_num, hipMemcpyDeviceToHost));
|
||||
// Cant use allof, incase of mismatch we need to show which value had a mismatch
|
||||
for (size_t i = 0; i < max_bf16_num; i++) {
|
||||
INFO("Comparing: " << f_in[i] << " for iter: " << i);
|
||||
REQUIRE(result[i] == 1);
|
||||
}
|
||||
HIP_CHECK(hipFree(in));
|
||||
HIP_CHECK(hipFree(out));
|
||||
}
|
||||
}
|
||||
@@ -8,6 +8,7 @@ set(TEST_SRC
|
||||
# AMD only tests
|
||||
set(AMD_TEST_SRC
|
||||
customOptions.cc
|
||||
linker.cc
|
||||
)
|
||||
|
||||
if(HIP_PLATFORM MATCHES "nvidia")
|
||||
|
||||
@@ -0,0 +1,128 @@
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
#include <hip/hiprtc.h>
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
|
||||
#include <cassert>
|
||||
#include <cstddef>
|
||||
#include <memory>
|
||||
#include <iostream>
|
||||
#include <iterator>
|
||||
#include <vector>
|
||||
|
||||
static constexpr auto NUM_THREADS{128};
|
||||
static constexpr auto NUM_BLOCKS{32};
|
||||
|
||||
static constexpr auto src{
|
||||
R"(
|
||||
extern "C"
|
||||
__global__
|
||||
void saxpy(float a, float* x, float* y, float* out, size_t n)
|
||||
{
|
||||
size_t tid = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (tid < n) {
|
||||
out[tid] = a * x[tid] + y[tid];
|
||||
}
|
||||
}
|
||||
)"};
|
||||
|
||||
TEST_CASE("Unit_RTC_LinkerAPI_Negative") {
|
||||
SECTION("get bitcode - nullptr size and code") {
|
||||
hiprtcProgram program;
|
||||
REQUIRE(hiprtcGetBitcodeSize(program, nullptr) == HIPRTC_ERROR_INVALID_INPUT);
|
||||
REQUIRE(hiprtcGetBitcode(program, nullptr) == HIPRTC_ERROR_INVALID_INPUT);
|
||||
}
|
||||
|
||||
SECTION("link create - nullptr image and input type") {
|
||||
hiprtcLinkState linkstate;
|
||||
REQUIRE(hiprtcLinkAddData(linkstate, HIPRTC_JIT_INPUT_LLVM_BITCODE, nullptr, 0, "code", 0,
|
||||
nullptr, nullptr) == HIPRTC_ERROR_INVALID_INPUT);
|
||||
REQUIRE(hiprtcLinkAddData(linkstate, HIPRTC_JIT_INPUT_CUBIN, &linkstate, 1, "code", 0, nullptr,
|
||||
nullptr) == HIPRTC_ERROR_INVALID_INPUT);
|
||||
}
|
||||
|
||||
SECTION("link complete - ") {
|
||||
hiprtcLinkState linkstate;
|
||||
REQUIRE(hiprtcLinkComplete(linkstate, nullptr, nullptr) == HIPRTC_ERROR_INVALID_INPUT);
|
||||
}
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_RTC_LinkerAPI") {
|
||||
hiprtcProgram program;
|
||||
HIPRTC_CHECK(hiprtcCreateProgram(&program, src, "saxpy", 0, nullptr, nullptr));
|
||||
|
||||
const char* options[]{"-fgpu-rdc"};
|
||||
HIPRTC_CHECK(hiprtcCompileProgram(program, 1, options));
|
||||
|
||||
size_t codesize = 0;
|
||||
HIPRTC_CHECK(hiprtcGetBitcodeSize(program, &codesize));
|
||||
|
||||
std::vector<char> code(codesize, '\0');
|
||||
HIPRTC_CHECK(hiprtcGetBitcode(program, &code[0]));
|
||||
|
||||
const char* isaopts[] = {"-mllvm", "-inline-threshold=1", "-mllvm", "-inlinehint-threshold=1"};
|
||||
std::vector<hiprtcJIT_option> jit_options = {HIPRTC_JIT_IR_TO_ISA_OPT_EXT,
|
||||
HIPRTC_JIT_IR_TO_ISA_OPT_COUNT_EXT};
|
||||
size_t isaoptssize = 4;
|
||||
const void* lopts[] = {(void*)isaopts, (void*)(isaoptssize)};
|
||||
hiprtcLinkState linkstate;
|
||||
HIPRTC_CHECK(hiprtcLinkCreate(jit_options.size(), jit_options.data(), (void**)lopts, &linkstate));
|
||||
HIPRTC_CHECK(hiprtcLinkAddData(linkstate, HIPRTC_JIT_INPUT_LLVM_BITCODE, code.data(), code.size(),
|
||||
"LinkISA", 0, nullptr, nullptr));
|
||||
|
||||
void* finaldata;
|
||||
size_t finalsize = 0;
|
||||
HIPRTC_CHECK(hiprtcLinkComplete(linkstate, &finaldata, &finalsize));
|
||||
|
||||
size_t n = 128 * 32;
|
||||
size_t bufferSize = n * sizeof(float);
|
||||
|
||||
float *dX, *dY, *dOut;
|
||||
HIP_CHECK(hipMalloc(&dX, bufferSize));
|
||||
HIP_CHECK(hipMalloc(&dY, bufferSize));
|
||||
HIP_CHECK(hipMalloc(&dOut, bufferSize));
|
||||
|
||||
hipModule_t module;
|
||||
hipFunction_t kernel;
|
||||
HIP_CHECK(hipModuleLoadData(&module, finaldata));
|
||||
HIP_CHECK(hipModuleGetFunction(&kernel, module, "saxpy"));
|
||||
|
||||
float a = 5.1f;
|
||||
std::unique_ptr<float[]> hX{new float[n]};
|
||||
std::unique_ptr<float[]> hY{new float[n]};
|
||||
std::unique_ptr<float[]> hOut{new float[n]};
|
||||
for (size_t i = 0; i < n; ++i) {
|
||||
hX[i] = static_cast<float>(i);
|
||||
hY[i] = static_cast<float>(i * 2);
|
||||
}
|
||||
|
||||
HIP_CHECK(hipMemcpy(dX, hX.get(), bufferSize, hipMemcpyHostToDevice));
|
||||
HIP_CHECK(hipMemcpy(dY, hY.get(), bufferSize, hipMemcpyHostToDevice));
|
||||
|
||||
struct {
|
||||
float a_;
|
||||
float* b_;
|
||||
float* c_;
|
||||
float* d_;
|
||||
size_t e_;
|
||||
} args{a, dX, dY, dOut, n};
|
||||
|
||||
auto size = sizeof(args);
|
||||
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
|
||||
HIP_LAUNCH_PARAM_END};
|
||||
|
||||
HIP_CHECK(hipModuleLaunchKernel(kernel, 32, 1, 1, 128, 1, 1, 0, nullptr, nullptr, config));
|
||||
|
||||
HIP_CHECK(hipMemcpy(hOut.get(), dOut, bufferSize, hipMemcpyDeviceToHost));
|
||||
|
||||
HIP_CHECK(hipFree(dX));
|
||||
HIP_CHECK(hipFree(dY));
|
||||
HIP_CHECK(hipFree(dOut));
|
||||
|
||||
HIP_CHECK(hipModuleUnload(module));
|
||||
|
||||
for (size_t i = 0; i < n; ++i) {
|
||||
REQUIRE(fabs(a * hX[i] + hY[i] - hOut[i]) <= fabs(hOut[i]) * 1e-6);
|
||||
}
|
||||
}
|
||||
@@ -32,49 +32,65 @@ texture<int4, hipTextureType2D, hipReadModeElementType> texInt4;
|
||||
texture<float4, hipTextureType2D, hipReadModeElementType> texFloat4;
|
||||
|
||||
extern "C" __global__ void tex2dKernelChar(char* outputData, int width, int height) {
|
||||
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
outputData[y * width + x] = tex2D(texChar, x, y);
|
||||
#endif
|
||||
}
|
||||
|
||||
extern "C" __global__ void tex2dKernelShort(short* outputData, int width, int height) {
|
||||
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
outputData[y * width + x] = tex2D(texShort, x, y);
|
||||
#endif
|
||||
}
|
||||
|
||||
extern "C" __global__ void tex2dKernelInt(int* outputData, int width, int height) {
|
||||
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
outputData[y * width + x] = tex2D(texInt, x, y);
|
||||
#endif
|
||||
}
|
||||
|
||||
extern "C" __global__ void tex2dKernelFloat(float* outputData, int width, int height) {
|
||||
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
outputData[y * width + x] = tex2D(texFloat, x, y);
|
||||
#endif
|
||||
}
|
||||
|
||||
extern "C" __global__ void tex2dKernelChar4(char4* outputData, int width, int height) {
|
||||
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
outputData[y * width + x] = tex2D(texChar4, x, y);
|
||||
#endif
|
||||
}
|
||||
|
||||
extern "C" __global__ void tex2dKernelShort4(short4* outputData, int width, int height) {
|
||||
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
outputData[y * width + x] = tex2D(texShort4, x, y);
|
||||
#endif
|
||||
}
|
||||
|
||||
extern "C" __global__ void tex2dKernelInt4(int4* outputData, int width, int height) {
|
||||
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
outputData[y * width + x] = tex2D(texInt4, x, y);
|
||||
#endif
|
||||
}
|
||||
|
||||
extern "C" __global__ void tex2dKernelFloat4(float4* outputData, int width, int height) {
|
||||
#if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
|
||||
int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
outputData[y * width + x] = tex2D(texFloat4, x, y);
|
||||
#endif
|
||||
}
|
||||
|
||||
Ссылка в новой задаче
Block a user