diff --git a/demos/CUDA/BlackScholes/BlackScholes_kernel.cuh b/demos/CUDA/BlackScholes/BlackScholes_kernel.cuh index 609eab932..bba537301 100644 --- a/demos/CUDA/BlackScholes/BlackScholes_kernel.cuh +++ b/demos/CUDA/BlackScholes/BlackScholes_kernel.cuh @@ -106,9 +106,7 @@ __global__ void BlackScholesGPU(float2* __restrict d_CallResult, BlackScholesBodyGPU(callResult2, putResult2, d_StockPrice[opt].y, d_OptionStrike[opt].y, d_OptionYears[opt].y, Riskfree, Volatility); - d_CallResult[opt].x = callResult1; - d_CallResult[opt].y = callResult2; - d_PutResult[opt].x = putResult1; - d_PutResult[opt].y = putResult2; + d_CallResult[opt] = make_float2(callResult1, callResult2); + d_PutResult[opt] = make_float2(putResult1, putResult2); } -} +} \ No newline at end of file diff --git a/include/clad/Differentiator/BuiltinDerivatives.h b/include/clad/Differentiator/BuiltinDerivatives.h index 8b8b4e09a..b2f6ba3e1 100644 --- a/include/clad/Differentiator/BuiltinDerivatives.h +++ b/include/clad/Differentiator/BuiltinDerivatives.h @@ -408,6 +408,15 @@ CUDA_HOST_DEVICE inline void sqrtf_pullback(float a, float d_y, float* d_a) { *d_a += (1.F / (2.F * sqrtf(a))) * d_y; } + +#ifdef __CUDACC__ +CUDA_HOST_DEVICE inline void make_float2_pullback(float a, float b, float2 d_y, + float* d_a, float* d_b) { + *d_a += d_y.x; + *d_b += d_y.y; +} +#endif + // These are required because C variants of mathematical functions are // defined in global namespace. using std::abs_pushforward; diff --git a/lib/Differentiator/ReverseModeVisitor.cpp b/lib/Differentiator/ReverseModeVisitor.cpp index 19b902ce1..e3c829ca7 100644 --- a/lib/Differentiator/ReverseModeVisitor.cpp +++ b/lib/Differentiator/ReverseModeVisitor.cpp @@ -1891,8 +1891,12 @@ Expr* getArraySizeExpr(const ArrayType* AT, ASTContext& context, e = CE->getNumArgs(); i != e; ++i) { const Expr* arg = CE->getArg(i); - const auto* PVD = FD->getParamDecl( - i - static_cast(isMethodOperatorCall)); + auto* PVD = const_cast(FD->getParamDecl( + i - static_cast(isMethodOperatorCall))); + if (PVD->getType()->isRValueReferenceType()) { + IdentifierInfo* RValueName = CreateUniqueIdentifier("_r"); + PVD->setDeclName(RValueName); + } StmtDiff argDiff{}; // We do not need to create result arg for arguments passed by reference // because the derivatives of arguments passed by reference are directly diff --git a/test/Gradient/STLCustomDerivatives.C b/test/Gradient/STLCustomDerivatives.C index 58ae5d64d..a32053e57 100644 --- a/test/Gradient/STLCustomDerivatives.C +++ b/test/Gradient/STLCustomDerivatives.C @@ -817,16 +817,16 @@ int main() { // CHECK-NEXT: std::vector _d_a({}); // CHECK-NEXT: std::vector a; // CHECK-NEXT: std::vector _t0 = a; -// CHECK-NEXT: {{.*}}push_back_reverse_forw(&a, 0{{.*}}, &_d_a, _r0); +// CHECK-NEXT: {{.*}}push_back_reverse_forw(&a, 0{{.*}}, &_d_a, _r1); // CHECK-NEXT: std::vector _t1 = a; -// CHECK-NEXT: {{.*}}ValueAndAdjoint _t2 = {{.*}}operator_subscript_reverse_forw(&a, 0, &_d_a, _r1); +// CHECK-NEXT: {{.*}}ValueAndAdjoint _t2 = {{.*}}operator_subscript_reverse_forw(&a, 0, &_d_a, _r2); // CHECK-NEXT: double _t3 = _t2.value; // CHECK-NEXT: _t2.value = x * x; // CHECK-NEXT: std::vector _t4 = a; -// CHECK-NEXT: {{.*}}ValueAndAdjoint _t5 = {{.*}}operator_subscript_reverse_forw(&a, 0, &_d_a, _r2); +// CHECK-NEXT: {{.*}}ValueAndAdjoint _t5 = {{.*}}operator_subscript_reverse_forw(&a, 0, &_d_a, _r3); // CHECK-NEXT: { -// CHECK-NEXT: {{.*}}size_type _r2 = 0{{.*}}; -// CHECK-NEXT: {{.*}}operator_subscript_pullback(&_t4, 0, 1, &_d_a, &_r2); +// CHECK-NEXT: {{.*}}size_type _r3 = 0{{.*}}; +// CHECK-NEXT: {{.*}}operator_subscript_pullback(&_t4, 0, 1, &_d_a, &_r3); // CHECK-NEXT: } // CHECK-NEXT: { // CHECK-NEXT: _t2.value = _t3; @@ -834,11 +834,11 @@ int main() { // CHECK-NEXT: _t2.adjoint = 0{{.*}}; // CHECK-NEXT: *_d_x += _r_d0 * x; // CHECK-NEXT: *_d_x += x * _r_d0; -// CHECK-NEXT: {{.*}}size_type _r1 = 0{{.*}}; -// CHECK-NEXT: {{.*}}operator_subscript_pullback(&_t1, 0, 0{{.*}}, &_d_a, &_r1); +// CHECK-NEXT: {{.*}}size_type _r2 = 0{{.*}}; +// CHECK-NEXT: {{.*}}operator_subscript_pullback(&_t1, 0, 0{{.*}}, &_d_a, &_r2); // CHECK-NEXT: } // CHECK-NEXT: { -// CHECK-NEXT: {{.*}}value_type _r0 = 0.; -// CHECK-NEXT: {{.*}}push_back_pullback(&_t0, 0{{.*}}, &_d_a, &_r0); +// CHECK-NEXT: {{.*}}value_type _r1 = 0.; +// CHECK-NEXT: {{.*}}push_back_pullback(&_t0, 0{{.*}}, &_d_a, &_r1); // CHECK-NEXT: } // CHECK-NEXT: }