已合并
修改stateless_drop_out_gen_mask算子+aclnn接口 #2299
huairuchen创建于 4月18日
修改stateless_drop_out_gen_mask算子+aclnn接口 #2299
已合并
从已删除 :pr_dropout合入到cann/ops-mathmaster
共 9 个文件变更+131-187
| @@ -107,7 +107,7 @@ ${INSTALL_DIR}表示CANN安装后文件路径；\$\{ops\_project\}表示算子 | |||
| 107 | | [aclnnDropoutDoMask](../../random/drop_out_do_mask/docs/aclnnDropoutDoMask.md) | 按照概率prob随机将输入中的元素置零，并将输出按照1/(1-prob)的比例放大。 | 默认确定性实现| 默认确定性实现 | | 107 | | [aclnnDropoutDoMask](../../random/drop_out_do_mask/docs/aclnnDropoutDoMask.md) | 按照概率prob随机将输入中的元素置零，并将输出按照1/(1-prob)的比例放大。 | 默认确定性实现| 默认确定性实现 | |
| 108 | | [aclnnDropoutGenMask](../../random/dsa_gen_bit_mask/docs/aclnnDropoutGenMask.md) | 训练过程中，按照概率prob生成mask，用于元素置零。 | 默认确定性实现| 默认确定性实现 | | 108 | | [aclnnDropoutGenMask](../../random/dsa_gen_bit_mask/docs/aclnnDropoutGenMask.md) | 训练过程中，按照概率prob生成mask，用于元素置零。 | 默认确定性实现| 默认确定性实现 | |
| 109 | | [aclnnDropoutGenMaskV2](../../random/dsa_gen_bit_mask/docs/aclnnDropoutGenMaskV2.md) | 训练过程中，按照概率prob生成mask，用于元素置零。 | 默认确定性实现| 默认确定性实现 | | 109 | | [aclnnDropoutGenMaskV2](../../random/dsa_gen_bit_mask/docs/aclnnDropoutGenMaskV2.md) | 训练过程中，按照概率prob生成mask，用于元素置零。 | 默认确定性实现| 默认确定性实现 | |
| 110 | -| [aclnnDropoutGenMaskV2Tensor](../../random/dsa_gen_bit_mask/docs/aclnnDropoutGenMaskV2Tensor.md) | 训练过程中，按照概率prob生成mask，用于元素置零。 | 默认确定性实现| - | | 110 | +| [aclnnDropoutGenMaskV2Tensor](../../random/dsa_gen_bit_mask/docs/aclnnDropoutGenMaskV2Tensor.md) | 训练过程中，按照概率prob生成mask，用于元素置零。 | 默认确定性实现| 默认确定性实现 | |
| 111 | | [aclnnDropoutV3](../../random/drop_out_v3/docs/aclnnDropoutV3.md) | 按照概率p随机将输入中的元素置零，并将输出按照1/(1-p)的比例缩放。 | 默认确定性实现| 默认确定性实现 | | 111 | | [aclnnDropoutV3](../../random/drop_out_v3/docs/aclnnDropoutV3.md) | 按照概率p随机将输入中的元素置零，并将输出按照1/(1-p)的比例缩放。 | 默认确定性实现| 默认确定性实现 | |
| 112 | | [aclnnEqScalar&aclnnInplaceEqScalar](../../math/equal/docs/aclnnEqScalar&aclnnInplaceEqScalar.md) | 计算self中的元素的值与other的值是否相等，将self每个元素与other的值的比较结果写入out中。 | 默认确定性实现| 默认确定性实现 | | 112 | | [aclnnEqScalar&aclnnInplaceEqScalar](../../math/equal/docs/aclnnEqScalar&aclnnInplaceEqScalar.md) | 计算self中的元素的值与other的值是否相等，将self每个元素与other的值的比较结果写入out中。 | 默认确定性实现| 默认确定性实现 | |
| 113 | | [aclnnEqTensor&aclnnInplaceEqTensor](../../math/equal/docs/aclnnEqTensor&aclnnInplaceEqTensor.md) | 计算两个Tensor中的元素是否相等，返回一个Tensor，self=other的为True(1.)，否则为False(0.)。 | 默认确定性实现| 默认确定性实现 | | 113 | | [aclnnEqTensor&aclnnInplaceEqTensor](../../math/equal/docs/aclnnEqTensor&aclnnInplaceEqTensor.md) | 计算两个Tensor中的元素是否相等，返回一个Tensor，self=other的为True(1.)，否则为False(0.)。 | 默认确定性实现| 默认确定性实现 | |
| @@ -6,7 +6,7 @@ | |||
| 6 | 6 | ||
| 7 | | 产品 | 是否支持 | | 7 | | 产品 | 是否支持 | |
| 8 | | :----------------------------------------------------------- | :------: | | 8 | | :----------------------------------------------------------- | :------: | |
| 9 | -| <term>Ascend 950PR/Ascend 950DT</term> | × | | 9 | +| <term>Ascend 950PR/Ascend 950DT</term> | √ | |
| 10 | | <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ | | 10 | | <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ | |
| 11 | | <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ | | 11 | | <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ | |
| 12 | | <term>Atlas 200I/500 A2 推理产品</term> | × | | 12 | | <term>Atlas 200I/500 A2 推理产品</term> | × | |
| @@ -1,5 +1,5 @@ | |||
| 1 | /** | 1 | /** |
| 2 | - * Copyright (c) 2025 Huawei Technologies Co., Ltd. | 2 | + * Copyright (c) 2026 Huawei Technologies Co., Ltd. |
| 3 | * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | 3 | * This program is free software, you can redistribute it and/or modify it under the terms and conditions of |
| 4 | * CANN Open Software License Agreement Version 2.0 (the "License"). | 4 | * CANN Open Software License Agreement Version 2.0 (the "License"). |
| 5 | * Please refer to the License for details. You may not use this file except in compliance with the License. | 5 | * Please refer to the License for details. You may not use this file except in compliance with the License. |
| @@ -24,6 +24,7 @@ | |||
| 24 | 24 | ||
| 25 | 25 | ||
| 26 | 26 | ||
| 27 | + | ||
| 27 | 28 | ||
| 28 | 29 | ||
| 29 | 30 | ||
| @@ -74,8 +75,8 @@ static inline bool CheckProbability(double prob) | |||
| 74 | 75 | ||
| 75 | static inline bool CheckSocVersionIsSupportDSA(void) | 76 | static inline bool CheckSocVersionIsSupportDSA(void) |
| 76 | { | 77 | { |
| 77 | - return GetCurrentPlatformInfo().GetSocVersion() >= SocVersion::ASCEND910B && | 78 | + auto curArch = GetCurrentPlatformInfo().GetCurNpuArch(); |
| 78 | - GetCurrentPlatformInfo().GetSocVersion() <= SocVersion::ASCEND910E; | 79 | + return curArch == NpuArch::DAV_2201 || IsRegBase(curArch); |
| 79 | } | 80 | } |
| 80 | 81 | ||
| 81 | static inline int64_t InferDSAOutShape(const aclIntArray* shape) | 82 | static inline int64_t InferDSAOutShape(const aclIntArray* shape) |
| @@ -164,11 +165,32 @@ static aclScalar* CreateDropout(float prob, op::DataType dtype, aclOpExecutor* e | |||
| 164 | } | 165 | } |
| 165 | } | 166 | } |
| 166 | 167 | ||
| 168 | +static const aclTensor* ComputeStatelessMask( | ||
| 169 | + const aclIntArray* shape, double prob, int64_t seed, int64_t offset, aclOpExecutor* executor) | ||
| 170 | +{ | ||
| 171 | + FVector<int64_t> seedVector = {seed}; | ||
| 172 | + aclIntArray* seedList = executor->AllocIntArray(seedVector.data(), seedVector.size()); | ||
| 173 | + auto seedTensor = executor->ConvertToTensor(seedList, op::DataType::DT_INT64); | ||
| 174 | + | ||
| 175 | + FVector<int64_t> seed1Vector = {0}; | ||
| 176 | + aclIntArray* seed1List = executor->AllocIntArray(seed1Vector.data(), seed1Vector.size()); | ||
| 177 | + auto seed1Tensor = executor->ConvertToTensor(seed1List, op::DataType::DT_INT64); | ||
| 178 | + | ||
| 179 | + FVector<int64_t> offsetVector = {0, offset}; | ||
| 180 | + aclIntArray* offsetList = executor->AllocIntArray(offsetVector.data(), offsetVector.size()); | ||
| 181 | + auto offsetTensor = executor->ConvertToTensor(offsetList, op::DataType::DT_INT64); | ||
| 182 | + | ||
| 183 | + auto probTensor = | ||
| 184 | + executor->ConvertToTensor(executor->AllocScalar(static_cast<float>(1 - prob)), op::DataType::DT_FLOAT); | ||
| 185 | + | ||
| 186 | + return l0op::StatelessDropoutGenMask(shape, probTensor, seedTensor, seed1Tensor, offsetTensor, executor); | ||
| 187 | +} | ||
| 188 | + | ||
| 167 | static const aclTensor* ComputeMask( | 189 | static const aclTensor* ComputeMask( |
| 168 | const aclIntArray* shape, double prob, int64_t seed, int64_t offset, aclTensor* out, aclOpExecutor* executor) | 190 | const aclIntArray* shape, double prob, int64_t seed, int64_t offset, aclTensor* out, aclOpExecutor* executor) |
| 169 | { | 191 | { |
| 170 | - auto socversion = GetCurrentPlatformInfo().GetSocVersion(); | 192 | + auto curArch = GetCurrentPlatformInfo().GetCurNpuArch(); |
| 171 | - if (socversion >= SocVersion::ASCEND910B && socversion <= SocVersion::ASCEND910_93) { | 193 | + if (curArch == NpuArch::DAV_2201) { |
| 172 | op::DataType dtype = out->GetDataType(); | 194 | op::DataType dtype = out->GetDataType(); |
| 173 | if (dtype == op::DataType::DT_UINT8) { | 195 | if (dtype == op::DataType::DT_UINT8) { |
| 174 | dtype = op::DataType::DT_FLOAT; | 196 | dtype = op::DataType::DT_FLOAT; |
| @@ -184,37 +206,21 @@ static const aclTensor* ComputeMask( | |||
| 184 | l0op::DSAGenBitMask(shapeSize * FLOAT_BIT_NUMBER, seed, offset, dropout, outTensorTemp, executor); | 206 | l0op::DSAGenBitMask(shapeSize * FLOAT_BIT_NUMBER, seed, offset, dropout, outTensorTemp, executor); |
| 185 | return out; | 207 | return out; |
| 186 | } else { | 208 | } else { |
| 187 | - FVector<int64_t> seedVector = {seed}; | 209 | + return ComputeStatelessMask(shape, prob, seed, offset, executor); |
| 188 | - auto seedTensor = executor->ConvertToTensor(seedVector.data(), seedVector.size(), op::DataType::DT_INT64); | ||
| 189 | - FVector<int64_t> seed1Vector = {0}; | ||
| 190 | - auto seed1Tensor = executor->ConvertToTensor(seed1Vector.data(), seed1Vector.size(), op::DataType::DT_INT64); | ||
| 191 | - FVector<int64_t> offsetVector = {0, offset}; | ||
| 192 | - auto offsetTensor = executor->ConvertToTensor(offsetVector.data(), offsetVector.size(), op::DataType::DT_INT64); | ||
| 193 | - FVector<float> probVector = {static_cast<float>(1 - prob)}; | ||
| 194 | - auto probTensor = executor->ConvertToTensor(probVector.data(), probVector.size(), DataType::DT_FLOAT); | ||
| 195 | - return l0op::StatelessDropoutGenMask(shape, probTensor, seedTensor, seed1Tensor, offsetTensor, executor); | ||
| 196 | } | 210 | } |
| 197 | } | 211 | } |
| 198 | 212 | ||
| 199 | static const aclTensor* ComputeMaskV2( | 213 | static const aclTensor* ComputeMaskV2( |
| 200 | const aclIntArray* shape, double prob, int64_t seed, int64_t offset, op::DataType dtype, aclOpExecutor* executor) | 214 | const aclIntArray* shape, double prob, int64_t seed, int64_t offset, op::DataType dtype, aclOpExecutor* executor) |
| 201 | { | 215 | { |
| 202 | - auto socversion = GetCurrentPlatformInfo().GetSocVersion(); | 216 | + auto curArch = GetCurrentPlatformInfo().GetCurNpuArch(); |
| 203 | - if (socversion >= SocVersion::ASCEND910B && socversion <= SocVersion::ASCEND910_93) { | 217 | + if (curArch == NpuArch::DAV_2201) { |
| 204 | int64_t shapeSize = InferDSAOutShapeV2(shape); | 218 | int64_t shapeSize = InferDSAOutShapeV2(shape); |
| 205 | auto dropout = CreateDropout(static_cast<float>(prob), dtype, executor); | 219 | auto dropout = CreateDropout(static_cast<float>(prob), dtype, executor); |
| 206 | CHECK_RET(dropout != nullptr, nullptr); | 220 | CHECK_RET(dropout != nullptr, nullptr); |
| 207 | return l0op::DSAGenBitMask(shapeSize * UINT8_BIT_NUMBER, seed, offset, dropout, executor); | 221 | return l0op::DSAGenBitMask(shapeSize * UINT8_BIT_NUMBER, seed, offset, dropout, executor); |
| 208 | } else { | 222 | } else { |
| 209 | - FVector<int64_t> seedVector = {seed}; | 223 | + return ComputeStatelessMask(shape, prob, seed, offset, executor); |
| 210 | - auto seedTensor = executor->ConvertToTensor(seedVector.data(), seedVector.size(), op::DataType::DT_INT64); | ||
| 211 | - FVector<int64_t> seed1Vector = {0}; | ||
| 212 | - auto seed1Tensor = executor->ConvertToTensor(seed1Vector.data(), seed1Vector.size(), op::DataType::DT_INT64); | ||
| 213 | - FVector<int64_t> offsetVector = {0, offset}; | ||
| 214 | - auto offsetTensor = executor->ConvertToTensor(offsetVector.data(), offsetVector.size(), op::DataType::DT_INT64); | ||
| 215 | - FVector<float> probVector = {static_cast<float>(1 - prob)}; | ||
| 216 | - auto probTensor = executor->ConvertToTensor(probVector.data(), probVector.size(), DataType::DT_FLOAT); | ||
| 217 | - return l0op::StatelessDropoutGenMask(shape, probTensor, seedTensor, seed1Tensor, offsetTensor, executor); | ||
| 218 | } | 224 | } |
| 219 | } | 225 | } |
| 220 | 226 | ||
| @@ -242,13 +248,28 @@ static const aclTensor* ComputeMaskV2Tensor( | |||
| 242 | const aclIntArray* shape, double prob, const aclTensor* seedTensor, const aclTensor* offsetTensor, int64_t offset, | 248 | const aclIntArray* shape, double prob, const aclTensor* seedTensor, const aclTensor* offsetTensor, int64_t offset, |
| 243 | op::DataType dtype, aclOpExecutor* executor) | 249 | op::DataType dtype, aclOpExecutor* executor) |
| 244 | { | 250 | { |
| 245 | - int64_t shapeSize = InferDSAOutShapeV2(shape); | 251 | + if (IsRegBase()) { |
| 246 | - auto dropout = CreateDropout(static_cast<float>(prob), dtype, executor); | 252 | + FVector<int64_t> seed1Vector = {0}; |
| 247 | - CHECK_RET(dropout != nullptr, nullptr); | 253 | + aclIntArray* seed1List = executor->AllocIntArray(seed1Vector.data(), seed1Vector.size()); |
| 248 | - auto concatTensor = ProcessOffsetTensor(offsetTensor, offset, executor); | 254 | + auto seed1Tensor = executor->ConvertToTensor(seed1List, op::DataType::DT_INT64); |
| 249 | - CHECK_RET(concatTensor != nullptr, nullptr); | ||
| 250 | 255 | ||
| 251 | - return l0op::DSAGenBitMaskTensor(shapeSize * UINT8_BIT_NUMBER, seedTensor, concatTensor, dropout, executor); | 256 | + FVector<int64_t> offsetVector{0, static_cast<int64_t>(offset)}; |
| 257 | + aclIntArray* offsetList = executor->AllocIntArray(offsetVector.data(), 2); | ||
| 258 | + auto tmpTensor = executor->ConvertToTensor(offsetList, op::DataType::DT_INT64); | ||
| 259 | + auto offset1Tensor = l0op::Add(offsetTensor, tmpTensor, executor); | ||
| 260 | + CHECK_RET(offset1Tensor != nullptr, nullptr); | ||
| 261 | + | ||
| 262 | + auto probTensor = executor->ConvertToTensor(executor->AllocScalar(static_cast<float>(1 - prob)), op::DataType::DT_FLOAT); | ||
| 263 | + | ||
| 264 | + return l0op::StatelessDropoutGenMask(shape, probTensor, seedTensor, seed1Tensor, offset1Tensor, executor); | ||
| 265 | + } else { | ||
| 266 | + int64_t shapeSize = InferDSAOutShapeV2(shape); | ||
| 267 | + auto dropout = CreateDropout(static_cast<float>(prob), dtype, executor); | ||
| 268 | + CHECK_RET(dropout != nullptr, nullptr); | ||
| 269 | + auto concatTensor = ProcessOffsetTensor(offsetTensor, offset, executor); | ||
| 270 | + CHECK_RET(concatTensor != nullptr, nullptr); | ||
| 271 | + return l0op::DSAGenBitMaskTensor(shapeSize * UINT8_BIT_NUMBER, seedTensor, concatTensor, dropout, executor); | ||
| 272 | + } | ||
| 252 | } | 273 | } |
| 253 | 274 | ||
| 254 | static bool IsDoubleEqual(double f1, double f2) | 275 | static bool IsDoubleEqual(double f1, double f2) |
| @@ -377,7 +398,6 @@ aclnnStatus aclnnDropoutGenMaskV2TensorGetWorkspaceSize( | |||
| 377 | auto uniqueExecutor = CREATE_EXECUTOR(); | 398 | auto uniqueExecutor = CREATE_EXECUTOR(); |
| 378 | CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR); | 399 | CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR); |
| 379 | 400 | ||
| 380 | - CHECK_RET(CheckSocVersionIsSupportDSA(), ACLNN_ERR_PARAM_INVALID); | ||
| 381 | // 固定写法，参数检查 | 401 | // 固定写法，参数检查 |
| 382 | auto ret = CheckParams(shape, prob, out); | 402 | auto ret = CheckParams(shape, prob, out); |
| 383 | CHECK_RET(ret == ACLNN_SUCCESS, ret); | 403 | CHECK_RET(ret == ACLNN_SUCCESS, ret); |
Mrandom/stateless_drop_out_gen_mask/op_host/arch35/stateless_drop_out_gen_mask_tiling_arch35.cpp+24-32
| @@ -1,5 +1,5 @@ | |||
| 1 | /** | 1 | /** |
| 2 | - * Copyright (c) 2025 Huawei Technologies Co., Ltd. | 2 | + * Copyright (c) 2026 Huawei Technologies Co., Ltd. |
| 3 | * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | 3 | * This program is free software, you can redistribute it and/or modify it under the terms and conditions of |
| 4 | * CANN Open Software License Agreement Version 2.0 (the "License"). | 4 | * CANN Open Software License Agreement Version 2.0 (the "License"). |
| 5 | * Please refer to the License for details. You may not use this file except in compliance with the License. | 5 | * Please refer to the License for details. You may not use this file except in compliance with the License. |
| @@ -77,37 +77,29 @@ OpTilingConfig StatelessDropOutGenMaskTiling::BuildOpConfig() | |||
| 77 | return ge::GRAPH_SUCCESS; | 77 | return ge::GRAPH_SUCCESS; |
| 78 | }; | 78 | }; |
| 79 | 79 | ||
| 80 | - // 获取key[2]：从attr1(seed) counter[4] attr(seed2) | 80 | + // key/counter 由 kernel 从 GM 直接读取，tiling 侧不再做 D2H memcpy |
| 81 | - config.getKeyAndCounter = [](gert::TilingContext* ctx, uint32_t key[2], uint32_t counter[4]) -> ge::graphStatus { | 81 | + // 通过 key[] 传递元数据给 kernel：key[0]=offsetElemCount, key[1]=seedByteSize |
| 82 | - //check offset | 82 | + config.getKeyAndCounter = []([[maybe_unused]] gert::TilingContext* ctx, uint32_t key[2], |
| 83 | + uint32_t counter[4]) -> ge::graphStatus { | ||
| 84 | + counter[0] = 0; counter[1] = 0; counter[2] = 0; counter[3] = 0; | ||
| 85 | + | ||
| 86 | + // offset element count (shape info, no D2H needed) | ||
| 83 | auto offsetTensor = ctx->GetRequiredInputTensor(IN_OFFSET_IDX); | 87 | auto offsetTensor = ctx->GetRequiredInputTensor(IN_OFFSET_IDX); |
| 84 | - OP_CHECK_NULL_WITH_CONTEXT(ctx, offsetTensor); | 88 | + uint32_t offsetElemCount = 2; |
| 85 | - auto offsetTensorSize = static_cast<int64_t>(offsetTensor->GetShapeSize()); | 89 | + if (offsetTensor != nullptr) { |
| 86 | - OP_CHECK_IF((offsetTensorSize != 1) && (offsetTensorSize != 2), | 90 | + offsetElemCount = static_cast<uint32_t>(offsetTensor->GetShapeSize()); |
| 87 | - OP_LOGE(ctx->GetNodeName(), "input offset shape_size should be 1 or 2, but got %ld.", offsetTensorSize), | ||
| 88 | - return ge::GRAPH_FAILED); | ||
| 89 | - // get input value of seed & offset. | ||
| 90 | - gert::Shape inputSeed_; | ||
| 91 | - gert::Shape inputOffset_; | ||
| 92 | - OP_CHECK_IF(ExtractTensorValue(ctx, IN_SEED_IDX, inputSeed_) != ge::GRAPH_SUCCESS, | ||
| 93 | - OP_LOGE(ctx->GetNodeName(), "get const shape of seed failed"), return ge::GRAPH_FAILED); | ||
| 94 | - OP_CHECK_IF(ExtractTensorValue(ctx, IN_OFFSET_IDX, inputOffset_) != ge::GRAPH_SUCCESS, | ||
| 95 | - OP_LOGE(ctx->GetNodeName(), "get const shape of offset failed"), return ge::GRAPH_FAILED); | ||
| 96 | - OP_LOGD(ctx->GetNodeName(), "const seed = %s, const offset = %s.", Ops::Base::ToString(inputSeed_).c_str(), | ||
| 97 | - Ops::Base::ToString(inputOffset_).c_str()); | ||
| 98 | - int64_t keyTemp = static_cast<int64_t>(inputSeed_[0]); | ||
| 99 | - std::vector<int64_t> counterTemp; | ||
| 100 | - if (offsetTensorSize == 1) { | ||
| 101 | - counterTemp = { 0, inputOffset_[0] }; | ||
| 102 | - } else { | ||
| 103 | - counterTemp = { inputOffset_[0], inputOffset_[1] }; | ||
| 104 | } | 91 | } |
| 105 | - key[0] = static_cast<int32_t>(keyTemp); | 92 | + key[0] = offsetElemCount; |
| 106 | - key[1] = static_cast<int32_t>(keyTemp >> RIGHT_SHIFT_NUM); // 32 for lower 32 bits | 93 | + |
| 107 | - counter[0] = static_cast<int32_t>(counterTemp[0]); | 94 | + // seed byte size: 4 for INT32, 8 for INT64 (dtype info, no D2H needed) |
| 108 | - counter[1] = static_cast<int32_t>(counterTemp[0] >> RIGHT_SHIFT_NUM); // 32 for lower 32 bits | 95 | + auto seedDesc = ctx->GetInputDesc(IN_SEED_IDX); |
| 109 | - counter[2] = static_cast<int32_t>(counterTemp[1]); | 96 | + uint32_t seedByteSize = 8; |
| 110 | - counter[3] = static_cast<int32_t>(counterTemp[1] >> RIGHT_SHIFT_NUM); // 32 for lower 32 bits | 97 | + if (seedDesc != nullptr) { |
| 98 | + auto seedDtype = seedDesc->GetDataType(); | ||
| 99 | + seedByteSize = (seedDtype == ge::DT_INT32) ? 4 : 8; | ||
| 100 | + } | ||
| 101 | + key[1] = seedByteSize; | ||
| 102 | + | ||
| 111 | return ge::GRAPH_SUCCESS; | 103 | return ge::GRAPH_SUCCESS; |
| 112 | }; | 104 | }; |
| 113 | 105 | ||
| @@ -142,5 +134,5 @@ static ge::graphStatus TilingStatelessDropOutGenMask(gert::TilingContext* tiling | |||
| 142 | IMPL_OP_OPTILING(StatelessDropOutGenMask) | 134 | IMPL_OP_OPTILING(StatelessDropOutGenMask) |
| 143 | .Tiling(TilingStatelessDropOutGenMask) | 135 | .Tiling(TilingStatelessDropOutGenMask) |
| 144 | .TilingParse<RandomOperatorCompileInfo>(TilingPrepare4StatelessDropOutGenMaskTiling) | 136 | .TilingParse<RandomOperatorCompileInfo>(TilingPrepare4StatelessDropOutGenMaskTiling) |
| 145 | - .TilingInputsDataDependency({ IN_SHAPE_IDX, IN_PROB_IDX, IN_SEED_IDX, IN_SEED1_IDX, IN_OFFSET_IDX }); | 137 | + .TilingInputsDataDependency({ IN_SHAPE_IDX, IN_PROB_IDX }); |
| 146 | -} // namespace optiling | 138 | +} // namespace optiling |
| @@ -1,5 +1,5 @@ | |||
| 1 | /** | 1 | /** |
| 2 | - * Copyright (c) 2025 Huawei Technologies Co., Ltd. | 2 | + * Copyright (c) 2026 Huawei Technologies Co., Ltd. |
| 3 | * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | 3 | * This program is free software, you can redistribute it and/or modify it under the terms and conditions of |
| 4 | * CANN Open Software License Agreement Version 2.0 (the "License"). | 4 | * CANN Open Software License Agreement Version 2.0 (the "License"). |
| 5 | * Please refer to the License for details. You may not use this file except in compliance with the License. | 5 | * Please refer to the License for details. You may not use this file except in compliance with the License. |
| @@ -37,7 +37,7 @@ template <typename T> | |||
| 37 | class StatelessDropOutGenMaskPt : public RandomKernelBaseOp{ | 37 | class StatelessDropOutGenMaskPt : public RandomKernelBaseOp{ |
| 38 | public: | 38 | public: |
| 39 | __aicore__ inline StatelessDropOutGenMaskPt(TPipe* pipe, const RandomUnifiedTilingDataStruct* __restrict tilingData) : RandomKernelBaseOp(tilingData),pipe_(pipe){}; | 39 | __aicore__ inline StatelessDropOutGenMaskPt(TPipe* pipe, const RandomUnifiedTilingDataStruct* __restrict tilingData) : RandomKernelBaseOp(tilingData),pipe_(pipe){}; |
| 40 | - __aicore__ inline void Init(GM_ADDR shape, GM_ADDR prob, GM_ADDR y); | 40 | + __aicore__ inline void Init(GM_ADDR shape, GM_ADDR prob, GM_ADDR seed, GM_ADDR offset, GM_ADDR y); |
| 41 | __aicore__ inline void Process(); | 41 | __aicore__ inline void Process(); |
| 42 | 42 | ||
| 43 | protected: | 43 | protected: |
| @@ -56,6 +56,7 @@ private: | |||
| 56 | TPipe* pipe_; | 56 | TPipe* pipe_; |
| 57 | GlobalTensor<T> probInputGm_; | 57 | GlobalTensor<T> probInputGm_; |
| 58 | GlobalTensor<uint8_t> outputGm_; | 58 | GlobalTensor<uint8_t> outputGm_; |
| 59 | + GlobalTensor<uint64_t> counterGm_; // offset tensor 的 GM 访问句柄 | ||
| 59 | 60 | ||
| 60 | TQue<QuePosition::VECOUT, BUFFER_NUM> outQueY_; | 61 | TQue<QuePosition::VECOUT, BUFFER_NUM> outQueY_; |
| 61 | TBuf<QuePosition::VECCALC> philoxQueBuf_; | 62 | TBuf<QuePosition::VECCALC> philoxQueBuf_; |
| @@ -74,10 +75,47 @@ private: | |||
| 74 | 75 | ||
| 75 | template <typename T> | 76 | template <typename T> |
| 76 | __aicore__ inline void StatelessDropOutGenMaskPt<T>::Init( | 77 | __aicore__ inline void StatelessDropOutGenMaskPt<T>::Init( |
| 77 | - GM_ADDR shape, GM_ADDR prob, GM_ADDR y) | 78 | + GM_ADDR shape, GM_ADDR prob, GM_ADDR seed, GM_ADDR offset, GM_ADDR y) |
| 78 | { | 79 | { |
| 79 | - //Parsetiling data | 80 | + // VarsInit 从 tiling_ 拷贝 key/counter（此时 tiling 中 key 存元数据，counter 已置零） |
| 80 | VarsInit(); | 81 | VarsInit(); |
| 82 | + | ||
| 83 | + // 读取 tiling 编码的元数据：key_[0]=offsetElemCount, key_[1]=seedByteSize | ||
| 84 | + uint32_t offsetElemCount = key_[0]; | ||
| 85 | + uint32_t seedByteSize = key_[1]; | ||
| 86 | + | ||
| 87 | + // 从 GM 直接读取 seed → key_[] | ||
| 88 | + constexpr uint32_t SHIFT_BITS = 32; | ||
| 89 | + int64_t keyVal = 0; | ||
| 90 | + if (seedByteSize == sizeof(int32_t)) { | ||
| 91 | + GlobalTensor<int32_t> seedGm; | ||
| 92 | + seedGm.SetGlobalBuffer(reinterpret_cast<__gm__ int32_t*>(seed), 1); | ||
| 93 | + keyVal = static_cast<int64_t>(seedGm(0)); | ||
| 94 | + } else { | ||
| 95 | + GlobalTensor<int64_t> seedGm; | ||
| 96 | + seedGm.SetGlobalBuffer(reinterpret_cast<__gm__ int64_t*>(seed), 1); | ||
| 97 | + keyVal = seedGm(0); | ||
| 98 | + } | ||
| 99 | + key_[0] = static_cast<uint32_t>(keyVal); // 截断取低32位，符号位不影响 | ||
| 100 | + // 先转 uint64_t 再右移：keyVal 是有符号 int64_t，负数直接右移是算术右移（高位补1）， | ||
| 101 | + // 转成无符号后右移才是逻辑右移（高位补0），保证拿到原始数据的高32位 | ||
| 102 | + key_[1] = static_cast<uint32_t>(static_cast<uint64_t>(keyVal) >> SHIFT_BITS); | ||
| 103 | + | ||
| 104 | + // 从 GM 直接读取 offset → counter_[] | ||
| 105 | + uint64_t counterVal0 = 0; | ||
| 106 | + uint64_t counterVal1 = 0; | ||
| 107 | + counterGm_.SetGlobalBuffer(reinterpret_cast<__gm__ uint64_t*>(offset), offsetElemCount); | ||
| 108 | + if (offsetElemCount == 1) { | ||
| 109 | + counterVal1 = counterGm_(0); | ||
| 110 | + } else { | ||
| 111 | + counterVal0 = counterGm_(0); | ||
| 112 | + counterVal1 = counterGm_(1); | ||
| 113 | + } | ||
| 114 | + counter_[0] = static_cast<uint32_t>(counterVal0); | ||
| 115 | + counter_[1] = static_cast<uint32_t>(counterVal0 >> SHIFT_BITS); | ||
| 116 | + counter_[2] = static_cast<uint32_t>(counterVal1); | ||
| 117 | + counter_[3] = static_cast<uint32_t>(counterVal1 >> SHIFT_BITS); | ||
| 118 | + | ||
| 81 | singleBufferProNum = tiling_->singleBufferSize; | 119 | singleBufferProNum = tiling_->singleBufferSize; |
| 82 | blockOffset_ = blockIdx_ * tiling_->normalCoreProNum; | 120 | blockOffset_ = blockIdx_ * tiling_->normalCoreProNum; |
| 83 | if (curCoreProNum_ < singleBufferProNum) { | 121 | if (curCoreProNum_ < singleBufferProNum) { |
| @@ -1,5 +1,5 @@ | |||
| 1 | /** | 1 | /** |
| 2 | - * Copyright (c) 2025 Huawei Technologies Co., Ltd. | 2 | + * Copyright (c) 2026 Huawei Technologies Co., Ltd. |
| 3 | * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | 3 | * This program is free software, you can redistribute it and/or modify it under the terms and conditions of |
| 4 | * CANN Open Software License Agreement Version 2.0 (the "License"). | 4 | * CANN Open Software License Agreement Version 2.0 (the "License"). |
| 5 | * Please refer to the License for details. You may not use this file except in compliance with the License. | 5 | * Please refer to the License for details. You may not use this file except in compliance with the License. |
| @@ -30,7 +30,7 @@ extern "C" __global__ __aicore__ void stateless_drop_out_gen_mask( | |||
| 30 | AscendC::TPipe pipe; | 30 | AscendC::TPipe pipe; |
| 31 | if (TILING_KEY_IS(STATELESS_DROP_OUT_GEN_MASK_DEFAULT_TILING_KEY)) { | 31 | if (TILING_KEY_IS(STATELESS_DROP_OUT_GEN_MASK_DEFAULT_TILING_KEY)) { |
| 32 | StatelessDropOutGenMask::StatelessDropOutGenMaskPt<DTYPE_PROB> op(&pipe, &tilingData); | 32 | StatelessDropOutGenMask::StatelessDropOutGenMaskPt<DTYPE_PROB> op(&pipe, &tilingData); |
| 33 | - op.Init(shape, prob, y); | 33 | + op.Init(shape, prob, seed, offset, y); |
| 34 | op.Process(); | 34 | op.Process(); |
| 35 | } | 35 | } |
| 36 | } | 36 | } |
Mrandom/stateless_drop_out_gen_mask/tests/ut/op_host/arch35/test_stateless_drop_out_gen_mask_tiling_arch35.cpp+9-9
| @@ -58,7 +58,7 @@ TEST_F(StatelessDropOutGenMaskTilingTest, stateless_drop_out_gen_mask_test_01) | |||
| 58 | }, | 58 | }, |
| 59 | &compileInfo); | 59 | &compileInfo); |
| 60 | uint64_t expectTilingKey = 100; | 60 | uint64_t expectTilingKey = 100; |
| 61 | - string expectTilingData = "64 256 256 13056 2 0 8 16384 0 0 0 "; | 61 | + string expectTilingData = "64 256 256 13056 34359738369 0 0 16384 0 0 0 0 "; |
| 62 | std::vector<size_t> expectWorkspaces = {0}; | 62 | std::vector<size_t> expectWorkspaces = {0}; |
| 63 | ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces); | 63 | ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces); |
| 64 | } | 64 | } |
| @@ -86,7 +86,7 @@ TEST_F(StatelessDropOutGenMaskTilingTest, stateless_drop_out_gen_mask_test_02) | |||
| 86 | }, | 86 | }, |
| 87 | &compileInfo); | 87 | &compileInfo); |
| 88 | uint64_t expectTilingKey = 100; | 88 | uint64_t expectTilingKey = 100; |
| 89 | - string expectTilingData = "54 768 256 13056 24 0 36 40960 0 0 0 "; | 89 | + string expectTilingData = "54 768 256 13056 17179869185 0 0 40960 0 0 0 0 "; |
| 90 | std::vector<size_t> expectWorkspaces = {0}; | 90 | std::vector<size_t> expectWorkspaces = {0}; |
| 91 | ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces); | 91 | ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces); |
| 92 | } | 92 | } |
| @@ -111,10 +111,10 @@ TEST_F(StatelessDropOutGenMaskTilingTest, stateless_drop_out_gen_mask_test_03) | |||
| 111 | }, | 111 | }, |
| 112 | { | 112 | { |
| 113 | {{{ 5120 }, { 5120 } },ge::DT_UINT8, ge::FORMAT_ND}, | 113 | {{{ 5120 }, { 5120 } },ge::DT_UINT8, ge::FORMAT_ND}, |
| 114 | - }, | 114 | + }, |
| 115 | &compileInfo); | 115 | &compileInfo); |
| 116 | uint64_t expectTilingKey = 100; | 116 | uint64_t expectTilingKey = 100; |
| 117 | - string expectTilingData = "54 768 256 13056 24 0 36 40960 0 0 0 "; | 117 | + string expectTilingData = "54 768 256 13056 17179869185 0 0 40960 0 0 0 0 "; |
| 118 | std::vector<size_t> expectWorkspaces = {0}; | 118 | std::vector<size_t> expectWorkspaces = {0}; |
| 119 | ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces); | 119 | ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces); |
| 120 | } | 120 | } |
| @@ -139,10 +139,10 @@ TEST_F(StatelessDropOutGenMaskTilingTest, stateless_drop_out_gen_mask_test_04) | |||
| 139 | }, | 139 | }, |
| 140 | { | 140 | { |
| 141 | {{{ 5120 }, { 5120 } },ge::DT_UINT8, ge::FORMAT_ND}, | 141 | {{{ 5120 }, { 5120 } },ge::DT_UINT8, ge::FORMAT_ND}, |
| 142 | - }, | 142 | + }, |
| 143 | &compileInfo); | 143 | &compileInfo); |
| 144 | uint64_t expectTilingKey = 100; | 144 | uint64_t expectTilingKey = 100; |
| 145 | - string expectTilingData = "54 768 256 13056 24 0 36 40960 0 0 0 "; | 145 | + string expectTilingData = "54 768 256 13056 17179869185 0 0 40960 0 0 0 0 "; |
| 146 | std::vector<size_t> expectWorkspaces = {0}; | 146 | std::vector<size_t> expectWorkspaces = {0}; |
| 147 | ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces); | 147 | ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces); |
| 148 | } | 148 | } |
| @@ -170,7 +170,7 @@ TEST_F(StatelessDropOutGenMaskTilingTest, stateless_drop_out_gen_mask_test_05) | |||
| 170 | }, | 170 | }, |
| 171 | &compileInfo); | 171 | &compileInfo); |
| 172 | uint64_t expectTilingKey = 100; | 172 | uint64_t expectTilingKey = 100; |
| 173 | - string expectTilingData = "64 12800 12800 13056 16 0 36 819200 0 0 0 "; | 173 | + string expectTilingData = "64 12800 12800 13056 34359738370 0 0 819200 0 0 0 0 "; |
| 174 | std::vector<size_t> expectWorkspaces = {0}; | 174 | std::vector<size_t> expectWorkspaces = {0}; |
| 175 | ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces); | 175 | ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces); |
| 176 | } | 176 | } |
| @@ -198,7 +198,7 @@ TEST_F(StatelessDropOutGenMaskTilingTest, stateless_drop_out_gen_mask_test_06) | |||
| 198 | }, | 198 | }, |
| 199 | &compileInfo); | 199 | &compileInfo); |
| 200 | uint64_t expectTilingKey = 100; | 200 | uint64_t expectTilingKey = 100; |
| 201 | - string expectTilingData = "40 512 512 13056 24 0 36 20480 0 0 0 "; | 201 | + string expectTilingData = "40 512 512 13056 34359738370 0 0 20480 0 0 0 0 "; |
| 202 | std::vector<size_t> expectWorkspaces = {0}; | 202 | std::vector<size_t> expectWorkspaces = {0}; |
| 203 | ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces); | 203 | ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces); |
| 204 | } | 204 | } |
| @@ -226,7 +226,7 @@ TEST_F(StatelessDropOutGenMaskTilingTest, stateless_drop_out_gen_mask_test_07) | |||
| 226 | }, | 226 | }, |
| 227 | &compileInfo); | 227 | &compileInfo); |
| 228 | uint64_t expectTilingKey = 100; | 228 | uint64_t expectTilingKey = 100; |
| 229 | - string expectTilingData = "48 512 512 13056 24 0 36 24576 0 0 0 "; | 229 | + string expectTilingData = "48 512 512 13056 34359738370 0 0 24576 0 0 0 0 "; |
| 230 | std::vector<size_t> expectWorkspaces = {0}; | 230 | std::vector<size_t> expectWorkspaces = {0}; |
| 231 | ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces); | 231 | ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces); |
| 232 | } | 232 | } |
Drandom/stateless_drop_out_gen_mask/tests/ut/op_kernel/CMakeLists.txt+0-26
| @@ -1,26 +0,0 @@ | |||
| 1 | -# ---------------------------------------------------------------------------- | ||
| 2 | -# Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | -# This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | -# CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | -# Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | -# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | -# See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | -# ---------------------------------------------------------------------------- | ||
| 10 | - | ||
| 11 | -if(UT_TEST_ALL OR OP_KERNEL_UT) | ||
| 12 | - set(KERNEL_STAGING_DIR ${CMAKE_CURRENT_BINARY_DIR}/kernel_dep_staging) | ||
| 13 | - file(MAKE_DIRECTORY ${KERNEL_STAGING_DIR}/stateless_drop_out_gen_mask/arch35) | ||
| 14 | - execute_process(COMMAND ${CMAKE_COMMAND} -E create_symlink | ||
| 15 | - ${PROJECT_SOURCE_DIR}/random/random_common/op_kernel | ||
| 16 | - ${KERNEL_STAGING_DIR}/random_common) | ||
| 17 | - | ||
| 18 | - set(stateless_drop_out_gen_mask_tiling_files | ||
| 19 | - ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/arch35/stateless_drop_out_gen_mask_tiling_arch35.cpp | ||
| 20 | - ${CMAKE_CURRENT_SOURCE_DIR}/../../../../random_common/op_host/arch35/random_tiling_arch35.cpp) | ||
| 21 | - AddOpTestCase( | ||
| 22 | - stateless_drop_out_gen_mask | ||
| 23 | - "ascend950" | ||
| 24 | - "-DDTYPE_PROB=float -DTestUtDefaultTilingStruct=RandomUnifiedTilingDataStruct -I${KERNEL_STAGING_DIR}/stateless_drop_out_gen_mask/arch35" | ||
| 25 | - "${stateless_drop_out_gen_mask_tiling_files}") | ||
| 26 | -endif() | ||
Drandom/stateless_drop_out_gen_mask/tests/ut/op_kernel/test_stateless_drop_out_gen_mask.cpp+0-80
| @@ -1,80 +0,0 @@ | |||
| 1 | -/** | ||
| 2 | - * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | - * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | - * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | - * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | - * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | - * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | - * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | - */ | ||
| 10 | - | ||
| 11 | - | ||
| 12 | - | ||
| 13 | - | ||
| 14 | - | ||
| 15 | - | ||
| 16 | - | ||
| 17 | -extern "C" __global__ __aicore__ void stateless_drop_out_gen_mask( | ||
| 18 | - GM_ADDR shape, GM_ADDR prob, GM_ADDR seed, GM_ADDR seed1, GM_ADDR offset, GM_ADDR y, GM_ADDR workspace, | ||
| 19 | - GM_ADDR tiling); | ||
| 20 | - | ||
| 21 | -namespace { | ||
| 22 | -constexpr uint32_t kNumBlocks = 1; | ||
| 23 | -constexpr uint64_t kTilingKey = 100; | ||
| 24 | -constexpr int64_t kElementCount = 256; | ||
| 25 | - | ||
| 26 | -inline size_t Align32(size_t size) | ||
| 27 | -{ | ||
| 28 | - return (size + 31U) / 32U * 32U; | ||
| 29 | -} | ||
| 30 | -} // namespace | ||
| 31 | - | ||
| 32 | -class StatelessDropOutGenMaskKernelTest : public testing::Test { | ||
| 33 | -}; | ||
| 34 | - | ||
| 35 | -TEST_F(StatelessDropOutGenMaskKernelTest, smoke_float) | ||
| 36 | -{ | ||
| 37 | - auto* shape = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(2 * sizeof(int32_t)))); | ||
| 38 | - auto* prob = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(float)))); | ||
| 39 | - auto* seed = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int64_t)))); | ||
| 40 | - auto* seed1 = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int64_t)))); | ||
| 41 | - auto* offset = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(int64_t)))); | ||
| 42 | - auto* y = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(kElementCount * sizeof(uint8_t)))); | ||
| 43 | - auto* workspace = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(1024 * 1024))); | ||
| 44 | - auto* tiling = static_cast<uint8_t*>(AscendC::GmAlloc(Align32(sizeof(RandomUnifiedTilingDataStruct)))); | ||
| 45 | - | ||
| 46 | - std::memset(y, 0, kElementCount * sizeof(uint8_t)); | ||
| 47 | - std::memset(tiling, 0, sizeof(RandomUnifiedTilingDataStruct)); | ||
| 48 | - reinterpret_cast<int32_t*>(shape)[0] = 16; | ||
| 49 | - reinterpret_cast<int32_t*>(shape)[1] = 16; | ||
| 50 | - *reinterpret_cast<float*>(prob) = 0.5f; | ||
| 51 | - *reinterpret_cast<int64_t*>(seed) = 42; | ||
| 52 | - *reinterpret_cast<int64_t*>(seed1) = 0; | ||
| 53 | - *reinterpret_cast<int64_t*>(offset) = 0; | ||
| 54 | - | ||
| 55 | - auto* tilingData = reinterpret_cast<RandomUnifiedTilingDataStruct*>(tiling); | ||
| 56 | - tilingData->usedCoreNum = kNumBlocks; | ||
| 57 | - tilingData->normalCoreProNum = kElementCount; | ||
| 58 | - tilingData->tailCoreProNum = kElementCount; | ||
| 59 | - tilingData->singleBufferSize = kElementCount; | ||
| 60 | - tilingData->key[0] = 42; | ||
| 61 | - tilingData->key[1] = 0; | ||
| 62 | - tilingData->counter[0] = 0; | ||
| 63 | - tilingData->counter[1] = 0; | ||
| 64 | - tilingData->counter[2] = 0; | ||
| 65 | - tilingData->counter[3] = 0; | ||
| 66 | - tilingData->outputSize = kElementCount; | ||
| 67 | - | ||
| 68 | - AscendC::SetKernelMode(KernelMode::AIV_MODE); | ||
| 69 | - ICPU_SET_TILING_KEY(kTilingKey); | ||
| 70 | - ICPU_RUN_KF(stateless_drop_out_gen_mask, kNumBlocks, shape, prob, seed, seed1, offset, y, workspace, tiling); | ||
| 71 | - | ||
| 72 | - AscendC::GmFree(shape); | ||
| 73 | - AscendC::GmFree(prob); | ||
| 74 | - AscendC::GmFree(seed); | ||
| 75 | - AscendC::GmFree(seed1); | ||
| 76 | - AscendC::GmFree(offset); | ||
| 77 | - AscendC::GmFree(y); | ||
| 78 | - AscendC::GmFree(workspace); | ||
| 79 | - AscendC::GmFree(tiling); | ||
| 80 | -} | ||