add static shared memory example
This commit is contained in:
parent
3875e179b4
commit
ba2c49abdd
|
@ -164,7 +164,8 @@ void ReplaceKernelLaunch(llvm::Module *M) {
|
||||||
FunctionType *ft = calledFunction->getFunctionType();
|
FunctionType *ft = calledFunction->getFunctionType();
|
||||||
DEBUG_INFO("Parent (Caller) Function Name: %s, "
|
DEBUG_INFO("Parent (Caller) Function Name: %s, "
|
||||||
"cudaLaunchKernel Function: %s, args : %d\n",
|
"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());
|
functionOperand->arg_size());
|
||||||
auto rep = kernels.find(functionOperand->getName().str());
|
auto rep = kernels.find(functionOperand->getName().str());
|
||||||
if (rep != kernels.end()) {
|
if (rep != kernels.end()) {
|
||||||
|
|
|
@ -0,0 +1,40 @@
|
||||||
|
// Get from: https://developer.nvidia.com/blog/using-shared-memory-cuda-cc/
|
||||||
|
#include <stdio.h>
|
||||||
|
#include <stdlib.h>
|
||||||
|
|
||||||
|
__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;
|
||||||
|
}
|
Loading…
Reference in New Issue