<p>So, im trying to write some code that utilizes Nvidia's CUDA architecture. I noticed that copying to and from the device was really hurting my overall performance, so now I am trying to move a large amount of data onto the device.</p>

<p>As this data is used in numerous functions, I would like it to be global. Yes, I can pass pointers around, but I would really like to know how to work with globals in this instance. </p>

<p>So, I have device functions that want to access a device allocated array.</p>

<p>Ideally, I could do something like:</p>

<pre><code>__device__ float* global_data;

main()
{
  cudaMalloc(global_data);
  kernel1&lt;&lt;&lt;blah&gt;&gt;&gt;(blah); //access global data
  kernel2&lt;&lt;&lt;blah&gt;&gt;&gt;(blah); //access global data again
}
</code></pre>

<p>However, I havent figured out how to create a dynamic array. I figured out a work around by declaring the array as follows:</p>

<pre><code>__device__ float global_data[REALLY_LARGE_NUMBER];
</code></pre>

<p>And while that doesn't require a cudaMalloc call, I would prefer the dynamic allocation approach.</p>

[{'answer_id': 79256, 'author': 'Community', 'author_id': -1, 'author_profile': 'https://Stackoverflow.com/users/-1', 'pm_score': 1, 'selected': False, 'text': '<p>Spend some time focusing on the copious documentation offered by NVIDIA.</p>\n\n<p>From the Programming Guide:</p>\n\n<pre><code>float* devPtr;\ncudaMalloc((void**)&amp;devPtr, 256 * sizeof(*devPtr));\ncudaMemset(devPtr, 0, 256 * sizeof(*devPtr));\n</code></pre>\n\n<p>That\'s a simple example of how to allocate memory.  Now, in your kernels, you should accept a pointer to a float like so:</p>\n\n<pre><code>__global__\nvoid kernel1(float *some_neat_data)\n{\n    some_neat_data[threadIdx.x]++;\n}\n\n__global__\nvoid kernel2(float *potentially_that_same_neat_data)\n{\n    potentially_that_same_neat_data[threadIdx.x] *= 0.3f;\n}\n</code></pre>\n\n<p>So now you can invoke them like so:</p>\n\n<pre><code>float* devPtr;\ncudaMalloc((void**)&amp;devPtr, 256 * sizeof(*devPtr));\ncudaMemset(devPtr, 0, 256 * sizeof(*devPtr));\n\nkernel1&lt;&lt;&lt;1,128&gt;&gt;&gt;(devPtr);\nkernel2&lt;&lt;&lt;1,128&gt;&gt;&gt;(devPtr);\n</code></pre>\n\n<blockquote>\n  <p>As this data is used in numerous\n  functions, I would like it to be\n  global.</p>\n</blockquote>\n\n<p>There are few good reasons to use globals.  This definitely is not one.  I\'ll leave it as an exercise to expand this example to include moving "devPtr" to a global scope.</p>\n\n<p>EDIT:</p>\n\n<p>Ok, the fundamental problem is this: your kernels can only access device memory and the only global-scope pointers that they can use are GPU ones.  When calling a kernel from your CPU, behind the scenes what happens is that the pointers and primitives get copied into GPU registers and/or shared memory before the kernel gets executed.</p>\n\n<p>So the closest I can suggest is this: use cudaMemcpyToSymbol() to achieve your goals.  But, in the background, consider that a different approach might be the Right Thing.</p>\n\n<pre><code>#include &lt;algorithm&gt;\n\n__constant__ float devPtr[1024];\n\n__global__\nvoid kernel1(float *some_neat_data)\n{\n    some_neat_data[threadIdx.x] = devPtr[0] * devPtr[1];\n}\n\n__global__\nvoid kernel2(float *potentially_that_same_neat_data)\n{\n    potentially_that_same_neat_data[threadIdx.x] *= devPtr[2];\n}\n\n\nint main(int argc, char *argv[])\n{\n    float some_data[256];\n    for (int i = 0; i &lt; sizeof(some_data) / sizeof(some_data[0]); i++)\n    {\n        some_data[i] = i * 2;\n    }\n    cudaMemcpyToSymbol(devPtr, some_data, std::min(sizeof(some_data), sizeof(devPtr) ));\n    float* otherDevPtr;\n    cudaMalloc((void**)&amp;otherDevPtr, 256 * sizeof(*otherDevPtr));\n    cudaMemset(otherDevPtr, 0, 256 * sizeof(*otherDevPtr));\n\n    kernel1&lt;&lt;&lt;1,128&gt;&gt;&gt;(otherDevPtr);\n    kernel2&lt;&lt;&lt;1,128&gt;&gt;&gt;(otherDevPtr);\n\n    return 0;\n}\n</code></pre>\n\n<p>Don\'t forget \'--host-compilation=c++\' for this example.</p>\n'}, {'answer_id': 79336, 'author': 'Voltaire', 'author_id': 14621, 'author_profile': 'https://Stackoverflow.com/users/14621', 'pm_score': 0, 'selected': False, 'text': "<p>Erm, it was exactly that problem of moving devPtr to global scope that was my problem.</p>\n\n<p>I have an implementation that does exactly that, with the two kernels having a pointer to data passed in. I explicitly don't want to pass in those pointers.</p>\n\n<p>I have read the documentation fairly closely, and hit up the nvidia forums (and google searched for an hour or so), but I haven't found an implementation of a global dynamic device array that actually runs (i have tried several that compile and then fail in new and interesting ways).</p>\n"}, {'answer_id': 79796, 'author': 'Community', 'author_id': -1, 'author_profile': 'https://Stackoverflow.com/users/-1', 'pm_score': 4, 'selected': True, 'text': '<p>Something like this should probably work.</p>\n\n<pre><code>#include &lt;algorithm&gt;\n\n#define NDEBUG\n#define CUT_CHECK_ERROR(errorMessage) do {                                 \\\n        cudaThreadSynchronize();                                           \\\n         cudaError_t err = cudaGetLastError();                             \\\n         if( cudaSuccess != err) {                                         \\\n                     fprintf(stderr, "Cuda error: %s in file \'%s\' in line %i : %s.\\n",    \\\n                                             errorMessage, __FILE__, __LINE__, cudaGetErrorString( err) );\\\n                     exit(EXIT_FAILURE);                                                  \\\n                 } } while (0)\n\n\n__device__ float *devPtr;\n\n__global__\nvoid kernel1(float *some_neat_data)\n{\n    devPtr = some_neat_data;\n}\n\n__global__\nvoid kernel2(void)\n{\n    devPtr[threadIdx.x] *= .3f;\n}\n\n\nint main(int argc, char *argv[])\n{\n    float* otherDevPtr;\n    cudaMalloc((void**)&amp;otherDevPtr, 256 * sizeof(*otherDevPtr));\n    cudaMemset(otherDevPtr, 0, 256 * sizeof(*otherDevPtr));\n\n    kernel1&lt;&lt;&lt;1,128&gt;&gt;&gt;(otherDevPtr);\n    CUT_CHECK_ERROR("kernel1");\n\n    kernel2&lt;&lt;&lt;1,128&gt;&gt;&gt;();\n\n    CUT_CHECK_ERROR("kernel2");\n\n    return 0;\n}\n</code></pre>\n\n<p>Give it a whirl.</p>\n'}, {'answer_id': 89477, 'author': 'Mark Borgerding', 'author_id': 3343, 'author_profile': 'https://Stackoverflow.com/users/3343', 'pm_score': 0, 'selected': False, 'text': '<p>check out the samples included with the SDK.  Many of those sample projects are a decent way to learn by example.</p>\n'}, {'answer_id': 94437, 'author': 'Voltaire', 'author_id': 14621, 'author_profile': 'https://Stackoverflow.com/users/14621', 'pm_score': 1, 'selected': False, 'text': '<p>I went ahead and tried the solution of allocating a temporary pointer and passing it to a simple global function similar to kernel1.</p>\n\n<p>The good news is that it does work :)</p>\n\n<p>However, I think it confuses the compiler as I now get "Advisory: Cannot tell what pointer points to, assuming global memory space" whenever I try to access the global data. Luckily, the assumption happens to be correct, but the warnings are annoying.</p>\n\n<p>Anyway, for the record - I have looked at many of the examples and did run through the nvidia exercises where the point is to get the output to say "Correct!". However, I haven\'t looked at <em>all</em> of them. If anyone knows of an sdk example where they do dynamic global device memory allocation, I would still like to know.</p>\n'}, {'answer_id': 3511490, 'author': 'Nils', 'author_id': 44232, 'author_profile': 'https://Stackoverflow.com/users/44232', 'pm_score': 0, 'selected': False, 'text': '<blockquote>\n  <p>As this data is used in numerous functions, I would like it to be global.</p>\n</blockquote>\n\n<p>-</p>\n\n<blockquote>\n  <p>There are few good reasons to use globals. This definitely is not one. I\'ll leave it as an \n  exercise to expand this example to include moving "devPtr" to a global scope.</p>\n</blockquote>\n\n<p>What if the kernel operates on a large const structure consisting of arrays? Using the so called constant memory is not an option, because it\'s very limited in size.. so then you have to put it in global memory..?</p>\n'}]
['https://Stackoverflow.com/questions/79121', 'https://Stackoverflow.com', 'https://Stackoverflow.com/users/14621/']