4
$\begingroup$

It seems that there are some errors in the documentation page Consider these functions in the documentation CUDA Programming. However, I'm not a expert in CUDA nor C, could you check whether these are indeed errors or I'm missing something?

error 1

__global__ void addTwo_kernel(mint * arry, mint len) {

    int index = threadIdx.x + blockIdx * blockDim.x;

    if (index >= len) return;

    arry[index] += 2;
}

Should blockIdx be blockIdx.x ?

error 2

void xAverage_cpu (mint * in, mint * out, mint width, mint height, mint pitch) {
      int xIndex, yIndex, idx;
      mint left, curr, right;

      for (yIndex = 0; yIndex < height; yIndex++) {
            for (xIndex = 0; xIndex < width; xIndex++) {
                  idx = x + y*pitch;

                  left = idx > 0 ? in[idx - 1] : in[yIndex*pitch];
                  curr = in[idx];
                  right = idx < width - 1 ? in[idx + 1] : in[width + yIndex*pitch];

                  out[x + y*pitch] = (left + curr + right)/3;
              }
        }
  }

Should x and y be xIndex and yIndex instead ?

error 3

void xAverage_cpu(mint * in, mint * out, mint width, mint height, mint pitch) {
  mint accum;
  mint left, curr, right;

  for (yIndex = 0; yIndex < height; yIndex++) {
      accum = 2*in[yIndex*pitch] + in[yIndex*pitch + 1]
      out[yIndex*pitch] = accum/3;

      accum -= in[yIndex*pitch];

      for (xIndex = 0; xIndex < width-1; xIndex++) {
          idx = x + y*pitch;

          accum += in[index+1];
          out[idx] = accum/3;
          accum -= in[idx-1];
      }
      accum += in[width - 2 + yIndex*pitch];
      out[width - 1 + yIndex*pitch] = accum/3;
  }
}

Should

  1. index in accum += in[index+1] be idx instead?
  2. a ";" is missing after accum = 2*in[yIndex*pitch] + in[yIndex*pitch + 1]?
  3. the declarations for xIndex, yIndex and idx are missing?
  4. xIndex = 0 in for (xIndex = 0; xIndex < width-1; xIndex++) should be xIndex = 1?

error 4

__global__ void xAverage_improved_kernel(mint * in, mint * out, mint width, mint height, mint pitch) {
  extern __shared__ smem[];

  int tx = threadIdx.x, dx = blockDim.x;
  int xIndex = dx*blockIdx.x + tx;
  int yIndex = blockDim.y*blockIdx.y + threadIdx.y;
  int index = xIndex + yIndex*pitch;

  if (yIndex >= height) return;

  smem[threadIdx.x+1] = index < width ? in[index] : in[yIndex*pitch + width - 1];

  if (tx == 0) {
      smem[0] = index > 0 ? in[index-1] : in[yIndex * pitch];
      smem[dx+1] = index+bx < width ? in[index+dx] : in[yIndex*pitch + width - 1];
  }

  __syncthreads();

  if (xIndex < width)
      out[index] = (smem[tx] + smem[tx+1] + smem[tx+2])/3;
}

Should

  1. bx be dx?
  2. extern __shared__ mint smem[] instead of extern __shared__ smem[]?
$\endgroup$
  • $\begingroup$ I don't really have experience here, but it seems that might be the case when you compare yours with the similar example under the heading "Second Kernel" a little further down that page, in which you indeed have blockIdx.x in an identical construction as the one you shown above. $\endgroup$ – MarcoB Sep 12 '15 at 0:23
  • $\begingroup$ Are you able to test these code snippets in their original form and with your proposed changes? I don't have access to CUDA capable hardware. $\endgroup$ – MarcoB Sep 12 '15 at 16:07
  • $\begingroup$ @MarcoB I'm able to test some, but these errors are all basic compiling errors. $\endgroup$ – xslittlegrass Sep 12 '15 at 16:09
  • 1
    $\begingroup$ If you have confirmed these as errors, you might want to include that in your question. Additionally, I'm afraid that I don't see what your question is here. It might be more expedient to report these documentation problems to Wolfram Support instead. $\endgroup$ – MarcoB Sep 12 '15 at 16:40
  • $\begingroup$ @MarcoB I'm not a expert on CUDA, so additional check from others may be needed. I'll report to Wolfram support shortly. I post here just to keep it in the bug tracking list the community is maintaining. $\endgroup$ – xslittlegrass Sep 12 '15 at 16:53

Your Answer

By clicking “Post Your Answer”, you agree to our terms of service, privacy policy and cookie policy

Browse other questions tagged or ask your own question.