Merged
update simt float math api doc && clarify SIMT VF array parameter error #2743
update simt float math api doc && clarify SIMT VF array parameter error #2743
Merged
wulinyuAdd on May 29
Total: 95 files changed+2356-785
@@ -55,4 +55,5 @@
55| ASCRT_REMQUO_BITS_F | 在remquof函数中提取商的最低3位有效位,即数值为3。 | math_constants.h |55| ASCRT_REMQUO_BITS_F | 在remquof函数中提取商的最低3位有效位,即数值为3。 | math_constants.h |
56| ASCRT_REMQUO_MASK_F | 用于提取remquof结果中低3位比特的位掩码。 | math_constants.h |56| ASCRT_REMQUO_MASK_F | 用于提取remquof结果中低3位比特的位掩码。 | math_constants.h |
57| ASCRT_TRIG_PLOSS_F | 105615.0F,表示三角函数计算中精度损失的临界阈值。 | math_constants.h |57| ASCRT_TRIG_PLOSS_F | 105615.0F,表示三角函数计算中精度损失的临界阈值。 | math_constants.h |
58- 58+| ASCRT_MIN_VAL_S | 定义int32类型的最小值。 | math_constants.h |
59+| ASCRT_MAX_VAL_S | 定义int32类型的最大值。 | math_constants.h |
@@ -55,4 +55,5 @@
55| ASCRT_REMQUO_BITS_F | 在remquof函数中提取商的最低3位有效位,即数值为3。 | math_constants.h |55| ASCRT_REMQUO_BITS_F | 在remquof函数中提取商的最低3位有效位,即数值为3。 | math_constants.h |
56| ASCRT_REMQUO_MASK_F | 用于提取remquof结果中低3位比特的位掩码。 | math_constants.h |56| ASCRT_REMQUO_MASK_F | 用于提取remquof结果中低3位比特的位掩码。 | math_constants.h |
57| ASCRT_TRIG_PLOSS_F | 105615.0F,表示三角函数计算中精度损失的临界阈值。 | math_constants.h |57| ASCRT_TRIG_PLOSS_F | 105615.0F,表示三角函数计算中精度损失的临界阈值。 | math_constants.h |
58- 58+| ASCRT_MIN_VAL_S | 定义int32类型的最小值。 | math_constants.h |
59+| ASCRT_MAX_VAL_S | 定义int32类型的最大值。 | math_constants.h |
@@ -54,23 +54,41 @@ float __fdividef(float x, float y)
54 54 
55## 调用示例55## 调用示例
56 56 
57-- SIMT编程场景:57+- SIMT编程场景:
58 58 
59 ```59 ```
60- __global__ __launch_bounds__(1024) void KernelFdividef(float* dst, float* x, float* y)60+ __global__ __launch_bounds__(256) void compute_fdivide_intrinsic(float *result, const float *x, const float *y, uint32_t count)
61 {61 {
62- int idx = threadIdx.x + blockIdx.x * blockDim.x;62+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
63- dst[idx] = __fdividef(x[idx], y[idx]);63+ if (idx >= count) {
64+ return;
65+ }
66+ result[idx] = __fdividef(x[idx], y[idx]);
64 }67 }
65 ```68 ```
66 69 
67-- SIMD与SIMT混合编程场景:70+- SIMD与SIMT混合编程场景:
68 71 
69 ```72 ```
70- __simt_vf__ __launch_bounds__(1024) inline void KernelFdividef(__gm__ float* dst, __gm__ float* x, __gm__ float* y)73+ __simt_vf__ __launch_bounds__(256) inline void compute_fdivide_intrinsic_vf(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, uint32_t count)
71 {74 {
72- int idx = threadIdx.x + blockIdx.x * blockDim.x;75+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
73- dst[idx] = __fdividef(x[idx], y[idx]);76+ if (idx >= count) {
77+ return;
78+ }
79+ result[idx] = __fdividef(x[idx], y[idx]);
80+ }
81+ 
82+ __global__ __vector__ void run_fdivide_intrinsic(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, uint32_t count)
83+ {
84+ asc_vf_call<compute_fdivide_intrinsic_vf>(dim3(256), result, x, y, count);
74 }85 }
75 ```86 ```
76 87 
88+输入输出示例如下:
89+ 
90+```
A
Aai_xinMay 30

缩进不对 下同

likedislike
91+x:0.25, 0.75, 1.25, 1.75
92+y:1.5, 2.5, 3.5, 4.5
93+result: 0.1666667 0.3 0.3571429 0.3888889
94+```
@@ -53,23 +53,40 @@ float __saturatef(float x)
53 53 
54## 调用示例54## 调用示例
55 55 
56-- SIMT编程场景:56+- SIMT编程场景:
57 57 
58 ```58 ```
59- __global__ __launch_bounds__(1024) void KernelSaturatef(float* dst, float* x)59+ __global__ __launch_bounds__(256) void compute_saturate_intrinsic(float *result, const float *x, uint32_t count)
60 {60 {
61- int idx = threadIdx.x + blockIdx.x * blockDim.x;61+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
62- dst[idx] = __saturatef(x[idx]);62+ if (idx >= count) {
63+ return;
64+ }
65+ result[idx] = __saturatef(x[idx]);
63 }66 }
64 ```67 ```
65 68 
66-- SIMD与SIMT混合编程场景:69+- SIMD与SIMT混合编程场景:
67 70 
68 ```71 ```
69- __simt_vf__ __launch_bounds__(1024) inline void KernelSaturatef(__gm__ float* dst, __gm__ float* x)72+ __simt_vf__ __launch_bounds__(256) inline void compute_saturate_intrinsic_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
70 {73 {
71- int idx = threadIdx.x + blockIdx.x * blockDim.x;74+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
72- dst[idx] = __saturatef(x[idx]);75+ if (idx >= count) {
76+ return;
77+ }
78+ result[idx] = __saturatef(x[idx]);
79+ }
80+ 
81+ __global__ __vector__ void run_saturate_intrinsic(__gm__ float *result, __gm__ const float *x, uint32_t count)
82+ {
83+ asc_vf_call<compute_saturate_intrinsic_vf>(dim3(256), result, x, count);
73 }84 }
74 ```85 ```
75 86 
87+输入输出示例如下:
88+ 
89+```
90+x:0.25, 0.75, 1.25, 1.75
91+result: 0.25 0.75 1 1
92+```
@@ -58,23 +58,40 @@ inline float acosf(float x)
58 58 
59## 调用示例59## 调用示例
60 60 
61-- SIMT编程场景:61+- SIMT编程场景:
62 62 
63 ```63 ```
64- __global__ __launch_bounds__(1024) void KernelAcos(float* dst, float* x)64+ __global__ __launch_bounds__(256) void compute_acosf(float *result, const float *x, uint32_t count)
65 {65 {
66- int idx = threadIdx.x + blockIdx.x * blockDim.x;66+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
67- dst[idx] = acosf(x[idx]); // 对src源地址的第idx个元素取反余弦值67+ if (idx >= count) {
68+ return;
69+ }
70+ result[idx] = acosf(x[idx]);
68 }71 }
69 ```72 ```
70 73 
71-- SIMD与SIMT混合编程场景:74+- SIMD与SIMT混合编程场景:
72 75 
73 ```76 ```
74- __simt_vf__ __launch_bounds__(1024) inline void KernelAcos(__gm__ float* dst, __gm__ float* x)77+ __simt_vf__ __launch_bounds__(256) inline void compute_acosf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
75 {78 {
76- int idx = threadIdx.x + blockIdx.x * blockDim.x;79+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
77- dst[idx] = acosf(x[idx]); // 对src源地址的第idx个元素取反余弦值80+ if (idx >= count) {
81+ return;
82+ }
83+ result[idx] = acosf(x[idx]);
84+ }
85+ 
86+ __global__ __vector__ void run_acosf(__gm__ float *result, __gm__ const float *x, uint32_t count)
87+ {
88+ asc_vf_call<compute_acosf_vf>(dim3(256), result, x, count);
78 }89 }
79 ```90 ```
80 91 
92+输入输出示例如下:
93+ 
94+```
95+x:-0.75, -0.25, 0.25, 0.75
96+result: 2.418875 1.823477 1.318116 0.7227179
97+```
@@ -56,23 +56,40 @@ inline float acoshf(float x)
56 56 
57## 调用示例57## 调用示例
58 58 
59-- SIMT编程场景:59+- SIMT编程场景:
60 60 
61 ```61 ```
62- __global__ __launch_bounds__(1024) void KernelAcosh(float* dst, float* x)62+ __global__ __launch_bounds__(256) void compute_acoshf(float *result, const float *x, uint32_t count)
63 {63 {
64- int idx = threadIdx.x + blockIdx.x * blockDim.x;64+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
65- dst[idx] = acoshf(x[idx]); // 对src源地址的第idx个元素取双曲反余弦值65+ if (idx >= count) {
66+ return;
67+ }
68+ result[idx] = acoshf(x[idx]);
66 }69 }
67 ```70 ```
68 71 
69-- SIMD与SIMT混合编程场景:72+- SIMD与SIMT混合编程场景:
70 73 
71 ```74 ```
72- __simt_vf__ __launch_bounds__(1024) inline void KernelAcosh(__gm__ float* dst, __gm__ float* x)75+ __simt_vf__ __launch_bounds__(256) inline void compute_acoshf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
73 {76 {
74- int idx = threadIdx.x + blockIdx.x * blockDim.x;77+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
75- dst[idx] = acoshf(x[idx]); // 对src源地址的第idx个元素取双曲反余弦值78+ if (idx >= count) {
79+ return;
80+ }
81+ result[idx] = acoshf(x[idx]);
82+ }
83+ 
84+ __global__ __vector__ void run_acoshf(__gm__ float *result, __gm__ const float *x, uint32_t count)
85+ {
86+ asc_vf_call<compute_acoshf_vf>(dim3(256), result, x, count);
76 }87 }
77 ```88 ```
78 89 
90+输入输出示例如下:
91+ 
92+```
93+x:1, 2, 3, 4
94+result: 0 1.316958 1.762747 2.063437
95+```
@@ -58,23 +58,40 @@ inline float asinf(float x)
58 58 
59## 调用示例59## 调用示例
60 60 
61-- SIMT编程场景:61+- SIMT编程场景:
62 62 
63 ```63 ```
64- __global__ __launch_bounds__(1024) void KernelAsin(float* dst, float* x)64+ __global__ __launch_bounds__(256) void compute_asinf(float *result, const float *x, uint32_t count)
65 {65 {
66- int idx = threadIdx.x + blockIdx.x * blockDim.x;66+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
67- dst[idx] = asinf(x[idx]); // 对src源地址的第idx个元素取反正弦值67+ if (idx >= count) {
68+ return;
69+ }
70+ result[idx] = asinf(x[idx]);
68 }71 }
69 ```72 ```
70 73 
71-- SIMD与SIMT混合编程场景:74+- SIMD与SIMT混合编程场景:
72 75 
73 ```76 ```
74- __simt_vf__ __launch_bounds__(1024) inline void KernelAsin(__gm__ float* dst, __gm__ float* x)77+ __simt_vf__ __launch_bounds__(256) inline void compute_asinf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
75 {78 {
76- int idx = threadIdx.x + blockIdx.x * blockDim.x;79+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
77- dst[idx] = asinf(x[idx]); // 对src源地址的第idx个元素取反正弦值80+ if (idx >= count) {
81+ return;
82+ }
83+ result[idx] = asinf(x[idx]);
84+ }
85+ 
86+ __global__ __vector__ void run_asinf(__gm__ float *result, __gm__ const float *x, uint32_t count)
87+ {
88+ asc_vf_call<compute_asinf_vf>(dim3(256), result, x, count);
78 }89 }
79 ```90 ```
80 91 
92+输入输出示例如下:
93+ 
94+```
95+x:-0.75, -0.25, 0.25, 0.75
96+result: -0.8480784 -0.2526802 0.2526802 0.8480784
97+```
@@ -55,23 +55,40 @@ inline float asinhf(float x)
55 55 
56## 调用示例56## 调用示例
57 57 
58-- SIMT编程场景:58+- SIMT编程场景:
59 59 
60 ```60 ```
61- __global__ __launch_bounds__(1024) void KernelAsinh(float* dst, float* x)61+ __global__ __launch_bounds__(256) void compute_asinhf(float *result, const float *x, uint32_t count)
62 {62 {
63- int idx = threadIdx.x + blockIdx.x * blockDim.x;63+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
64- dst[idx] = asinhf(x[idx]); // 对src源地址的第idx个元素取双曲反正弦值64+ if (idx >= count) {
65+ return;
66+ }
67+ result[idx] = asinhf(x[idx]);
65 }68 }
66 ```69 ```
67 70 
68-- SIMD与SIMT混合编程场景:71+- SIMD与SIMT混合编程场景:
69 72 
70 ```73 ```
71- __simt_vf__ __launch_bounds__(1024) inline void KernelAsinh(__gm__ float* dst, __gm__ float* x)74+ __simt_vf__ __launch_bounds__(256) inline void compute_asinhf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
72 {75 {
73- int idx = threadIdx.x + blockIdx.x * blockDim.x;76+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
74- dst[idx] = asinhf(x[idx]); // 对src源地址的第idx个元素取双曲反正弦值77+ if (idx >= count) {
78+ return;
79+ }
80+ result[idx] = asinhf(x[idx]);
81+ }
82+ 
83+ __global__ __vector__ void run_asinhf(__gm__ float *result, __gm__ const float *x, uint32_t count)
84+ {
85+ asc_vf_call<compute_asinhf_vf>(dim3(256), result, x, count);
75 }86 }
76 ```87 ```
77 88 
89+输入输出示例如下:
90+ 
91+```
92+x:0.25, 0.75, 1.25, 1.75
93+result: 0.2474664 0.6931472 1.047593 1.325898
94+```
@@ -37,16 +37,24 @@ inline float atan2f(float y, float x)
37 37 
38y/x的反正切值。38y/x的反正切值。
39 39 
40-- 当x,y任意为nan时,返回值为nan。40+特殊取值场景下的返回值如下表所示。
41-- 当x=0,y=0时,返回值为nan。41+ 
42-- 当y=inf,x=inf时,返回值为π/4。42+| 输入值x | 输入值y | 返回值 |
43-- 当y=-inf,x=inf时,返回值为-π/4。43+| --- | --- | --- |
44-- 当y=1,x=inf时,返回值为0.0。44+| 任意值 | nan | nan |
45-- 当y=inf,x=-inf时,返回值为π3/4。45+| nan | 任意值 | nan |
46-- 当y=-inf,x=-inf时,返回值为-π3/4。46+| 正值(含+0) | +0 | +0 |
47-- 当y=1,x=-inf时,返回值为π。47+| 正值(含+0) | -0 | -0 |
48-- 当y=inf,x=1时,返回值为π/2。48+| 负值(含-0) | +0 | π |
49-- 当y=-inf,x=1时,返回值为-π/2。49+| 负值(含-0) | -0 | -π |
50+| inf | inf | π/4 |
51+| inf | -inf | -π/4 |
52+| inf | 1 | 0.0 |
53+| -inf | inf | 3π/4 |
54+| -inf | -inf | -3π/4 |
55+| -inf | 1 | π |
56+| 1 | inf | π/2 |
57+| 1 | -inf | -π/2 |
50 58 
51## 约束说明59## 约束说明
52 60 
@@ -62,23 +70,41 @@ y/x的反正切值。
62 70 
63## 调用示例71## 调用示例
64 72 
65-- SIMT编程场景:73+- SIMT编程场景:
66 74 
67 ```75 ```
68- __global__ __launch_bounds__(1024) void KernelAtan2(float* dst, float* y, float* x)76+ __global__ __launch_bounds__(256) void compute_atan2f(float *result, const float *x, const float *y, uint32_t count)
69 {77 {
70- int idx = threadIdx.x + blockIdx.x * blockDim.x;78+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
71- dst[idx] = atan2f(y[idx], x[idx]);79+ if (idx >= count) {
80+ return;
81+ }
82+ result[idx] = atan2f(y[idx], x[idx]);
72 }83 }
73 ```84 ```
74 85 
75-- SIMD与SIMT混合编程场景:86+- SIMD与SIMT混合编程场景:
76 87 
77 ```88 ```
78- __simt_vf__ __launch_bounds__(1024) inline void KernelAtan2(__gm__ float* dst, __gm__ float* y, __gm__ float* x)89+ __simt_vf__ __launch_bounds__(256) inline void compute_atan2f_vf(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, uint32_t count)
79 {90 {
80- int idx = threadIdx.x + blockIdx.x * blockDim.x;91+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
81- dst[idx] = atan2f(y[idx], x[idx]);92+ if (idx >= count) {
93+ return;
94+ }
95+ result[idx] = atan2f(y[idx], x[idx]);
96+ }
97+ 
98+ __global__ __vector__ void run_atan2f(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, uint32_t count)
99+ {
100+ asc_vf_call<compute_atan2f_vf>(dim3(256), result, x, y, count);
82 }101 }
83 ```102 ```
84 103 
104+输入输出示例如下:
105+ 
106+```
107+y:1.5, 2.5, 3.5, 4.5
108+x:0.25, 0.75, 1.25, 1.75
109+result: 1.405648 1.27934 1.227772 1.199905
110+```
@@ -54,23 +54,40 @@ inline float atanf(float x)
54 54 
55## 调用示例55## 调用示例
56 56 
57-- SIMT编程场景:57+- SIMT编程场景:
58 58 
59 ```59 ```
60- __global__ __launch_bounds__(1024) void KernelAtan(float* dst, float* x)60+ __global__ __launch_bounds__(256) void compute_atanf(float *result, const float *x, uint32_t count)
61 {61 {
62- int idx = threadIdx.x + blockIdx.x * blockDim.x;62+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
63- dst[idx] = atanf(x[idx]);63+ if (idx >= count) {
64+ return;
65+ }
66+ result[idx] = atanf(x[idx]);
64 }67 }
65 ```68 ```
66 69 
67-- SIMD与SIMT混合编程场景:70+- SIMD与SIMT混合编程场景:
68 71 
69 ```72 ```
70- __simt_vf__ __launch_bounds__(1024) inline void KernelAtan(__gm__ float* dst, __gm__ float* x)73+ __simt_vf__ __launch_bounds__(256) inline void compute_atanf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
71 {74 {
72- int idx = threadIdx.x + blockIdx.x * blockDim.x;75+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
73- dst[idx] = atanf(x[idx]);76+ if (idx >= count) {
77+ return;
78+ }
79+ result[idx] = atanf(x[idx]);
80+ }
81+ 
82+ __global__ __vector__ void run_atanf(__gm__ float *result, __gm__ const float *x, uint32_t count)
83+ {
84+ asc_vf_call<compute_atanf_vf>(dim3(256), result, x, count);
74 }85 }
75 ```86 ```
76 87 
88+输入输出示例如下:
89+ 
90+```
91+x:0.25, 0.75, 1.25, 1.75
92+result: 0.2449787 0.6435012 0.8960554 1.05165
93+```
@@ -54,23 +54,40 @@ inline float atanhf(float x)
54 54 
55## 调用示例55## 调用示例
56 56 
57-- SIMT编程场景:57+- SIMT编程场景:
58 58 
59 ```59 ```
60- __global__ __launch_bounds__(1024) void KernelAtanh(float* dst, float* x)60+ __global__ __launch_bounds__(256) void compute_atanhf(float *result, const float *x, uint32_t count)
61 {61 {
62- int idx = threadIdx.x + blockIdx.x * blockDim.x;62+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
63- dst[idx] = atanhf(x[idx]);63+ if (idx >= count) {
64+ return;
65+ }
66+ result[idx] = atanhf(x[idx]);
64 }67 }
65 ```68 ```
66 69 
67-- SIMD与SIMT混合编程场景:70+- SIMD与SIMT混合编程场景:
68 71 
69 ```72 ```
70- __simt_vf__ __launch_bounds__(1024) inline void KernelAtanh(__gm__ float* dst, __gm__ float* x)73+ __simt_vf__ __launch_bounds__(256) inline void compute_atanhf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
71 {74 {
72- int idx = threadIdx.x + blockIdx.x * blockDim.x;75+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
73- dst[idx] = atanhf(x[idx]);76+ if (idx >= count) {
77+ return;
78+ }
79+ result[idx] = atanhf(x[idx]);
80+ }
81+ 
82+ __global__ __vector__ void run_atanhf(__gm__ float *result, __gm__ const float *x, uint32_t count)
83+ {
84+ asc_vf_call<compute_atanhf_vf>(dim3(256), result, x, count);
74 }85 }
75 ```86 ```
76 87 
88+输入输出示例如下:
89+ 
90+```
91+x:-0.75, -0.25, 0.25, 0.75
92+result: -0.972955 -0.2554128 0.2554128 0.9729551
93+```
@@ -55,23 +55,40 @@ x的立方根。
55 55 
56## 调用示例56## 调用示例
57 57 
58-- SIMT编程场景:58+- SIMT编程场景:
59 59 
60 ```60 ```
61- __global__ __launch_bounds__(1024) void KernelCbrt(float* x, float* y)61+ __global__ __launch_bounds__(256) void compute_cbrtf(float *result, const float *x, uint32_t count)
62 {62 {
63- int idx = threadIdx.x + blockIdx.x * blockDim.x;63+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
64- y[idx] = cbrtf(x[idx]);64+ if (idx >= count) {
65+ return;
66+ }
67+ result[idx] = cbrtf(x[idx]);
65 }68 }
66 ```69 ```
67 70 
68-- SIMD与SIMT混合编程场景:71+- SIMD与SIMT混合编程场景:
69 72 
70 ```73 ```
71- __simt_vf__ __launch_bounds__(1024) inline void KernelCbrt(__gm__ float* x, __gm__ float* y)74+ __simt_vf__ __launch_bounds__(256) inline void compute_cbrtf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
72 {75 {
73- int idx = threadIdx.x + blockIdx.x * blockDim.x;76+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
74- y[idx] = cbrtf(x[idx]);77+ if (idx >= count) {
78+ return;
79+ }
80+ result[idx] = cbrtf(x[idx]);
81+ }
82+ 
83+ __global__ __vector__ void run_cbrtf(__gm__ float *result, __gm__ const float *x, uint32_t count)
84+ {
85+ asc_vf_call<compute_cbrtf_vf>(dim3(256), result, x, count);
75 }86 }
76 ```87 ```
77 88 
89+输入输出示例如下:
90+ 
91+```
92+x:0.25, 0.75, 1.25, 1.75
93+result: 0.6299605 0.9085603 1.077217 1.205071
94+```
@@ -52,23 +52,40 @@ inline float ceilf(float x)
52 52 
53## 调用示例53## 调用示例
54 54 
55-- SIMT编程场景:55+- SIMT编程场景:
56 56 
57 ```57 ```
58- __global__ __launch_bounds__(1024) void KernelCeil(float* dst, float* x)58+ __global__ __launch_bounds__(256) void compute_ceilf(float *result, const float *x, uint32_t count)
59 {59 {
60- int idx = threadIdx.x + blockIdx.x * blockDim.x;60+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
61- dst[idx] = ceilf(x[idx]);61+ if (idx >= count) {
62+ return;
63+ }
64+ result[idx] = ceilf(x[idx]);
62 }65 }
63 ```66 ```
64 67 
65-- SIMD与SIMT混合编程场景:68+- SIMD与SIMT混合编程场景:
66 69 
67 ```70 ```
68- __simt_vf__ __launch_bounds__(1024) inline void KernelCeil(__gm__ float* dst, __gm__ float* x)71+ __simt_vf__ __launch_bounds__(256) inline void compute_ceilf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
69 {72 {
70- int idx = threadIdx.x + blockIdx.x * blockDim.x;73+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
71- dst[idx] = ceilf(x[idx]);74+ if (idx >= count) {
75+ return;
76+ }
77+ result[idx] = ceilf(x[idx]);
78+ }
79+ 
80+ __global__ __vector__ void run_ceilf(__gm__ float *result, __gm__ const float *x, uint32_t count)
81+ {
82+ asc_vf_call<compute_ceilf_vf>(dim3(256), result, x, count);
72 }83 }
73 ```84 ```
74 85 
86+输入输出示例如下:
87+ 
88+```
89+x:0.25, 0.75, 1.25, 1.75
90+result: 1 1 2 2
91+```
@@ -55,23 +55,41 @@ inline float copysignf(float x, float y)
55 55 
56## 调用示例56## 调用示例
57 57 
58-- SIMT编程场景:58+- SIMT编程场景:
59 59 
60 ```60 ```
61- __global__ __launch_bounds__(1024) void KerneCopySign(float* dst, float* x, float* y)61+ __global__ __launch_bounds__(256) void compute_copysignf(float *result, const float *x, const float *y, uint32_t count)
62 {62 {
63- int idx = threadIdx.x + blockIdx.x * blockDim.x;63+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
64- dst[idx] = copysignf(x[idx], y[idx]);64+ if (idx >= count) {
65+ return;
66+ }
67+ result[idx] = copysignf(x[idx], y[idx]);
65 }68 }
66 ```69 ```
67 70 
68-- SIMD与SIMT混合编程场景:71+- SIMD与SIMT混合编程场景:
69 72 
70 ```73 ```
71- __simt_vf__ __launch_bounds__(1024) inline void KerneCopySign(__gm__ float* dst, __gm__ float* x, __gm__ float* y)74+ __simt_vf__ __launch_bounds__(256) inline void compute_copysignf_vf(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, uint32_t count)
72 {75 {
73- int idx = threadIdx.x + blockIdx.x * blockDim.x;76+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
74- dst[idx] = copysignf(x[idx], y[idx]);77+ if (idx >= count) {
78+ return;
79+ }
80+ result[idx] = copysignf(x[idx], y[idx]);
81+ }
82+ 
83+ __global__ __vector__ void run_copysignf(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, uint32_t count)
84+ {
85+ asc_vf_call<compute_copysignf_vf>(dim3(256), result, x, y, count);
75 }86 }
76 ```87 ```
77 88 
89+输入输出示例如下:
90+ 
91+```
92+x:0.25, 0.75, 1.25, 1.75
93+y:1.5, 2.5, 3.5, 4.5
94+result: 0.25 0.75 1.25 1.75
95+```
@@ -54,23 +54,40 @@ inline float cosf(float x)
54 54 
55## 调用示例55## 调用示例
56 56 
57-- SIMT编程场景:57+- SIMT编程场景:
58 58 
59 ```59 ```
60- __global__ __launch_bounds__(1024) void KernelCos(float* dst, float* x)60+ __global__ __launch_bounds__(256) void compute_cosf(float *result, const float *x, uint32_t count)
61 {61 {
62- int idx = threadIdx.x + blockIdx.x * blockDim.x;62+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
63- dst[idx] = cosf(x[idx]); // 对src源地址的第idx个元素取三角函数余弦值63+ if (idx >= count) {
64+ return;
65+ }
66+ result[idx] = cosf(x[idx]);
64 }67 }
65 ```68 ```
66 69 
67-- SIMD与SIMT混合编程场景:70+- SIMD与SIMT混合编程场景:
68 71 
69 ```72 ```
70- __simt_vf__ __launch_bounds__(1024) inline void KernelCos(__gm__ float* dst, __gm__ float* x)73+ __simt_vf__ __launch_bounds__(256) inline void compute_cosf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
71 {74 {
72- int idx = threadIdx.x + blockIdx.x * blockDim.x;75+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
73- dst[idx] = cosf(x[idx]); // 对src源地址的第idx个元素取三角函数余弦值76+ if (idx >= count) {
77+ return;
78+ }
79+ result[idx] = cosf(x[idx]);
80+ }
81+ 
82+ __global__ __vector__ void run_cosf(__gm__ float *result, __gm__ const float *x, uint32_t count)
83+ {
84+ asc_vf_call<compute_cosf_vf>(dim3(256), result, x, count);
74 }85 }
75 ```86 ```
76 87 
88+输入输出示例如下:
89+ 
90+```
91+x:0.25, 0.75, 1.25, 1.75
92+result: 0.9689124 0.7316889 0.3153224 -0.1782461
93+```
@@ -55,23 +55,40 @@ inline float coshf(float x)
55 55 
56## 调用示例56## 调用示例
57 57 
58-- SIMT编程场景:58+- SIMT编程场景:
59 59 
60 ```60 ```
61- __global__ __launch_bounds__(1024) void KernelCosh(float* dst, float* x)61+ __global__ __launch_bounds__(256) void compute_coshf(float *result, const float *x, uint32_t count)
62 {62 {
63- int idx = threadIdx.x + blockIdx.x * blockDim.x;63+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
64- dst[idx] = coshf(x[idx]); // 对src源地址的第idx个元素取双曲余弦值64+ if (idx >= count) {
65+ return;
66+ }
67+ result[idx] = coshf(x[idx]);
65 }68 }
66 ```69 ```
67 70 
68-- SIMD与SIMT混合编程场景:71+- SIMD与SIMT混合编程场景:
69 72 
70 ```73 ```
71- __simt_vf__ __launch_bounds__(1024) inline void KernelCosh(__gm__ float* dst, __gm__ float* x)74+ __simt_vf__ __launch_bounds__(256) inline void compute_coshf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
72 {75 {
73- int idx = threadIdx.x + blockIdx.x * blockDim.x;76+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
74- dst[idx] = coshf(x[idx]); // 对src源地址的第idx个元素取双曲余弦值77+ if (idx >= count) {
78+ return;
79+ }
80+ result[idx] = coshf(x[idx]);
81+ }
82+ 
83+ __global__ __vector__ void run_coshf(__gm__ float *result, __gm__ const float *x, uint32_t count)
84+ {
85+ asc_vf_call<compute_coshf_vf>(dim3(256), result, x, count);
75 }86 }
76 ```87 ```
77 88 
89+输入输出示例如下:
90+ 
91+```
92+x:0.25, 0.75, 1.25, 1.75
93+result: 1.031413 1.294683 1.888424 2.964188
94+```
@@ -56,23 +56,40 @@ inline float cospif(float x)
56 56 
57## 调用示例57## 调用示例
58 58 
59-- SIMT编程场景:59+- SIMT编程场景:
60 60 
61 ```61 ```
62- __global__ __launch_bounds__(1024) void KernelCospi(float* dst, float* x)62+ __global__ __launch_bounds__(256) void compute_cospif(float *result, const float *x, uint32_t count)
63 {63 {
64- int idx = threadIdx.x + blockIdx.x * blockDim.x;64+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
65- dst[idx] = cospif(x[idx]); // 对src源地址的第idx个元素取与π相乘的余弦值65+ if (idx >= count) {
66+ return;
67+ }
68+ result[idx] = cospif(x[idx]);
66 }69 }
67 ```70 ```
68 71 
69-- SIMD与SIMT混合编程场景:72+- SIMD与SIMT混合编程场景:
70 73 
71 ```74 ```
72- __simt_vf__ __launch_bounds__(1024) inline void KernelCospi(__gm__ float* dst, __gm__ float* x)75+ __simt_vf__ __launch_bounds__(256) inline void compute_cospif_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
73 {76 {
74- int idx = threadIdx.x + blockIdx.x * blockDim.x;77+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
75- dst[idx] = cospif(x[idx]); // 对src源地址的第idx个元素取与π相乘的余弦值78+ if (idx >= count) {
79+ return;
80+ }
81+ result[idx] = cospif(x[idx]);
82+ }
83+ 
84+ __global__ __vector__ void run_cospif(__gm__ float *result, __gm__ const float *x, uint32_t count)
85+ {
86+ asc_vf_call<compute_cospif_vf>(dim3(256), result, x, count);
76 }87 }
77 ```88 ```
78 89 
90+输入输出示例如下:
91+ 
92+```
93+x:0.25, 0.75, 1.25, 1.75
94+result: 0.7071068 -0.7071068 -0.7071066 0.707107
95+```
@@ -56,23 +56,40 @@ inline float cyl_bessel_i0f(float x)
56 56 
57## 调用示例57## 调用示例
58 58 
59-- SIMT编程场景:59+- SIMT编程场景:
60 60 
61 ```61 ```
62- __global__ __launch_bounds__(1024) void KernelCylBesselI0(float* dst, float* x)62+ __global__ __launch_bounds__(256) void compute_cyl_bessel_i0f(float *result, const float *x, uint32_t count)
63 {63 {
64- int idx = threadIdx.x + blockIdx.x * blockDim.x;64+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
65- dst[idx] = cyl_bessel_i0f(x[idx]);65+ if (idx >= count) {
66+ return;
67+ }
68+ result[idx] = cyl_bessel_i0f(x[idx]);
66 }69 }
67 ```70 ```
68 71 
69-- SIMD与SIMT混合编程场景:72+- SIMD与SIMT混合编程场景:
70 73 
71 ```74 ```
72- __simt_vf__ __launch_bounds__(1024) inline void KernelCylBesselI0(__gm__ float* dst, __gm__ float* x)75+ __simt_vf__ __launch_bounds__(256) inline void compute_cyl_bessel_i0f_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
73 {76 {
74- int idx = threadIdx.x + blockIdx.x * blockDim.x;77+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
75- dst[idx] = cyl_bessel_i0f(x[idx]);78+ if (idx >= count) {
79+ return;
80+ }
81+ result[idx] = cyl_bessel_i0f(x[idx]);
82+ }
83+ 
84+ __global__ __vector__ void run_cyl_bessel_i0f(__gm__ float *result, __gm__ const float *x, uint32_t count)
85+ {
86+ asc_vf_call<compute_cyl_bessel_i0f_vf>(dim3(256), result, x, count);
76 }87 }
77 ```88 ```
78 89 
90+输入输出示例如下:
91+ 
92+```
93+x:0.25, 0.75, 1.25, 1.75
94+result: 1.015686 1.145647 1.430469 1.925252
95+```
@@ -57,23 +57,40 @@ inline float cyl_bessel_i1f(float x)
57 57 
58## 调用示例58## 调用示例
59 59 
60-- SIMT编程场景:60+- SIMT编程场景:
61 61 
62 ```62 ```
63- __global__ __launch_bounds__(1024) void KernelCylBesselI1(float* dst, float* x)63+ __global__ __launch_bounds__(256) void compute_cyl_bessel_i1f(float *result, const float *x, uint32_t count)
64 {64 {
65- int idx = threadIdx.x + blockIdx.x * blockDim.x;65+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
66- dst[idx] = cyl_bessel_i1f(x[idx]);66+ if (idx >= count) {
67+ return;
68+ }
69+ result[idx] = cyl_bessel_i1f(x[idx]);
67 }70 }
68 ```71 ```
69 72 
70-- SIMD与SIMT混合编程场景:73+- SIMD与SIMT混合编程场景:
71 74 
72 ```75 ```
73- __simt_vf__ __launch_bounds__(1024) inline void KernelCylBesselI1(__gm__ float* dst, __gm__ float* x)76+ __simt_vf__ __launch_bounds__(256) inline void compute_cyl_bessel_i1f_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
74 {77 {
75- int idx = threadIdx.x + blockIdx.x * blockDim.x;78+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
76- dst[idx] = cyl_bessel_i1f(x[idx]);79+ if (idx >= count) {
80+ return;
81+ }
82+ result[idx] = cyl_bessel_i1f(x[idx]);
83+ }
84+ 
85+ __global__ __vector__ void run_cyl_bessel_i1f(__gm__ float *result, __gm__ const float *x, uint32_t count)
86+ {
87+ asc_vf_call<compute_cyl_bessel_i1f_vf>(dim3(256), result, x, count);
77 }88 }
78 ```89 ```
79 90 
91+输入输出示例如下:
92+ 
93+```
94+x:0.25, 0.75, 1.25, 1.75
95+result: 0.1259791 0.4019924 0.7552814 1.255537
96+```
@@ -52,23 +52,40 @@ inline float erfcf(float x)
52 52 
53## 调用示例53## 调用示例
54 54 
55-- SIMT编程场景:55+- SIMT编程场景:
56 56 
57 ```57 ```
58- __global__ __launch_bounds__(1024) void KernelErfc(float* x, float* y)58+ __global__ __launch_bounds__(256) void compute_erfcf(float *result, const float *x, uint32_t count)
59 {59 {
60- int idx = threadIdx.x + blockIdx.x * blockDim.x;60+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
61- y[idx] = erfcf(x[idx]);61+ if (idx >= count) {
62+ return;
63+ }
64+ result[idx] = erfcf(x[idx]);
62 }65 }
63 ```66 ```
64 67 
65-- SIMD与SIMT混合编程场景:68+- SIMD与SIMT混合编程场景:
66 69 
67 ```70 ```
68- __simt_vf__ __launch_bounds__(1024) inline void KernelErfc(__gm__ float* x, __gm__ float* y)71+ __simt_vf__ __launch_bounds__(256) inline void compute_erfcf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
69 {72 {
70- int idx = threadIdx.x + blockIdx.x * blockDim.x;73+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
71- y[idx] = erfcf(x[idx]);74+ if (idx >= count) {
75+ return;
76+ }
77+ result[idx] = erfcf(x[idx]);
78+ }
79+ 
80+ __global__ __vector__ void run_erfcf(__gm__ float *result, __gm__ const float *x, uint32_t count)
81+ {
82+ asc_vf_call<compute_erfcf_vf>(dim3(256), result, x, count);
72 }83 }
73 ```84 ```
74 85 
86+输入输出示例如下:
87+ 
88+```
89+x:0.25, 0.75, 1.25, 1.75
90+result: 0.7236736 0.2888444 0.07709987 0.01332833
91+```
@@ -53,23 +53,40 @@ inline float erfcinvf(float x)
53 53 
54## 调用示例54## 调用示例
55 55 
56-- SIMT编程场景:56+- SIMT编程场景:
57 57 
58 ```58 ```
59- __global__ __launch_bounds__(1024) void KernelErfc(float* x, float* y)59+ __global__ __launch_bounds__(256) void compute_erfcinvf(float *result, const float *x, uint32_t count)
60 {60 {
61- int idx = threadIdx.x + blockIdx.x * blockDim.x;61+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
62- y[idx] = erfcinvf(x[idx]);62+ if (idx >= count) {
63+ return;
64+ }
65+ result[idx] = erfcinvf(x[idx]);
63 }66 }
64 ```67 ```
65 68 
66-- SIMD与SIMT混合编程场景:69+- SIMD与SIMT混合编程场景:
67 70 
68 ```71 ```
69- __simt_vf__ __launch_bounds__(1024) inline void KernelErfc(__gm__ float* x, __gm__ float* y)72+ __simt_vf__ __launch_bounds__(256) inline void compute_erfcinvf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
70 {73 {
71- int idx = threadIdx.x + blockIdx.x * blockDim.x;74+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
72- y[idx] = erfcinvf(x[idx]);75+ if (idx >= count) {
76+ return;
77+ }
78+ result[idx] = erfcinvf(x[idx]);
79+ }
80+ 
81+ __global__ __vector__ void run_erfcinvf(__gm__ float *result, __gm__ const float *x, uint32_t count)
82+ {
83+ asc_vf_call<compute_erfcinvf_vf>(dim3(256), result, x, count);
73 }84 }
74 ```85 ```
75 86 
87+输入输出示例如下:
88+ 
89+```
90+x:0.25, 0.5, 0.75, 1
91+result: 0.8134198 0.4769363 0.2253121 0
92+```
@@ -53,23 +53,40 @@ inline float erfcxf(float x)
53 53 
54## 调用示例54## 调用示例
55 55 
56-- SIMT编程场景:56+- SIMT编程场景:
57 57 
58 ```58 ```
59- __global__ __launch_bounds__(1024) void KernelErfc(float* x, float* y)59+ __global__ __launch_bounds__(256) void compute_erfcxf(float *result, const float *x, uint32_t count)
60 {60 {
61- int idx = threadIdx.x + blockIdx.x * blockDim.x;61+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
62- y[idx] = erfcxf(x[idx]);62+ if (idx >= count) {
63+ return;
64+ }
65+ result[idx] = erfcxf(x[idx]);
63 }66 }
64 ```67 ```
65 68 
66-- SIMD与SIMT混合编程场景:69+- SIMD与SIMT混合编程场景:
67 70 
68 ```71 ```
69- __simt_vf__ __launch_bounds__(1024) inline void KernelErfc(__gm__ float* x, __gm__ float* y)72+ __simt_vf__ __launch_bounds__(256) inline void compute_erfcxf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
70 {73 {
71- int idx = threadIdx.x + blockIdx.x * blockDim.x;74+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
72- y[idx] = erfcxf(x[idx]);75+ if (idx >= count) {
76+ return;
77+ }
78+ result[idx] = erfcxf(x[idx]);
79+ }
80+ 
81+ __global__ __vector__ void run_erfcxf(__gm__ float *result, __gm__ const float *x, uint32_t count)
82+ {
83+ asc_vf_call<compute_erfcxf_vf>(dim3(256), result, x, count);
73 }84 }
74 ```85 ```
75 86 
87+输入输出示例如下:
88+ 
89+```
90+x:0.25, 0.75, 1.25, 1.75
91+result: 0.7703465 0.5069377 0.3678229 0.2849723
92+```
@@ -53,23 +53,40 @@ inline float erff(float x)
53 53 
54## 调用示例54## 调用示例
55 55 
56-- SIMT编程场景:56+- SIMT编程场景:
57 57 
58 ```58 ```
59- __global__ __launch_bounds__(1024) void KernelErf(float* x, float* y)59+ __global__ __launch_bounds__(256) void compute_erff(float *result, const float *x, uint32_t count)
60 {60 {
61- int idx = threadIdx.x + blockIdx.x * blockDim.x;61+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
62- y[idx] = erff(x[idx]);62+ if (idx >= count) {
63+ return;
64+ }
65+ result[idx] = erff(x[idx]);
63 }66 }
64 ```67 ```
65 68 
66-- SIMD与SIMT混合编程场景:69+- SIMD与SIMT混合编程场景:
67 70 
68 ```71 ```
69- __simt_vf__ __launch_bounds__(1024) inline void KernelErf(__gm__ float* x, __gm__ float* y)72+ __simt_vf__ __launch_bounds__(256) inline void compute_erff_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
70 {73 {
71- int idx = threadIdx.x + blockIdx.x * blockDim.x;74+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
72- y[idx] = erff(x[idx]);75+ if (idx >= count) {
76+ return;
77+ }
78+ result[idx] = erff(x[idx]);
79+ }
80+ 
81+ __global__ __vector__ void run_erff(__gm__ float *result, __gm__ const float *x, uint32_t count)
82+ {
83+ asc_vf_call<compute_erff_vf>(dim3(256), result, x, count);
73 }84 }
74 ```85 ```
75 86 
87+输入输出示例如下:
88+ 
89+```
90+x:0.25, 0.75, 1.25, 1.75
91+result: 0.2763264 0.7111557 0.9229001 0.9866717
92+```
@@ -54,23 +54,40 @@ inline float erfinvf(float x)
54 54 
55## 调用示例55## 调用示例
56 56 
57-- SIMT编程场景:57+- SIMT编程场景:
58 58 
59 ```59 ```
60- __global__ __launch_bounds__(1024) void KernelErfc(float* x, float* y)60+ __global__ __launch_bounds__(256) void compute_erfinvf(float *result, const float *x, uint32_t count)
61 {61 {
62- int idx = threadIdx.x + blockIdx.x * blockDim.x;62+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
63- y[idx] = erfinvf(x[idx]);63+ if (idx >= count) {
64+ return;
65+ }
66+ result[idx] = erfinvf(x[idx]);
64 }67 }
65 ```68 ```
66 69 
67-- SIMD与SIMT混合编程场景:70+- SIMD与SIMT混合编程场景:
68 71 
69 ```72 ```
70- __simt_vf__ __launch_bounds__(1024) inline void KernelErfc(__gm__ float* x, __gm__ float* y)73+ __simt_vf__ __launch_bounds__(256) inline void compute_erfinvf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
71 {74 {
72- int idx = threadIdx.x + blockIdx.x * blockDim.x;75+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
73- y[idx] = erfinvf(x[idx]);76+ if (idx >= count) {
77+ return;
78+ }
79+ result[idx] = erfinvf(x[idx]);
80+ }
81+ 
82+ __global__ __vector__ void run_erfinvf(__gm__ float *result, __gm__ const float *x, uint32_t count)
83+ {
84+ asc_vf_call<compute_erfinvf_vf>(dim3(256), result, x, count);
74 }85 }
75 ```86 ```
76 87 
88+输入输出示例如下:
89+ 
90+```
91+x:-0.75, -0.25, 0.25, 0.75
92+result: -0.8134198 -0.2253121 0.2253121 0.8134198
93+```
@@ -54,23 +54,40 @@ inline float exp10f(float x)
54 54 
55## 调用示例55## 调用示例
56 56 
57-- SIMT编程场景:57+- SIMT编程场景:
58 58 
59 ```59 ```
60- __global__ __launch_bounds__(1024) void KernelExp10(float* dst, float* x)60+ __global__ __launch_bounds__(256) void compute_exp10f(float *result, const float *x, uint32_t count)
61 {61 {
62- int idx = threadIdx.x + blockIdx.x * blockDim.x;62+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
63- dst[idx] = exp10f(x[idx]);63+ if (idx >= count) {
64+ return;
65+ }
66+ result[idx] = exp10f(x[idx]);
64 }67 }
65 ```68 ```
66 69 
67-- SIMD与SIMT混合编程场景:70+- SIMD与SIMT混合编程场景:
68 71 
69 ```72 ```
70- __simt_vf__ __launch_bounds__(1024) inline void KernelExp10(__gm__ float* dst, __gm__ float* x)73+ __simt_vf__ __launch_bounds__(256) inline void compute_exp10f_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
71 {74 {
72- int idx = threadIdx.x + blockIdx.x * blockDim.x;75+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
73- dst[idx] = exp10f(x[idx]);76+ if (idx >= count) {
77+ return;
78+ }
79+ result[idx] = exp10f(x[idx]);
80+ }
81+ 
82+ __global__ __vector__ void run_exp10f(__gm__ float *result, __gm__ const float *x, uint32_t count)
83+ {
84+ asc_vf_call<compute_exp10f_vf>(dim3(256), result, x, count);
74 }85 }
75 ```86 ```
76 87 
88+输入输出示例如下:
89+ 
90+```
91+x:0.25, 0.75, 1.25, 1.75
92+result: 1.778279 5.623414 17.7828 56.23413
93+```
@@ -59,23 +59,40 @@ inline float exp2f(float x)
59 59 
60## 调用示例60## 调用示例
61 61 
62-- SIMT编程场景:62+- SIMT编程场景:
63 63 
64 ```64 ```
65- __global__ __launch_bounds__(1024) void KernelExp2(float* dst, float* x)65+ __global__ __launch_bounds__(256) void compute_exp2f(float *result, const float *x, uint32_t count)
66 {66 {
67- int idx = threadIdx.x + blockIdx.x * blockDim.x;67+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
68- dst[idx] = exp2f(x[idx]);68+ if (idx >= count) {
69+ return;
70+ }
71+ result[idx] = exp2f(x[idx]);
69 }72 }
70 ```73 ```
71 74 
72-- SIMD与SIMT混合编程场景:75+- SIMD与SIMT混合编程场景:
73 76 
74 ```77 ```
75- __simt_vf__ __launch_bounds__(1024) inline void KernelExp2(__gm__ float* dst, __gm__ float* x)78+ __simt_vf__ __launch_bounds__(256) inline void compute_exp2f_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
76 {79 {
77- int idx = threadIdx.x + blockIdx.x * blockDim.x;80+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
78- dst[idx] = exp2f(x[idx]);81+ if (idx >= count) {
82+ return;
83+ }
84+ result[idx] = exp2f(x[idx]);
85+ }
86+ 
87+ __global__ __vector__ void run_exp2f(__gm__ float *result, __gm__ const float *x, uint32_t count)
88+ {
89+ asc_vf_call<compute_exp2f_vf>(dim3(256), result, x, count);
79 }90 }
80 ```91 ```
81 92 
93+输入输出示例如下:
94+ 
95+```
96+x:0.25, 0.75, 1.25, 1.75
97+result: 1.189207 1.681793 2.378414 3.363586
98+```
@@ -54,23 +54,40 @@ e的x次方。
54 54 
55## 调用示例55## 调用示例
56 56 
57-- SIMT编程场景:57+- SIMT编程场景:
58 58 
59 ```59 ```
60- __global__ __launch_bounds__(1024) void KernelExp(float* dst, float* x)60+ __global__ __launch_bounds__(256) void compute_expf(float *result, const float *x, uint32_t count)
61 {61 {
62- int idx = threadIdx.x + blockIdx.x * blockDim.x;62+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
63- dst[idx] = expf(x[idx]);63+ if (idx >= count) {
64+ return;
65+ }
66+ result[idx] = expf(x[idx]);
64 }67 }
65 ```68 ```
66 69 
67-- SIMD与SIMT混合编程场景:70+- SIMD与SIMT混合编程场景:
68 71 
69 ```72 ```
70- __simt_vf__ __launch_bounds__(1024) inline void KernelExp(__gm__ float* dst, __gm__ float* x)73+ __simt_vf__ __launch_bounds__(256) inline void compute_expf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
71 {74 {
72- int idx = threadIdx.x + blockIdx.x * blockDim.x;75+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
73- dst[idx] = expf(x[idx]);76+ if (idx >= count) {
77+ return;
78+ }
79+ result[idx] = expf(x[idx]);
80+ }
81+ 
82+ __global__ __vector__ void run_expf(__gm__ float *result, __gm__ const float *x, uint32_t count)
83+ {
84+ asc_vf_call<compute_expf_vf>(dim3(256), result, x, count);
74 }85 }
75 ```86 ```
76 87 
88+输入输出示例如下:
89+ 
90+```
91+x:0.25, 0.75, 1.25, 1.75
92+result: 1.284025 2.117 3.490343 5.754602
93+```
@@ -54,23 +54,40 @@ e的x次方减1。
54 54 
55## 调用示例55## 调用示例
56 56 
57-- SIMT编程场景:57+- SIMT编程场景:
58 58 
59 ```59 ```
60- __global__ __launch_bounds__(1024) void KernelExpm1(float* dst, float* x)60+ __global__ __launch_bounds__(256) void compute_expm1f(float *result, const float *x, uint32_t count)
61 {61 {
62- int idx = threadIdx.x + blockIdx.x * blockDim.x;62+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
63- dst[idx] = expm1f(x[idx]);63+ if (idx >= count) {
64+ return;
65+ }
66+ result[idx] = expm1f(x[idx]);
64 }67 }
65 ```68 ```
66 69 
67-- SIMD与SIMT混合编程场景:70+- SIMD与SIMT混合编程场景:
68 71 
69 ```72 ```
70- __simt_vf__ __launch_bounds__(1024) inline void KernelExpm1(__gm__ float* dst, __gm__ float* x)73+ __simt_vf__ __launch_bounds__(256) inline void compute_expm1f_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
71 {74 {
72- int idx = threadIdx.x + blockIdx.x * blockDim.x;75+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
73- dst[idx] = expm1f(x[idx]);76+ if (idx >= count) {
77+ return;
78+ }
79+ result[idx] = expm1f(x[idx]);
80+ }
81+ 
82+ __global__ __vector__ void run_expm1f(__gm__ float *result, __gm__ const float *x, uint32_t count)
83+ {
84+ asc_vf_call<compute_expm1f_vf>(dim3(256), result, x, count);
74 }85 }
75 ```86 ```
76 87 
88+输入输出示例如下:
89+ 
90+```
91+x:0.25, 0.75, 1.25, 1.75
92+result: 0.2840254 1.117 2.490343 4.754602
93+```
@@ -55,23 +55,40 @@ inline float fabsf(float x)
55 55 
56## 调用示例56## 调用示例
57 57 
58-- SIMT编程场景:58+- SIMT编程场景:
59 59 
60 ```60 ```
61- __global__ __launch_bounds__(1024) void KernelAbs(float* dst, float* x)61+ __global__ __launch_bounds__(256) void compute_fabsf(float *result, const float *x, uint32_t count)
62 {62 {
63- int idx = threadIdx.x + blockIdx.x * blockDim.x;63+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
64- dst[idx] = fabsf(x[idx]);64+ if (idx >= count) {
65+ return;
66+ }
67+ result[idx] = fabsf(x[idx]);
65 }68 }
66 ```69 ```
67 70 
68-- SIMD与SIMT混合编程场景:71+- SIMD与SIMT混合编程场景:
69 72 
70 ```73 ```
71- __simt_vf__ __launch_bounds__(1024) inline void KernelAbs(__gm__ float* dst, __gm__ float* x)74+ __simt_vf__ __launch_bounds__(256) inline void compute_fabsf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
72 {75 {
73- int idx = threadIdx.x + blockIdx.x * blockDim.x;76+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
74- dst[idx] = fabsf(x[idx]);77+ if (idx >= count) {
78+ return;
79+ }
80+ result[idx] = fabsf(x[idx]);
81+ }
82+ 
83+ __global__ __vector__ void run_fabsf(__gm__ float *result, __gm__ const float *x, uint32_t count)
84+ {
85+ asc_vf_call<compute_fabsf_vf>(dim3(256), result, x, count);
75 }86 }
76 ```87 ```
77 88 
89+输入输出示例如下:
90+ 
91+```
92+x:0.25, 0.75, 1.25, 1.75
93+result: 0.25 0.75 1.25 1.75
94+```
@@ -56,23 +56,41 @@ inline float fdimf(float x, float y)
56 56 
57## 调用示例57## 调用示例
58 58 
59-- SIMT编程场景:59+- SIMT编程场景:
60 60 
61 ```61 ```
62- __global__ __launch_bounds__(1024) void KernelDim(float* dst, float* x, float* y)62+ __global__ __launch_bounds__(256) void compute_fdimf(float *result, const float *x, const float *y, uint32_t count)
63 {63 {
64- int idx = threadIdx.x + blockIdx.x * blockDim.x;64+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
65- dst[idx] = fdimf(x[idx], y[idx]);65+ if (idx >= count) {
66+ return;
67+ }
68+ result[idx] = fdimf(x[idx], y[idx]);
66 }69 }
67 ```70 ```
68 71 
69-- SIMD与SIMT混合编程场景:72+- SIMD与SIMT混合编程场景:
70 73 
71 ```74 ```
72- __simt_vf__ __launch_bounds__(1024) inline void KernelDim(__gm__ float* dst, __gm__ float* x, __gm__ float* y)75+ __simt_vf__ __launch_bounds__(256) inline void compute_fdimf_vf(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, uint32_t count)
73 {76 {
74- int idx = threadIdx.x + blockIdx.x * blockDim.x;77+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
75- dst[idx] = fdimf(x[idx], y[idx]);78+ if (idx >= count) {
79+ return;
80+ }
81+ result[idx] = fdimf(x[idx], y[idx]);
82+ }
83+ 
84+ __global__ __vector__ void run_fdimf(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, uint32_t count)
85+ {
86+ asc_vf_call<compute_fdimf_vf>(dim3(256), result, x, y, count);
76 }87 }
77 ```88 ```
78 89 
90+输入输出示例如下:
91+ 
92+```
93+x:0.25, 0.75, 1.25, 1.75
94+y:1.5, 2.5, 3.5, 4.5
95+result: 0 0 0 0
96+```
@@ -54,22 +54,41 @@ inline float fdividef(float x, float y)
54 54 
55## 调用示例55## 调用示例
56 56 
57-- SIMT编程场景:57+- SIMT编程场景:
58 58 
59 ```59 ```
60- __global__ __launch_bounds__(1024) void KernelFdividef(float* dst, float* x, float* y)60+ __global__ __launch_bounds__(256) void compute_fdividef(float *result, const float *x, const float *y, uint32_t count)
61 {61 {
62- int idx = threadIdx.x + blockIdx.x * blockDim.x;62+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
63- dst[idx] = fdividef(x[idx], y[idx]);63+ if (idx >= count) {
64+ return;
65+ }
66+ result[idx] = fdividef(x[idx], y[idx]);
64 }67 }
65 ```68 ```
66 69 
67-- SIMD与SIMT混合编程场景:70+- SIMD与SIMT混合编程场景:
68 71 
69 ```72 ```
70- __simt_vf__ __launch_bounds__(1024) inline void KernelFdividef(__gm__ float* dst, __gm__ float* x, __gm__ float* y)73+ __simt_vf__ __launch_bounds__(256) inline void compute_fdividef_vf(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, uint32_t count)
71 {74 {
72- int idx = threadIdx.x + blockIdx.x * blockDim.x;75+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
73- dst[idx] = fdividef(x[idx], y[idx]);76+ if (idx >= count) {
77+ return;
78+ }
79+ result[idx] = fdividef(x[idx], y[idx]);
80+ }
81+ 
82+ __global__ __vector__ void run_fdividef(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, uint32_t count)
83+ {
84+ asc_vf_call<compute_fdividef_vf>(dim3(256), result, x, y, count);
74 }85 }
75 ```86 ```
87+ 
88+输入输出示例如下:
89+ 
90+```
91+x:0.25, 0.75, 1.25, 1.75
92+y:1.5, 2.5, 3.5, 4.5
93+result: 0.1666667 0.3 0.3571429 0.3888889
94+```
@@ -52,23 +52,40 @@ inline float floorf(float x)
52 52 
53## 调用示例53## 调用示例
54 54 
55-- SIMT编程场景:55+- SIMT编程场景:
56 56 
57 ```57 ```
58- __global__ __launch_bounds__(1024) void KernelIsFinite(float* dst, float* x)58+ __global__ __launch_bounds__(256) void compute_floorf(float *result, const float *x, uint32_t count)
59 {59 {
60- int idx = threadIdx.x + blockIdx.x * blockDim.x;60+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
61- dst[idx] = floorf(x[idx]);61+ if (idx >= count) {
62+ return;
63+ }
64+ result[idx] = floorf(x[idx]);
62 }65 }
63 ```66 ```
64 67 
65-- SIMD与SIMT混合编程场景:68+- SIMD与SIMT混合编程场景:
66 69 
67 ```70 ```
68- __simt_vf__ __launch_bounds__(1024) inline void KernelIsFinite(__gm__ float* dst, __gm__ float* x)71+ __simt_vf__ __launch_bounds__(256) inline void compute_floorf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
69 {72 {
70- int idx = threadIdx.x + blockIdx.x * blockDim.x;73+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
71- dst[idx] = floorf(x[idx]);74+ if (idx >= count) {
75+ return;
76+ }
77+ result[idx] = floorf(x[idx]);
78+ }
79+ 
80+ __global__ __vector__ void run_floorf(__gm__ float *result, __gm__ const float *x, uint32_t count)
81+ {
82+ asc_vf_call<compute_floorf_vf>(dim3(256), result, x, count);
72 }83 }
73 ```84 ```
74 85 
86+输入输出示例如下:
87+ 
88+```
89+x:0.25, 0.75, 1.25, 1.75
90+result: 0 0 1 1
91+```
@@ -60,21 +60,42 @@ x \* y + z的值。
60 60 
61## 调用示例61## 调用示例
62 62 
63-- SIMT编程场景:63+- SIMT编程场景:
64 64 
65 ```65 ```
66- __global__ __launch_bounds__(1024) void KernelFma(float* dst, float* x, float* y, float* z){66+ __global__ __launch_bounds__(256) void compute_fmaf(float *result, const float *x, const float *y, const float *z, uint32_t count)
67- int idx = threadIdx.x + blockIdx.x * blockDim.x;67+ {
68- dst[idx] = fmaf(x[idx], y[idx], z[idx]);68+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
69+ if (idx >= count) {
70+ return;
71+ }
72+ result[idx] = fmaf(x[idx], y[idx], z[idx]);
69 }73 }
70 ```74 ```
71 75 
72-- SIMD与SIMT混合编程场景:76+- SIMD与SIMT混合编程场景:
73 77 
74 ```78 ```
75- __simt_vf__ __launch_bounds__(1024) inline void KernelFma(__gm__ float* dst, __gm__ float* x, __gm__ float* y, __gm__ float* z){79+ __simt_vf__ __launch_bounds__(256) inline void compute_fmaf_vf(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, __gm__ const float *z, uint32_t count)
76- int idx = threadIdx.x + blockIdx.x * blockDim.x;80+ {
77- dst[idx] = fmaf(x[idx], y[idx], z[idx]);81+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
82+ if (idx >= count) {
83+ return;
84+ }
85+ result[idx] = fmaf(x[idx], y[idx], z[idx]);
86+ }
87+ 
88+ __global__ __vector__ void run_fmaf(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, __gm__ const float *z, uint32_t count)
89+ {
90+ asc_vf_call<compute_fmaf_vf>(dim3(256), result, x, y, z, count);
78 }91 }
79 ```92 ```
80 93 
94+输入输出示例如下:
95+ 
96+```
97+x:0.25, 0.75, 1.25, 1.75
98+y:1.5, 2.5, 3.5, 4.5
99+z:-0.5, 0.5, 1.5, 2.5
100+result: -0.125 2.375 5.875 10.375
101+```
@@ -60,23 +60,41 @@ inline float fmaxf(float x, float y)
60 60 
61## 调用示例61## 调用示例
62 62 
63-- SIMT编程场景:63+- SIMT编程场景:
64 64 
65 ```65 ```
66- __global__ __launch_bounds__(1024) void KernelMax(float* dst, float* x, float* y)66+ __global__ __launch_bounds__(256) void compute_fmaxf(float *result, const float *x, const float *y, uint32_t count)
67 {67 {
68- int idx = threadIdx.x + blockIdx.x * blockDim.x;68+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
69- dst[idx] = fmaxf(x[idx], y[idx]);69+ if (idx >= count) {
70+ return;
71+ }
72+ result[idx] = fmaxf(x[idx], y[idx]);
70 }73 }
71 ```74 ```
72 75 
73-- SIMD与SIMT混合编程场景:76+- SIMD与SIMT混合编程场景:
74 77 
75 ```78 ```
76- __simt_vf__ __launch_bounds__(1024) inline void KernelMax(__gm__ float* dst, __gm__ float* x, __gm__ float* y)79+ __simt_vf__ __launch_bounds__(256) inline void compute_fmaxf_vf(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, uint32_t count)
77 {80 {
78- int idx = threadIdx.x + blockIdx.x * blockDim.x;81+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
79- dst[idx] = fmaxf(x[idx], y[idx]);82+ if (idx >= count) {
83+ return;
84+ }
85+ result[idx] = fmaxf(x[idx], y[idx]);
86+ }
87+ 
88+ __global__ __vector__ void run_fmaxf(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, uint32_t count)
89+ {
90+ asc_vf_call<compute_fmaxf_vf>(dim3(256), result, x, y, count);
80 }91 }
81 ```92 ```
82 93 
94+输入输出示例如下:
95+ 
96+```
97+x:0.25, 0.75, 1.25, 1.75
98+y:1.5, 2.5, 3.5, 4.5
99+result: 1.5 2.5 3.5 4.5
100+```
@@ -60,23 +60,41 @@ inline float fminf(float x, float y)
60 60 
61## 调用示例61## 调用示例
62 62 
63-- SIMT编程场景:63+- SIMT编程场景:
64 64 
65 ```65 ```
66- __global__ __launch_bounds__(1024) void KernelMin(float* dst, float* x, float* y)66+ __global__ __launch_bounds__(256) void compute_fminf(float *result, const float *x, const float *y, uint32_t count)
67 {67 {
68- int idx = threadIdx.x + blockIdx.x * blockDim.x;68+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
69- dst[idx] = fminf(x[idx], y[idx]);69+ if (idx >= count) {
70+ return;
71+ }
72+ result[idx] = fminf(x[idx], y[idx]);
70 }73 }
71 ```74 ```
72 75 
73-- SIMD与SIMT混合编程场景:76+- SIMD与SIMT混合编程场景:
74 77 
75 ```78 ```
76- __simt_vf__ __launch_bounds__(1024) inline void KernelMin(__gm__ float* dst, __gm__ float* x, __gm__ float* y)79+ __simt_vf__ __launch_bounds__(256) inline void compute_fminf_vf(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, uint32_t count)
77 {80 {
78- int idx = threadIdx.x + blockIdx.x * blockDim.x;81+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
79- dst[idx] = fminf(x[idx], y[idx]);82+ if (idx >= count) {
83+ return;
84+ }
85+ result[idx] = fminf(x[idx], y[idx]);
86+ }
87+ 
88+ __global__ __vector__ void run_fminf(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, uint32_t count)
89+ {
90+ asc_vf_call<compute_fminf_vf>(dim3(256), result, x, y, count);
80 }91 }
81 ```92 ```
82 93 
94+输入输出示例如下:
95+ 
96+```
97+x:0.25, 0.75, 1.25, 1.75
98+y:1.5, 2.5, 3.5, 4.5
99+result: 0.25 0.75 1.25 1.75
100+```
@@ -54,23 +54,41 @@ inline float fmodf(float x, float y)
54 54 
55## 调用示例55## 调用示例
56 56 
57-- SIMT编程场景:57+- SIMT编程场景:
58 58 
59 ```59 ```
60- __global__ __launch_bounds__(1024) void KernelMod(float* dst, float* x, float* y)60+ __global__ __launch_bounds__(256) void compute_fmodf(float *result, const float *x, const float *y, uint32_t count)
61 {61 {
62- int idx = threadIdx.x + blockIdx.x * blockDim.x;62+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
63- dst[idx] = fmodf(x[idx], y[idx]);63+ if (idx >= count) {
64+ return;
65+ }
66+ result[idx] = fmodf(x[idx], y[idx]);
64 }67 }
65 ```68 ```
66 69 
67-- SIMD与SIMT混合编程场景:70+- SIMD与SIMT混合编程场景:
68 71 
69 ```72 ```
70- __simt_vf__ __launch_bounds__(1024) inline void KernelMod(__gm__ float* dst, __gm__ float* x, __gm__ float* y)73+ __simt_vf__ __launch_bounds__(256) inline void compute_fmodf_vf(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, uint32_t count)
71 {74 {
72- int idx = threadIdx.x + blockIdx.x * blockDim.x;75+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
73- dst[idx] = fmodf(x[idx], y[idx]);76+ if (idx >= count) {
77+ return;
78+ }
79+ result[idx] = fmodf(x[idx], y[idx]);
80+ }
81+ 
82+ __global__ __vector__ void run_fmodf(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, uint32_t count)
83+ {
84+ asc_vf_call<compute_fmodf_vf>(dim3(256), result, x, y, count);
74 }85 }
75 ```86 ```
76 87 
88+输入输出示例如下:
89+ 
90+```
91+x:0.25, 0.75, 1.25, 1.75
92+y:1.25, 2.25, 3.25, 4.25
93+result: 0.25 0.75 1.25 1.75
94+```
@@ -38,7 +38,7 @@ inline float frexpf(float x, int *exp)
38将x转换为归一化\[1/2, 1\)的有符号数乘以2的积分幂时,返回该有符号数。38将x转换为归一化\[1/2, 1\)的有符号数乘以2的积分幂时,返回该有符号数。
39 39 
40- 当x为inf时,返回值为inf,exp为0。40- 当x为inf时,返回值为inf,exp为0。
41-- 当x为-inf时,返回值为inf,exp为0。41+- 当x为-inf时,返回值为-inf,exp为0。
42- 当x为nan时,返回值为nan,exp为0。42- 当x为nan时,返回值为nan,exp为0。
43 43 
44## 约束说明44## 约束说明
@@ -55,24 +55,41 @@ inline float frexpf(float x, int *exp)
55 55 
56## 调用示例56## 调用示例
57 57 
58-- SIMT编程场景:58+- SIMT编程场景:
59 59 
60 ```60 ```
61- __global__ __launch_bounds__(1024) void KernelFrexp(float* dst1, int* dst2, float* x)61+ __global__ __launch_bounds__(256) void compute_frexpf(float *result, int *exp_out, const float *x, uint32_t count)
62 {62 {
63- int idx = threadIdx.x + blockIdx.x * blockDim.x;63+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
64- dst1[idx] = frexpf(x[idx], dst2 + idx);64+ if (idx >= count) {
65+ return;
66+ }
67+ result[idx] = frexpf(x[idx], exp_out + idx);
65 }68 }
66 ```69 ```
67 70 
68-- SIMD与SIMT混合编程场景:71+- SIMD与SIMT混合编程场景:
69- 
70- SIMD与SIMT混合编程场景,需要显式使用地址空间限定符表示地址空间:\_\_gm\_\_表示Global Memory内存空间,\_\_ubuf\_\_表示Unified Buffer内存空间,栈空间无需添加地址空间限定符。
71 72 
72 ```73 ```
73- __simt_vf__ __launch_bounds__(1024) inline void KernelFrexp(__gm__ float* dst1, __gm__ int* dst2, __gm__ float* x)74+ __simt_vf__ __launch_bounds__(256) inline void compute_frexpf_vf(__gm__ float *result, __gm__ int *exp_out, __gm__ const float *x, uint32_t count)
74 {75 {
75- int idx = threadIdx.x + blockIdx.x * blockDim.x;76+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
76- dst1[idx] = frexpf(x[idx], dst2 + idx);77+ if (idx >= count) {
78+ return;
79+ }
80+ result[idx] = frexpf(x[idx], exp_out + idx);
81+ }
82+ 
83+ __global__ __vector__ void run_frexpf(__gm__ float *result, __gm__ int *exp_out, __gm__ const float *x, uint32_t count)
84+ {
85+ asc_vf_call<compute_frexpf_vf>(dim3(256), result, exp_out, x, count);
77 }86 }
78 ```87 ```
88+ 
89+输入输出示例如下:
90+ 
91+```
92+x:0.25, 0.75, 1.25, 1.75
93+result: 0.5 0.75 0.625 0.875
94+exp_out: -1 0 1 1
95+```
@@ -55,23 +55,41 @@ x^2 + y^2的平方根。
55 55 
56## 调用示例56## 调用示例
57 57 
58-- SIMT编程场景:58+- SIMT编程场景:
59 59 
60 ```60 ```
61- __global__ __launch_bounds__(1024) void KernelHypot(float* dst, float* x, float* y)61+ __global__ __launch_bounds__(256) void compute_hypotf(float *result, const float *x, const float *y, uint32_t count)
62 {62 {
63- int idx = threadIdx.x + blockIdx.x * blockDim.x;63+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
64- dst[idx] = hypotf(x[idx], y[idx]);64+ if (idx >= count) {
65+ return;
66+ }
67+ result[idx] = hypotf(x[idx], y[idx]);
65 }68 }
66 ```69 ```
67 70 
68-- SIMD与SIMT混合编程场景:71+- SIMD与SIMT混合编程场景:
69 72 
70 ```73 ```
71- __simt_vf__ __launch_bounds__(1024) inline void KernelHypot(__gm__ float* dst, __gm__ float* x, __gm__ float* y)74+ __simt_vf__ __launch_bounds__(256) inline void compute_hypotf_vf(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, uint32_t count)
72 {75 {
73- int idx = threadIdx.x + blockIdx.x * blockDim.x;76+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
74- dst[idx] = hypotf(x[idx], y[idx]);77+ if (idx >= count) {
78+ return;
79+ }
80+ result[idx] = hypotf(x[idx], y[idx]);
81+ }
82+ 
83+ __global__ __vector__ void run_hypotf(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, uint32_t count)
84+ {
85+ asc_vf_call<compute_hypotf_vf>(dim3(256), result, x, y, count);
75 }86 }
76 ```87 ```
77 88 
89+输入输出示例如下:
90+ 
91+```
92+x:0.25, 0.75, 1.25, 1.75
93+y:1.5, 2.5, 3.5, 4.5
94+result: 1.520691 2.610077 3.716517 4.828302
95+```
@@ -34,12 +34,17 @@ inline int ilogbf(float x)
34 34 
35## 返回值说明35## 返回值说明
36 36 
37-以2为底的x的对数,并向下取整后的值。37+对于有限非零值,返回以2为底的x的绝对值的对数,并向下取整后的整数值。
38 38 
39-- 当x为inf时,返回值为2147483647。39+特殊取值场景下的返回值如下表所示。
40-- 当x为-inf时,返回值为2147483647。40+ 
41-- 当x为nan时,返回值为-2147483648。41+| 输入值x | 返回值 |
42-- 当x为0时,返回值为-2147483648。42+| --- | --- |
43+| +0 | ASCRT_MIN_VAL_S |
44+| -0 | ASCRT_MIN_VAL_S |
45+| inf | ASCRT_MAX_VAL_S |
46+| -inf | ASCRT_MAX_VAL_S |
47+| nan | ASCRT_MIN_VAL_S |
43 48 
44## 约束说明49## 约束说明
45 50 
@@ -55,23 +60,40 @@ inline int ilogbf(float x)
55 60 
56## 调用示例61## 调用示例
57 62 
58-- SIMT编程场景:63+- SIMT编程场景:
59 64 
60 ```65 ```
61- __global__ __launch_bounds__(1024) void KernelILogb(int* dst, float* x)66+ __global__ __launch_bounds__(256) void compute_ilogbf(float *result, const float *x, uint32_t count)
62 {67 {
63- int idx = threadIdx.x + blockIdx.x * blockDim.x;68+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
64- dst[idx] = ilogbf(x[idx]);69+ if (idx >= count) {
70+ return;
71+ }
72+ result[idx] = static_cast<float>(ilogbf(x[idx]));
65 }73 }
66 ```74 ```
67 75 
68-- SIMD与SIMT混合编程场景:76+- SIMD与SIMT混合编程场景:
69 77 
70 ```78 ```
71- __simt_vf__ __launch_bounds__(1024) inline void KernelILogb(__gm__ int* dst, __gm__ float* x)79+ __simt_vf__ __launch_bounds__(256) inline void compute_ilogbf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
72 {80 {
73- int idx = threadIdx.x + blockIdx.x * blockDim.x;81+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
74- dst[idx] = ilogbf(x[idx]);82+ if (idx >= count) {
83+ return;
84+ }
85+ result[idx] = static_cast<float>(ilogbf(x[idx]));
86+ }
87+ 
88+ __global__ __vector__ void run_ilogbf(__gm__ float *result, __gm__ const float *x, uint32_t count)
89+ {
90+ asc_vf_call<compute_ilogbf_vf>(dim3(256), result, x, count);
75 }91 }
76 ```92 ```
77 93 
94+输入输出示例如下:
95+ 
96+```
97+x:0.25, 0.75, 1.25, 1.75
98+result: -2 -1 0 0
99+```
@@ -49,23 +49,40 @@ inline bool isfinite(float x)
49 49 
50## 调用示例50## 调用示例
51 51 
52-- SIMT编程场景:52+- SIMT编程场景:
53 53 
54 ```54 ```
55- __global__ __launch_bounds__(1024) void KernelIsFinite(bool* dst, float* x)55+ __global__ __launch_bounds__(256) void compute_isfinite(float *result, const float *x, uint32_t count)
56 {56 {
57- int idx = threadIdx.x + blockIdx.x * blockDim.x;57+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
58- dst[idx] = isfinite(x[idx]);58+ if (idx >= count) {
59+ return;
60+ }
61+ result[idx] = isfinite(x[idx]) ? 1.0F : 0.0F;
59 }62 }
60 ```63 ```
61 64 
62-- SIMD与SIMT混合编程场景:65+- SIMD与SIMT混合编程场景:
63 66 
64 ```67 ```
65- __simt_vf__ __launch_bounds__(1024) inline void KernelIsFinite(__gm__ bool* dst, __gm__ float* x)68+ __simt_vf__ __launch_bounds__(256) inline void compute_isfinite_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
66 {69 {
67- int idx = threadIdx.x + blockIdx.x * blockDim.x;70+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
68- dst[idx] = isfinite(x[idx]);71+ if (idx >= count) {
72+ return;
73+ }
74+ result[idx] = isfinite(x[idx]) ? 1.0F : 0.0F;
75+ }
76+ 
77+ __global__ __vector__ void run_isfinite(__gm__ float *result, __gm__ const float *x, uint32_t count)
78+ {
79+ asc_vf_call<compute_isfinite_vf>(dim3(256), result, x, count);
69 }80 }
70 ```81 ```
71 82 
83+输入输出示例如下:
84+ 
85+```
86+x:0.25, 0.75, 1.25, 1.75
87+result: 1 1 1 1
88+```
@@ -49,23 +49,40 @@ inline bool isinf(float x)
49 49 
50## 调用示例50## 调用示例
51 51 
52-- SIMT编程场景:52+- SIMT编程场景:
53 53 
54 ```54 ```
55- __global__ __launch_bounds__(1024) void KernelIsInf(bool* dst, float* x)55+ __global__ __launch_bounds__(256) void compute_isinf(float *result, const float *x, uint32_t count)
56 {56 {
57- int idx = threadIdx.x + blockIdx.x * blockDim.x;57+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
58- dst[idx] = isinf(x[idx]);58+ if (idx >= count) {
59+ return;
60+ }
61+ result[idx] = isinf(x[idx]) ? 1.0F : 0.0F;
59 }62 }
60 ```63 ```
61 64 
62-- SIMD与SIMT混合编程场景:65+- SIMD与SIMT混合编程场景:
63 66 
64 ```67 ```
65- __simt_vf__ __launch_bounds__(1024) inline void KernelIsInf(__gm__ bool* dst, __gm__ float* x)68+ __simt_vf__ __launch_bounds__(256) inline void compute_isinf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
66 {69 {
67- int idx = threadIdx.x + blockIdx.x * blockDim.x;70+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
68- dst[idx] = isinf(x[idx]);71+ if (idx >= count) {
72+ return;
73+ }
74+ result[idx] = isinf(x[idx]) ? 1.0F : 0.0F;
75+ }
76+ 
77+ __global__ __vector__ void run_isinf(__gm__ float *result, __gm__ const float *x, uint32_t count)
78+ {
79+ asc_vf_call<compute_isinf_vf>(dim3(256), result, x, count);
69 }80 }
70 ```81 ```
71 82 
83+输入输出示例如下:
84+ 
85+```
86+x:0.25, 0.75, 1.25, 1.75
87+result: 0 0 0 0
88+```
@@ -49,23 +49,40 @@ inline bool isnan(float x)
49 49 
50## 调用示例50## 调用示例
51 51 
52-- SIMT编程场景:52+- SIMT编程场景:
53 53 
54 ```54 ```
55- __global__ __launch_bounds__(1024) void KernelIsNan(bool* dst, float* x)55+ __global__ __launch_bounds__(256) void compute_isnan(float *result, const float *x, uint32_t count)
56 {56 {
57- int idx = threadIdx.x + blockIdx.x * blockDim.x;57+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
58- dst[idx] = isnan(x[idx]);58+ if (idx >= count) {
59+ return;
60+ }
61+ result[idx] = isnan(x[idx]) ? 1.0F : 0.0F;
59 }62 }
60 ```63 ```
61 64 
62-- SIMD与SIMT混合编程场景:65+- SIMD与SIMT混合编程场景:
63 66 
64 ```67 ```
65- __simt_vf__ __launch_bounds__(1024) inline void KernelIsNan(__gm__ bool* dst, __gm__ float* x)68+ __simt_vf__ __launch_bounds__(256) inline void compute_isnan_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
66 {69 {
67- int idx = threadIdx.x + blockIdx.x * blockDim.x;70+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
68- dst[idx] = isnan(x[idx]);71+ if (idx >= count) {
72+ return;
73+ }
74+ result[idx] = isnan(x[idx]) ? 1.0F : 0.0F;
75+ }
76+ 
77+ __global__ __vector__ void run_isnan(__gm__ float *result, __gm__ const float *x, uint32_t count)
78+ {
79+ asc_vf_call<compute_isnan_vf>(dim3(256), result, x, count);
69 }80 }
70 ```81 ```
71 82 
83+输入输出示例如下:
84+ 
85+```
86+x:0.25, 0.75, 1.25, 1.75
87+result: 0 0 0 0
88+```
@@ -53,23 +53,40 @@ inline float j0f(float x)
53 53 
54## 调用示例54## 调用示例
55 55 
56-- SIMT编程场景:56+- SIMT编程场景:
57 57 
58 ```58 ```
59- __global__ __launch_bounds__(256) void KernelJ0(float* dst, float* x)59+ __global__ __launch_bounds__(256) void compute_j0f(float *result, const float *x, uint32_t count)
60 {60 {
61- int idx = threadIdx.x + blockIdx.x * blockDim.x;61+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
62- dst[idx] = j0f(x[idx]);62+ if (idx >= count) {
63+ return;
64+ }
65+ result[idx] = j0f(x[idx]);
63 }66 }
64 ```67 ```
65 68 
66-- SIMD与SIMT混合编程场景:69+- SIMD与SIMT混合编程场景:
67 70 
68 ```71 ```
69- __simt_vf__ __launch_bounds__(256) inline void KernelJ0(__gm__ float* dst, __gm__ float* x)72+ __simt_vf__ __launch_bounds__(256) inline void compute_j0f_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
70 {73 {
71- int idx = threadIdx.x + blockIdx.x * blockDim.x;74+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
72- dst[idx] = j0f(x[idx]);75+ if (idx >= count) {
76+ return;
77+ }
78+ result[idx] = j0f(x[idx]);
79+ }
80+ 
81+ __global__ __vector__ void run_j0f(__gm__ float *result, __gm__ const float *x, uint32_t count)
82+ {
83+ asc_vf_call<compute_j0f_vf>(dim3(256), result, x, count);
73 }84 }
74 ```85 ```
75 86 
87+输入输出示例如下:
88+ 
89+```
90+x:0.25, 0.75, 1.25, 1.75
91+result: 0.9844361 0.8642423 0.6459062 0.3690326
92+```
@@ -54,23 +54,40 @@ inline float j1f(float x)
54 54 
55## 调用示例55## 调用示例
56 56 
57-- SIMT编程场景:57+- SIMT编程场景:
58 58 
59 ```59 ```
60- __global__ __launch_bounds__(256) void KernelJ1(float* dst, float* x)60+ __global__ __launch_bounds__(256) void compute_j1f(float *result, const float *x, uint32_t count)
61 {61 {
62- int idx = threadIdx.x + blockIdx.x * blockDim.x;62+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
63- dst[idx] = j1f(x[idx]);63+ if (idx >= count) {
64+ return;
65+ }
66+ result[idx] = j1f(x[idx]);
64 }67 }
65 ```68 ```
66 69 
67-- SIMD与SIMT混合编程场景:70+- SIMD与SIMT混合编程场景:
68 71 
69 ```72 ```
70- __simt_vf__ __launch_bounds__(256) inline void KernelJ1(__gm__ float* dst, __gm__ float* x)73+ __simt_vf__ __launch_bounds__(256) inline void compute_j1f_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
71 {74 {
72- int idx = threadIdx.x + blockIdx.x * blockDim.x;75+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
73- dst[idx] = j1f(x[idx]);76+ if (idx >= count) {
77+ return;
78+ }
79+ result[idx] = j1f(x[idx]);
80+ }
81+ 
82+ __global__ __vector__ void run_j1f(__gm__ float *result, __gm__ const float *x, uint32_t count)
83+ {
84+ asc_vf_call<compute_j1f_vf>(dim3(256), result, x, count);
74 }85 }
75 ```86 ```
76 87 
88+输入输出示例如下:
89+ 
90+```
91+x:0.25, 0.75, 1.25, 1.75
92+result: 0.124026 0.3492436 0.5106233 0.5801562
93+```
@@ -56,23 +56,41 @@ inline float jnf(int n, float x)
56 56 
57## 调用示例57## 调用示例
58 58 
59-- SIMT编程场景:59+- SIMT编程场景:
60 60 
61 ```61 ```
62- __global__ __launch_bounds__(256) void KernelJn(float* dst, int* n, float* x)62+ __global__ __launch_bounds__(256) void compute_jnf(float *result, const int *n, const float *x, uint32_t count)
63 {63 {
64- int idx = threadIdx.x + blockIdx.x * blockDim.x;64+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
65- dst[idx] = jnf(n[idx], x[idx]);65+ if (idx >= count) {
66+ return;
67+ }
68+ result[idx] = jnf(n[idx], x[idx]);
66 }69 }
67 ```70 ```
68 71 
69-- SIMD与SIMT混合编程场景:72+- SIMD与SIMT混合编程场景:
70 73 
71 ```74 ```
72- __simt_vf__ __launch_bounds__(256) inline void KernelJn(__gm__ float* dst, __gm__ int* n, __gm__ float* x)75+ __simt_vf__ __launch_bounds__(256) inline void compute_jnf_vf(__gm__ float *result, __gm__ const int *n, __gm__ const float *x, uint32_t count)
73 {76 {
74- int idx = threadIdx.x + blockIdx.x * blockDim.x;77+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
75- dst[idx] = jnf(n[idx], x[idx]);78+ if (idx >= count) {
79+ return;
80+ }
81+ result[idx] = jnf(n[idx], x[idx]);
82+ }
83+ 
84+ __global__ __vector__ void run_jnf(__gm__ float *result, __gm__ const int *n, __gm__ const float *x, uint32_t count)
85+ {
86+ asc_vf_call<compute_jnf_vf>(dim3(256), result, n, x, count);
76 }87 }
77 ```88 ```
78 89 
90+输入输出示例如下:
91+ 
92+```
93+n:1, 2, 3, 1
94+x:0.25, 0.75, 1.25, 1.75
95+result: 0.124026 0.06707399 0.03686836 0.5801562
96+```
@@ -55,23 +55,41 @@ inline float ldexpf(float x, int exp)
55 55 
56## 调用示例56## 调用示例
57 57 
58-- SIMT编程场景:58+- SIMT编程场景:
59 59 
60 ```60 ```
61- __global__ __launch_bounds__(1024) void KernelLdexp(float* dst, float* x, int* exp)61+ __global__ __launch_bounds__(256) void compute_ldexpf(float *result, const int *n, const float *x, uint32_t count)
62 {62 {
63- int idx = threadIdx.x + blockIdx.x * blockDim.x;63+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
64- dst[idx] = ldexpf(x[idx], exp[idx]);64+ if (idx >= count) {
65+ return;
66+ }
67+ result[idx] = ldexpf(x[idx], n[idx]);
65 }68 }
66 ```69 ```
67 70 
68-- SIMD与SIMT混合编程场景:71+- SIMD与SIMT混合编程场景:
69 72 
70 ```73 ```
71- __simt_vf__ __launch_bounds__(1024) inline void KernelLdexp(__gm__ float* dst, __gm__ float* x, __gm__ int* exp)74+ __simt_vf__ __launch_bounds__(256) inline void compute_ldexpf_vf(__gm__ float *result, __gm__ const int *n, __gm__ const float *x, uint32_t count)
72 {75 {
73- int idx = threadIdx.x + blockIdx.x * blockDim.x;76+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
74- dst[idx] = ldexpf(x[idx], exp[idx]);77+ if (idx >= count) {
78+ return;
79+ }
80+ result[idx] = ldexpf(x[idx], n[idx]);
81+ }
82+ 
83+ __global__ __vector__ void run_ldexpf(__gm__ float *result, __gm__ const int *n, __gm__ const float *x, uint32_t count)
84+ {
85+ asc_vf_call<compute_ldexpf_vf>(dim3(256), result, n, x, count);
75 }86 }
76 ```87 ```
77 88 
89+输入输出示例如下:
90+ 
91+```
92+n:1, 2, 3, 1
93+x:0.25, 0.75, 1.25, 1.75
94+result: 0.5 3 10 3.5
95+```
@@ -57,23 +57,40 @@ inline float lgammaf(float x)
57 57 
58## 调用示例58## 调用示例
59 59 
60-- SIMT编程场景:60+- SIMT编程场景:
61 61 
62 ```62 ```
63- __global__ __launch_bounds__(1024) void KernelLgamma(float* dst, float* x)63+ __global__ __launch_bounds__(256) void compute_lgammaf(float *result, const float *x, uint32_t count)
64 {64 {
65- int idx = threadIdx.x + blockIdx.x * blockDim.x;65+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
66- dst[idx] = lgammaf(x[idx]);66+ if (idx >= count) {
67+ return;
68+ }
69+ result[idx] = lgammaf(x[idx]);
67 }70 }
68 ```71 ```
69 72 
70-- SIMD与SIMT混合编程场景:73+- SIMD与SIMT混合编程场景:
71 74 
72 ```75 ```
73- __simt_vf__ __launch_bounds__(1024) inline void KernelLgamma(__gm__ float* dst, __gm__ float* x)76+ __simt_vf__ __launch_bounds__(256) inline void compute_lgammaf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
74 {77 {
75- int idx = threadIdx.x + blockIdx.x * blockDim.x;78+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
76- dst[idx] = lgammaf(x[idx]);79+ if (idx >= count) {
80+ return;
81+ }
82+ result[idx] = lgammaf(x[idx]);
83+ }
84+ 
85+ __global__ __vector__ void run_lgammaf(__gm__ float *result, __gm__ const float *x, uint32_t count)
86+ {
87+ asc_vf_call<compute_lgammaf_vf>(dim3(256), result, x, count);
77 }88 }
78 ```89 ```
79 90 
91+输入输出示例如下:
92+ 
93+```
94+x:0.25, 0.75, 1.25, 1.75
95+result: 1.288023 0.203281 -0.09827184 -0.08440112
96+```
@@ -53,23 +53,40 @@ inline long long int llrintf(float x)
53 53 
54## 调用示例54## 调用示例
55 55 
56-- SIMT编程场景:56+- SIMT编程场景:
57 57 
58 ```58 ```
59- __global__ __launch_bounds__(1024) void KernelRint(long long int* dst, float* x)59+ __global__ __launch_bounds__(256) void compute_llrintf(float *result, const float *x, uint32_t count)
60 {60 {
61- int idx = threadIdx.x + blockIdx.x * blockDim.x;61+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
62- dst[idx] = llrintf(x[idx]);62+ if (idx >= count) {
63+ return;
64+ }
65+ result[idx] = static_cast<float>(llrintf(x[idx]));
63 }66 }
64 ```67 ```
65 68 
66-- SIMD与SIMT混合编程场景:69+- SIMD与SIMT混合编程场景:
67 70 
68 ```71 ```
69- __simt_vf__ __launch_bounds__(1024) inline void KernelRint(__gm__ long long int* dst, __gm__ float* x)72+ __simt_vf__ __launch_bounds__(256) inline void compute_llrintf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
70 {73 {
71- int idx = threadIdx.x + blockIdx.x * blockDim.x;74+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
72- dst[idx] = llrintf(x[idx]);75+ if (idx >= count) {
76+ return;
77+ }
78+ result[idx] = static_cast<float>(llrintf(x[idx]));
79+ }
80+ 
81+ __global__ __vector__ void run_llrintf(__gm__ float *result, __gm__ const float *x, uint32_t count)
82+ {
83+ asc_vf_call<compute_llrintf_vf>(dim3(256), result, x, count);
73 }84 }
74 ```85 ```
75 86 
87+输入输出示例如下:
88+ 
89+```
90+x:0.25, 0.75, 1.25, 1.75
91+result: 0 1 1 2
92+```
@@ -50,23 +50,40 @@ inline long long int llroundf(float x)
50 50 
51## 调用示例51## 调用示例
52 52 
53-- SIMT编程场景:53+- SIMT编程场景:
54 54 
55 ```55 ```
56- __global__ __launch_bounds__(1024) void KernelRound(long long int* dst, float* x)56+ __global__ __launch_bounds__(256) void compute_llroundf(float *result, const float *x, uint32_t count)
57 {57 {
58- int idx = threadIdx.x + blockIdx.x * blockDim.x;58+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
59- dst[idx] = llroundf(x[idx]);59+ if (idx >= count) {
60+ return;
61+ }
62+ result[idx] = static_cast<float>(llroundf(x[idx]));
60 }63 }
61 ```64 ```
62 65 
63-- SIMD与SIMT混合编程场景:66+- SIMD与SIMT混合编程场景:
64 67 
65 ```68 ```
66- __simt_vf__ __launch_bounds__(1024) inline void KernelRound(__gm__ long long int* dst, __gm__ float* x)69+ __simt_vf__ __launch_bounds__(256) inline void compute_llroundf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
67 {70 {
68- int idx = threadIdx.x + blockIdx.x * blockDim.x;71+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
69- dst[idx] = llroundf(x[idx]);72+ if (idx >= count) {
73+ return;
74+ }
75+ result[idx] = static_cast<float>(llroundf(x[idx]));
76+ }
77+ 
78+ __global__ __vector__ void run_llroundf(__gm__ float *result, __gm__ const float *x, uint32_t count)
79+ {
80+ asc_vf_call<compute_llroundf_vf>(dim3(256), result, x, count);
70 }81 }
71 ```82 ```
72 83 
84+输入输出示例如下:
85+ 
86+```
87+x:0.25, 0.75, 1.25, 1.75
88+result: 0 1 1 2
89+```
@@ -54,23 +54,40 @@ inline float log10f(float x)
54 54 
55## 调用示例55## 调用示例
56 56 
57-- SIMT编程场景:57+- SIMT编程场景:
58 58 
59 ```59 ```
60- __global__ __launch_bounds__(1024) void KernelLog10(float* dst, float* x)60+ __global__ __launch_bounds__(256) void compute_log10f(float *result, const float *x, uint32_t count)
61 {61 {
62- int idx = threadIdx.x + blockIdx.x * blockDim.x;62+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
63- dst[idx] = log10f(x[idx]);63+ if (idx >= count) {
64+ return;
65+ }
66+ result[idx] = log10f(x[idx]);
64 }67 }
65 ```68 ```
66 69 
67-- SIMD与SIMT混合编程场景:70+- SIMD与SIMT混合编程场景:
68 71 
69 ```72 ```
70- __simt_vf__ __launch_bounds__(1024) inline void KernelLog10(__gm__ float* dst, __gm__ float* x)73+ __simt_vf__ __launch_bounds__(256) inline void compute_log10f_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
71 {74 {
72- int idx = threadIdx.x + blockIdx.x * blockDim.x;75+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
73- dst[idx] = log10f(x[idx]);76+ if (idx >= count) {
77+ return;
78+ }
79+ result[idx] = log10f(x[idx]);
80+ }
81+ 
82+ __global__ __vector__ void run_log10f(__gm__ float *result, __gm__ const float *x, uint32_t count)
83+ {
84+ asc_vf_call<compute_log10f_vf>(dim3(256), result, x, count);
74 }85 }
75 ```86 ```
76 87 
88+输入输出示例如下:
89+ 
90+```
91+x:1, 2, 3, 4
92+result: 0 0.30103 0.4771213 0.60206
93+```
@@ -54,23 +54,40 @@ inline float log1pf(float x)
54 54 
55## 调用示例55## 调用示例
56 56 
57-- SIMT编程场景:57+- SIMT编程场景:
58 58 
59 ```59 ```
60- __global__ __launch_bounds__(1024) void KernelLog1p(float* dst, float* x)60+ __global__ __launch_bounds__(256) void compute_log1pf(float *result, const float *x, uint32_t count)
61 {61 {
62- int idx = threadIdx.x + blockIdx.x * blockDim.x;62+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
63- dst[idx] = log1pf(x[idx]);63+ if (idx >= count) {
64+ return;
65+ }
66+ result[idx] = log1pf(x[idx]);
64 }67 }
65 ```68 ```
66 69 
67-- SIMD与SIMT混合编程场景:70+- SIMD与SIMT混合编程场景:
68 71 
69 ```72 ```
70- __simt_vf__ __launch_bounds__(1024) inline void KernelLog1p(__gm__ float* dst, __gm__ float* x)73+ __simt_vf__ __launch_bounds__(256) inline void compute_log1pf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
71 {74 {
72- int idx = threadIdx.x + blockIdx.x * blockDim.x;75+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
73- dst[idx] = log1pf(x[idx]);76+ if (idx >= count) {
77+ return;
78+ }
79+ result[idx] = log1pf(x[idx]);
80+ }
81+ 
82+ __global__ __vector__ void run_log1pf(__gm__ float *result, __gm__ const float *x, uint32_t count)
83+ {
84+ asc_vf_call<compute_log1pf_vf>(dim3(256), result, x, count);
74 }85 }
75 ```86 ```
76 87 
88+输入输出示例如下:
89+ 
90+```
91+x:0.25, 1.25, 2.25, 3.25
92+result: 0.2231435 0.8109302 1.178655 1.446919
93+```
@@ -54,22 +54,40 @@ inline float log2f(float x)
54 54 
55## 调用示例55## 调用示例
56 56 
57-- SIMT编程场景:57+- SIMT编程场景:
58 58 
59 ```59 ```
60- __global__ __launch_bounds__(1024) void KernelLog2(float* dst, float* x)60+ __global__ __launch_bounds__(256) void compute_log2f(float *result, const float *x, uint32_t count)
61 {61 {
62- int idx = threadIdx.x + blockIdx.x * blockDim.x;62+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
63- dst[idx] = log2f(x[idx]);63+ if (idx >= count) {
64+ return;
65+ }
66+ result[idx] = log2f(x[idx]);
64 }67 }
65 ```68 ```
66 69 
67-- SIMD与SIMT混合编程场景:70+- SIMD与SIMT混合编程场景:
68 71 
69 ```72 ```
70- __simt_vf__ __launch_bounds__(1024) inline void KernelLog2(__gm__ float* dst, __gm__ float* x)73+ __simt_vf__ __launch_bounds__(256) inline void compute_log2f_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
71 {74 {
72- int idx = threadIdx.x + blockIdx.x * blockDim.x;75+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
73- dst[idx] = log2f(x[idx]);76+ if (idx >= count) {
77+ return;
78+ }
79+ result[idx] = log2f(x[idx]);
80+ }
81+ 
82+ __global__ __vector__ void run_log2f(__gm__ float *result, __gm__ const float *x, uint32_t count)
83+ {
84+ asc_vf_call<compute_log2f_vf>(dim3(256), result, x, count);
74 }85 }
75 ```86 ```
87+ 
88+输入输出示例如下:
89+ 
90+```
91+x:1, 2, 3, 4
92+result: 0 1 1.584962 2
93+```
@@ -55,23 +55,40 @@ inline float logbf(float x)
55 55 
56## 调用示例56## 调用示例
57 57 
58-- SIMT编程场景:58+- SIMT编程场景:
59 59 
60 ```60 ```
61- __global__ __launch_bounds__(1024) void KernelLogb(float* dst, float* x)61+ __global__ __launch_bounds__(256) void compute_logbf(float *result, const float *x, uint32_t count)
62 {62 {
63- int idx = threadIdx.x + blockIdx.x * blockDim.x;63+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
64- dst[idx] = logbf(x[idx]);64+ if (idx >= count) {
65+ return;
66+ }
67+ result[idx] = logbf(x[idx]);
65 }68 }
66 ```69 ```
67 70 
68-- SIMD与SIMT混合编程场景:71+- SIMD与SIMT混合编程场景:
69 72 
70 ```73 ```
71- __simt_vf__ __launch_bounds__(1024) inline void KernelLogb(__gm__ float* dst, __gm__ float* x)74+ __simt_vf__ __launch_bounds__(256) inline void compute_logbf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
72 {75 {
73- int idx = threadIdx.x + blockIdx.x * blockDim.x;76+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
74- dst[idx] = logbf(x[idx]);77+ if (idx >= count) {
78+ return;
79+ }
80+ result[idx] = logbf(x[idx]);
81+ }
82+ 
83+ __global__ __vector__ void run_logbf(__gm__ float *result, __gm__ const float *x, uint32_t count)
84+ {
85+ asc_vf_call<compute_logbf_vf>(dim3(256), result, x, count);
75 }86 }
76 ```87 ```
77 88 
89+输入输出示例如下:
90+ 
91+```
92+x:1, 2, 3, 4
93+result: 0 1 1 2
94+```
@@ -54,23 +54,40 @@ inline float logf(float x)
54 54 
55## 调用示例55## 调用示例
56 56 
57-- SIMT编程场景:57+- SIMT编程场景:
58 58 
59 ```59 ```
60- __launch_bounds__(1024) void KernelLog(float* dst, float* x)60+ __global__ __launch_bounds__(256) void compute_logf(float *result, const float *x, uint32_t count)
61 {61 {
62- int idx = threadIdx.x + blockIdx.x * blockDim.x;62+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
63- dst[idx] = logf(x[idx]);63+ if (idx >= count) {
64+ return;
65+ }
66+ result[idx] = logf(x[idx]);
64 }67 }
65 ```68 ```
66 69 
67-- SIMD与SIMT混合编程场景:70+- SIMD与SIMT混合编程场景:
68 71 
69 ```72 ```
70- __simt_vf__ __launch_bounds__(1024) inline void KernelLog(__gm__ float* dst, __gm__ float* x)73+ __simt_vf__ __launch_bounds__(256) inline void compute_logf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
71 {74 {
72- int idx = threadIdx.x + blockIdx.x * blockDim.x;75+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
73- dst[idx] = logf(x[idx]);76+ if (idx >= count) {
77+ return;
78+ }
79+ result[idx] = logf(x[idx]);
80+ }
81+ 
82+ __global__ __vector__ void run_logf(__gm__ float *result, __gm__ const float *x, uint32_t count)
83+ {
84+ asc_vf_call<compute_logf_vf>(dim3(256), result, x, count);
74 }85 }
75 ```86 ```
76 87 
88+输入输出示例如下:
89+ 
90+```
91+x:1, 2, 3, 4
92+result: 0 0.6931472 1.098612 1.386294
93+```
@@ -53,23 +53,40 @@ inline long int lrintf(float x)
53 53 
54## 调用示例54## 调用示例
55 55 
56-- SIMT编程场景:56+- SIMT编程场景:
57 57 
58 ```58 ```
59- __global__ __launch_bounds__(1024) void KernelRint(long int* dst, float* x)59+ __global__ __launch_bounds__(256) void compute_lrintf(float *result, const float *x, uint32_t count)
60 {60 {
61- int idx = threadIdx.x + blockIdx.x * blockDim.x;61+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
62- dst[idx] = lrintf(x[idx]);62+ if (idx >= count) {
63+ return;
64+ }
65+ result[idx] = static_cast<float>(lrintf(x[idx]));
63 }66 }
64 ```67 ```
65 68 
66-- SIMD与SIMT混合编程场景:69+- SIMD与SIMT混合编程场景:
67 70 
68 ```71 ```
69- __simt_vf__ __launch_bounds__(1024) inline void KernelRint(__gm__ long int* dst, __gm__ float* x)72+ __simt_vf__ __launch_bounds__(256) inline void compute_lrintf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
70 {73 {
71- int idx = threadIdx.x + blockIdx.x * blockDim.x;74+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
72- dst[idx] = lrintf(x[idx]);75+ if (idx >= count) {
76+ return;
77+ }
78+ result[idx] = static_cast<float>(lrintf(x[idx]));
79+ }
80+ 
81+ __global__ __vector__ void run_lrintf(__gm__ float *result, __gm__ const float *x, uint32_t count)
82+ {
83+ asc_vf_call<compute_lrintf_vf>(dim3(256), result, x, count);
73 }84 }
74 ```85 ```
75 86 
87+输入输出示例如下:
88+ 
89+```
90+x:0.25, 0.75, 1.25, 1.75
91+result: 0 1 1 2
92+```
@@ -50,23 +50,40 @@ inline long int lroundf(float x)
50 50 
51## 调用示例51## 调用示例
52 52 
53-- SIMT编程场景:53+- SIMT编程场景:
54 54 
55 ```55 ```
56- __global__ __launch_bounds__(1024) void KernelRound(long int* dst, float* x)56+ __global__ __launch_bounds__(256) void compute_lroundf(float *result, const float *x, uint32_t count)
57 {57 {
58- int idx = threadIdx.x + blockIdx.x * blockDim.x;58+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
59- dst[idx] = lroundf(x[idx]);59+ if (idx >= count) {
60+ return;
61+ }
62+ result[idx] = static_cast<float>(lroundf(x[idx]));
60 }63 }
61 ```64 ```
62 65 
63-- SIMD与SIMT混合编程场景:66+- SIMD与SIMT混合编程场景:
64 67 
65 ```68 ```
66- __simt_vf__ __launch_bounds__(1024) inline void KernelRound(__gm__ long int* dst, __gm__ float* x)69+ __simt_vf__ __launch_bounds__(256) inline void compute_lroundf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
67 {70 {
68- int idx = threadIdx.x + blockIdx.x * blockDim.x;71+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
69- dst[idx] = lroundf(x[idx]);72+ if (idx >= count) {
73+ return;
74+ }
75+ result[idx] = static_cast<float>(lroundf(x[idx]));
76+ }
77+ 
78+ __global__ __vector__ void run_lroundf(__gm__ float *result, __gm__ const float *x, uint32_t count)
79+ {
80+ asc_vf_call<compute_lroundf_vf>(dim3(256), result, x, count);
70 }81 }
71 ```82 ```
72 83 
84+输入输出示例如下:
85+ 
86+```
87+x:0.25, 0.75, 1.25, 1.75
88+result: 0 1 1 2
89+```
@@ -53,25 +53,41 @@ inline float modff(float x, float *n)
53 53 
54## 调用示例54## 调用示例
55 55 
56-- SIMT编程场景:56+- SIMT编程场景:
57 57 
58 ```58 ```
59- __global__ __launch_bounds__(1024) void KernelModff(float* dst, float* out, float* x)59+ __global__ __launch_bounds__(256) void compute_modff(float *result, float *integer_part, const float *x, uint32_t count)
60 {60 {
61- int idx = threadIdx.x + blockIdx.x * blockDim.x;61+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
62- dst[idx] = modff(x[idx], out + idx);62+ if (idx >= count) {
63+ return;
64+ }
65+ result[idx] = modff(x[idx], integer_part + idx);
63 }66 }
64 ```67 ```
65 68 
66-- SIMD与SIMT混合编程场景:69+- SIMD与SIMT混合编程场景:
67- 
68- SIMD与SIMT混合编程场景,需要显式使用地址空间限定符表示地址空间:\_\_gm\_\_表示Global Memory内存空间,\_\_ubuf\_\_表示Unified Buffer内存空间,栈空间无需添加地址空间限定符。
69 70 
70 ```71 ```
71- __simt_vf__ __launch_bounds__(1024) inline void KernelModff(__gm__ float* dst, __gm__ float* out, __gm__ float* x)72+ __simt_vf__ __launch_bounds__(256) inline void compute_modff_vf(__gm__ float *result, __gm__ float *integer_part, __gm__ const float *x, uint32_t count)
72 {73 {
73- int idx = threadIdx.x + blockIdx.x * blockDim.x;74+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
74- dst[idx] = modff(x[idx], out + idx);75+ if (idx >= count) {
76+ return;
77+ }
78+ result[idx] = modff(x[idx], integer_part + idx);
79+ }
80+ 
81+ __global__ __vector__ void run_modff(__gm__ float *result, __gm__ float *integer_part, __gm__ const float *x, uint32_t count)
82+ {
83+ asc_vf_call<compute_modff_vf>(dim3(256), result, integer_part, x, count);
75 }84 }
76 ```85 ```
77 86 
87+输入输出示例如下:
88+ 
89+```
90+x:0.25, 0.75, 1.25, 1.75
91+result: 0.25 0.75 0.25 0.75
92+integer_part: 0 0 1 1
93+```
@@ -52,23 +52,40 @@ inline float nearbyintf(float x)
52 52 
53## 调用示例53## 调用示例
54 54 
55-- SIMT编程场景:55+- SIMT编程场景:
56 56 
57 ```57 ```
58- __global__ __launch_bounds__(1024) void KernelNearByInt(float* dst, float* x)58+ __global__ __launch_bounds__(256) void compute_nearbyintf(float *result, const float *x, uint32_t count)
59 {59 {
60- int idx = threadIdx.x + blockIdx.x * blockDim.x;60+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
61- dst[idx] = nearbyintf(x[idx]);61+ if (idx >= count) {
62+ return;
63+ }
64+ result[idx] = nearbyintf(x[idx]);
62 }65 }
63 ```66 ```
64 67 
65-- SIMD与SIMT混合编程场景:68+- SIMD与SIMT混合编程场景:
66 69 
67 ```70 ```
68- __simt_vf__ __launch_bounds__(1024) inline void KernelNearByInt(__gm__ float* dst, __gm__ float* x)71+ __simt_vf__ __launch_bounds__(256) inline void compute_nearbyintf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
69 {72 {
70- int idx = threadIdx.x + blockIdx.x * blockDim.x;73+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
71- dst[idx] = nearbyintf(x[idx]);74+ if (idx >= count) {
75+ return;
76+ }
77+ result[idx] = nearbyintf(x[idx]);
78+ }
79+ 
80+ __global__ __vector__ void run_nearbyintf(__gm__ float *result, __gm__ const float *x, uint32_t count)
81+ {
82+ asc_vf_call<compute_nearbyintf_vf>(dim3(256), result, x, count);
72 }83 }
73 ```84 ```
74 85 
86+输入输出示例如下:
87+ 
88+```
89+x:0.25, 0.75, 1.25, 1.75
90+result: 0 1 1 2
91+```
@@ -59,40 +59,41 @@ inline float nextafterf(float x, float y)
59 59 
60## 调用示例60## 调用示例
61 61 
62-- SIMT编程场景:62+- SIMT编程场景:
63 63 
64 ```64 ```
65- __simt_vf__ __launch_bounds__(1024) inline void KernelNextAfter(__gm__ float* dst, __gm__ float* x, __gm__ float* y)65+ __global__ __launch_bounds__(256) void compute_nextafterf(float *result, const float *x, const float *y, uint32_t count)
66 {66 {
67- int idx = threadIdx.x + blockIdx.x * blockDim.x;67+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
68- dst[idx] = nextafterf(x[idx], y[idx]);68+ if (idx >= count) {
69+ return;
70+ }
71+ result[idx] = nextafterf(x[idx], y[idx]);
69 }72 }
70 ```73 ```
71 74 
72-- SIMD与SIMT混合编程场景:75+- SIMD与SIMT混合编程场景:
73 76 
74 ```77 ```
75- __simt_vf__ __launch_bounds__(1024) inline void KernelNextAfter(__gm__ float* dst, __gm__ float* x, __gm__ float* y)78+ __simt_vf__ __launch_bounds__(256) inline void compute_nextafterf_vf(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, uint32_t count)
76 {79 {
77- int idx = threadIdx.x + blockIdx.x * blockDim.x;80+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
78- dst[idx] = nextafterf(x[idx], y[idx]);81+ if (idx >= count) {
79- printf("x[%d] = %f, y[%d] = %f, dst[%d] = %f, 输入x的16进制表示为: 0x%x, 输出dst的16进制表示为: 0x%x\n", idx, x[idx], idx, y[idx], idx, dst[idx], x[idx], dst[idx]);82+ return;
83+ }
84+ result[idx] = nextafterf(x[idx], y[idx]);
85+ }
86+ 
87+ __global__ __vector__ void run_nextafterf(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, uint32_t count)
88+ {
89+ asc_vf_call<compute_nextafterf_vf>(dim3(256), result, x, y, count);
80 }90 }
81 ```91 ```
82 92 
83- 程序输入为:93+输入输出示例如下:
84- 
85- ```
86- x:[10.0f 10.0f -10.0f -10.0f ...]
87- y:[20.0f -20.0f 20.0f -20.0f ...]
88- ```
89- 
90- 程序前4个线程的运行时打印效果如下:
91- 
92- ```
93- x[0] = 10.000000, y[0] = 20.000000, dst[0] = 10.000001, 输入x的16进制表示为: 0x41200000, 输出dst的16进制表示为: 0x41200001
94- x[1] = 10.000000, y[1] = -20.000000, dst[1] = 9.999999, 输入x的16进制表示为: 0x41200000, 输出dst的16进制表示为: 0x411fffff
95- x[2] = -10.000000, y[2] = 20.000000, dst[2] = -9.999999, 输入x的16进制表示为: 0xc1200000, 输出dst的16进制表示为: 0xc11fffff
96- x[3] = -10.000000, y[3] = -20.000000, dst[3] = -10.000001, 输入x的16进制表示为: 0xc1200000, 输出dst的16进制表示为: 0xc1200001
97- ```
98 94 
95+```
96+x:0.25, 0.75, 1.25, 1.75
97+y:1.5, 2.5, 3.5, 4.5
98+result: 0.25 0.7500001 1.25 1.75
99+```
@@ -58,23 +58,42 @@ a^2 + b^2 + c^2的平方根。
58 58 
59## 调用示例59## 调用示例
60 60 
61-- SIMT编程场景:61+- SIMT编程场景:
62 62 
63 ```63 ```
64- __global__ __launch_bounds__(1024) void KernelNorm3d(float* dst, float* a, float* b, float* c)64+ __global__ __launch_bounds__(256) void compute_norm3df(float *result, const float *x, const float *y, const float *z, uint32_t count)
65 {65 {
66- int idx = threadIdx.x + blockIdx.x * blockDim.x;66+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
67- dst[idx] = norm3df(a[idx], b[idx], c[idx]);67+ if (idx >= count) {
68+ return;
69+ }
70+ result[idx] = norm3df(x[idx], y[idx], z[idx]);
68 }71 }
69 ```72 ```
70 73 
71-- SIMD与SIMT混合编程场景:74+- SIMD与SIMT混合编程场景:
72 75 
73 ```76 ```
74- __simt_vf__ __launch_bounds__(1024) inline void KernelNorm3d(__gm__ float* dst, __gm__ float* a, __gm__ float* b, __gm__ float* c)77+ __simt_vf__ __launch_bounds__(256) inline void compute_norm3df_vf(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, __gm__ const float *z, uint32_t count)
75 {78 {
76- int idx = threadIdx.x + blockIdx.x * blockDim.x;79+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
77- dst[idx] = norm3df(a[idx], b[idx], c[idx]);80+ if (idx >= count) {
81+ return;
82+ }
83+ result[idx] = norm3df(x[idx], y[idx], z[idx]);
84+ }
85+ 
86+ __global__ __vector__ void run_norm3df(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, __gm__ const float *z, uint32_t count)
87+ {
88+ asc_vf_call<compute_norm3df_vf>(dim3(256), result, x, y, z, count);
78 }89 }
79 ```90 ```
80 91 
92+输入输出示例如下:
93+ 
94+```
95+x:0.25, 0.75, 1.25, 1.75
96+y:1.5, 2.5, 3.5, 4.5
97+z:-0.5, 0.5, 1.5, 2.5
98+result: 1.600781 2.657537 4.007805 5.437141
99+```
@@ -57,23 +57,43 @@ a^2 + b^2 + c^2 + d^2的平方根。
57 57 
58## 调用示例58## 调用示例
59 59 
60-- SIMT编程场景:60+- SIMT编程场景:
61 61 
62 ```62 ```
63- __global__ __launch_bounds__(1024) void KernelNorm4d(float* dst, float* a, float* b, float* c, float* d)63+ __global__ __launch_bounds__(256) void compute_norm4df(float *result, const float *x, const float *y, const float *z, const float *w, uint32_t count)
64 {64 {
65- int idx = threadIdx.x + blockIdx.x * blockDim.x;65+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
66- dst[idx] = norm4df(a[idx], b[idx], c[idx], d[idx]);66+ if (idx >= count) {
67+ return;
68+ }
69+ result[idx] = norm4df(x[idx], y[idx], z[idx], w[idx]);
67 }70 }
68 ```71 ```
69 72 
70-- SIMD与SIMT混合编程场景:73+- SIMD与SIMT混合编程场景:
71 74 
72 ```75 ```
73- __simt_vf__ __launch_bounds__(1024) inline void KernelNorm4d(__gm__ float* dst, __gm__ float* a, __gm__ float* b, __gm__ float* c, __gm__ float* d)76+ __simt_vf__ __launch_bounds__(256) inline void compute_norm4df_vf(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, __gm__ const float *z, __gm__ const float *w, uint32_t count)
74 {77 {
75- int idx = threadIdx.x + blockIdx.x * blockDim.x;78+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
76- dst[idx] = norm4df(a[idx], b[idx], c[idx], d[idx]);79+ if (idx >= count) {
80+ return;
81+ }
82+ result[idx] = norm4df(x[idx], y[idx], z[idx], w[idx]);
83+ }
84+ 
85+ __global__ __vector__ void run_norm4df(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, __gm__ const float *z, __gm__ const float *w, uint32_t count)
86+ {
87+ asc_vf_call<compute_norm4df_vf>(dim3(256), result, x, y, z, w, count);
77 }88 }
78 ```89 ```
79 90 
91+输入输出示例如下:
92+ 
93+```
94+x:0.25, 0.75, 1.25, 1.75
95+y:1.5, 2.5, 3.5, 4.5
96+z:-0.5, 0.5, 1.5, 2.5
97+w:2, 3, 4, 5
98+result: 2.561738 4.007805 5.662376 7.386643
99+```
@@ -54,23 +54,40 @@ inline float normcdff(float x)
54 54 
55## 调用示例55## 调用示例
56 56 
57-- SIMT编程场景:57+- SIMT编程场景:
58 58 
59 ```59 ```
60- __global__ __launch_bounds__(1024) void KernelNormcdf(float* dst, float* x)60+ __global__ __launch_bounds__(256) void compute_normcdff(float *result, const float *x, uint32_t count)
61 {61 {
62- int idx = threadIdx.x + blockIdx.x * blockDim.x;62+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
63- dst[idx] = normcdff(x[idx]);63+ if (idx >= count) {
64+ return;
65+ }
66+ result[idx] = normcdff(x[idx]);
64 }67 }
65 ```68 ```
66 69 
67-- SIMD与SIMT混合编程场景:70+- SIMD与SIMT混合编程场景:
68 71 
69 ```72 ```
70- __simt_vf__ __launch_bounds__(1024) inline void KernelNormcdf(__gm__ float* dst, __gm__ float* x)73+ __simt_vf__ __launch_bounds__(256) inline void compute_normcdff_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
71 {74 {
72- int idx = threadIdx.x + blockIdx.x * blockDim.x;75+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
73- dst[idx] = normcdff(x[idx]);76+ if (idx >= count) {
77+ return;
78+ }
79+ result[idx] = normcdff(x[idx]);
80+ }
81+ 
82+ __global__ __vector__ void run_normcdff(__gm__ float *result, __gm__ const float *x, uint32_t count)
83+ {
84+ asc_vf_call<compute_normcdff_vf>(dim3(256), result, x, count);
74 }85 }
75 ```86 ```
76 87 
88+输入输出示例如下:
89+ 
90+```
91+x:0.25, 0.75, 1.25, 1.75
92+result: 0.5987064 0.7733727 0.8943502 0.9599409
93+```
@@ -55,23 +55,40 @@ inline float normcdfinvf(float x)
55 55 
56## 调用示例56## 调用示例
57 57 
58-- SIMT编程场景:58+- SIMT编程场景:
59 59 
60 ```60 ```
61- __global__ __launch_bounds__(1024) void KernelNormcdfinv(float* dst, float* x)61+ __global__ __launch_bounds__(256) void compute_normcdfinvf(float *result, const float *x, uint32_t count)
62 {62 {
63- int idx = threadIdx.x + blockIdx.x * blockDim.x;63+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
64- dst[idx] = normcdfinvf(x[idx]);64+ if (idx >= count) {
65+ return;
66+ }
67+ result[idx] = normcdfinvf(x[idx]);
65 }68 }
66 ```69 ```
67 70 
68-- SIMD与SIMT混合编程场景:71+- SIMD与SIMT混合编程场景:
69 72 
70 ```73 ```
71- __simt_vf__ __launch_bounds__(1024) inline void KernelNormcdfinv(__gm__ float* dst, __gm__ float* x)74+ __simt_vf__ __launch_bounds__(256) inline void compute_normcdfinvf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
72 {75 {
73- int idx = threadIdx.x + blockIdx.x * blockDim.x;76+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
74- dst[idx] = normcdfinvf(x[idx]);77+ if (idx >= count) {
78+ return;
79+ }
80+ result[idx] = normcdfinvf(x[idx]);
81+ }
82+ 
83+ __global__ __vector__ void run_normcdfinvf(__gm__ float *result, __gm__ const float *x, uint32_t count)
84+ {
85+ asc_vf_call<compute_normcdfinvf_vf>(dim3(256), result, x, count);
75 }86 }
76 ```87 ```
77 88 
89+输入输出示例如下:
90+ 
91+```
92+x:-0.75, -0.25, 0.25, 0.75
93+result: -0.6744897 0.6744897 nan nan
94+```
@@ -57,25 +57,41 @@ a\[0\]^2 + a\[1\]^2 + ...+ a\[n-1\]^2的平方根。
57 57 
58## 调用示例58## 调用示例
59 59 
60-- SIMT编程场景:60+- SIMT编程场景:
61 61 
62 ```62 ```
63- __global__ __launch_bounds__(1024) void KernelNorm(float* dst, int* n, float* a)63+ __global__ __launch_bounds__(256) void compute_normf(float *result, const int *n, float *vector_data, uint32_t count)
64 {64 {
65- int idx = threadIdx.x + blockIdx.x * blockDim.x;65+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
66- dst[idx] = normf(n[idx], a);66+ if (idx >= count) {
67+ return;
68+ }
69+ result[idx] = normf(n[idx], vector_data + idx * 4);
67 }70 }
68 ```71 ```
69 72 
70-- SIMD与SIMT混合编程场景:73+- SIMD与SIMT混合编程场景:
71- 
72- SIMD与SIMT混合编程场景,需要显式使用地址空间限定符表示地址空间:\_\_gm\_\_表示Global Memory内存空间,\_\_ubuf\_\_表示Unified Buffer内存空间,栈空间无需添加地址空间限定符。
73 74 
74 ```75 ```
75- __simt_vf__ __launch_bounds__(1024) inline void KernelNorm(__gm__ float* dst, __gm__ int* n, __gm__ float* a)76+ __simt_vf__ __launch_bounds__(256) inline void compute_normf_vf(__gm__ float *result, __gm__ const int *n, __gm__ float *vector_data, uint32_t count)
76 {77 {
77- int idx = threadIdx.x + blockIdx.x * blockDim.x;78+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
78- dst[idx] = normf(n[idx], a);79+ if (idx >= count) {
80+ return;
81+ }
82+ result[idx] = normf(n[idx], vector_data + idx * 4);
83+ }
84+ 
85+ __global__ __vector__ void run_normf(__gm__ float *result, __gm__ const int *n, __gm__ float *vector_data, uint32_t count)
86+ {
87+ asc_vf_call<compute_normf_vf>(dim3(256), result, n, vector_data, count);
79 }88 }
80 ```89 ```
81 90 
91+输入输出示例如下:
92+ 
93+```
94+n:1, 2, 3, 1
95+vector_data:[[1, 2, 3, 4], [2, 3, 4, 5], [3, 4, 5, 6], [4, 5, 6, 7]]
96+result: 1 3.605551 7.071068 4
97+```
@@ -75,23 +75,41 @@ x的y次幂的结果。
75 75 
76## 调用示例76## 调用示例
77 77 
78-- SIMT编程场景:78+- SIMT编程场景:
79 79 
80 ```80 ```
81- __global__ __launch_bounds__(1024) void KernelPow(float* dst, float* x, float* y)81+ __global__ __launch_bounds__(256) void compute_powf(float *result, const float *x, const float *y, uint32_t count)
82 {82 {
83- int idx = threadIdx.x + blockIdx.x * blockDim.x;83+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
84- dst[idx] = powf(x[idx], y[idx]);84+ if (idx >= count) {
85+ return;
86+ }
87+ result[idx] = powf(x[idx], y[idx]);
85 }88 }
86 ```89 ```
87 90 
88-- SIMD与SIMT混合编程场景:91+- SIMD与SIMT混合编程场景:
89 92 
90 ```93 ```
91- __simt_vf__ __launch_bounds__(1024) inline void KernelPow(__gm__ float* dst, __gm__ float* x, __gm__ float* y)94+ __simt_vf__ __launch_bounds__(256) inline void compute_powf_vf(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, uint32_t count)
92 {95 {
93- int idx = threadIdx.x + blockIdx.x * blockDim.x;96+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
94- dst[idx] = powf(x[idx], y[idx]);97+ if (idx >= count) {
98+ return;
99+ }
100+ result[idx] = powf(x[idx], y[idx]);
101+ }
102+ 
103+ __global__ __vector__ void run_powf(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, uint32_t count)
104+ {
105+ asc_vf_call<compute_powf_vf>(dim3(256), result, x, y, count);
95 }106 }
96 ```107 ```
97 108 
109+输入输出示例如下:
110+ 
111+```
112+x:0.25, 0.75, 1.25, 1.75
113+y:1.25, 2.25, 3.25, 4.25
114+result: 0.25 0.5625 1.953125 9.378906
115+```
@@ -55,23 +55,40 @@ x的立方根的倒数。
55 55 
56## 调用示例56## 调用示例
57 57 
58-- SIMT编程场景:58+- SIMT编程场景:
59 59 
60 ```60 ```
61- __global__ __launch_bounds__(1024) void KernelRcbrt(float* x, float* y)61+ __global__ __launch_bounds__(256) void compute_rcbrtf(float *result, const float *x, uint32_t count)
62 {62 {
63- int idx = threadIdx.x + blockIdx.x * blockDim.x;63+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
64- y[idx] = rcbrtf(x[idx]);64+ if (idx >= count) {
65+ return;
66+ }
67+ result[idx] = rcbrtf(x[idx]);
65 }68 }
66 ```69 ```
67 70 
68-- SIMD与SIMT混合编程场景:71+- SIMD与SIMT混合编程场景:
69 72 
70 ```73 ```
71- __simt_vf__ __launch_bounds__(1024) inline void KernelRcbrt(__gm__ float* x, __gm__ float* y)74+ __simt_vf__ __launch_bounds__(256) inline void compute_rcbrtf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
72 {75 {
73- int idx = threadIdx.x + blockIdx.x * blockDim.x;76+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
74- y[idx] = rcbrtf(x[idx]);77+ if (idx >= count) {
78+ return;
79+ }
80+ result[idx] = rcbrtf(x[idx]);
81+ }
82+ 
83+ __global__ __vector__ void run_rcbrtf(__gm__ float *result, __gm__ const float *x, uint32_t count)
84+ {
85+ asc_vf_call<compute_rcbrtf_vf>(dim3(256), result, x, count);
75 }86 }
76 ```87 ```
77 88 
89+输入输出示例如下:
90+ 
91+```
92+x:0.25, 0.75, 1.25, 1.75
93+result: 1.587401 1.100642 0.9283177 0.8298265
94+```
@@ -54,23 +54,41 @@ inline float remainderf(float x, float y)
54 54 
55## 调用示例55## 调用示例
56 56 
57-- SIMT编程场景:57+- SIMT编程场景:
58 58 
59 ```59 ```
60- __global__ __launch_bounds__(1024) void KernelRemainder(float* dst, float* x, float* y)60+ __global__ __launch_bounds__(256) void compute_remainderf(float *result, const float *x, const float *y, uint32_t count)
61 {61 {
62- int idx = threadIdx.x + blockIdx.x * blockDim.x;62+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
63- dst[idx] = remainderf(x[idx], y[idx]);63+ if (idx >= count) {
64+ return;
65+ }
66+ result[idx] = remainderf(x[idx], y[idx]);
64 }67 }
65 ```68 ```
66 69 
67-- SIMD与SIMT混合编程场景:70+- SIMD与SIMT混合编程场景:
68 71 
69 ```72 ```
70- __simt_vf__ __launch_bounds__(1024) inline void KernelRemainder(__gm__ float* dst, __gm__ float* x, __gm__ float* y)73+ __simt_vf__ __launch_bounds__(256) inline void compute_remainderf_vf(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, uint32_t count)
71 {74 {
72- int idx = threadIdx.x + blockIdx.x * blockDim.x;75+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
73- dst[idx] = remainderf(x[idx], y[idx]);76+ if (idx >= count) {
77+ return;
78+ }
79+ result[idx] = remainderf(x[idx], y[idx]);
80+ }
81+ 
82+ __global__ __vector__ void run_remainderf(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, uint32_t count)
83+ {
84+ asc_vf_call<compute_remainderf_vf>(dim3(256), result, x, y, count);
74 }85 }
75 ```86 ```
76 87 
88+输入输出示例如下:
89+ 
90+```
91+x:0.25, 0.75, 1.25, 1.75
92+y:1.25, 2.25, 3.25, 4.25
93+result: 0.25 0.75 1.25 1.75
94+```
@@ -55,25 +55,42 @@ inline float remquof(float x, float y, int *quo)
55 55 
56## 调用示例56## 调用示例
57 57 
58-- SIMT编程场景:58+- SIMT编程场景:
59 59 
60 ```60 ```
61- __global__ __launch_bounds__(1024) void KernelRemQuo(float* dst, float* x, float* y, int* quo)61+ __global__ __launch_bounds__(256) void compute_remquof(float *result, int *quo, const float *x, const float *y, uint32_t count)
62 {62 {
63- int idx = threadIdx.x + blockIdx.x * blockDim.x;63+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
64- dst[idx] = remquof(x[idx], y[idx], quo + idx);64+ if (idx >= count) {
65- }65+ return;
66+ }
67+ result[idx] = remquof(x[idx], y[idx], quo + idx);
68+ }
66 ```69 ```
67 70 
68-- SIMD与SIMT混合编程场景:71+- SIMD与SIMT混合编程场景:
69- 
70- SIMD与SIMT混合编程场景,需要显式使用地址空间限定符表示地址空间:\_\_gm\_\_表示Global Memory内存空间,\_\_ubuf\_\_表示Unified Buffer内存空间,栈空间无需添加地址空间限定符。
71 72 
72 ```73 ```
73- __simt_vf__ __launch_bounds__(1024) inline void KernelRemQuo(__gm__ float* dst, __gm__ float* x, __gm__ float* y, __gm__ int* quo)74+ __simt_vf__ __launch_bounds__(256) inline void compute_remquof_vf(__gm__ float *result, __gm__ int *quo, __gm__ const float *x, __gm__ const float *y, uint32_t count)
74 {75 {
75- int idx = threadIdx.x + blockIdx.x * blockDim.x;76+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
76- dst[idx] = remquof(x[idx], y[idx], quo + idx);77+ if (idx >= count) {
77- }78+ return;
79+ }
80+ result[idx] = remquof(x[idx], y[idx], quo + idx);
81+ }
82+ 
83+ __global__ __vector__ void run_remquof(__gm__ float *result, __gm__ int *quo, __gm__ const float *x, __gm__ const float *y, uint32_t count)
84+ {
85+ asc_vf_call<compute_remquof_vf>(dim3(256), result, quo, x, y, count);
86+ }
78 ```87 ```
79 88 
89+输入输出示例如下:
90+ 
91+```
92+x:0.25, 0.75, 1.25, 1.75
93+y:1.25, 2.25, 3.25, 4.25
94+result: 0.25 0.75 1.25 1.75
95+quo: 0 0 0 0
96+```
@@ -57,23 +57,41 @@ x^2 + y^2的平方根的倒数。
57 57 
58## 调用示例58## 调用示例
59 59 
60-- SIMT编程场景:60+- SIMT编程场景:
61 61 
62 ```62 ```
63- __global__ __launch_bounds__(1024) void KernelPow(float* dst, float* x, float* y)63+ __global__ __launch_bounds__(256) void compute_rhypotf(float *result, const float *x, const float *y, uint32_t count)
64 {64 {
65- int idx = threadIdx.x + blockIdx.x * blockDim.x;65+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
66- dst[idx] = rhypotf(x[idx], y[idx]);66+ if (idx >= count) {
67+ return;
68+ }
69+ result[idx] = rhypotf(x[idx], y[idx]);
67 }70 }
68 ```71 ```
69 72 
70-- SIMD与SIMT混合编程场景:73+- SIMD与SIMT混合编程场景:
71 74 
72 ```75 ```
73- __simt_vf__ __launch_bounds__(1024) inline void KernelPow(__gm__ float* dst, __gm__ float* x, __gm__ float* y)76+ __simt_vf__ __launch_bounds__(256) inline void compute_rhypotf_vf(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, uint32_t count)
74 {77 {
75- int idx = threadIdx.x + blockIdx.x * blockDim.x;78+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
76- dst[idx] = rhypotf(x[idx], y[idx]);79+ if (idx >= count) {
80+ return;
81+ }
82+ result[idx] = rhypotf(x[idx], y[idx]);
83+ }
84+ 
85+ __global__ __vector__ void run_rhypotf(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, uint32_t count)
86+ {
87+ asc_vf_call<compute_rhypotf_vf>(dim3(256), result, x, y, count);
77 }88 }
78 ```89 ```
79 90 
91+输入输出示例如下:
92+ 
93+```
94+x:0.25, 0.75, 1.25, 1.75
95+y:1.5, 2.5, 3.5, 4.5
96+result: 0.7844645 0.421637 0.2871833 0.2175713
97+```
@@ -55,23 +55,40 @@ inline float rintf(float x)
55 55 
56## 调用示例56## 调用示例
57 57 
58-- SIMT编程场景:58+- SIMT编程场景:
59 59 
60 ```60 ```
61- __global__ __launch_bounds__(1024) void KernelRint(float* dst, float* x)61+ __global__ __launch_bounds__(256) void compute_rintf(float *result, const float *x, uint32_t count)
62 {62 {
63- int idx = threadIdx.x + blockIdx.x * blockDim.x;63+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
64- dst[idx] = rintf(x[idx]);64+ if (idx >= count) {
65+ return;
66+ }
67+ result[idx] = rintf(x[idx]);
65 }68 }
66 ```69 ```
67 70 
68-- SIMD与SIMT混合编程场景:71+- SIMD与SIMT混合编程场景:
69 72 
70 ```73 ```
71- __simt_vf__ __launch_bounds__(1024) inline void KernelRint(__gm__ float* dst, __gm__ float* x)74+ __simt_vf__ __launch_bounds__(256) inline void compute_rintf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
72 {75 {
73- int idx = threadIdx.x + blockIdx.x * blockDim.x;76+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
74- dst[idx] = rintf(x[idx]);77+ if (idx >= count) {
78+ return;
79+ }
80+ result[idx] = rintf(x[idx]);
81+ }
82+ 
83+ __global__ __vector__ void run_rintf(__gm__ float *result, __gm__ const float *x, uint32_t count)
84+ {
85+ asc_vf_call<compute_rintf_vf>(dim3(256), result, x, count);
75 }86 }
76 ```87 ```
77 88 
89+输入输出示例如下:
90+ 
91+```
92+x:0.25, 0.75, 1.25, 1.75
93+result: 0 1 1 2
94+```
@@ -62,23 +62,42 @@ a^2 + b^2 + c^2的平方根的倒数。
62 62 
63## 调用示例63## 调用示例
64 64 
65-- SIMT编程场景:65+- SIMT编程场景:
66 66 
67 ```67 ```
68- __global__ __launch_bounds__(1024) void KernelRnorm3d(float* dst, float* a, float* b, float* c)68+ __global__ __launch_bounds__(256) void compute_rnorm3df(float *result, const float *x, const float *y, const float *z, uint32_t count)
69 {69 {
70- int idx = threadIdx.x + blockIdx.x * blockDim.x;70+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
71- dst[idx] = rnorm3df(a[idx], b[idx], c[idx]);71+ if (idx >= count) {
72+ return;
73+ }
74+ result[idx] = rnorm3df(x[idx], y[idx], z[idx]);
72 }75 }
73 ```76 ```
74 77 
75-- SIMD与SIMT混合编程场景:78+- SIMD与SIMT混合编程场景:
76 79 
77 ```80 ```
78- __simt_vf__ __launch_bounds__(1024) inline void KernelRnorm3d(__gm__ float* dst, __gm__ float* a, __gm__ float* b, __gm__ float* c)81+ __simt_vf__ __launch_bounds__(256) inline void compute_rnorm3df_vf(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, __gm__ const float *z, uint32_t count)
79 {82 {
80- int idx = threadIdx.x + blockIdx.x * blockDim.x;83+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
81- dst[idx] = rnorm3df(a[idx], b[idx], c[idx]);84+ if (idx >= count) {
85+ return;
86+ }
87+ result[idx] = rnorm3df(x[idx], y[idx], z[idx]);
88+ }
89+ 
90+ __global__ __vector__ void run_rnorm3df(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, __gm__ const float *z, uint32_t count)
91+ {
92+ asc_vf_call<compute_rnorm3df_vf>(dim3(256), result, x, y, z, count);
82 }93 }
83 ```94 ```
84 95 
96+输入输出示例如下:
97+ 
98+```
99+x:0.25, 0.75, 1.25, 1.75
100+y:1.5, 2.5, 3.5, 4.5
101+z:-0.5, 0.5, 1.5, 2.5
102+result: 0.624695 0.3762884 0.2495131 0.1839202
103+```
@@ -59,23 +59,43 @@ a^2 + b^2 + c^2 + d^2的平方根的倒数。
59 59 
60## 调用示例60## 调用示例
61 61 
62-- SIMT编程场景:62+- SIMT编程场景:
63 63 
64 ```64 ```
65- __global__ __launch_bounds__(1024) void KernelRnorm4d(float* dst, float* a, float* b, float* c, float* d)65+ __global__ __launch_bounds__(256) void compute_rnorm4df(float *result, const float *x, const float *y, const float *z, const float *w, uint32_t count)
66 {66 {
67- int idx = threadIdx.x + blockIdx.x * blockDim.x;67+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
68- dst[idx] = rnorm4df(a[idx], b[idx], c[idx], d[idx]);68+ if (idx >= count) {
69+ return;
70+ }
71+ result[idx] = rnorm4df(x[idx], y[idx], z[idx], w[idx]);
69 }72 }
70 ```73 ```
71 74 
72-- SIMD与SIMT混合编程场景:75+- SIMD与SIMT混合编程场景:
73 76 
74 ```77 ```
75- __simt_vf__ __launch_bounds__(1024) inline void KernelRnorm4d(__gm__ float* dst, __gm__ float* a, __gm__ float* b, __gm__ float* c, __gm__ float* d)78+ __simt_vf__ __launch_bounds__(256) inline void compute_rnorm4df_vf(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, __gm__ const float *z, __gm__ const float *w, uint32_t count)
76 {79 {
77- int idx = threadIdx.x + blockIdx.x * blockDim.x;80+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
78- dst[idx] = rnorm4df(a[idx], b[idx], c[idx], d[idx]);81+ if (idx >= count) {
82+ return;
83+ }
84+ result[idx] = rnorm4df(x[idx], y[idx], z[idx], w[idx]);
85+ }
86+ 
87+ __global__ __vector__ void run_rnorm4df(__gm__ float *result, __gm__ const float *x, __gm__ const float *y, __gm__ const float *z, __gm__ const float *w, uint32_t count)
88+ {
89+ asc_vf_call<compute_rnorm4df_vf>(dim3(256), result, x, y, z, w, count);
79 }90 }
80 ```91 ```
81 92 
93+输入输出示例如下:
94+ 
95+```
96+x:0.25, 0.75, 1.25, 1.75
97+y:1.5, 2.5, 3.5, 4.5
98+z:-0.5, 0.5, 1.5, 2.5
99+w:2, 3, 4, 5
100+result: 0.39036 0.2495131 0.1766043 0.1353795
101+```
@@ -40,7 +40,7 @@ a\[0\]^2 + a\[1\]^2 + ...+ a\[n-1\]^2的平方根的倒数。
40- 若a\[0\]^2 + a\[1\]^2 + ...+ a\[n-1\]^2的平方根超出float最大范围,返回值为0。40- 若a\[0\]^2 + a\[1\]^2 + ...+ a\[n-1\]^2的平方根超出float最大范围,返回值为0。
41- 若a\[0\]^2 + a\[1\]^2 + ...+ a\[n-1\]^2平方根的倒数超出float最大范围,返回值为inf。41- 若a\[0\]^2 + a\[1\]^2 + ...+ a\[n-1\]^2平方根的倒数超出float最大范围,返回值为inf。
42- 若a\[0\]、a\[1\] 、... 、a\[n-1\]都为0,返回值为inf。42- 若a\[0\]、a\[1\] 、... 、a\[n-1\]都为0,返回值为inf。
43-- 若a\[0\]、a\[1\] 、... 、a\[n-1\]\]任意一个或多个为±inf,返回值为0。43+- 若a\[0\]、a\[1\] 、... 、a\[n-1\]任意一个或多个为±inf,返回值为0。
44- 若a\[0\]、a\[1\] 、... 、a\[n-1\]任意一个或多个为nan同时不是±inf,返回值为nan。44- 若a\[0\]、a\[1\] 、... 、a\[n-1\]任意一个或多个为nan同时不是±inf,返回值为nan。
45- 若n小于1,返回a\[0\]绝对值的倒数,若a\[0\]为0,返回值为inf。45- 若n小于1,返回a\[0\]绝对值的倒数,若a\[0\]为0,返回值为inf。
46 46 
@@ -59,25 +59,41 @@ a\[0\]^2 + a\[1\]^2 + ...+ a\[n-1\]^2的平方根的倒数。
59 59 
60## 调用示例60## 调用示例
61 61 
62-- SIMT编程场景:62+- SIMT编程场景:
63 63 
64 ```64 ```
65- __global__ __launch_bounds__(1024) void KernelRnorm(float* dst, int* n, float* a)65+ __global__ __launch_bounds__(256) void compute_rnormf(float *result, const int *n, float *vector_data, uint32_t count)
66 {66 {
67- int idx = threadIdx.x + blockIdx.x * blockDim.x;67+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
68- dst[idx] = rnormf(n[idx], a);68+ if (idx >= count) {
69+ return;
70+ }
71+ result[idx] = rnormf(n[idx], vector_data + idx * 4);
69 }72 }
70 ```73 ```
71 74 
72-- SIMD与SIMT混合编程场景:75+- SIMD与SIMT混合编程场景:
73- 
74- SIMD与SIMT混合编程场景,需要显式使用地址空间限定符表示地址空间:\_\_gm\_\_表示Global Memory内存空间,\_\_ubuf\_\_表示Unified Buffer内存空间,栈空间无需添加地址空间限定符。
75 76 
76 ```77 ```
77- __simt_vf__ __launch_bounds__(1024) inline void KernelRnorm(__gm__ float* dst, __gm__ int* n, __gm__ float* a)78+ __simt_vf__ __launch_bounds__(256) inline void compute_rnormf_vf(__gm__ float *result, __gm__ const int *n, __gm__ float *vector_data, uint32_t count)
78 {79 {
79- int idx = threadIdx.x + blockIdx.x * blockDim.x;80+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
80- dst[idx] = rnormf(n[idx], a);81+ if (idx >= count) {
82+ return;
83+ }
84+ result[idx] = rnormf(n[idx], vector_data + idx * 4);
85+ }
86+ 
87+ __global__ __vector__ void run_rnormf(__gm__ float *result, __gm__ const int *n, __gm__ float *vector_data, uint32_t count)
88+ {
89+ asc_vf_call<compute_rnormf_vf>(dim3(256), result, n, vector_data, count);
81 }90 }
82 ```91 ```
83 92 
93+输入输出示例如下:
94+ 
95+```
96+n:1, 2, 3, 1
97+vector_data:[[1, 2, 3, 4], [2, 3, 4, 5], [3, 4, 5, 6], [4, 5, 6, 7]]
98+result: 1 0.2773501 0.1414214 0.25
99+```
@@ -52,23 +52,40 @@ inline float roundf(float x)
52 52 
53## 调用示例53## 调用示例
54 54 
55-- SIMT编程场景:55+- SIMT编程场景:
56 56 
57 ```57 ```
58- __global__ __launch_bounds__(1024) void KernelRound(float* dst, float* x)58+ __global__ __launch_bounds__(256) void compute_roundf(float *result, const float *x, uint32_t count)
59 {59 {
60- int idx = threadIdx.x + blockIdx.x * blockDim.x;60+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
61- dst[idx] = roundf(x[idx]);61+ if (idx >= count) {
62+ return;
63+ }
64+ result[idx] = roundf(x[idx]);
62 }65 }
63 ```66 ```
64 67 
65-- SIMD与SIMT混合编程场景:68+- SIMD与SIMT混合编程场景:
66 69 
67 ```70 ```
68- __simt_vf__ __launch_bounds__(1024) inline void KernelRound(__gm__ float* dst, __gm__ float* x)71+ __simt_vf__ __launch_bounds__(256) inline void compute_roundf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
69 {72 {
70- int idx = threadIdx.x + blockIdx.x * blockDim.x;73+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
71- dst[idx] = roundf(x[idx]);74+ if (idx >= count) {
75+ return;
76+ }
77+ result[idx] = roundf(x[idx]);
78+ }
79+ 
80+ __global__ __vector__ void run_roundf(__gm__ float *result, __gm__ const float *x, uint32_t count)
81+ {
82+ asc_vf_call<compute_roundf_vf>(dim3(256), result, x, count);
72 }83 }
73 ```84 ```
74 85 
86+输入输出示例如下:
87+ 
88+```
89+x:0.25, 0.75, 1.25, 1.75
90+result: 0 1 1 2
91+```
@@ -55,23 +55,40 @@ inline float rsqrtf(float x)
55 55 
56## 调用示例56## 调用示例
57 57 
58-- SIMT编程场景:58+- SIMT编程场景:
59 59 
60 ```60 ```
61- __global__ __launch_bounds__(1024) void KernelRsqrt(float* dst, float* x)61+ __global__ __launch_bounds__(256) void compute_rsqrtf(float *result, const float *x, uint32_t count)
62 {62 {
63- int idx = threadIdx.x + blockIdx.x * blockDim.x;63+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
64- dst[idx] = rsqrtf(x[idx]);64+ if (idx >= count) {
65+ return;
66+ }
67+ result[idx] = rsqrtf(x[idx]);
65 }68 }
66 ```69 ```
67 70 
68-- SIMD与SIMT混合编程场景:71+- SIMD与SIMT混合编程场景:
69 72 
70 ```73 ```
71- __simt_vf__ __launch_bounds__(1024) inline void KernelRsqrt(__gm__ float* dst, __gm__ float* x)74+ __simt_vf__ __launch_bounds__(256) inline void compute_rsqrtf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
72 {75 {
73- int idx = threadIdx.x + blockIdx.x * blockDim.x;76+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
74- dst[idx] = rsqrtf(x[idx]);77+ if (idx >= count) {
78+ return;
79+ }
80+ result[idx] = rsqrtf(x[idx]);
81+ }
82+ 
83+ __global__ __vector__ void run_rsqrtf(__gm__ float *result, __gm__ const float *x, uint32_t count)
84+ {
85+ asc_vf_call<compute_rsqrtf_vf>(dim3(256), result, x, count);
75 }86 }
76 ```87 ```
77 88 
89+输入输出示例如下:
90+ 
91+```
92+x:0.25, 1.25, 2.25, 3.25
93+result: 2 0.8944272 0.6666667 0.5547002
94+```
@@ -55,22 +55,41 @@ inline float scalblnf(float x, int64_t n)
55 55 
56## 调用示例56## 调用示例
57 57 
58-- SIMT编程场景:58+- SIMT编程场景:
59 59 
60 ```60 ```
61- __global__ __launch_bounds__(1024) void KernelScalblnf(float* dst, float* x, int64_t* n)61+ __global__ __launch_bounds__(256) void compute_scalblnf(float *result, const int *n, const float *x, uint32_t count)
62 {62 {
63- int idx = threadIdx.x + blockIdx.x * blockDim.x;63+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
64- dst[idx] = scalblnf(x[idx], n[idx]);64+ if (idx >= count) {
65+ return;
66+ }
67+ result[idx] = scalblnf(x[idx], static_cast<int64_t>(n[idx]));
65 }68 }
66 ```69 ```
67 70 
68-- SIMD与SIMT混合编程场景:71+- SIMD与SIMT混合编程场景:
69 72 
70 ```73 ```
71- __simt_vf__ __launch_bounds__(1024) inline void KernelScalblnf(__gm__ float* dst, __gm__ float* x, __gm__ int64_t* n)74+ __simt_vf__ __launch_bounds__(256) inline void compute_scalblnf_vf(__gm__ float *result, __gm__ const int *n, __gm__ const float *x, uint32_t count)
72 {75 {
73- int idx = threadIdx.x + blockIdx.x * blockDim.x;76+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
74- dst[idx] = scalblnf(x[idx], n[idx]);77+ if (idx >= count) {
78+ return;
79+ }
80+ result[idx] = scalblnf(x[idx], static_cast<int64_t>(n[idx]));
81+ }
82+ 
83+ __global__ __vector__ void run_scalblnf(__gm__ float *result, __gm__ const int *n, __gm__ const float *x, uint32_t count)
84+ {
85+ asc_vf_call<compute_scalblnf_vf>(dim3(256), result, n, x, count);
75 }86 }
76 ```87 ```
88+ 
89+输入输出示例如下:
90+ 
91+```
92+n:1, 2, 3, 1
93+x:0.25, 0.75, 1.25, 1.75
94+result: 0.5 3 10 3.5
95+```
@@ -55,23 +55,41 @@ inline float scalbnf(float x, int32_t n)
55 55 
56## 调用示例56## 调用示例
57 57 
58-- SIMT编程场景:58+- SIMT编程场景:
59 59 
60 ```60 ```
61- __global__ __launch_bounds__(1024) void KernelScalbn(float* dst, float* x, int* n)61+ __global__ __launch_bounds__(256) void compute_scalbnf(float *result, const int *n, const float *x, uint32_t count)
62 {62 {
63- int idx = threadIdx.x + blockIdx.x * blockDim.x;63+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
64- dst[idx] = scalbnf(x[idx], n[idx]);64+ if (idx >= count) {
65+ return;
66+ }
67+ result[idx] = scalbnf(x[idx], n[idx]);
65 }68 }
66 ```69 ```
67 70 
68-- SIMD与SIMT混合编程场景:71+- SIMD与SIMT混合编程场景:
69 72 
70 ```73 ```
71- __simt_vf__ __launch_bounds__(1024) inline void KernelScalbn(__gm__ float* dst, __gm__ float* x, __gm__ int* n)74+ __simt_vf__ __launch_bounds__(256) inline void compute_scalbnf_vf(__gm__ float *result, __gm__ const int *n, __gm__ const float *x, uint32_t count)
72 {75 {
73- int idx = threadIdx.x + blockIdx.x * blockDim.x;76+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
74- dst[idx] = scalbnf(x[idx], n[idx]);77+ if (idx >= count) {
78+ return;
79+ }
80+ result[idx] = scalbnf(x[idx], n[idx]);
81+ }
82+ 
83+ __global__ __vector__ void run_scalbnf(__gm__ float *result, __gm__ const int *n, __gm__ const float *x, uint32_t count)
84+ {
85+ asc_vf_call<compute_scalbnf_vf>(dim3(256), result, n, x, count);
75 }86 }
76 ```87 ```
77 88 
89+输入输出示例如下:
90+ 
91+```
92+n:1, 2, 3, 1
93+x:0.25, 0.75, 1.25, 1.75
94+result: 0.5 3 10 3.5
95+```
@@ -34,9 +34,12 @@ inline int signbit(float x)
34 34 
35返回输入数据的符号位。35返回输入数据的符号位。
36 36 
37-- 当x为nan时,返回值为0。37+- 当x为+nan时,返回值为0。
38+- 当x为-nan时,返回值为1。
38- 当x为inf时,返回值为0。39- 当x为inf时,返回值为0。
39- 当x为-inf时,返回值为1。40- 当x为-inf时,返回值为1。
41+- 当x为+0时,返回值为0。
42+- 当x为-0时,返回值为1。
40 43 
41## 约束说明44## 约束说明
42 45 
@@ -52,23 +55,40 @@ inline int signbit(float x)
52 55 
53## 调用示例56## 调用示例
54 57 
55-- SIMT编程场景:58+- SIMT编程场景:
56 59 
57 ```60 ```
58- __global__ __launch_bounds__(1024) void KernelSignbit(int* dst, float* x)61+ __global__ __launch_bounds__(256) void compute_signbit(float *result, const float *x, uint32_t count)
59 {62 {
60- int idx = threadIdx.x + blockIdx.x * blockDim.x;63+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
61- dst[idx] = signbit(x[idx]);64+ if (idx >= count) {
65+ return;
66+ }
67+ result[idx] = static_cast<float>(signbit(x[idx]));
62 }68 }
63 ```69 ```
64 70 
65-- SIMD与SIMT混合编程场景:71+- SIMD与SIMT混合编程场景:
66 72 
67 ```73 ```
68- __simt_vf__ __launch_bounds__(1024) inline void KernelSignbit(__gm__ int* dst, __gm__ float* x)74+ __simt_vf__ __launch_bounds__(256) inline void compute_signbit_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
69 {75 {
70- int idx = threadIdx.x + blockIdx.x * blockDim.x;76+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
71- dst[idx] = signbit(x[idx]);77+ if (idx >= count) {
78+ return;
79+ }
80+ result[idx] = static_cast<float>(signbit(x[idx]));
81+ }
82+ 
83+ __global__ __vector__ void run_signbit(__gm__ float *result, __gm__ const float *x, uint32_t count)
84+ {
85+ asc_vf_call<compute_signbit_vf>(dim3(256), result, x, count);
72 }86 }
73 ```87 ```
74 88 
89+输入输出示例如下:
90+ 
91+```
92+x:0.25, 0.75, 1.25, 1.75
93+result: 0 0 0 0
94+```
@@ -52,25 +52,41 @@ inline void sincosf(float x, float *s, float *c)
52 52 
53## 调用示例53## 调用示例
54 54 
55-- SIMT编程场景:55+- SIMT编程场景:
56 56 
57 ```57 ```
58- __global__ __launch_bounds__(1024) void KernelSinCos(float* s, float* c, float* x)58+ __global__ __launch_bounds__(256) void compute_sincosf(float *sin_result, float *cos_result, const float *x, uint32_t count)
59 {59 {
60- int idx = threadIdx.x + blockIdx.x * blockDim.x;60+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
61- sincosf(x[idx], s + idx, c + idx); // 对源地址的第idx个元素取三角函数正弦值和余弦值61+ if (idx >= count) {
62+ return;
63+ }
64+ sincosf(x[idx], sin_result + idx, cos_result + idx);
62 }65 }
63 ```66 ```
64 67 
65-- SIMD与SIMT混合编程场景:68+- SIMD与SIMT混合编程场景:
66- 
67- SIMD与SIMT混合编程场景,需要显式使用地址空间限定符表示地址空间:\_\_gm\_\_表示Global Memory内存空间,\_\_ubuf\_\_表示Unified Buffer内存空间,栈空间无需添加地址空间限定符。
68 69 
69 ```70 ```
70- __simt_vf__ __launch_bounds__(1024) inline void KernelSinCos(__gm__ float* s, __gm__ float* c, __gm__ float* x)71+ __simt_vf__ __launch_bounds__(256) inline void compute_sincosf_vf(__gm__ float *sin_result, __gm__ float *cos_result, __gm__ const float *x, uint32_t count)
71 {72 {
72- int idx = threadIdx.x + blockIdx.x * blockDim.x;73+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
73- sincosf(x[idx], s + idx, c + idx); // 对源地址的第idx个元素取三角函数正弦值和余弦值74+ if (idx >= count) {
75+ return;
76+ }
77+ sincosf(x[idx], sin_result + idx, cos_result + idx);
78+ }
79+ 
80+ __global__ __vector__ void run_sincosf(__gm__ float *sin_result, __gm__ float *cos_result, __gm__ const float *x, uint32_t count)
81+ {
82+ asc_vf_call<compute_sincosf_vf>(dim3(256), sin_result, cos_result, x, count);
74 }83 }
75 ```84 ```
76 85 
86+输入输出示例如下:
87+ 
88+```
89+x:0.25, 0.75, 1.25, 1.75
90+sin_result: 0.247404 0.6816388 0.9489846 0.983986
91+cos_result: 0.9689124 0.7316889 0.3153224 -0.1782461
92+```
@@ -53,25 +53,41 @@ inline void sincospif(float x, float *s, float *c)
53 53 
54## 调用示例54## 调用示例
55 55 
56-- SIMT编程场景:56+- SIMT编程场景:
57 57 
58 ```58 ```
59- __global__ __launch_bounds__(1024) void KernelSinCospi(float* s, float* c, float* x)59+ __global__ __launch_bounds__(256) void compute_sincospif(float *sin_result, float *cos_result, const float *x, uint32_t count)
60 {60 {
61- int idx = threadIdx.x + blockIdx.x * blockDim.x;61+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
62- sincospif(x[idx], s + idx, c + idx); // 对源地址的第idx个元素取与π相乘后的三角函数正弦值和余弦值62+ if (idx >= count) {
63+ return;
64+ }
65+ sincospif(x[idx], sin_result + idx, cos_result + idx);
63 }66 }
64 ```67 ```
65 68 
66-- SIMD与SIMT混合编程场景:69+- SIMD与SIMT混合编程场景:
67- 
68- SIMD与SIMT混合编程场景,需要显式使用地址空间限定符表示地址空间:\_\_gm\_\_表示Global Memory内存空间,\_\_ubuf\_\_表示Unified Buffer内存空间,栈空间无需添加地址空间限定符。
69 70 
70 ```71 ```
71- __simt_vf__ __launch_bounds__(1024) inline void KernelSinCospi(__gm__ float* s, __gm__ float* c, __gm__ float* x)72+ __simt_vf__ __launch_bounds__(256) inline void compute_sincospif_vf(__gm__ float *sin_result, __gm__ float *cos_result, __gm__ const float *x, uint32_t count)
72 {73 {
73- int idx = threadIdx.x + blockIdx.x * blockDim.x;74+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
74- sincospif(x[idx], s + idx, c + idx); // 对源地址的第idx个元素取与π相乘后的三角函数正弦值和余弦值75+ if (idx >= count) {
76+ return;
77+ }
78+ sincospif(x[idx], sin_result + idx, cos_result + idx);
79+ }
80+ 
81+ __global__ __vector__ void run_sincospif(__gm__ float *sin_result, __gm__ float *cos_result, __gm__ const float *x, uint32_t count)
82+ {
83+ asc_vf_call<compute_sincospif_vf>(dim3(256), sin_result, cos_result, x, count);
75 }84 }
76 ```85 ```
77 86 
87+输入输出示例如下:
88+ 
89+```
90+x:0.25, 0.75, 1.25, 1.75
91+sin_result: 0.7071068 0.7071068 -0.7071069 -0.7071065
92+cos_result: 0.7071068 -0.7071068 -0.7071066 0.707107
93+```
@@ -54,23 +54,40 @@ inline float sinf(float x)
54 54 
55## 调用示例55## 调用示例
56 56 
57-- SIMT编程场景:57+- SIMT编程场景:
58 58 
59 ```59 ```
60- __global__ __launch_bounds__(1024) void KernelSin(float* dst, float* x)60+ __global__ __launch_bounds__(256) void compute_sinf(float *result, const float *x, uint32_t count)
61 {61 {
62- int idx = threadIdx.x + blockIdx.x * blockDim.x;62+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
63- dst[idx] = sinf(x[idx]); // 对src源地址的第idx个元素取三角函数正弦值63+ if (idx >= count) {
64+ return;
65+ }
66+ result[idx] = sinf(x[idx]);
64 }67 }
65 ```68 ```
66 69 
67-- SIMD与SIMT混合编程场景:70+- SIMD与SIMT混合编程场景:
68 71 
69 ```72 ```
70- __simt_vf__ __launch_bounds__(1024) inline void KernelSin(__gm__ float* dst, __gm__ float* x)73+ __simt_vf__ __launch_bounds__(256) inline void compute_sinf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
71 {74 {
72- int idx = threadIdx.x + blockIdx.x * blockDim.x;75+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
73- dst[idx] = sinf(x[idx]); // 对src源地址的第idx个元素取三角函数正弦值76+ if (idx >= count) {
77+ return;
78+ }
79+ result[idx] = sinf(x[idx]);
80+ }
81+ 
82+ __global__ __vector__ void run_sinf(__gm__ float *result, __gm__ const float *x, uint32_t count)
83+ {
84+ asc_vf_call<compute_sinf_vf>(dim3(256), result, x, count);
74 }85 }
75 ```86 ```
76 87 
88+输入输出示例如下:
89+ 
90+```
91+x:0.25, 0.75, 1.25, 1.75
92+result: 0.247404 0.6816388 0.9489846 0.983986
93+```
@@ -55,23 +55,40 @@ inline float sinhf(float x)
55 55 
56## 调用示例56## 调用示例
57 57 
58-- SIMT编程场景:58+- SIMT编程场景:
59 59 
60 ```60 ```
61- __global__ __launch_bounds__(1024) void KernelSinh(float* dst, float* x)61+ __global__ __launch_bounds__(256) void compute_sinhf(float *result, const float *x, uint32_t count)
62 {62 {
63- int idx = threadIdx.x + blockIdx.x * blockDim.x;63+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
64- dst[idx] = sinhf(x[idx]); // 对src源地址的第idx个元素取双曲正弦值64+ if (idx >= count) {
65+ return;
66+ }
67+ result[idx] = sinhf(x[idx]);
65 }68 }
66 ```69 ```
67 70 
68-- SIMD与SIMT混合编程场景:71+- SIMD与SIMT混合编程场景:
69 72 
70 ```73 ```
71- __simt_vf__ __launch_bounds__(1024) inline void KernelSinh(__gm__ float* dst, __gm__ float* x)74+ __simt_vf__ __launch_bounds__(256) inline void compute_sinhf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
72 {75 {
73- int idx = threadIdx.x + blockIdx.x * blockDim.x;76+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
74- dst[idx] = sinhf(x[idx]); // 对src源地址的第idx个元素取双曲正弦值77+ if (idx >= count) {
78+ return;
79+ }
80+ result[idx] = sinhf(x[idx]);
81+ }
82+ 
83+ __global__ __vector__ void run_sinhf(__gm__ float *result, __gm__ const float *x, uint32_t count)
84+ {
85+ asc_vf_call<compute_sinhf_vf>(dim3(256), result, x, count);
75 }86 }
76 ```87 ```
77 88 
89+输入输出示例如下:
90+ 
91+```
92+x:0.25, 0.75, 1.25, 1.75
93+result: 0.2526123 0.8223166 1.601919 2.790414
94+```
@@ -56,23 +56,40 @@ inline float sinpif(float x)
56 56 
57## 调用示例57## 调用示例
58 58 
59-- SIMT编程场景:59+- SIMT编程场景:
60 60 
61 ```61 ```
62- __global__ __launch_bounds__(1024) void KernelSinpi(float* dst, float* x)62+ __global__ __launch_bounds__(256) void compute_sinpif(float *result, const float *x, uint32_t count)
63 {63 {
64- int idx = threadIdx.x + blockIdx.x * blockDim.x;64+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
65- dst[idx] = sinpif(x[idx]); // 对src源地址的第idx个元素取与π相乘的正弦值65+ if (idx >= count) {
66+ return;
67+ }
68+ result[idx] = sinpif(x[idx]);
66 }69 }
67 ```70 ```
68 71 
69-- SIMD与SIMT混合编程场景:72+- SIMD与SIMT混合编程场景:
70 73 
71 ```74 ```
72- __simt_vf__ __launch_bounds__(1024) inline void KernelSinpi(__gm__ float* dst, __gm__ float* x)75+ __simt_vf__ __launch_bounds__(256) inline void compute_sinpif_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
73 {76 {
74- int idx = threadIdx.x + blockIdx.x * blockDim.x;77+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
75- dst[idx] = sinpif(x[idx]); // 对src源地址的第idx个元素取与π相乘的正弦值78+ if (idx >= count) {
79+ return;
80+ }
81+ result[idx] = sinpif(x[idx]);
82+ }
83+ 
84+ __global__ __vector__ void run_sinpif(__gm__ float *result, __gm__ const float *x, uint32_t count)
85+ {
86+ asc_vf_call<compute_sinpif_vf>(dim3(256), result, x, count);
76 }87 }
77 ```88 ```
78 89 
90+输入输出示例如下:
91+ 
92+```
93+x:0.25, 0.75, 1.25, 1.75
94+result: 0.7071068 0.7071068 -0.7071069 -0.7071065
95+```
@@ -55,23 +55,40 @@ inline float sqrtf(float x)
55 55 
56## 调用示例56## 调用示例
57 57 
58-- SIMT编程场景:58+- SIMT编程场景:
59 59 
60 ```60 ```
61- __global__ __launch_bounds__(1024) void KernelSqrt(float* dst, float* x)61+ __global__ __launch_bounds__(256) void compute_sqrtf(float *result, const float *x, uint32_t count)
62 {62 {
63- int idx = threadIdx.x + blockIdx.x * blockDim.x;63+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
64- dst[idx] = sqrtf(x[idx]);64+ if (idx >= count) {
65+ return;
66+ }
67+ result[idx] = sqrtf(x[idx]);
65 }68 }
66 ```69 ```
67 70 
68-- SIMD与SIMT混合编程场景:71+- SIMD与SIMT混合编程场景:
69 72 
70 ```73 ```
71- __simt_vf__ __launch_bounds__(1024) inline void KernelSqrt(__gm__ float* dst, __gm__ float* x)74+ __simt_vf__ __launch_bounds__(256) inline void compute_sqrtf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
72 {75 {
73- int idx = threadIdx.x + blockIdx.x * blockDim.x;76+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
74- dst[idx] = sqrtf(x[idx]);77+ if (idx >= count) {
78+ return;
79+ }
80+ result[idx] = sqrtf(x[idx]);
81+ }
82+ 
83+ __global__ __vector__ void run_sqrtf(__gm__ float *result, __gm__ const float *x, uint32_t count)
84+ {
85+ asc_vf_call<compute_sqrtf_vf>(dim3(256), result, x, count);
75 }86 }
76 ```87 ```
77 88 
89+输入输出示例如下:
90+ 
91+```
92+x:0.25, 1.25, 2.25, 3.25
93+result: 0.5 1.118034 1.5 1.802776
94+```
@@ -54,23 +54,40 @@ inline float tanf(float x)
54 54 
55## 调用示例55## 调用示例
56 56 
57-- SIMT编程场景:57+- SIMT编程场景:
58 58 
59 ```59 ```
60- __global__ __launch_bounds__(1024) void KernelTan(float* dst, float* x)60+ __global__ __launch_bounds__(256) void compute_tanf(float *result, const float *x, uint32_t count)
61 {61 {
62- int idx = threadIdx.x + blockIdx.x * blockDim.x;62+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
63- dst[idx] = tanf(x[idx]);63+ if (idx >= count) {
64+ return;
65+ }
66+ result[idx] = tanf(x[idx]);
64 }67 }
65 ```68 ```
66 69 
67-- SIMD与SIMT混合编程场景:70+- SIMD与SIMT混合编程场景:
68 71 
69 ```72 ```
70- __simt_vf__ __launch_bounds__(1024) inline void KernelTan(__gm__ float* dst, __gm__ float* x)73+ __simt_vf__ __launch_bounds__(256) inline void compute_tanf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
71 {74 {
72- int idx = threadIdx.x + blockIdx.x * blockDim.x;75+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
73- dst[idx] = tanf(x[idx]);76+ if (idx >= count) {
77+ return;
78+ }
79+ result[idx] = tanf(x[idx]);
80+ }
81+ 
82+ __global__ __vector__ void run_tanf(__gm__ float *result, __gm__ const float *x, uint32_t count)
83+ {
84+ asc_vf_call<compute_tanf_vf>(dim3(256), result, x, count);
74 }85 }
75 ```86 ```
76 87 
88+输入输出示例如下:
89+ 
90+```
91+x:0.25, 0.75, 1.25, 1.75
92+result: 0.2553419 0.9315965 3.00957 -5.52038
93+```
@@ -54,23 +54,40 @@ inline float tanhf(float x)
54 54 
55## 调用示例55## 调用示例
56 56 
57-- SIMT编程场景:57+- SIMT编程场景:
58 58 
59 ```59 ```
60- __global__ __launch_bounds__(1024) void KernelTanh(float* dst, float* x)60+ __global__ __launch_bounds__(256) void compute_tanhf(float *result, const float *x, uint32_t count)
61 {61 {
62- int idx = threadIdx.x + blockIdx.x * blockDim.x;62+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
63- dst[idx] = tanhf(x[idx]);63+ if (idx >= count) {
64+ return;
65+ }
66+ result[idx] = tanhf(x[idx]);
64 }67 }
65 ```68 ```
66 69 
67-- SIMD与SIMT混合编程场景:70+- SIMD与SIMT混合编程场景:
68 71 
69 ```72 ```
70- __simt_vf__ __launch_bounds__(1024) inline void KernelTanh(__gm__ float* dst, __gm__ float* x)73+ __simt_vf__ __launch_bounds__(256) inline void compute_tanhf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
71 {74 {
72- int idx = threadIdx.x + blockIdx.x * blockDim.x;75+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
73- dst[idx] = tanhf(x[idx]);76+ if (idx >= count) {
77+ return;
78+ }
79+ result[idx] = tanhf(x[idx]);
80+ }
81+ 
82+ __global__ __vector__ void run_tanhf(__gm__ float *result, __gm__ const float *x, uint32_t count)
83+ {
84+ asc_vf_call<compute_tanhf_vf>(dim3(256), result, x, count);
74 }85 }
75 ```86 ```
76 87 
88+输入输出示例如下:
89+ 
90+```
91+x:0.25, 0.75, 1.25, 1.75
92+result: 0.2449186 0.6351489 0.8482836 0.9413756
93+```
@@ -56,23 +56,40 @@ inline float tanpif(float x)
56 56 
57## 调用示例57## 调用示例
58 58 
59-- SIMT编程场景:59+- SIMT编程场景:
60 60 
61 ```61 ```
62- __global__ __launch_bounds__(1024) void KernelTanPi(float* dst, float* x)62+ __global__ __launch_bounds__(256) void compute_tanpif(float *result, const float *x, uint32_t count)
63 {63 {
64- int idx = threadIdx.x + blockIdx.x * blockDim.x;64+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
65- dst[idx] = tanpif(x[idx]);65+ if (idx >= count) {
66+ return;
67+ }
68+ result[idx] = tanpif(x[idx]);
66 }69 }
67 ```70 ```
68 71 
69-- SIMD与SIMT混合编程场景:72+- SIMD与SIMT混合编程场景:
70 73 
71 ```74 ```
72- __simt_vf__ __launch_bounds__(1024) inline void KernelTanPi(__gm__ float* dst, __gm__ float* x)75+ __simt_vf__ __launch_bounds__(256) inline void compute_tanpif_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
73 {76 {
74- int idx = threadIdx.x + blockIdx.x * blockDim.x;77+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
75- dst[idx] = tanpif(x[idx]);78+ if (idx >= count) {
79+ return;
80+ }
81+ result[idx] = tanpif(x[idx]);
82+ }
83+ 
84+ __global__ __vector__ void run_tanpif(__gm__ float *result, __gm__ const float *x, uint32_t count)
85+ {
86+ asc_vf_call<compute_tanpif_vf>(dim3(256), result, x, count);
76 }87 }
77 ```88 ```
78 89 
90+输入输出示例如下:
91+ 
92+```
93+x:0.25, 0.75, 1.25, 1.75
94+result: 1 -0.9999999 1 -0.9999993
95+```
@@ -58,23 +58,40 @@ inline float tgammaf(float x)
58 58 
59## 调用示例59## 调用示例
60 60 
61-- SIMT编程场景:61+- SIMT编程场景:
62 62 
63 ```63 ```
64- __global__ __launch_bounds__(1024) void KernelTgamma(float* dst, float* x)64+ __global__ __launch_bounds__(256) void compute_tgammaf(float *result, const float *x, uint32_t count)
65 {65 {
66- int idx = threadIdx.x + blockIdx.x * blockDim.x;66+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
67- dst[idx] = tgammaf(x[idx]);67+ if (idx >= count) {
68+ return;
69+ }
70+ result[idx] = tgammaf(x[idx]);
68 }71 }
69 ```72 ```
70 73 
71-- SIMD与SIMT混合编程场景:74+- SIMD与SIMT混合编程场景:
72 75 
73 ```76 ```
74- __simt_vf__ __launch_bounds__(1024) inline void KernelTgamma(__gm__ float* dst, __gm__ float* x)77+ __simt_vf__ __launch_bounds__(256) inline void compute_tgammaf_vf(__gm__ float *result, __gm__ const float *x, uint32_t count)
75 {78 {
76- int idx = threadIdx.x + blockIdx.x * blockDim.x;79+ const uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
77- dst[idx] = tgammaf(x[idx]);80+ if (idx >= count) {
81+ return;
82+ }
83+ result[idx] = tgammaf(x[idx]);
84+ }
85+ 
86+ __global__ __vector__ void run_tgammaf(__gm__ float *result, __gm__ const float *x, uint32_t count)
87+ {
88+ asc_vf_call<compute_tgammaf_vf>(dim3(256), result, x, count);
78 }89 }
79 ```90 ```
80 91 
92+输入输出示例如下:
93+ 
94+```
95+x:0.25, 0.75, 1.25, 1.75
96+result: 3.62561 1.225417 0.9064025 0.9190626
97+```