From f57047a49b0fc6734ae3caa3fdcdc04790e930d6 Mon Sep 17 00:00:00 2001 From: yukun-hh Date: Sat, 26 Sep 2026 22:21:54 +0800 Subject: [PATCH] =?UTF-8?q?=E5=90=91=E9=87=8F=E5=8A=A0=E6=B3=95?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- 1(vectoradd)/main1.cu | 13 +++++++++++++ 1(vectoradd)/main2.cu | 27 +++++++++++++++++++++++++++ 1(vectoradd)/main3.cu | 34 ++++++++++++++++++++++++++++++++++ 3 files changed, 74 insertions(+) create mode 100644 1(vectoradd)/main1.cu create mode 100644 1(vectoradd)/main2.cu create mode 100644 1(vectoradd)/main3.cu diff --git a/1(vectoradd)/main1.cu b/1(vectoradd)/main1.cu new file mode 100644 index 0000000..70ed5cb --- /dev/null +++ b/1(vectoradd)/main1.cu @@ -0,0 +1,13 @@ +#include +#include +__global__ void vector_add(__half* A, const __half* B, int N) { + int idx=blockDim.x*blockIdx.x+threadIdx.x; + if(idx>>(A, B, numel); + cudaDeviceSynchronize(); +} diff --git a/1(vectoradd)/main2.cu b/1(vectoradd)/main2.cu new file mode 100644 index 0000000..7157da6 --- /dev/null +++ b/1(vectoradd)/main2.cu @@ -0,0 +1,27 @@ +#include +#include +__global__ void vector_add(__half* A, const __half* B, int numel) { + int idx=blockDim.x*blockIdx.x+threadIdx.x; + __half2* A2 = reinterpret_cast< __half2* >(A); + const __half2* B2 = reinterpret_cast(B); + int base = blockIdx.x * blockDim.x * 4 + threadIdx.x; + int i0 = base; + int i1 = base + blockDim.x; + int i2 = base + blockDim.x * 2; + int i3 = base + blockDim.x * 3; + int n2 = numel/2; + if(i0 < n2) A2[i0] = __hadd2(A2[i0],B2[i0]); + if(i1 < n2) A2[i1] = __hadd2(A2[i1],B2[i1]); + if(i2 < n2) A2[i2] = __hadd2(A2[i2],B2[i2]); + if(i3 < n2) A2[i3] = __hadd2(A2[i3],B2[i3]); + if(idx == 0 && (numel&1)){ + A[numel-1]= __hadd(A[numel-1],B[numel-1]); + } +} +extern "C" void run_kernel(__half* A, const __half* B, int64_t numel){ + int threads = 1024; + int numel2=numel/2; + int works=(numel2+3)/4; + int blocks = (works + threads-1) / threads; + vector_add<<>>(A,B,numel); +} diff --git a/1(vectoradd)/main3.cu b/1(vectoradd)/main3.cu new file mode 100644 index 0000000..d6c3425 --- /dev/null +++ b/1(vectoradd)/main3.cu @@ -0,0 +1,34 @@ +#include +#include +struct __align__(16) half4 { + __half2 a, b, c, d; +}; +__global__ void __launch_bounds__(256, 2) fp16_add_kernel(__half* __restrict__ A, + const __half* __restrict__ B, + int64_t numel) { + int64_t idx = ((int64_t)blockIdx.x * blockDim.x + threadIdx.x) << 3; + + if (idx + 7 < numel) { + half4 a4 = *reinterpret_cast(A + idx); + half4 b4 = *reinterpret_cast(B + idx); + + a4.a = __hadd2(a4.a, b4.a); + a4.b = __hadd2(a4.b, b4.b); + a4.c = __hadd2(a4.c, b4.c); + a4.d = __hadd2(a4.d, b4.d); + + *reinterpret_cast(A + idx) = a4; + } else if (idx < numel) { + #pragma unroll + for (int i = 0; i < 8 && idx + i < numel; ++i) + A[idx + i] = __hadd(A[idx + i], B[idx + i]); + } +} + +extern "C" void run_kernel(__half* A, const __half* B, int64_t numel) { + + const int threads = 256; + int64_t blocks = (numel + ((int64_t)threads << 3) - 1) >> 11; + + fp16_add_kernel<<>>(A, B, numel); +}