From 0e8792865cc74f635f0ca79b7f791dfee04277bc Mon Sep 17 00:00:00 2001 From: kchristin Date: Wed, 20 Nov 2024 23:46:59 +0200 Subject: [PATCH 1/2] Clone printf and fprintf calls --- lib/Differentiator/ReverseModeVisitor.cpp | 3 +++ test/CUDA/GradientKernels.cu | 2 ++ 2 files changed, 5 insertions(+) diff --git a/lib/Differentiator/ReverseModeVisitor.cpp b/lib/Differentiator/ReverseModeVisitor.cpp index b90dacb99..1018a4436 100644 --- a/lib/Differentiator/ReverseModeVisitor.cpp +++ b/lib/Differentiator/ReverseModeVisitor.cpp @@ -1456,6 +1456,9 @@ Expr* getArraySizeExpr(const ArrayType* AT, ASTContext& context, return StmtDiff(Clone(CE)); } + if (FD->getNameAsString() == "printf" || FD->getNameAsString() == "fprintf") + return StmtDiff(Clone(CE)); + Expr* CUDAExecConfig = nullptr; if (const auto* KCE = dyn_cast(CE)) CUDAExecConfig = Clone(KCE->getConfig()); diff --git a/test/CUDA/GradientKernels.cu b/test/CUDA/GradientKernels.cu index 328a1f50d..6afd6b633 100644 --- a/test/CUDA/GradientKernels.cu +++ b/test/CUDA/GradientKernels.cu @@ -504,6 +504,7 @@ double fn_memory(double *out, double *in) { //CHECK-NEXT:} void launch_add_kernel_4(int *out, int *in, const int N) { + printf("Launching add_kernel_4 for size: %d\n", N); int *in_dev = nullptr; cudaMalloc(&in_dev, N * sizeof(int)); cudaMemcpy(in_dev, in, N * sizeof(int), cudaMemcpyHostToDevice); @@ -520,6 +521,7 @@ void launch_add_kernel_4(int *out, int *in, const int N) { // CHECK: void launch_add_kernel_4_grad_0_1(int *out, int *in, const int N, int *_d_out, int *_d_in) { //CHECK-NEXT: int _d_N = 0; +//CHECK-NEXT: printf("Launching add_kernel_4 for size: %d\n", N); //CHECK-NEXT: int *_d_in_dev = nullptr; //CHECK-NEXT: int *in_dev = nullptr; //CHECK-NEXT: cudaMalloc(&_d_in_dev, N * sizeof(int)); From 1ef0d2375c7ac823689a5e2a9bc8861a7ab93f3b Mon Sep 17 00:00:00 2001 From: kchristin Date: Fri, 22 Nov 2024 14:38:59 +0200 Subject: [PATCH 2/2] Add FIXME and print loop index --- lib/Differentiator/ReverseModeVisitor.cpp | 1 + test/CUDA/GradientKernels.cu | 4 ++-- 2 files changed, 3 insertions(+), 2 deletions(-) diff --git a/lib/Differentiator/ReverseModeVisitor.cpp b/lib/Differentiator/ReverseModeVisitor.cpp index 1018a4436..b0c6c001c 100644 --- a/lib/Differentiator/ReverseModeVisitor.cpp +++ b/lib/Differentiator/ReverseModeVisitor.cpp @@ -1456,6 +1456,7 @@ Expr* getArraySizeExpr(const ArrayType* AT, ASTContext& context, return StmtDiff(Clone(CE)); } + // FIXME: Revisit this when variadic functions are supported. if (FD->getNameAsString() == "printf" || FD->getNameAsString() == "fprintf") return StmtDiff(Clone(CE)); diff --git a/test/CUDA/GradientKernels.cu b/test/CUDA/GradientKernels.cu index 6afd6b633..969a1b1c1 100644 --- a/test/CUDA/GradientKernels.cu +++ b/test/CUDA/GradientKernels.cu @@ -436,6 +436,7 @@ double fn_memory(double *out, double *in) { cudaMemcpy(out_host, out, 10 * sizeof(double), cudaMemcpyDeviceToHost); double res = 0; for (int i=0; i < 10; ++i) { + printf("Writing result of out[%d]\n", i); res += out_host[i]; } free(out_host); @@ -469,6 +470,7 @@ double fn_memory(double *out, double *in) { //CHECK-NEXT: break; //CHECK-NEXT: } //CHECK-NEXT: _t0++; +//CHECK-NEXT: printf("Writing result of out[%d]\n", i); //CHECK-NEXT: clad::push(_t1, res); //CHECK-NEXT: res += out_host[i]; //CHECK-NEXT: } @@ -504,7 +506,6 @@ double fn_memory(double *out, double *in) { //CHECK-NEXT:} void launch_add_kernel_4(int *out, int *in, const int N) { - printf("Launching add_kernel_4 for size: %d\n", N); int *in_dev = nullptr; cudaMalloc(&in_dev, N * sizeof(int)); cudaMemcpy(in_dev, in, N * sizeof(int), cudaMemcpyHostToDevice); @@ -521,7 +522,6 @@ void launch_add_kernel_4(int *out, int *in, const int N) { // CHECK: void launch_add_kernel_4_grad_0_1(int *out, int *in, const int N, int *_d_out, int *_d_in) { //CHECK-NEXT: int _d_N = 0; -//CHECK-NEXT: printf("Launching add_kernel_4 for size: %d\n", N); //CHECK-NEXT: int *_d_in_dev = nullptr; //CHECK-NEXT: int *in_dev = nullptr; //CHECK-NEXT: cudaMalloc(&_d_in_dev, N * sizeof(int));