diff --git a/compilation/HostTranslation/src/x86/ReplaceCudaBuiltin.cpp b/compilation/HostTranslation/src/x86/ReplaceCudaBuiltin.cpp index b7d0030..5c1ea5d 100644 --- a/compilation/HostTranslation/src/x86/ReplaceCudaBuiltin.cpp +++ b/compilation/HostTranslation/src/x86/ReplaceCudaBuiltin.cpp @@ -164,7 +164,8 @@ void ReplaceKernelLaunch(llvm::Module *M) { FunctionType *ft = calledFunction->getFunctionType(); DEBUG_INFO("Parent (Caller) Function Name: %s, " "cudaLaunchKernel Function: %s, args : %d\n", - func_name, functionOperand->getName().str().c_str(), + func_name.c_str(), + functionOperand->getName().str().c_str(), functionOperand->arg_size()); auto rep = kernels.find(functionOperand->getName().str()); if (rep != kernels.end()) { diff --git a/examples/sharedMemory/reverse.cu b/examples/sharedMemory/reverse.cu new file mode 100644 index 0000000..a990835 --- /dev/null +++ b/examples/sharedMemory/reverse.cu @@ -0,0 +1,40 @@ +// Get from: https://developer.nvidia.com/blog/using-shared-memory-cuda-cc/ +#include +#include + +__global__ void staticReverse(int *d, int n) +{ + __shared__ int s[64]; + int t = threadIdx.x; + int tr = n-t-1; + s[t] = d[t]; + __syncthreads(); + d[t] = s[tr]; +} + +int main() +{ + const int n = 64; + int a[n], r[n], d[n]; + + for (int i = 0; i < n; i++) { + a[i] = i; + r[i] = n-i-1; + d[i] = 0; + } + + int *d_d; + cudaMalloc(&d_d, n * sizeof(int)); + + // run version with static shared memory + cudaMemcpy(d_d, a, n*sizeof(int), cudaMemcpyHostToDevice); + staticReverse<<<1,n>>>(d_d, n); + cudaMemcpy(d, d_d, n*sizeof(int), cudaMemcpyDeviceToHost); + for (int i = 0; i < n; i++) + if (d[i] != r[i]) { + printf("Error: d[%d]!=r[%d] (%d, %d)n", i, i, d[i], r[i]); + return 1; + } + printf("PASS\n"); + return 0; +}