diff --git a/catch/hipTestMain/config/config_amd_linux_common.json b/catch/hipTestMain/config/config_amd_linux_common.json index f9799a723c..d54eb6b9e0 100644 --- a/catch/hipTestMain/config/config_amd_linux_common.json +++ b/catch/hipTestMain/config/config_amd_linux_common.json @@ -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", diff --git a/catch/unit/deviceLib/CMakeLists.txt b/catch/unit/deviceLib/CMakeLists.txt index 006b0767ed..60fafe7d99 100644 --- a/catch/unit/deviceLib/CMakeLists.txt +++ b/catch/unit/deviceLib/CMakeLists.txt @@ -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 diff --git a/catch/unit/deviceLib/bfloat16.cc b/catch/unit/deviceLib/bfloat16.cc new file mode 100644 index 0000000000..677b5b09e2 --- /dev/null +++ b/catch/unit/deviceLib/bfloat16.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 +#include + +#include +#include +#include +#include + +// 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 getAllBF16() { + constexpr unsigned char max_mantissa = std::numeric_limits::max() >> 1; // 7 bits + const size_t max_bf16_num = + 2 /*sign*/ * std::pow(2, 8) /*exponent*/ * std::pow(2, 7) /*mantissa*/; + + std::vector 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::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(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(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(&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(size); + auto h_c = std::make_unique(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(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 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 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 res_cmp = {0x3f80, 0x3f00, 0x3eab, 0x7f7e, 0x7f80 /*Inf*/}; + std::vector 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 in = { + std::numeric_limits::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 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 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 f_res_s(max_bf16_num, 0.0f); + std::vector 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 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 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 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)); + } +} diff --git a/catch/unit/rtc/CMakeLists.txt b/catch/unit/rtc/CMakeLists.txt index 534289455b..7d4b4893ac 100644 --- a/catch/unit/rtc/CMakeLists.txt +++ b/catch/unit/rtc/CMakeLists.txt @@ -8,6 +8,7 @@ set(TEST_SRC # AMD only tests set(AMD_TEST_SRC customOptions.cc + linker.cc ) if(HIP_PLATFORM MATCHES "nvidia") diff --git a/catch/unit/rtc/linker.cc b/catch/unit/rtc/linker.cc new file mode 100644 index 0000000000..8733487ccf --- /dev/null +++ b/catch/unit/rtc/linker.cc @@ -0,0 +1,128 @@ +#include + +#include +#include + + +#include +#include +#include +#include +#include +#include + +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 code(codesize, '\0'); + HIPRTC_CHECK(hiprtcGetBitcode(program, &code[0])); + + const char* isaopts[] = {"-mllvm", "-inline-threshold=1", "-mllvm", "-inlinehint-threshold=1"}; + std::vector 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 hX{new float[n]}; + std::unique_ptr hY{new float[n]}; + std::unique_ptr hOut{new float[n]}; + for (size_t i = 0; i < n; ++i) { + hX[i] = static_cast(i); + hY[i] = static_cast(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); + } +} \ No newline at end of file diff --git a/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp b/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp index a1d3985de5..d20b1d0864 100644 --- a/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp +++ b/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp @@ -32,49 +32,65 @@ texture texInt4; texture 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 }