init comiit

This commit is contained in:
RubenCGomes 2025-10-06 11:35:12 +01:00
commit 6586f9d4c2
No known key found for this signature in database
GPG Key ID: 7791B4374E875525
26 changed files with 3179 additions and 0 deletions

8
.idea/.gitignore vendored Normal file
View File

@ -0,0 +1,8 @@
# Default ignored files
/shelf/
/workspace.xml
# Editor-based HTTP Client requests
/httpRequests/
# Datasource local storage ignored files
/dataSources/
/dataSources.local.xml

View File

@ -0,0 +1,6 @@
<?xml version="1.0" encoding="UTF-8"?>
<project version="4">
<component name="AgentMigrationStateService">
<option name="migrationStatus" value="COMPLETED" />
</component>
</project>

View File

@ -0,0 +1,6 @@
<?xml version="1.0" encoding="UTF-8"?>
<project version="4">
<component name="AskMigrationStateService">
<option name="migrationStatus" value="COMPLETED" />
</component>
</project>

View File

@ -0,0 +1,6 @@
<?xml version="1.0" encoding="UTF-8"?>
<project version="4">
<component name="Ask2AgentMigrationStateService">
<option name="migrationStatus" value="COMPLETED" />
</component>
</project>

View File

@ -0,0 +1,6 @@
<?xml version="1.0" encoding="UTF-8"?>
<project version="4">
<component name="EditMigrationStateService">
<option name="migrationStatus" value="COMPLETED" />
</component>
</project>

248
.idea/editor.xml Normal file
View File

@ -0,0 +1,248 @@
<?xml version="1.0" encoding="UTF-8"?>
<project version="4">
<component name="BackendCodeEditorSettings">
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CDeclarationWithImplicitIntType/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CommentTypo/@EntryIndexedValue" value="DO_NOT_SHOW" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=ConstevalIfIsAlwaysConstant/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppAbstractClassWithoutSpecifier/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppAbstractFinalClass/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppAbstractVirtualFunctionCallInCtor/@EntryIndexedValue" value="ERROR" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppAccessSpecifierWithNoDeclarations/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppAwaiterTypeIsNotClass/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppBooleanIncrementExpression/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppBoostFormatBadCode/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppBoostFormatLegacyCode/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppBoostFormatMixedArgs/@EntryIndexedValue" value="ERROR" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppBoostFormatTooFewArgs/@EntryIndexedValue" value="ERROR" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppBoostFormatTooManyArgs/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppCStyleCast/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppCVQualifierCanNotBeAppliedToReference/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppClassCanBeFinal/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppClassIsIncomplete/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppClassNeedsConstructorBecauseOfUninitializedMember/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppClassNeverUsed/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppCompileTimeConstantCanBeReplacedWithBooleanConstant/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppConceptNeverUsed/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppConditionalExpressionCanBeSimplified/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppConstParameterInDeclaration/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppConstValueFunctionReturnType/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppCoroutineCallResolveError/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDFAArrayIndexOutOfBounds/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDFAConstantConditions/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDFAConstantFunctionResult/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDFAConstantParameter/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDFADeletedPointer/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDFAEndlessLoop/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDFAInfiniteRecursion/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDFAInvalidatedMemory/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDFALocalValueEscapesFunction/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDFALocalValueEscapesScope/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDFALoopConditionNotUpdated/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDFAMemoryLeak/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDFANotInitializedField/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDFANullDereference/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDFATimeOver/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDFAUnreachableCode/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDFAUnreachableFunctionCall/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDFAUnreadVariable/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDFAUnusedValue/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDeclarationHidesLocal/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDeclarationHidesUncapturedLocal/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDeclarationSpecifierWithoutDeclarators/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDeclaratorDisambiguatedAsFunction/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDeclaratorNeverUsed/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDeclaratorUsedBeforeInitialization/@EntryIndexedValue" value="ERROR" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDefaultCaseNotHandledInSwitchStatement/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDefaultInitializationWithNoUserConstructor/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDefaultIsUsedAsIdentifier/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDefaultedSpecialMemberFunctionIsImplicitlyDeleted/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDeletingVoidPointer/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDependentTemplateWithoutTemplateKeyword/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDependentTypeWithoutTypenameKeyword/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDeprecatedEntity/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDeprecatedOverridenMethod/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDeprecatedRegisterStorageClassSpecifier/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDereferenceOperatorLimitExceeded/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDiscardedPostfixOperatorResult/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDoxygenSyntaxError/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDoxygenUndocumentedParameter/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppDoxygenUnresolvedReference/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppEmptyDeclaration/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppEnforceCVQualifiersOrder/@EntryIndexedValue" value="DO_NOT_SHOW" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppEnforceCVQualifiersPlacement/@EntryIndexedValue" value="DO_NOT_SHOW" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppEnforceDoStatementBraces/@EntryIndexedValue" value="DO_NOT_SHOW" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppEnforceForStatementBraces/@EntryIndexedValue" value="DO_NOT_SHOW" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppEnforceFunctionDeclarationStyle/@EntryIndexedValue" value="DO_NOT_SHOW" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppEnforceIfStatementBraces/@EntryIndexedValue" value="DO_NOT_SHOW" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppEnforceNestedNamespacesStyle/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppEnforceOverridingDestructorStyle/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppEnforceOverridingFunctionStyle/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppEnforceTypeAliasCodeStyle/@EntryIndexedValue" value="DO_NOT_SHOW" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppEnforceWhileStatementBraces/@EntryIndexedValue" value="DO_NOT_SHOW" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppEntityAssignedButNoRead/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppEntityUsedOnlyInUnevaluatedContext/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppEnumeratorNeverUsed/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppEqualOperandsInBinaryExpression/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppEvaluationFailure/@EntryIndexedValue" value="ERROR" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppExplicitSpecializationInNonNamespaceScope/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppExpressionWithoutSideEffects/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppFinalFunctionInFinalClass/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppFinalNonOverridingVirtualFunction/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppForLoopCanBeReplacedWithWhile/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppForwardEnumDeclarationWithoutUnderlyingType/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppFunctionDoesntReturnValue/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppFunctionIsNotImplemented/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppFunctionResultShouldBeUsed/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppFunctionalStyleCast/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppHeaderHasBeenAlreadyIncluded/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppHiddenFunction/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppHidingFunction/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppIdenticalOperandsInBinaryExpression/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppIfCanBeReplacedByConstexprIf/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppImplicitDefaultConstructorNotAvailable/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppIncompatiblePointerConversion/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppIncompleteSwitchStatement/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppInconsistentNaming/@EntryIndexedValue" value="DO_NOT_SHOW" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppIntegralToPointerConversion/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppInvalidLineContinuation/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppJoinDeclarationAndAssignment/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppLambdaCaptureNeverUsed/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppLocalVariableMayBeConst/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppLocalVariableMightNotBeInitialized/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppLocalVariableWithNonTrivialDtorIsNeverUsed/@EntryIndexedValue" value="DO_NOT_SHOW" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppLongFloat/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppMemberFunctionMayBeConst/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppMemberFunctionMayBeStatic/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppMemberInitializersOrder/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppMismatchedClassTags/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppMissingIncludeGuard/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppMissingKeywordThrow/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppModulePartitionWithSeveralPartitionUnits/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppMsExtAddressOfClassRValue/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppMsExtBindingRValueToLvalueReference/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppMsExtCopyElisionInCopyInitDeclarator/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppMsExtDoubleUserConversionInCopyInit/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppMsExtNotInitializedStaticConstLocalVar/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppMsExtReinterpretCastFromNullptr/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppMultiCharacterLiteral/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppMultiCharacterWideLiteral/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppMustBePublicVirtualToImplementInterface/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppMutableSpecifierOnReferenceMember/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppNoDiscardExpression/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppNodiscardFunctionWithoutReturnValue/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppNonExceptionSafeResourceAcquisition/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppNonExplicitConversionOperator/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppNonExplicitConvertingConstructor/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppNonInlineFunctionDefinitionInHeaderFile/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppNonInlineVariableDefinitionInHeaderFile/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppNotAllPathsReturnValue/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppObjectMemberMightNotBeInitialized/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppOutParameterMustBeWritten/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppParameterMayBeConst/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppParameterMayBeConstPtrOrRef/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppParameterNamesMismatch/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppParameterNeverUsed/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppPassValueParameterByConstReference/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppPointerConversionDropsQualifiers/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppPointerToIntegralConversion/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppPolymorphicClassWithNonVirtualPublicDestructor/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppPossiblyErroneousEmptyStatements/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppPossiblyUninitializedMember/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppPossiblyUnintendedObjectSlicing/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppPrecompiledHeaderIsNotIncluded/@EntryIndexedValue" value="ERROR" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppPrecompiledHeaderNotFound/@EntryIndexedValue" value="ERROR" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppPrintfBadFormat/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppPrintfExtraArg/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppPrintfMissedArg/@EntryIndexedValue" value="ERROR" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppPrintfRiskyFormat/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppPrivateSpecialMemberFunctionIsNotImplemented/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRangeBasedForIncompatibleReference/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedefinitionOfDefaultArgumentInOverrideFunction/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedundantAccessSpecifier/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedundantBaseClassAccessSpecifier/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedundantBaseClassInitializer/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedundantBooleanExpressionArgument/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedundantCastExpression/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedundantComplexityInComparison/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedundantConditionalExpression/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedundantConstSpecifier/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedundantControlFlowJump/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedundantDereferencingAndTakingAddress/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedundantElaboratedTypeSpecifier/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedundantElseKeyword/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedundantElseKeywordInsideCompoundStatement/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedundantEmptyDeclaration/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedundantEmptyStatement/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedundantExportKeyword/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedundantFwdClassOrEnumSpecifier/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedundantInlineSpecifier/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedundantLambdaParameterList/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedundantMemberInitializer/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedundantNamespaceDefinition/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedundantParentheses/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedundantQualifier/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedundantQualifierADL/@EntryIndexedValue" value="DO_NOT_SHOW" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedundantStaticSpecifierOnMemberAllocationFunction/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedundantStaticSpecifierOnThreadLocalLocalVariable/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedundantTemplateArguments/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedundantTemplateKeyword/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedundantTypenameKeyword/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedundantVoidArgumentList/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRedundantZeroInitializerInAggregateInitialization/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppReinterpretCastFromVoidPtr/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppRemoveRedundantBraces/@EntryIndexedValue" value="DO_NOT_SHOW" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppReplaceMemsetWithZeroInitialization/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppReplaceTieWithStructuredBinding/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppReturnNoValueInNonVoidFunction/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppSmartPointerVsMakeFunction/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppSomeObjectMembersMightNotBeInitialized/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppSpecialFunctionWithoutNoexceptSpecification/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppStaticAssertFailure/@EntryIndexedValue" value="ERROR" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppStaticDataMemberInUnnamedStruct/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppStaticSpecifierOnAnonymousNamespaceMember/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppStringLiteralToCharPointerConversion/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppTabsAreDisallowed/@EntryIndexedValue" value="DO_NOT_SHOW" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppTemplateArgumentsCanBeDeduced/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppTemplateParameterNeverUsed/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppTemplateParameterShadowing/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppThrowExpressionCanBeReplacedWithRethrow/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppTooWideScope/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppTooWideScopeInitStatement/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppTypeAliasNeverUsed/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppUninitializedDependentBaseClass/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppUninitializedNonStaticDataMember/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppUnionMemberOfReferenceType/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppUnmatchedPragmaEndRegionDirective/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppUnmatchedPragmaRegionDirective/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppUnnamedNamespaceInHeaderFile/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppUnnecessaryWhitespace/@EntryIndexedValue" value="DO_NOT_SHOW" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppUnsignedZeroComparison/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppUnusedIncludeDirective/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppUseAlgorithmWithCount/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppUseAssociativeContains/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppUseAuto/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppUseAutoForNumeric/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppUseElementsView/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppUseEraseAlgorithm/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppUseFamiliarTemplateSyntaxForGenericLambdas/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppUseRangeAlgorithm/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppUseStdSize/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppUseStructuredBinding/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppUseTypeTraitAlias/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppUserDefinedLiteralSuffixDoesNotStartWithUnderscore/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppUsingResultOfAssignmentAsCondition/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppVariableCanBeMadeConstexpr/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppVirtualFunctionCallInsideCtor/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppVirtualFunctionInFinalClass/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppVolatileParameterInDeclaration/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppWarningDirective/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppWrongIncludesOrder/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppWrongSlashesInIncludeDirective/@EntryIndexedValue" value="HINT" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppZeroConstantCanBeReplacedWithNullptr/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=CppZeroValuedExpressionUsedAsNullPointer/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=IdentifierTypo/@EntryIndexedValue" value="DO_NOT_SHOW" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=IfStdIsConstantEvaluatedCanBeReplaced/@EntryIndexedValue" value="SUGGESTION" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=StdIsConstantEvaluatedWillAlwaysEvaluateToConstant/@EntryIndexedValue" value="WARNING" type="string" />
<option name="/Default/CodeInspection/Highlighting/InspectionSeverities/=StringLiteralTypo/@EntryIndexedValue" value="DO_NOT_SHOW" type="string" />
</component>
</project>

