Geek Answers Handbook

Cuda doesn't give the right answer when array size exceeds 1,000,000

I wrote simple code to reduce the amount, which seems to work fine until I increase the size of the array to 1 million, which can be a problem.

#define BLOCK_SIZE 128 #define ARRAY_SIZE 10000 cudaError_t addWithCuda(const long *input, long *output, int totalBlocks, size_t size); __global__ void sumKernel(const long *input, long *output) { int tid = threadIdx.x; int bid = blockDim.x * blockIdx.x; __shared__ long data[BLOCK_SIZE]; if(bid+tid < ARRAY_SIZE) data[tid] = input[bid+tid]; else data[tid] = 0; __syncthreads(); for(int i = BLOCK_SIZE/2; i >= 1; i >>= 1) { if(tid < i) data[tid] += data[tid + i]; __syncthreads(); } if(tid == 0) output[blockIdx.x] = data[0]; } int main() { int totalBlocks = ARRAY_SIZE/BLOCK_SIZE; if(ARRAY_SIZE % BLOCK_SIZE != 0) totalBlocks++; long *input = (long*) malloc(ARRAY_SIZE * sizeof(long) ); long *output = (long*) malloc(totalBlocks * sizeof(long) ); for(int i=0; i<ARRAY_SIZE; i++) { input[i] = i+1 ; } // Add vectors in parallel. cudaError_t cudaStatus = addWithCuda(input, output, totalBlocks, ARRAY_SIZE); if (cudaStatus != cudaSuccess) { fprintf(stderr, "addWithCuda failed!"); return 1; } long ans = 0; for(int i =0 ; i < totalBlocks ;i++) { ans = ans + output[i]; } printf("Final Ans : %ld",ans); // cudaDeviceReset must be called before exiting in order for profiling and // tracing tools such as Nsight and Visual Profiler to show complete traces. cudaStatus = cudaDeviceReset(); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaDeviceReset failed!"); return 1; } getchar(); return 0; } // Helper function for using CUDA to add vectors in parallel. cudaError_t addWithCuda(const long *input, long *output, int totalBlocks, size_t size) { long *dev_input = 0; long *dev_output = 0; cudaError_t cudaStatus; // Choose which GPU to run on, change this on a multi-GPU system. cudaStatus = cudaSetDevice(0); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaSetDevice failed! Do you have a CUDA-capable GPU installed?"); goto Error; } // Allocate GPU buffers for two vectors (one input, one output) . cudaStatus = cudaMalloc((void**)&dev_input, size * sizeof(long)); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMalloc failed!"); goto Error; } cudaStatus = cudaMalloc((void**)&dev_output, totalBlocks * sizeof(long)); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMalloc failed!"); goto Error; } // Copy input vectors from host memory to GPU buffers. cudaStatus = cudaMemcpy(dev_input, input, size * sizeof(long), cudaMemcpyHostToDevice); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMemcpy failed!"); goto Error; } cudaStatus = cudaMemcpy(dev_output, output, (totalBlocks) * sizeof(long), cudaMemcpyHostToDevice); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMemcpy failed!"); goto Error; } // Launch a kernel on the GPU with one thread for each element. sumKernel<<<totalBlocks, BLOCK_SIZE>>>(dev_input, dev_output); // cudaDeviceSynchronize waits for the kernel to finish, and returns // any errors encountered during the launch. cudaStatus = cudaDeviceSynchronize(); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus); goto Error; } // Copy output vector from GPU buffer to host memory. cudaStatus = cudaMemcpy(output, dev_output, totalBlocks * sizeof(long), cudaMemcpyDeviceToHost); if (cudaStatus != cudaSuccess) { fprintf(stderr, "cudaMemcpy failed!"); goto Error; } Error: cudaFree(dev_input); cudaFree(dev_output); return cudaStatus; }

and just for reference, if he needs to do something with my GPU device, my GPU is a GTXX 650ti. and here is the GPU info:

Maximum number of threads per multiprocessor: 2048

Maximum number of threads per block: 1024

Maximum dimensions of each block dimension: 1024 x 1024 x 64

Maximum dimensions of each grid size: 2147483647 x 65535 x 65535

Maximum memory step: 2147483647 bytes

Texture Alignment: 512 Bytes

+4
source share
3 answers

Actually, the answer = could not fit into the long one , so after using long double for data types this problem was solved. Thanks everyone!

+2
source

One of the problems with your code is that your last cudaMemcpy is not configured correctly:

 cudaMemcpy(output, dev_output, totalBlocks * sizeof(int), cudaMemcpyDeviceToHost);

All your data is long data, so you should copy using sizeof(long) not sizeof(int)

Another problem in your code is the wrong printf format identifier for a long data type:

 printf("\n %d \n",output[i]);

use something like this:

 printf("\n %ld \n",output[i]);

You may also have a problem with a large number of blocks if you do not compile for sm_30 architecture. In this case, proper cuda error checking identifies the problem.

+1
source

You do not check for errors after sumKernel<<<totalBlocks, BLOCK_SIZE>>>(dev_input, dev_output); . Usually, if you check the last error that occurred, it should give an invalid configuration argument error. Try adding the following after the sumKernel line.

 cudaStatus = cudaGetLastError(); if (cudaStatus != cudaSuccess) { printf(stderr, "sumKernel failed: %s\n", cudaGetErrorString(cudaStatus)); goto Error; }

See this question for more details about the error.

0
source

More articles:


All Articles