Merged
update simt float math api doc && clarify SIMT VF array parameter error #2743
wulinyuAdd on May 29
update simt float math api doc && clarify SIMT VF array parameter error #2743
Merged
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 | |||
| 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 | ||
| 38 | y/x的反正切值。 | 38 | y/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 | +``` | ||
缩进不对 下同