View File

@ -0,0 +1,48 @@
<component name="InspectionProjectProfileManager">
<profile version="1.0">
<option name="myName" value="Project Default" />
<inspection_tool class="DuplicatedCode" enabled="false" level="WEAK WARNING" enabled_by_default="false" />
<inspection_tool class="HtmlUnknownAttribute" enabled="true" level="WARNING" enabled_by_default="true">
<option name="myValues">
<value>
<list size="2">
<item index="0" class="java.lang.String" itemvalue="x-data" />
<item index="1" class="java.lang.String" itemvalue="x-init" />
</list>
</value>
</option>
<option name="myCustomValuesEnabled" value="true" />
</inspection_tool>
<inspection_tool class="HtmlUnknownTag" enabled="true" level="WARNING" enabled_by_default="true">
<option name="myValues">
<value>
<list size="8">
<item index="0" class="java.lang.String" itemvalue="nobr" />
<item index="1" class="java.lang.String" itemvalue="noembed" />
<item index="2" class="java.lang.String" itemvalue="comment" />
<item index="3" class="java.lang.String" itemvalue="noscript" />
<item index="4" class="java.lang.String" itemvalue="embed" />
<item index="5" class="java.lang.String" itemvalue="script" />
<item index="6" class="java.lang.String" itemvalue="p" />
<item index="7" class="java.lang.String" itemvalue="include" />
</list>
</value>
</option>
<option name="myCustomValuesEnabled" value="true" />
</inspection_tool>
<inspection_tool class="PyPep8NamingInspection" enabled="true" level="WEAK WARNING" enabled_by_default="true">
<option name="ignoredErrors">
<list>
<option value="N802" />
<option value="N806" />
<option value="N803" />
</list>
</option>
</inspection_tool>
<inspection_tool class="SpellCheckingInspection" enabled="false" level="TYPO" enabled_by_default="false">
<option name="processCode" value="true" />
<option name="processLiterals" value="true" />
<option name="processComments" value="true" />
</inspection_tool>
</profile>
</component>

21
.idea/misc.xml Normal file
View File

@ -0,0 +1,21 @@
<?xml version="1.0" encoding="UTF-8"?>
<project version="4">
<component name="CMakePythonSetting">
<option name="pythonIntegrationState" value="YES" />
</component>
<component name="ExternalStorageConfigurationManager" enabled="true" />
<component name="MakefileSettings">
<option name="linkedExternalProjectsSettings">
<MakefileProjectSettings>
<option name="externalProjectPath" value="$PROJECT_DIR$" />
<option name="modules">
<set>
<option value="$PROJECT_DIR$" />
</set>
</option>
<option name="version" value="2" />
</MakefileProjectSettings>
</option>
</component>
<component name="MakefileWorkspace" PROJECT_DIR="$PROJECT_DIR$" />
</project>

6
.idea/vcs.xml Normal file
View File

@ -0,0 +1,6 @@
<?xml version="1.0" encoding="UTF-8"?>
<project version="4">
<component name="VcsDirectoryMappings">
<mapping directory="$PROJECT_DIR$" vcs="Git" />
</component>
</project>

BIN
AAD_A1.pdf Normal file

Binary file not shown.

270
aad_cuda_utilities.h Normal file
View File

