title: Cuda Create 3d texture and cudaArray(3d) from device memory
body: <p>im trying to create a texture 3d from a part of a device array.</p>
  
  <p>To do this, these are my steps:</p>
  
  <ol>
  <li>malloc Device Array</li>
  <li>Write Device Array</li>
  <li>Create CudaArray (3D)</li>
  <li>Bind Texture to CudaArray</li>
  </ol>
  
  <p>The way im doing it it creates no compiler errors, but when i run cuda-memcheck it's failing when im trying to fetch data from the Texture.</p>
  
  <blockquote>
    <p>Invalid <strong>global</strong> read of size 8 .. Address 0x10dfaf3a0 is out of bounds</p>
  </blockquote>
  
  <p>Thats why i'm guessing i declared the texture Array wrong.
  here is how i access the texture:</p>
  
  <blockquote>
    <p>tex3D(NoiseTextures[i],x,y,z)</p>
  </blockquote>
  
  <p>The way im doing the steps mentioned above:</p>
  
  <p>1.Malloc Device Array</p>
  
  <pre><code>cudaMalloc((void **)&amp;d_Noise, sqrSizeNoise*nNoise*sizeof(float));
  </code></pre>
  
  <p>2.Write Device Array</p>
  
  <pre><code>curandCreateGenerator(&amp;gen,CURAND_RNG_PSEUDO_DEFAULT);
  curandSetPseudoRandomGeneratorSeed(gen,Seed);
  curandGenerateUniform(gen, d_Noise, sqrSizeNoise*nNoise);
  curandDestroyGenerator(gen);
  </code></pre>
  
  <p>3+4.Creating the Cuda Array and binding it to the texture (Im guessing the mistake is here)</p>
  
  <pre><code>cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc&lt;float&gt;();//cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
  cudaArray *d_cuArr;
  cudaMalloc3DArray(&amp;d_cuArr, &amp;channelDesc, make_cudaExtent(SizeNoise,SizeNoise,SizeNoise), 0);
  cudaMemcpy3DParms copyParams = {0};
  
  //Loop for every separated Noise Texture (nNoise = 4)
  for(int i = 0; i &lt; nNoise; i++){
  
      //initialize the textures
      NoiseTextures[i] = texture&lt;float, 3, cudaReadModeElementType&gt;(1,cudaFilterModeLinear,cudaAddressModeWrap,channelDesc);
  
      //Array creation
      //+(sqrSizeNoise*i) is to separate the created Noise Array into smaller parts with the size of SizeNoise^3
      copyParams.srcPtr   = make_cudaPitchedPtr(d_Noise+(sqrSizeNoise*i), SizeNoise*sizeof(float), SizeNoise, SizeNoise);
      copyParams.dstArray = d_cuArr;
      copyParams.extent   = make_cudaExtent(SizeNoise,SizeNoise,SizeNoise);
      copyParams.kind     = cudaMemcpyDeviceToDevice;
      checkCudaErrors(cudaMemcpy3D(&amp;copyParams));
      //Array creation End
  
      //new Bind
      // set texture parameters
      NoiseTextures[i].normalized = true;                      // access with normalized texture coordinates
      NoiseTextures[i].filterMode = cudaFilterModeLinear;      // linear interpolation
      NoiseTextures[i].addressMode[0] = cudaAddressModeWrap;   // wrap texture coordinates
      NoiseTextures[i].addressMode[1] = cudaAddressModeWrap;
      NoiseTextures[i].addressMode[2] = cudaAddressModeWrap;
  
      // bind array to 3D texture
      checkCudaErrors(cudaBindTextureToArray(NoiseTextures[i], d_cuArr, channelDesc));
      //end Bind
  }
  cudaFreeArray(d_cuArr);
  </code></pre>
  
  <p>I've Pasted this code snippet to Pastebin so its easier to look at with colors etc.
  <a href="http://pastebin.com/SM3dYd38" rel="nofollow">http://pastebin.com/SM3dYd38</a></p>
  
  <p>I hope I clearly described my problem. If not pls comment!</p>
  
  <p>Can you help me with this?
  Thanks for reading,</p>
  
  <p>Cery</p>
  
  <p>Edit:
  Here is a complete code so you can try it on your own machine:</p>
  
  <pre><code>#include &lt;helper_cuda.h&gt;  
  #include &lt;helper_functions.h&gt;
  #include &lt;helper_cuda_gl.h&gt;
  #include &lt;texture_types.h&gt;
  #include &lt;cuda_runtime.h&gt;
  #include &lt;curand.h&gt;
  
  static texture&lt;float, 3, cudaReadModeElementType&gt; NoiseTextures[4];//texture Array
  float *d_NoiseTest;//Device Array with random floats
  int SizeNoiseTest = 32;
  int sqrSizeNoiseTest = 32768;
  
  void CreateTexture();
  
  __global__ void AccesTexture(texture&lt;float, 3, cudaReadModeElementType&gt;* NoiseTextures)
  {
          int test = tex3D(NoiseTextures[0],threadIdx.x,threadIdx.y,threadIdx.z);//by using this the error occurs
  }
  
  int
  main(int argc, char **argv)
  {
          CreateTexture();
  }
  void CreateTexture()
  {
          //curand Random Generator (needs compiler link -lcurand)
          curandGenerator_t gen;
          cudaMalloc((void **)&amp;d_NoiseTest, sqrSizeNoiseTest*4*sizeof(float));//Allocation of device Array
          curandCreateGenerator(&amp;gen,CURAND_RNG_PSEUDO_DEFAULT);
          curandSetPseudoRandomGeneratorSeed(gen,1234ULL);
          curandGenerateUniform(gen, d_NoiseTest, sqrSizeNoiseTest*4);//writing data to d_NoiseTest
          curandDestroyGenerator(gen);
  
          //cudaArray Descriptor
          cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc&lt;float&gt;();
          //cuda Array
          cudaArray *d_cuArr;
          cudaMalloc3DArray(&amp;d_cuArr, &amp;channelDesc, make_cudaExtent(SizeNoiseTest*sizeof(float),SizeNoiseTest,SizeNoiseTest), 0);
          cudaMemcpy3DParms copyParams = {0};
  
          //Loop for every separated Noise Texture (4 = 4)
          for(int i = 0; i &lt; 4; i++){
  
                  //initialize the textures
                  NoiseTextures[i] = texture&lt;float, 3, cudaReadModeElementType&gt;(1,cudaFilterModeLinear,cudaAddressModeWrap,channelDesc);
  
                  //Array creation
                  //+(sqrSizeNoise*i) is to separate the created Noise Array into smaller parts with the size of SizeNoise^3
                  copyParams.srcPtr   = make_cudaPitchedPtr(d_NoiseTest+(sqrSizeNoiseTest*i), SizeNoiseTest*sizeof(float), SizeNoiseTest, SizeNoiseTest);
                  copyParams.dstArray = d_cuArr;
                  copyParams.extent   = make_cudaExtent(SizeNoiseTest*sizeof(float),SizeNoiseTest,SizeNoiseTest);
                  copyParams.kind     = cudaMemcpyDeviceToDevice;
                  checkCudaErrors(cudaMemcpy3D(&amp;copyParams));
                  //Array creation End
  
                  //new Bind
                  // set texture parameters
                  NoiseTextures[i].normalized = true;                      // access with normalized texture coordinates
                  NoiseTextures[i].filterMode = cudaFilterModeLinear;      // linear interpolation
                  NoiseTextures[i].addressMode[0] = cudaAddressModeWrap;   // wrap texture coordinates
                  NoiseTextures[i].addressMode[1] = cudaAddressModeWrap;
                  NoiseTextures[i].addressMode[2] = cudaAddressModeWrap;
  
                  // bind array to 3D texture
                  checkCudaErrors(cudaBindTextureToArray(NoiseTextures[i], d_cuArr, channelDesc));
                  //end Bind
          }
          cudaFreeArray(d_cuArr);
  
          AccesTexture&lt;&lt;&lt;1,dim3(4,4,4)&gt;&gt;&gt;(NoiseTextures);
  }
  </code></pre>
  
  <p>You need to link -lcurand though. And include CUDA-6.0/samples/common/inc</p>
  
  <p>Im now getting a different error in this code </p>
  
  <blockquote>
    <p>code=11(cudaErrorInvalidValue) "cudaMemcpy3D(&amp;copyParams)"</p>
  </blockquote>
  
  <p>Even though it's the exact same code then my original. - Im starting to get completely confused.
  Thank you for your help</p>