@ -0,0 +1,270 @@
//
// Tomás Oliveira e Silva, September 2025
//
// Arquiteturas de Alto Desempenho 2025/2026
//
// CUDA driver API stuff
//
#ifndef AAD_CUDA_UTILITIES
#define AAD_CUDA_UTILITIES
#include <cuda.h>
//
// data type used to store all CUDA related stuff
//
#define MAX_N_ARGUMENTS 4
typedef struct
{
// input data
int device_number; // number of the device to initialize
char *cubin_file_name; // name of the cubin file to load (NULL if not needed)
char *kernel_name; // name of the CUDA kernel to load (NULL if not needed)
u32_t data_size[2]; // the number of bytes of the two data arrays to allocate on the host and on the device (0 if not needed)
// persistent data
CUdevice cu_device; // the device yhandle
char device_name[256]; // the device name
CUcontext cu_context; // the device context
CUmodule cu_module; // the loaded cubin file contents
CUfunction cu_kernel; // the pointer to the CUDA kernel
CUstream cu_stream; // the command stream
void *host_data[2]; // the pointers to the host data
CUdeviceptr device_data[2]; // the pointers to the device data
// launch kernel data
unsigned int grid_dim_x; // the number of grid blocks (in the X dimension, the only one we will use here)
unsigned int block_dim_x; // the number of threads in a block (in the X dimension, the only one we will use here, should be equal to RECOMENDED_CUDA_BLOCK_SIZE)
int n_kernel_arguments; // number of kernel arguments
void *arg[MAX_N_ARGUMENTS]; // pointers to the kernel argument data
}
cuda_data_t;
//
// CU_CALL --- macro that should be used to call a CUDA driver API function and to test its return value
//
// it should be used to test the return value of calls such as
// cuInit(device_number);
// cuDeviceGet(&cu_device,device_number);
//
// in these cases, f_name is, respectively, cuInit and cuDeviceGet, and args is, respectively,
// (device_number) and (&cu_device,device_number)
//
#define CU_CALL(f_name,args) \
do \
{ \
CUresult e = f_name args; \
if(e != CUDA_SUCCESS) \
{ /* the call failed, terminate the program */ \
fprintf(stderr,"" # f_name "() returned %s (file %s, line %d)\n",cu_error_string(e),__FILE__,__LINE__); \
exit(1); \
} \
} \
while(0)
//
// terse description of the CUDA error codes (replacement of the error code number by its name)
//
static const char *cu_error_string(CUresult e)
{
static char error_string[64];
# define CASE(error_code) case error_code: return "" # error_code;
switch((int)e)
{ // list of error codes extracted from cuda.h (TODO: /usr/local/cuda-10.2/targets/x86_64-linux/include/CL)
default: sprintf(error_string,"unknown error code (%d)",(int)e); return(error_string);
CASE(CUDA_SUCCESS );
CASE(CUDA_ERROR_INVALID_VALUE );
CASE(CUDA_ERROR_OUT_OF_MEMORY );
CASE(CUDA_ERROR_NOT_INITIALIZED );
CASE(CUDA_ERROR_DEINITIALIZED );
CASE(CUDA_ERROR_PROFILER_DISABLED );
CASE(CUDA_ERROR_PROFILER_NOT_INITIALIZED );
CASE(CUDA_ERROR_PROFILER_ALREADY_STARTED );
CASE(CUDA_ERROR_PROFILER_ALREADY_STOPPED );
CASE(CUDA_ERROR_NO_DEVICE );
CASE(CUDA_ERROR_INVALID_DEVICE );
CASE(CUDA_ERROR_INVALID_IMAGE );
CASE(CUDA_ERROR_INVALID_CONTEXT );
CASE(CUDA_ERROR_CONTEXT_ALREADY_CURRENT );
CASE(CUDA_ERROR_MAP_FAILED );
CASE(CUDA_ERROR_UNMAP_FAILED );
CASE(CUDA_ERROR_ARRAY_IS_MAPPED );
CASE(CUDA_ERROR_ALREADY_MAPPED );
CASE(CUDA_ERROR_NO_BINARY_FOR_GPU );
CASE(CUDA_ERROR_ALREADY_ACQUIRED );
CASE(CUDA_ERROR_NOT_MAPPED );
CASE(CUDA_ERROR_NOT_MAPPED_AS_ARRAY );
CASE(CUDA_ERROR_NOT_MAPPED_AS_POINTER );
CASE(CUDA_ERROR_ECC_UNCORRECTABLE );
CASE(CUDA_ERROR_UNSUPPORTED_LIMIT );
CASE(CUDA_ERROR_CONTEXT_ALREADY_IN_USE );
CASE(CUDA_ERROR_PEER_ACCESS_UNSUPPORTED );
CASE(CUDA_ERROR_INVALID_PTX );
CASE(CUDA_ERROR_INVALID_GRAPHICS_CONTEXT );
CASE(CUDA_ERROR_NVLINK_UNCORRECTABLE );
CASE(CUDA_ERROR_INVALID_SOURCE );
CASE(CUDA_ERROR_FILE_NOT_FOUND );
CASE(CUDA_ERROR_SHARED_OBJECT_SYMBOL_NOT_FOUND);
CASE(CUDA_ERROR_SHARED_OBJECT_INIT_FAILED );
CASE(CUDA_ERROR_OPERATING_SYSTEM );
CASE(CUDA_ERROR_INVALID_HANDLE );
CASE(CUDA_ERROR_NOT_FOUND );
CASE(CUDA_ERROR_NOT_READY );
CASE(CUDA_ERROR_ILLEGAL_ADDRESS );
CASE(CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES );
CASE(CUDA_ERROR_LAUNCH_TIMEOUT );
CASE(CUDA_ERROR_LAUNCH_INCOMPATIBLE_TEXTURING );
CASE(CUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED );
CASE(CUDA_ERROR_PEER_ACCESS_NOT_ENABLED );
CASE(CUDA_ERROR_PRIMARY_CONTEXT_ACTIVE );
CASE(CUDA_ERROR_CONTEXT_IS_DESTROYED );
CASE(CUDA_ERROR_ASSERT );
CASE(CUDA_ERROR_TOO_MANY_PEERS );
CASE(CUDA_ERROR_HOST_MEMORY_ALREADY_REGISTERED);
CASE(CUDA_ERROR_HOST_MEMORY_NOT_REGISTERED );
CASE(CUDA_ERROR_HARDWARE_STACK_ERROR );
CASE(CUDA_ERROR_ILLEGAL_INSTRUCTION );
CASE(CUDA_ERROR_MISALIGNED_ADDRESS );
CASE(CUDA_ERROR_INVALID_ADDRESS_SPACE );
CASE(CUDA_ERROR_INVALID_PC );
CASE(CUDA_ERROR_LAUNCH_FAILED );
CASE(CUDA_ERROR_NOT_PERMITTED );
CASE(CUDA_ERROR_NOT_SUPPORTED );
CASE(CUDA_ERROR_UNKNOWN );
};
# undef CASE
}
//
// synchonize the stream command buffer
//
static void synchronize_cuda(cuda_data_t *cd)
{
CU_CALL( cuStreamSynchronize , (cd->cu_stream) );
}
//
// initialize the CUDA driver API interface
//
// load a single cubin file, with a single CUDA kernel
// allocate up to two storage areas both on the host and on the device
//
static void initialize_cuda(cuda_data_t *cd)
{
//
// initialize the driver API interface
//
CU_CALL( cuInit , (0) );
//
// open the CUDA device
//
CU_CALL( cuDeviceGet , (&cd->cu_device,cd->device_number) );
//
// get information about the CUDA device
//
CU_CALL( cuDeviceGetName , (cd->device_name,(int)sizeof(cd->device_name) - 1,cd->cu_device) );
printf("initialize_cuda(): CUDA code running on a %s (device %d, CUDA %u.%u.%u)\n",cd->device_name,cd->device_number,CUDA_VERSION / 1000,(CUDA_VERSION / 10) % 100,CUDA_VERSION % 10);
//
// create a context
//
CU_CALL( cuCtxCreate , (&cd->cu_context,CU_CTX_SCHED_YIELD,cd->cu_device) );
CU_CALL( cuCtxSetCacheConfig , (CU_FUNC_CACHE_PREFER_L1) );
//
// load precompiled modules
//
CU_CALL( cuModuleLoad , (&cd->cu_module,cd->cubin_file_name) );
//
// get the kernel function pointers
//
CU_CALL( cuModuleGetFunction, (&cd->cu_kernel,cd->cu_module,cd->kernel_name) );
//
// create a command stream (we could have used the default stream)
//
CU_CALL( cuStreamCreate, (&cd->cu_stream,CU_STREAM_NON_BLOCKING) );
//
// allocate host and device memory
//
for(int i = 0;i < 2;i++)
if(cd->data_size[i] > 0u)
{
CU_CALL( cuMemAllocHost , ((void **)&cd->host_data[i] ,(size_t)cd->data_size[i]) );
CU_CALL( cuMemAlloc , (&cd->device_data[i],(size_t)cd->data_size[i]) );
}
else
cd->host_data[i] = NULL;
//
// catch any lingering errors
//
synchronize_cuda(cd);
}
//
// terminate the CUDA driver API interface
//
static void terminate_cuda(cuda_data_t *cd)
{
CU_CALL( cuStreamDestroy, (cd->cu_stream) );
for(int i = 0;i < 2;i++)
if(cd->data_size[i] > 0u)
{
CU_CALL( cuMemFreeHost , (cd->host_data[i]) );
CU_CALL( cuMemFree , (cd->device_data[i]) );
}
CU_CALL( cuModuleUnload , (cd->cu_module) );
CU_CALL( cuCtxDestroy , (cd->cu_context) );
}
//
// copy data from the host to the device and from the device to the host
//
static void host_to_device_copy(cuda_data_t *cd,int idx)
{
if(idx < 0 || idx > 1 || cd->data_size[idx] == 0u)
{
fprintf(stderr,"host_to_device_copy(): bad idx\n");
exit(1);
}
CU_CALL( cuMemcpyHtoD , (cd->device_data[idx],(void *)cd->host_data[idx],(size_t)cd->data_size[idx]) );
synchronize_cuda(cd);
}
static void device_to_host_copy(cuda_data_t *cd,int idx)
{
if(idx < 0 || idx > 1 || cd->data_size[idx] == 0u)
{
fprintf(stderr,"device_to_host_copy(): bad idx\n");
exit(1);
}
CU_CALL( cuMemcpyDtoH , ((void *)cd->host_data[idx],cd->device_data[idx],(size_t)cd->data_size[idx]) );
synchronize_cuda(cd);
}
//
// launch a CUDA kernel (with 0 bytes of shared memory and no extra options)
//
static void lauch_kernel(cuda_data_t *cd)
{
if(cd->block_dim_x != (unsigned int)RECOMENDED_CUDA_BLOCK_SIZE)
fprintf(stderr,"lauch_kernel(): block_dim_x should be equal to %d\n",RECOMENDED_CUDA_BLOCK_SIZE);
CU_CALL( cuLaunchKernel , (cd->cu_kernel,cd->grid_dim_x,1u,1u,cd->block_dim_x,1u,1u,0u,cd->cu_stream,&cd->arg[0],NULL) );
synchronize_cuda(cd);
}
#endif

48
aad_data_types.h Normal file
View File

@ -0,0 +1,48 @@
//
// Tomás Oliveira e Silva, September 2025
//
// Arquiteturas de Alto Desempenho 2025/2026
//
// integer data types
//
#ifndef AAD_DATA_TYPES
#define AAD_DATA_TYPES
//
// scalar data types (for a typical 64-bit processor)
//
typedef signed char s08_t; // 8-bit signed integer
typedef unsigned char u08_t; // 8-bit unsigned integer
typedef signed short s16_t; // 16-bit signed integer
typedef unsigned short u16_t; // 16-bit unsigned integer
typedef signed int s32_t; // 32-bit signed integer
typedef unsigned int u32_t; // 32-bit unsigned integer
typedef signed long s64_t; // 64-bit signed integer
typedef unsigned long u64_t; // 64-bit unsigned integer
//
// vector data types (this probably will only work on the gcc compiler)
//
#if defined(__AVX__)
typedef int v4si __attribute__((vector_size(16))) __attribute__((aligned(16)));
#endif
#if defined(__AVX2__)
typedef int v8si __attribute__((vector_size(32))) __attribute__((aligned(32)));
#endif
#if defined(__AVX512F__)
typedef int v16si __attribute__((vector_size(64))) __attribute__((aligned(64)));
#endif
#if defined(__ARM_NEON)
# include <arm_neon.h>
#endif
//
// the end!
//
#endif

240
aad_sha1.h Normal file
View File

@ -0,0 +1,240 @@
//
// Tomás Oliveira e Silva, September 2025
//
// Arquiteturas de Alto Desempenho 2025/2026
//
// template for the computation of the SHA1 secure hash
//
//
// compute the SHA1 secure hash of a custom message with exactly 55 bytes
//
// the general SHA1 secure hash algorithm ingests data in chunks of 64 bytes; at the end there must
// be one byte of padding, with value 0x80, zero or more bytes of zeros, also for padding, appended
// until the last chunk has 56 bytes, and a final 8 byte integer holding the number of bits of the
// entire message
//
// by restricting the data to have 55 bytes or less the algorithm can be simplified, and only one
// chunk is needed; that is what is done below
//
#ifndef AAD_SHA1
#define AAD_SHA1
//
// number of threads in each CUDA block
//
// we place this here to simplify things (aad_sha1_cuda_kernel.cu includes this file...)
//
#define RECOMENDED_CUDA_BLOCK_SIZE 128
//
// each custom message has exactly 55 bytes, and must be followed by an additional byte with the
// value 0x80
// these 55+1=56 bytes must be stored in a 32-bit integer array with 14 elements as illustrated in
// the test code below; the secure hash has to be interpreted in the same way --- don't blame the
// teacher for this; that is how the SHA1 secure hash is described in the 3174 request for comments
// (https://datatracker.ietf.org/doc/html/rfc3174)
//
// the SHA1 secure hash of the 55 bytes message is computed using a macro called CUSTOM_SHA1_CODE
// it must be customized using the following additional macros:
// T --- the data type
// C(c) --- how to expand the constant c
// ROTATE(x,n) --- how to rotate x left by n bits
// DATA(idx) --- how to access the data at index idx, 0 <= idx <= 13
// HASH(idx) --- how to access the hash at index idx, 0 <= idx <= 4
// see aad_sha1_cpu.h for examples
//
// each custom message is stored in the locations
// DATA(0), DATA(1), ..., DATA(13)
// each SHA1 secure hash is stored in the locations
// HASH(0), HASH(1), ..., HASH(4)
//
//
// first group of 20 iterations (0 <= t <= 19)
//
#define SHA1_F1(x,y,z) ((x & y) | (~x & z))
#define SHA1_K1 0x5A827999u
//
// second group of 20 iterations (20 <= t <= 39)
//
#define SHA1_F2(x,y,z) (x ^ y ^ z)
#define SHA1_K2 0x6ED9EBA1u
//
// third group of 20 iterations (40 <= t <= 59)
//
#define SHA1_F3(x,y,z) ((x & y) | (x & z) | (y & z))
#define SHA1_K3 0x8F1BBCDCu
//
// fourth group of 20 iterations (60 <= t <= 79)
//
#define SHA1_F4(x,y,z) (x ^ y ^ z)
#define SHA1_K4 0xCA62C1D6u
//
// data mixing function
//
#define SHA1_D(t) \
do \
{ \
T tmp = w[((t) - 3) & 15] ^ w[((t) - 8) & 15] ^ w[((t) - 14) & 15] ^ w[((t) - 16) & 15]; \
w[(t) & 15] = ROTATE(tmp,1); \
} \
while(0)
//
// state mixing function
//
#define SHA1_S(F,t,K) \
do \
{ \
T tmp = ROTATE(a,5) + F(b,c,d) + e + w[(t) & 15] + C(K); \
e = d; \
d = c; \
c = ROTATE(b,30); \
b = a; \
a = tmp; \
} \
while(0)
//
// the CUSTOM_SHA1_CODE macro, for a little-endian processor
//
// everything is loop unrolled to make sure all indices are static integers, so the compiler
// has no excuse to produce sub-optimal code (the w[16] array can even become 16 separate
// integer variables, the CUDA compiler actually does this)
//
#define CUSTOM_SHA1_CODE() \
do \
{ \
/* local variables */ \
T a,b,c,d,e,w[16]; \
/* initial state */ \
a = C(0x67452301u); \
b = C(0xEFCDAB89u); \
c = C(0x98BADCFEu); \
d = C(0x10325476u); \
e = C(0xC3D2E1F0u); \
/* copy data to the internal buffer */ \
w[ 0] = DATA( 0); \
w[ 1] = DATA( 1); \
w[ 2] = DATA( 2); \
w[ 3] = DATA( 3); \
w[ 4] = DATA( 4); \
w[ 5] = DATA( 5); \
w[ 6] = DATA( 6); \
w[ 7] = DATA( 7); \
w[ 8] = DATA( 8); \
w[ 9] = DATA( 9); \
w[10] = DATA(10); \
w[11] = DATA(11); \
w[12] = DATA(12); \
w[13] = DATA(13); /* WARNING: DATA(13) & 0xFF must be 0x80 (SHA1 padding) */ \
w[14] = C(0); \
w[15] = C(440); /* the message has 55*8 bits */ \
/* first group of 20 iterations (0 <= t <= 19) */ \
SHA1_S(SHA1_F1, 0,SHA1_K1); \
SHA1_S(SHA1_F1, 1,SHA1_K1); \
SHA1_S(SHA1_F1, 2,SHA1_K1); \
SHA1_S(SHA1_F1, 3,SHA1_K1); \
SHA1_S(SHA1_F1, 4,SHA1_K1); \
SHA1_S(SHA1_F1, 5,SHA1_K1); \
SHA1_S(SHA1_F1, 6,SHA1_K1); \
SHA1_S(SHA1_F1, 7,SHA1_K1); \
SHA1_S(SHA1_F1, 8,SHA1_K1); \
SHA1_S(SHA1_F1, 9,SHA1_K1); \
SHA1_S(SHA1_F1,10,SHA1_K1); \
SHA1_S(SHA1_F1,11,SHA1_K1); \
SHA1_S(SHA1_F1,12,SHA1_K1); \
SHA1_S(SHA1_F1,13,SHA1_K1); \
SHA1_S(SHA1_F1,14,SHA1_K1); \
SHA1_S(SHA1_F1,15,SHA1_K1); \
SHA1_D(16); SHA1_S(SHA1_F1,16,SHA1_K1); \
SHA1_D(17); SHA1_S(SHA1_F1,17,SHA1_K1); \
SHA1_D(18); SHA1_S(SHA1_F1,18,SHA1_K1); \
SHA1_D(19); SHA1_S(SHA1_F1,19,SHA1_K1); \
/* second group of 20 iterations (20 <= t <= 39) */ \
SHA1_D(20); SHA1_S(SHA1_F2,20,SHA1_K2); \
SHA1_D(21); SHA1_S(SHA1_F2,21,SHA1_K2); \
SHA1_D(22); SHA1_S(SHA1_F2,22,SHA1_K2); \
SHA1_D(23); SHA1_S(SHA1_F2,23,SHA1_K2); \
SHA1_D(24); SHA1_S(SHA1_F2,24,SHA1_K2); \
SHA1_D(25); SHA1_S(SHA1_F2,25,SHA1_K2); \
SHA1_D(26); SHA1_S(SHA1_F2,26,SHA1_K2); \
SHA1_D(27); SHA1_S(SHA1_F2,27,SHA1_K2); \
SHA1_D(28); SHA1_S(SHA1_F2,28,SHA1_K2); \
SHA1_D(29); SHA1_S(SHA1_F2,29,SHA1_K2); \
SHA1_D(30); SHA1_S(SHA1_F2,30,SHA1_K2); \
SHA1_D(31); SHA1_S(SHA1_F2,31,SHA1_K2); \
SHA1_D(32); SHA1_S(SHA1_F2,32,SHA1_K2); \
SHA1_D(33); SHA1_S(SHA1_F2,33,SHA1_K2); \
SHA1_D(34); SHA1_S(SHA1_F2,34,SHA1_K2); \
SHA1_D(35); SHA1_S(SHA1_F2,35,SHA1_K2); \
SHA1_D(36); SHA1_S(SHA1_F2,36,SHA1_K2); \
SHA1_D(37); SHA1_S(SHA1_F2,37,SHA1_K2); \
SHA1_D(38); SHA1_S(SHA1_F2,38,SHA1_K2); \
SHA1_D(39); SHA1_S(SHA1_F2,39,SHA1_K2); \
/* third group of 20 iterations (40 <= t <= 59) */ \
SHA1_D(40); SHA1_S(SHA1_F3,40,SHA1_K3); \
SHA1_D(41); SHA1_S(SHA1_F3,41,SHA1_K3); \
SHA1_D(42); SHA1_S(SHA1_F3,42,SHA1_K3); \
SHA1_D(43); SHA1_S(SHA1_F3,43,SHA1_K3); \
SHA1_D(44); SHA1_S(SHA1_F3,44,SHA1_K3); \
SHA1_D(45); SHA1_S(SHA1_F3,45,SHA1_K3); \
SHA1_D(46); SHA1_S(SHA1_F3,46,SHA1_K3); \
SHA1_D(47); SHA1_S(SHA1_F3,47,SHA1_K3); \
SHA1_D(48); SHA1_S(SHA1_F3,48,SHA1_K3); \
SHA1_D(49); SHA1_S(SHA1_F3,49,SHA1_K3); \
SHA1_D(50); SHA1_S(SHA1_F3,50,SHA1_K3); \
SHA1_D(51); SHA1_S(SHA1_F3,51,SHA1_K3); \
SHA1_D(52); SHA1_S(SHA1_F3,52,SHA1_K3); \
SHA1_D(53); SHA1_S(SHA1_F3,53,SHA1_K3); \
SHA1_D(54); SHA1_S(SHA1_F3,54,SHA1_K3); \
SHA1_D(55); SHA1_S(SHA1_F3,55,SHA1_K3); \
SHA1_D(56); SHA1_S(SHA1_F3,56,SHA1_K3); \
SHA1_D(57); SHA1_S(SHA1_F3,57,SHA1_K3); \
SHA1_D(58); SHA1_S(SHA1_F3,58,SHA1_K3); \
SHA1_D(59); SHA1_S(SHA1_F3,59,SHA1_K3); \
/* fourth group of 20 iterations (60 <= t <= 79) */ \
SHA1_D(60); SHA1_S(SHA1_F4,60,SHA1_K4); \
SHA1_D(61); SHA1_S(SHA1_F4,61,SHA1_K4); \
SHA1_D(62); SHA1_S(SHA1_F4,62,SHA1_K4); \
SHA1_D(63); SHA1_S(SHA1_F4,63,SHA1_K4); \
SHA1_D(64); SHA1_S(SHA1_F4,64,SHA1_K4); \
SHA1_D(65); SHA1_S(SHA1_F4,65,SHA1_K4); \
SHA1_D(66); SHA1_S(SHA1_F4,66,SHA1_K4); \
SHA1_D(67); SHA1_S(SHA1_F4,67,SHA1_K4); \
SHA1_D(68); SHA1_S(SHA1_F4,68,SHA1_K4); \
SHA1_D(69); SHA1_S(SHA1_F4,69,SHA1_K4); \
SHA1_D(70); SHA1_S(SHA1_F4,70,SHA1_K4); \
SHA1_D(71); SHA1_S(SHA1_F4,71,SHA1_K4); \
SHA1_D(72); SHA1_S(SHA1_F4,72,SHA1_K4); \
SHA1_D(73); SHA1_S(SHA1_F4,73,SHA1_K4); \
SHA1_D(74); SHA1_S(SHA1_F4,74,SHA1_K4); \
SHA1_D(75); SHA1_S(SHA1_F4,75,SHA1_K4); \
SHA1_D(76); SHA1_S(SHA1_F4,76,SHA1_K4); \
SHA1_D(77); SHA1_S(SHA1_F4,77,SHA1_K4); \
SHA1_D(78); SHA1_S(SHA1_F4,78,SHA1_K4); \
SHA1_D(79); SHA1_S(SHA1_F4,79,SHA1_K4); \
/* update state (in this special case, finish) */ \
HASH(0) = a + C(0x67452301u); \
HASH(1) = b + C(0xEFCDAB89u); \
HASH(2) = c + C(0x98BADCFEu); \
HASH(3) = d + C(0x10325476u); \
HASH(4) = e + C(0xC3D2E1F0u); \
} \
while(0)
//
// the end!
//
#endif

141
aad_sha1_cpu.h Normal file
View File

@ -0,0 +1,141 @@
//
// Tomás Oliveira e Silva, September 2025
//
// Arquiteturas de Alto Desempenho 2025/2026
//
// SHA1 secure hash implementations on the CPU
//
#ifndef AAD_SHA1_CPU
#define AAD_SHA1_CPU
#include "aad_sha1.h"
#define FOUR(c) (int)(c),(int)(c),(int)(c),(int)(c)
//
// reference implementation (no SIMD instructions)
//
__attribute__((unused))
static void sha1(u32_t *data,u32_t *hash)
{ // one message -> one SHA1 hash
# define T u32_t
# define C(c) (c)
# define ROTATE(x,n) (((x) << (n)) | ((x) >> (32 - (n))))
# define DATA(idx) data[idx]
# define HASH(idx) hash[idx]
CUSTOM_SHA1_CODE();
# undef T
# undef C
# undef ROTATE
# undef DATA
# undef HASH
}
//
// implementation using avx instructions (Intel/AMD)
//
#if defined(__AVX__)
__attribute__((unused))
static void sha1_avx(v4si *interleaved4_data,v4si *interleaved4_hash)
{ // four interleaved messages -> four interleaved SHA1 secure hashes
# define T v4si
# define C(c) (v4si){ FOUR(c) }
# define ROTATE(x,n) (__builtin_ia32_pslldi128(x,n) | __builtin_ia32_psrldi128(x,32 - (n)))
# define DATA(idx) interleaved4_data[idx]
# define HASH(idx) interleaved4_hash[idx]
CUSTOM_SHA1_CODE();
# undef T
# undef C
# undef ROTATE
# undef DATA
# undef HASH
}
#endif
//
// implementation using avx2 instructions (Intel/AMD)
//
#if defined(__AVX2__)
__attribute__((unused))
static void sha1_avx2(v8si *interleaved8_data,v8si *interleaved8_hash)
{ // eight interleaved messages -> eight interleaved SHA1 secure hashes
# define T v8si
# define C(c) (v8si){ FOUR(c),FOUR(c) }
# define ROTATE(x,n) (__builtin_ia32_pslldi256(x,n) | __builtin_ia32_psrldi256(x,32 - (n)))
# define DATA(idx) interleaved8_data[idx]
# define HASH(idx) interleaved8_hash[idx]
CUSTOM_SHA1_CODE();
# undef T
# undef C
# undef ROTATE
# undef DATA
# undef HASH
}
#endif
//
// implementation using avx512f instructions (Intel/AMD)
//
#if defined(__AVX512F__)
__attribute__((unused))
static void sha1_avx512f(v16si *interleaved16_data,v16si *interleaved16_hash)
{ // sixteen interleaved messages -> sixteen interleaved SHA1 secure hashes
# define T v16si
# define C(c) (v16si){ FOUR(c),FOUR(c),FOUR(c),FOUR(c) }
# define ROTATE(x,n) __builtin_ia32_prold512_mask(x,n,x,0xFFFF)
# define DATA(idx) interleaved16_data[idx]
# define HASH(idx) interleaved16_hash[idx]
CUSTOM_SHA1_CODE();
# undef T
# undef C
# undef ROTATE
# undef DATA
# undef HASH
}
#endif
//
// implementation using neon instructions (ARM)
//
#if defined(__ARM_NEON)
__attribute__((unused))
static void sha1_neon(uint32x4_t *interleaved4_data,uint32x4_t *interleaved4_hash)
{ // four interleaved messages -> four interleaved SHA1 secure hashes
# define T uint32x4_t
# define C(c) (uint32x4_t){ FOUR(c) }
# define ROTATE(x,n) (vshlq_n_u32(x,n) | vshrq_n_u32(x,32 - (n)))
# define DATA(idx) interleaved4_data[idx]
# define HASH(idx) interleaved4_hash[idx]
CUSTOM_SHA1_CODE();
# undef T
# undef C
# undef ROTATE
# undef DATA
# undef HASH
}
#endif
//
// the end!
//
#endif

416
aad_sha1_cpu_tests.c Normal file
View File

@ -0,0 +1,416 @@
//
// Tomás Oliveira e Silva, September 2025
//
// Arquiteturas de Alto Desempenho 2025/2026
//
#include <time.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include "aad_data_types.h"
#include "aad_utilities.h"
#include "aad_sha1_cpu.h"
//
// test the reference implementation
//
static void test_sha1(int n_tests,int n_measurements)
{
static union { u08_t c[14 * 4]; u32_t i[14]; } data; // the data as bytes and as 32-bit integers
static union { u08_t c[ 5 * 4]; u32_t i[ 5]; } hash; // the hash as bytes and as 32-bit integers
char command[320]; // 320 is more than enough
char response[64]; // 64 is more than enough
char computed[64]; // 64 is more than enough
double hashes_per_second;
int n,i,idx;
u32_t sum;
FILE *fp;
// test
response[40] = '\0';
for(n = 0;n < n_tests;n++)
{
// create random data (55 bytes)
for(i = 0;i < 55;i++)
data.c[i ^ 3] = random_byte();
// append padding (a SHA1 thing...)
data.c[55 ^ 3] = 0x80;
// compute its SHA1 secure hash
sha1(&data.i[0],&hash.i[0]);
// convert the secure hash into a string
idx = 0;
for(i = 0;i < 20;i++)
idx += sprintf(&computed[idx],"%02x",(int)hash.c[i ^ 3] & 0xFF);
if(idx >= (int)sizeof(computed))
{
fprintf(stderr,"computed[] is too small\n");
exit(1);
}
// construct the command to test the SHA1 secure hash
idx = sprintf(&command[0],"/bin/echo -en '"); // do not rely on the bash echo builtin command
for(i = 0;i < 55;i++)
idx += sprintf(&command[idx],"\\x%02x",data.c[i ^ 3]);
idx += sprintf(&command[idx],"' | sha1sum");
if(idx >= (int)sizeof(command))
{
fprintf(stderr,"command[] is too small\n");
exit(1);
}
// run it and get its output
fp = popen(command,"r");
if(fp == NULL)
{
fprintf(stderr,"popen() failed\n");
exit(1);
}
if(fread((void *)&response[0],sizeof(char),(size_t)40,fp) != (size_t)40)
{
fprintf(stderr,"fread() failed\n");
exit(1);
}
pclose(fp);
// compare them
if(memcmp((void *)response,(void *)computed,(size_t)40) != 0)
{ // print everything
fprintf(stderr,"sha1() failure for n=%d:\n",n);
for(i = 0;i < 55;i++)
fprintf(stderr," message[%2d] = %02x\n",i,(int)data.c[i ^ 3] & 0xFF);
for(i = 0;i < 20;i++)
fprintf(stderr," hash[%2d] = %02x\n",i,(int)hash.c[i ^ 3] % 0xFF);
fprintf(stderr," sha1sum output: %s\n",response);
fprintf(stderr," sha1() output: %s\n",computed);
for(i = 0;i < 40 && response[i] == computed[i];i++)
;
fprintf(stderr," mismatch at %d\n",i);
exit(1);
}
}
// warmup (turbo boost...)
for(i = n = 0;i < 1000000;i++)
n += (int)random_byte();
if(n == 0)
fprintf(stderr,"sha1(): this should not be possible, n=0\n");
// measure
time_measurement();
sum = 0u;
for(n = 0;n < n_measurements;n++)
{
data.i[0]++;
sha1(&data.i[0],&hash.i[0]);
sum += hash.i[4];
}
time_measurement();
if(sum == 0u)
fprintf(stderr,"sha1(): what a coincidence, sum=0\n");
hashes_per_second = (double)n_measurements / cpu_time_delta();
// report
printf("sha1() passed (%d test%s, %.0f secure hashes per second)\n",n_tests,(n_tests == 1) ? "" : "s",hashes_per_second);
}
//
// test the avx implementation
//
#if defined(__AVX__)
static void test_sha1_avx(int n_tests,int n_measurements)
{
#define N_LANES 4
static union { u08_t c[14 * 4]; u32_t i[14]; } data[N_LANES]; // the data as bytes and as 32-bit integers
static union { u08_t c[ 5 * 4]; u32_t i[ 5]; } hash[N_LANES]; // the hash as bytes and as 32-bit integers
static u32_t interleaved_data[14][N_LANES] __attribute__((aligned(16)));
static u32_t interleaved_hash[5][N_LANES] __attribute__((aligned(16)));
double hashes_per_second;
int n,i,lane;
u32_t sum;
// test
for(n = 0;n < n_tests;n++)
{
// the data and the secure hash for the reference implementation
for(lane = 0;lane < N_LANES;lane++)
{
// create random data (55 bytes)
for(i = 0;i < 55;i++)
data[lane].c[i ^ 3] = random_byte();
// append padding (a SHA1 thing...)
data[lane].c[55 ^ 3] = 0x80;
// compute its SHA1 secure hash
sha1(&data[lane].i[0],&hash[lane].i[0]);
}
// interleave (transpose) the data for the avx implementation
for(lane = 0;lane < N_LANES;lane++)
for(i = 0;i < 14;i++)
interleaved_data[i][lane] = data[lane].i[i];
// compute the four secure hashes in one go
sha1_avx((v4si *)&interleaved_data[0],(v4si *)&interleaved_hash[0]);
// test
for(lane = 0;lane < N_LANES;lane++)
for(i = 0;i < 5;i++)
if(interleaved_hash[i][lane] != hash[lane].i[i])
{
fprintf(stderr,"sha1_avx() failure for n=%d (bad/good):\n",n);
for(i = 0;i < 5;i++)
for(lane = 0;lane < N_LANES;lane++)
fprintf(stderr,"%s%08X/%08X%s",(lane == 0) ? " " : " ",interleaved_hash[i][lane] ,hash[lane].i[i],(lane == N_LANES - 1) ? "\n" : "");
exit(1);
}
}
// measure
time_measurement();
sum = 0u;
for(n = 0;n < n_measurements;n++)
{
interleaved_data[0][0]++;
sha1(&data[lane].i[0],&hash[lane].i[0]);
sum += interleaved_hash[4][0];
}
time_measurement();
if(sum == 0u)
fprintf(stderr,"sha1_avx(): what a coincidence, sum=0\n");
hashes_per_second = (double)n_measurements * (double)N_LANES / cpu_time_delta();
// report
printf("sha1_avx() passed (%d test%s, %.0f secure hashes per second)\n",n_tests,(n_tests == 1) ? "" : "s",hashes_per_second);
# undef N_LANES
}
#endif
//
// test the avx2 implementation
//
#if defined(__AVX2__)
static void test_sha1_avx2(int n_tests,int n_measurements)
{
#define N_LANES 8
static union { u08_t c[14 * 4]; u32_t i[14]; } data[N_LANES]; // the data as bytes and as 32-bit integers
static union { u08_t c[ 5 * 4]; u32_t i[ 5]; } hash[N_LANES]; // the hash as bytes and as 32-bit integers
static u32_t interleaved_data[14][N_LANES] __attribute__((aligned(32)));
static u32_t interleaved_hash[5][N_LANES] __attribute__((aligned(32)));
double hashes_per_second;
int n,i,lane;
u32_t sum;
// test
for(n = 0;n < n_tests;n++)
{
// the data and the secure hash for the reference implementation
for(lane = 0;lane < N_LANES;lane++)
{
// create random data (55 bytes)
for(i = 0;i < 55;i++)
data[lane].c[i ^ 3] = random_byte();
// append padding (a SHA1 thing...)
data[lane].c[55 ^ 3] = 0x80;
// compute its SHA1 secure hash
sha1(&data[lane].i[0],&hash[lane].i[0]);
}
// interleave (transpose) the data for the avx2 implementation
for(lane = 0;lane < N_LANES;lane++)
for(i = 0;i < 14;i++)
interleaved_data[i][lane] = data[lane].i[i];
// compute the eight secure hashes in one go
sha1_avx2((v8si *)&interleaved_data[0],(v8si *)&interleaved_hash[0]);
// test
for(lane = 0;lane < N_LANES;lane++)
for(i = 0;i < 5;i++)
if(interleaved_hash[i][lane] != hash[lane].i[i])
{
fprintf(stderr,"sha1_avx2() failure for n=%d (bad/good):\n",n);
for(i = 0;i < 5;i++)
for(lane = 0;lane < N_LANES;lane++)
fprintf(stderr,"%s%08X/%08X%s",(lane == 0) ? " " : " ",interleaved_hash[i][lane] ,hash[lane].i[i],(lane == N_LANES - 1) ? "\n" : "");
exit(1);
}
}
// measure
time_measurement();
sum = 0u;
for(n = 0;n < n_measurements;n++)
{
interleaved_data[0][0]++;
sha1(&data[lane].i[0],&hash[lane].i[0]);
sum += interleaved_hash[4][0];
}
time_measurement();
if(sum == 0u)
fprintf(stderr,"sha1_avx2(): what a coincidence, sum=0\n");
hashes_per_second = (double)n_measurements * (double)N_LANES / cpu_time_delta();
// report
printf("sha1_avx2() passed (%d test%s, %.0f secure hashes per second)\n",n_tests,(n_tests == 1) ? "" : "s",hashes_per_second);
# undef N_LANES
}
#endif
//
// test the avx512f implementation
//
#if defined(__AVX512F__)
static void test_sha1_avx512f(int n_tests,int n_measurements)
{
#define N_LANES 16
static union { u08_t c[14 * 4]; u32_t i[14]; } data[N_LANES]; // the data as bytes and as 32-bit integers
static union { u08_t c[ 5 * 4]; u32_t i[ 5]; } hash[N_LANES]; // the hash as bytes and as 32-bit integers
static u32_t interleaved_data[14][N_LANES] __attribute__((aligned(64)));
static u32_t interleaved_hash[5][N_LANES] __attribute__((aligned(64)));
double hashes_per_second;
int n,i,lane;
u32_t sum;
// test
for(n = 0;n < n_tests;n++)
{
// the data and the secure hash for the reference implementation
for(lane = 0;lane < N_LANES;lane++)
{
// create random data (55 bytes)
for(i = 0;i < 55;i++)
data[lane].c[i ^ 3] = random_byte();
// append padding (a SHA1 thing...)
data[lane].c[55 ^ 3] = 0x80;
// compute its SHA1 secure hash
sha1(&data[lane].i[0],&hash[lane].i[0]);
}
// interleave (transpose) the data for the avx512f implementation
for(lane = 0;lane < N_LANES;lane++)
for(i = 0;i < 14;i++)
interleaved_data[i][lane] = data[lane].i[i];
// compute the sixteen secure hashes in one go
sha1_avx512f((v16si *)&interleaved_data[0],(v16si *)&interleaved_hash[0]);
// test
for(lane = 0;lane < N_LANES;lane++)
for(i = 0;i < 5;i++)
if(interleaved_hash[i][lane] != hash[lane].i[i])
{
fprintf(stderr,"sha1_avx512f() failure for n=%d (bad/good):\n",n);
for(i = 0;i < 5;i++)
for(lane = 0;lane < N_LANES;lane++)
fprintf(stderr,"%s%08X/%08X%s",(lane == 0) ? " " : " ",interleaved_hash[i][lane] ,hash[lane].i[i],(lane == N_LANES - 1) ? "\n" : "");
exit(1);
}
}
// measure
time_measurement();
sum = 0u;
for(n = 0;n < n_measurements;n++)
{
interleaved_data[0][0]++;
sha1(&data[lane].i[0],&hash[lane].i[0]);
sum += interleaved_hash[4][0];
}
time_measurement();
if(sum == 0u)
fprintf(stderr,"sha1_avx512f(): what a coincidence, sum=0\n");
hashes_per_second = (double)n_measurements * (double)N_LANES / cpu_time_delta();
// report
printf("sha1_avx512f() passed (%d test%s, %.0f secure hashes per second)\n",n_tests,(n_tests == 1) ? "" : "s",hashes_per_second);
# undef N_LANES
}
#endif
//
// test the neon implementation
//
#if defined(__ARM_NEON)
static void test_sha1_neon(int n_tests,int n_measurements)
{
#define N_LANES 4
static union { u08_t c[14 * 4]; u32_t i[14]; } data[N_LANES]; // the data as bytes and as 32-bit integers
static union { u08_t c[ 5 * 4]; u32_t i[ 5]; } hash[N_LANES]; // the hash as bytes and as 32-bit integers
static u32_t interleaved_data[14][N_LANES] __attribute__((aligned(16)));
static u32_t interleaved_hash[5][N_LANES] __attribute__((aligned(16)));
double hashes_per_second;
int n,i,lane;
u32_t sum;
// test
for(n = 0;n < n_tests;n++)
{
// the data and the secure hash for the reference implementation
for(lane = 0;lane < N_LANES;lane++)
{
// create random data (55 bytes)
for(i = 0;i < 55;i++)
data[lane].c[i ^ 3] = random_byte();
// append padding (a SHA1 thing...)
data[lane].c[55 ^ 3] = 0x80;
// compute its SHA1 secure hash
sha1(&data[lane].i[0],&hash[lane].i[0]);
}
// interleave (transpose) the data for the neon implementation
for(lane = 0;lane < N_LANES;lane++)
for(i = 0;i < 14;i++)
interleaved_data[i][lane] = data[lane].i[i];
// compute the four secure hashes in one go
sha1_neon((uint32x4_t *)&interleaved_data[0],(uint32x4_t *)&interleaved_hash[0]);
// test
for(lane = 0;lane < N_LANES;lane++)
for(i = 0;i < 5;i++)
if(interleaved_hash[i][lane] != hash[lane].i[i])
{
fprintf(stderr,"sha1_neon() failure for n=%d (bad/good):\n",n);
for(i = 0;i < 5;i++)
for(lane = 0;lane < N_LANES;lane++)
fprintf(stderr,"%s%08X/%08X%s",(lane == 0) ? " " : " ",interleaved_hash[i][lane] ,hash[lane].i[i],(lane == N_LANES - 1) ? "\n" : "");
exit(1);
}
}
// measure
time_measurement();
sum = 0u;
for(n = 0;n < n_measurements;n++)
{
interleaved_data[0][0]++;
sha1(&data[lane].i[0],&hash[lane].i[0]);
sum += interleaved_hash[4][0];
}
time_measurement();
if(sum == 0u)
fprintf(stderr,"sha1_neon(): what a coincidence, sum=0\n");
hashes_per_second = (double)n_measurements * (double)N_LANES / cpu_time_delta();
// report
printf("sha1_neon() passed (%d test%s, %.0f secure hashes per second)\n",n_tests,(n_tests == 1) ? "" : "s",hashes_per_second);
# undef N_LANES
}
#endif
//
// main program
//
int main(void)
{
int n_tests = 1000;
int n_measurements = 10000000;
test_sha1(n_tests,n_measurements);
#if defined(__AVX__)
test_sha1_avx(n_tests,n_measurements);
#endif
#if defined(__AVX2__)
test_sha1_avx2(n_tests,n_measurements);
#endif
#if defined(__AVX512F__)
test_sha1_avx512f(n_tests,n_measurements);
#endif
#if defined(__ARM_NEON)
test_sha1_neon(n_tests,n_measurements);
#endif
return 0;
}

84
aad_sha1_cuda_kernel.cu Normal file
View File

@ -0,0 +1,84 @@
//
// Tomás Oliveira e Silva, September 2025
//
// Arquiteturas de Alto Desempenho 2025/2026
//
//
// sha1_cuda_kernel() --- each CUDA thread computes the SHA1 secure hash of one message
//
// this kernel should only be used to validate the SHA1 secure hash code in CUDA
//
#include "aad_sha1.h"
typedef unsigned int u32_t;
//
// the nvcc compiler stores w[] in registers (constant indices!)
//
// global thread number: n = threadIdx.x + blockDim.x * blockIdx.x
// global warp number: n >> 5
// warp thread number: n & 31 -- the lane
//
extern "C" __global__ __launch_bounds__(RECOMENDED_CUDA_BLOCK_SIZE,1)
void sha1_cuda_kernel(u32_t *interleaved32_data,u32_t *interleaved32_hash)
{
u32_t n;
//
// get the global thread number (to make things easier, only the x dimension is used)
//
n = (u32_t)threadIdx.x + (u32_t)blockDim.x * (u32_t)blockIdx.x;
//
// adjust data and hash pointers; together with the DATA and HASH macros below, these pointer adjustments ensure that
// the 32 threads of a warp access consecutive memory addresses; for one warp addresses grow from the left to the
// right, and then from top to bottom
// +----------------------+----------------------+- ... -+----------------------+----------------------+
// | | | | | |
// | data[ 0] for lane 0 | data[ 0] for lane 1 | ..... | data[ 0] for lane 30 | data[ 0] for lane 31 |
// | | | | | |
// +----------------------+----------------------+- ... -+----------------------+----------------------+
// | | | | | |
// | data[ 1] for lane 0 | data[ 1] for lane 1 | ..... | data[ 1] for lane 30 | data[ 1] for lane 31 |
// | | | | | |
// +----------------------+----------------------+- ... -+----------------------+----------------------+
// ...
// +----------------------+----------------------+- ... -+----------------------+----------------------+
// | | | | | |
// | data[13] for lane 0 | data[13] for lane 1 | ..... | data[13] for lane 30 | data[13] for lane 31 |
// | | | | | |
// +----------------------+----------------------+- ... -+----------------------+----------------------+
// this is followed by the data for the next warp
// +----------------------+----------------------+- ... -+----------------------+----------------------+
// | | | | | |
// | data[ 0] for lane 0 | data[ 0] for lane 1 | ..... | data[ 0] for lane 30 | data[ 0] for lane 31 |
// | | | | | |
// +----------------------+----------------------+- ... -+----------------------+----------------------+
// ...
// And so on. The interleaved32_data is CONCEPTUALLY organized in the following way
// interleaved32_data[number_of_warps] [14] [32]
// [warp_number] [idx] [lane]
// for the same warp number and the same idx, the data for the 32 lanes (warp thread number) are in consecutive addresses
//
// the same happens for the interleaved32_hash, but the indices go only from 0 to 4
//
interleaved32_data = &interleaved32_data[(n >> 5u) * (32u * 14u) + (n & 31u)];
interleaved32_hash = &interleaved32_hash[(n >> 5u) * (32u * 5u) + (n & 31u)];
//
// compute the SHA1 secure hash
//
# define T u32_t
# define C(c) (c)
# define ROTATE(x,n) (((x) << (n)) | ((x) >> (32 - (n))))
# define DATA(idx) interleaved32_data[32u * (idx)]
# define HASH(idx) interleaved32_hash[32u * (idx)]
CUSTOM_SHA1_CODE();
# undef T
# undef C
# undef ROTATE
# undef DATA
# undef HASH
}

105
aad_sha1_cuda_test.c Normal file
View File

@ -0,0 +1,105 @@
//
// Tomás Oliveira e Silva, September 2025
//
// Arquiteturas de Alto Desempenho 2025/2026
//
#include <time.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include "aad_data_types.h"
#include "aad_utilities.h"
#include "aad_sha1_cpu.h"
#include "aad_cuda_utilities.h"
static void test_sha1_cuda(int n_tests)
{
u32_t n,*interleaved32_data,*interleaved32_hash,data[14],hash[5],good_hash[5];
double host_to_device_time,kernel_time,device_to_host_time,hashes_per_second;
cuda_data_t cd;
if(n_tests <= 0 || n_tests > (1 << 24) || n_tests % RECOMENDED_CUDA_BLOCK_SIZE != 0)
{
fprintf(stderr,"test_sha1_cuda(): bad number of tests\n");
exit(1);
}
// initialize
cd.device_number = 0; // first device
cd.cubin_file_name = "sha1_cuda_kernel.cubin";
cd.kernel_name = "sha1_cuda_kernel";
cd.data_size[0] = (u32_t)n_tests * (u32_t)14 * (u32_t)sizeof(u32_t); // size of the data array
cd.data_size[1] = (u32_t)n_tests * (u32_t) 5 * (u32_t)sizeof(u32_t); // size of the hash array
fprintf(stderr,"test_sha1_cuda(): %.3f MiB bytes for the interleaved32_data[] array\n",(double)cd.data_size[0] / (double)(1 << 20));
fprintf(stderr,"test_sha1_cuda(): %.3f MiB bytes for the interleaved32_hash[] array\n",(double)cd.data_size[1] / (double)(1 << 20));
initialize_cuda(&cd);
interleaved32_data = (u32_t *)cd.host_data[0];
interleaved32_hash = (u32_t *)cd.host_data[1];
// random interleaved32_data
n = cd.data_size[0];
while(n != 0u)
((u08_t *)interleaved32_data)[--n] = random_byte();
// run SHA1 in the CUDA device
time_measurement();
host_to_device_copy(&cd,0); // idx=0 means that the interleaved32_data is copied to the CUDA device
time_measurement();
host_to_device_time = wall_time_delta();
cd.grid_dim_x = (u32_t)n_tests / (u32_t)RECOMENDED_CUDA_BLOCK_SIZE;
cd.block_dim_x = (u32_t)RECOMENDED_CUDA_BLOCK_SIZE;
cd.n_kernel_arguments = 2;
cd.arg[0] = &cd.device_data[0]; // interleaved32_data
cd.arg[1] = &cd.device_data[1]; // interleaved32_hash
time_measurement();
lauch_kernel(&cd);
time_measurement();
kernel_time = wall_time_delta();
time_measurement();
device_to_host_copy(&cd,1); // idx=1 means that the interleaved32_hash is copied to the host
time_measurement();
device_to_host_time = wall_time_delta();
// test
for(n = 0;n < n_tests;n++)
{
// deinterleave the data and the hash
// on the CUDA side, the data for each warp is clustered together; what follows must match what is in the CUDA kernel
// each warp has 32 threads
int warp_number = n / 32;
int lane = n % 32;
for(int idx = 0;idx < 14;idx++)
data[idx] = interleaved32_data[32 * 14 * warp_number + 32 * idx + lane];
for(int idx = 0;idx < 5;idx++)
hash[idx] = interleaved32_hash[32 * 5 * warp_number + 32 * idx + lane];
// compute the SHA1 secure hahs on the cpu
sha1(&data[0],&good_hash[0]);
// compare them
for(int idx = 0;idx < 5;idx++)
if(hash[idx] != good_hash[idx])
{
fprintf(stderr,"test_sha1_cuda() failed for n=%d\n",n);
for(idx = 0;idx < 14;idx++)
fprintf(stderr,"%2d 0x%08X\n",idx,data[idx]);
fprintf(stderr,"---\n");
for(idx = 0;idx < 5;idx++)
fprintf(stderr,"%2d 0x%08X 0x%08X\n",idx,good_hash[idx],hash[idx]);
exit(1);
}
}
// cleanup
terminate_cuda(&cd);
hashes_per_second = (double)n_tests / kernel_time;
printf("sha1_cuda_kernel() passed (%d test%s, %.0f secure hashes per second)\n",n_tests,(n_tests == 1) ? "" : "s",hashes_per_second);
printf(" host -> device --- %.6f seconds\n",host_to_device_time);
printf(" kernel ----------- %.6f seconds\n",kernel_time);
printf(" device -> host --- %.6f seconds\n",device_to_host_time);
}
//
// main program
//
int main(void)
{
test_sha1_cuda(128 * 65536);
return 0;
}

65
aad_utilities.h Normal file
View File

@ -0,0 +1,65 @@
//
// Tomás Oliveira e Silva, September 2025
//
// Arquiteturas de Alto Desempenho 2025/2026
//
// useful functions (all functions are marked with the unused attribute)
// the compiler will not complain if they are not actually used in the code
//
#ifndef AAD_UTILITIES
#define AAD_UTILITIES
//
// measure elapsed and wall times --- requires <time.h>
//
// warning: Linux and macOS only, if clock_gettime() is not available, consider using clock()
//
static struct timespec measured_cpu_time[2],measured_wall_time[2];
__attribute__((unused))
static void time_measurement(void)
{
measured_cpu_time[0] = measured_cpu_time[1];
(void)clock_gettime(CLOCK_PROCESS_CPUTIME_ID,&measured_cpu_time[1]);
measured_wall_time[0] = measured_wall_time[1];
(void)clock_gettime(CLOCK_MONOTONIC_RAW,&measured_wall_time[1]);
}
__attribute__((unused))
static double cpu_time_delta(void)
{
return ((double)measured_cpu_time[1].tv_sec - (double)measured_cpu_time[0].tv_sec)
+ 1.0e-9 * ((double)measured_cpu_time[1].tv_nsec - (double)measured_cpu_time[0].tv_nsec);
}
__attribute__((unused))
static double wall_time_delta(void)
{
return ((double)measured_wall_time[1].tv_sec - (double)measured_wall_time[0].tv_sec)
+ 1.0e-9 * ((double)measured_wall_time[1].tv_nsec - (double)measured_wall_time[0].tv_nsec);
}
//
// linear congruential pseudo-random number generator with period 2^32
// see, for example, https://en.wikipedia.org/wiki/Linear_congruential_generator
// not good for cryptographic applications, but good enough to generate test data
//
__attribute__((unused))
u08_t random_byte(void)
{
static u32_t x = 0u;
x = 3134521u * x + 1u;
return (u08_t)x;
}
//
// the end!
//
#endif

145
aad_vault.h Normal file
View File

@ -0,0 +1,145 @@
//
// Tomás Oliveira e Silva, September 2025
//
// Arquiteturas de Alto Desempenho 2025/2026
//
// implements a vault for all found DETI coins
//
#ifndef AAD_VAULT
#define AAD_VAULT
static void save_coin(u32_t coin[14])
{
# define VAULT_FILE_NAME "deti_coins_v2_vault.txt"
# define MAX_SAVED_COINS 65536u
static u08_t saved_coins[MAX_SAVED_COINS][4 + 55];
static u32_t n_saved_coins = 0u;
static u08_t deti_coin_v2_template[56u] =
{ // non-zero entries are mandatory, the others are arbitrary
[ 0u] = (u08_t)'D',
[ 1u] = (u08_t)'E',
[ 2u] = (u08_t)'T',
[ 3u] = (u08_t)'I',
[ 4u] = (u08_t)' ',
[ 5u] = (u08_t)'c',
[ 6u] = (u08_t)'o',
[ 7u] = (u08_t)'i',
[ 8u] = (u08_t)'n',
[ 9u] = (u08_t)' ',
[10u] = (u08_t)'2',
[11u] = (u08_t)' ',
[54u] = (u08_t)'\n',
[55u] = (u08_t)0x80
};
static int error_tolerance_count = 4; // number of errors to tolerate before bailing out
u32_t idx,n,hash[5];
char *reason;
u08_t *s;
//
// handle a NULL argument (meaning: save all stored DETI coins) or an already full buffer
//
if(coin == NULL || n_saved_coins == MAX_SAVED_COINS)
{
if(n_saved_coins > 0u)
{
FILE *fp = fopen(VAULT_FILE_NAME,"a");
if(fp == NULL ||
fwrite((void *)&saved_coins[0][0],(size_t)(4 + 55),(size_t)n_saved_coins,fp) != (size_t)n_saved_coins ||
fflush(fp) != 0 ||
fclose(fp) != 0)
{
fprintf(stderr,"save_coin(): error while updating file \"" VAULT_FILE_NAME "\"\n");
exit(1);
}
}
n_saved_coins = 0u;
}
if(coin == NULL)
return;
//
// compute the SHA1 secure hash
//
sha1(coin,hash);
//
// make sure that the coin has the appropriate format
//
for(idx = 0u;idx < 56u;idx++)
if((deti_coin_v2_template[idx] != (u08_t)0 && deti_coin_v2_template[idx] != ((u08_t *)coin)[idx ^ 3]) || (idx >= 12u && idx <= 53u && ((char *)coin)[idx ^ 3] == '\n'))
{
reason = "coin does not match the template";
error:
fprintf(stderr,"save_coin(): bad DETI coin v2 format (%s)\n",reason);
fprintf(stderr," coin contents\n");
fprintf(stderr," idx template coin\n");
fprintf(stderr," --- --------- ---------\n");
for(idx = 0u;idx < 56u;idx++)
{
u08_t t = deti_coin_v2_template[idx];
u08_t c = ((u08_t *)coin)[idx ^ 3];
fprintf(stderr," %3u",idx);
if(t == (u08_t)0)
fprintf(stderr," arbitrary");
else if(t == '\n')
fprintf(stderr," 0x%02X '\\n'",(int)t);
else if(t >= 32 && t <= 126)
fprintf(stderr," 0x%02X '%c'",(int)t,t);
else
fprintf(stderr," 0x%02X ",(int)t);
fprintf(stderr," 0x%02X",(int)c);
if(c == '\n')
fprintf(stderr," '\\n'");
else if(c == '\b')
fprintf(stderr," '\\b'");
else if(c >= 32 && c <= 126)
fprintf(stderr," '%c'",c);
else
fprintf(stderr," ");
if((t != (u08_t)0 && t != c) || (idx >= 12u && idx <= 53u && c == '\n'))
fprintf(stderr," error");
fprintf(stderr,"\n");
}
fprintf(stderr," --- --------- ---------\n");
fprintf(stderr," SHA1 secure hash\n");
fprintf(stderr," idx value\n");
fprintf(stderr," --- ----------\n");
for(idx = 0u;idx < 5u;idx++)
fprintf(stderr," %u 0x%08X%s\n",idx,hash[idx],(idx == 0u && hash[idx] != 0xAAD20250u) ? " error" : "");
fprintf(stderr," --- ----------\n\n");
if(--error_tolerance_count < 0)
exit(1); // too many errors, exit
return; // ignore this coin
}
//
// chech the DETI coin v2 signature
//
if(hash[0] != 0xAAD20250u)
{
reason = "bad coin signature";
goto error;
}
//
// count the number of leading zeros bits of the last 4 32-bit words of the SHA1 secure hash
//
for(n = 0u;n < 128u;n++)
if((hash[1u + n / 32u] >> (31u - n % 32u)) % 2u != 0u)
break;
//
// save the coin in the buffer
// format of each line: "Vuv:" "coin_data" where u and v are ascii digits that encode, in base 10, the reported power of the coin
//
if(n > 99u)
n = 99u;
s = &saved_coins[n_saved_coins++][0];
*s++ = (u08_t)'V';
*s++ = (u08_t)('0' + n / 10u);
*s++ = (u08_t)('0' + n % 10u);
*s++ = (u08_t)':';
for(idx = 0u;idx < 55u;idx++)
*s++ = ((u08_t *)coin)[idx ^ 3];
# undef VAULT_FILE_NAME
# undef MAX_SAVED_COINS
}
#endif

1
deti_coin_example.txt Normal file
View File

@ -0,0 +1 @@
DETI coin 2 000004408026884AAD é fixe.

BIN
logo.png Normal file

Binary file not shown.

After

Width:  |  Height:  |  Size: 255 KiB

66
makefile Normal file
View File

@ -0,0 +1,66 @@
#
# Arquiteturas de Alto Desempenho 2025/2026
#
# makefile for the first practical assignment (A1)
#
# makefile automatic variables:
# $@ is the name of the target
# $< is the name of the first prerequisite
# $^ is the list of names of all prerequisites (without duplicates)
#
#
# CUDA installation directory --- /usr/local/cuda or $(CUDA_HOME)
#
CUDA_DIR = /opt/cuda
#
# OpenCL installation directory (for a NVidia graphics card, sama as CUDA)
#
OPENCL_DIR = $(CUDA_DIR)
#
# CUDA device architecture
#
# GeForce GTX 1660 Ti --- sm_75
# RTX A2000 Ada --------- sm_86
# RTX A6000 Ada --------- sm_86
# RTX 4070 -------------- sm_89
#
CUDA_ARCH = sm_75
#
# clean up
#
clean:
rm -f sha1_tests
rm -f sha1_cuda_test sha1_cuda_kernel.cubin
rm -f a.out
#
# test the CUSTOM_SHA1_CODE macro
#
sha1_tests: aad_sha1_cpu_tests.c aad_sha1.h aad_data_types.h aad_utilities.h makefile
cc -march=native -Wall -Wshadow -Werror -O3 $< -o $@
sha1_cuda_test: aad_sha1_cuda_test.c sha1_cuda_kernel.cubin aad_sha1.h aad_data_types.h aad_utilities.h aad_cuda_utilities.h makefile
cc -march=native -Wall -Wshadow -Werror -O3 -I$(CUDA_DIR)/include $< -o $@ -lcuda
#
# compile the CUDA kernels
#
sha1_cuda_kernel.cubin: aad_sha1_cuda_kernel.cu aad_sha1.h makefile
nvcc -arch=$(CUDA_ARCH) --compiler-options -O2,-Wall -I$(CUDA_DIR)/include --cubin $< -o $@
all: sha1_tests sha1_cuda_test sha1_cuda_kernel.cubin

1235
rfc3174.txt Normal file

File diff suppressed because it is too large Load Diff

BIN
sha1_cuda_kernel.cubin Normal file

Binary file not shown.

BIN
sha1_tests Executable file

Binary file not shown.

8
test_vault.bash Normal file
View File

@ -0,0 +1,8 @@
#! /bin/bash
set -e
vault_file=deti_coins_v2_vault.txt
while IFS= read -r line; do
hash=$(echo "$line" | cut -b 5- | sha1sum | cut -b 1-40)
echo "$line --- $hash"
done < ${vault_file}