<p dir="ltr">hi everyone!</p>
<p dir="ltr">I tried to compile a CUDA program using clang-3.4. This program is taken<br>
from NVIDIA_CUDA-5.0_samples collection and it's a very simple program adding<br>
two vectors.</p>
<p dir="ltr">A few modifications to the original code were done, including<br>
1. I substituted the __global__ CUDA C keyword with __attribute__((global))<br>
in order to use clang as the compiler.<br>
2. <stdlib.h> <math.h> were added.<br>
3. declarations of blockDim, blockIdx, threadIdx were added.</p>
<p dir="ltr">// ==================== code begin ========================<br>
/**</p>
<p dir="ltr"> * Vector addition: C = A + B.<br>
*<br>
* This sample is a very basic sample that implements element by element<br>
* vector addition. It is the same as the sample illustrating Chapter 2<br>
* of the programming guide with some additions like error checking.<br>
*/</p>
<p dir="ltr">#include <stdio.h></p>
<p dir="ltr">// For the CUDA runtime routines (prefixed with "cuda_")<br>
#include <cuda_runtime.h></p>
<p dir="ltr">#include <malloc.h><br>
#include <stdlib.h><br>
#include <math.h><br>
extern dim3 blockDim, blockIdx, threadIdx;</p>
<p dir="ltr">/**<br>
* CUDA Kernel Device code<br>
*<br>
* Computes the vector addition of A and B into C. The 3 vectors have the same<br>
* number of elements numElements.<br>
*/<br>
__attribute__((global)) void<br>
vectorAdd(const float *A, const float *B, float *C, int numElements)<br>
{<br>
int i = blockDim.x * blockIdx.x + threadIdx.x;</p>
<p dir="ltr"> if (i < numElements)<br>
{<br>
C[i] = A[i] + B[i];<br>
}<br>
}</p>
<p dir="ltr">int main(void)<br>
{<br>
... ...<br>
return 0;<br>
}</p>
<p dir="ltr">//==================== code end ========================</p>
<p dir="ltr">$ clang -std=cuda -I/usr/local/cuda-5.0/include -o vectorAdd1 vectorAdd1.cu -L/usr/local/cuda-5.0/lib64 -lcudart</p>
<p dir="ltr">The compiling was successful, but running the program gives:</p>
<p dir="ltr"> [Vector addition of 50000 elements]<br>
Copy input data from the host memory to the CUDA device<br>
CUDA kernel launch with 196 blocks of 256 threads<br>
Failed to launch vectorAdd kernel (error code invalid device function )!</p>
<p dir="ltr">while the expected output is</p>
<p dir="ltr"> [Vector addition of 50000 elements]<br>
Copy input data from the host memory to the CUDA device<br>
CUDA kernel launch with 196 blocks of 256 threads<br>
Copy output data from the CUDA device to the host memory<br>
Done</p>
<p dir="ltr">The result said that the vectorAdd function was not a valid kernel. To see what happend,<br>
I compiled the program to LLVM IR. Only the IR of function vectorAdd is shown below.</p>
<p dir="ltr">$ clang -S -emit-llvm -std=cuda -I/usr/local/cuda-5.0/include -o vectorAdd1.ll vectorAdd1.cu</p>
<p dir="ltr">//==================== code begin ==========================<br>
define void @_Z9vectorAddPKfS0_Pfi(float* %A, float* %B, float* %C, i32 %numElements) #0 {<br>
%1 = alloca float*, align 8<br>
%2 = alloca float*, align 8<br>
%3 = alloca float*, align 8<br>
%4 = alloca i32, align 4<br>
store float* %A, float** %1, align 8<br>
store float* %B, float** %2, align 8<br>
store float* %C, float** %3, align 8<br>
store i32 %numElements, i32* %4, align 4<br>
%5 = bitcast float** %1 to i8*<br>
%6 = call i32 @cudaSetupArgument(i8* %5, i64 ptrtoint (i1** getelementptr (i1** null, i32 1) to i64), i64 0)<br>
%7 = icmp eq i32 %6, 0<br>
br i1 %7, label %8, label %22</p>
<p dir="ltr">; <label>:8 ; preds = %0<br>
%9 = bitcast float** %2 to i8*<br>
%10 = call i32 @cudaSetupArgument(i8* %9, i64 ptrtoint (i1** getelementptr (i1** null, i32 1) to i64), i64 ptrtoint (float** getelementptr ({ float*, float*, float*, i32 }* null, i64 0, i32 1) to i64))<br>
%11 = icmp eq i32 %10, 0<br>
br i1 %11, label %12, label %22</p>
<p dir="ltr">; <label>:12 ; preds = %8<br>
%13 = bitcast float** %3 to i8*<br>
%14 = call i32 @cudaSetupArgument(i8* %13, i64 ptrtoint (i1** getelementptr (i1** null, i32 1) to i64), i64 ptrtoint (float** getelementptr ({ float*, float*, float*, i32 }* null, i64 0, i32 2) to i64))<br>
%15 = icmp eq i32 %14, 0<br>
br i1 %15, label %16, label %22</p>
<p dir="ltr">; <label>:16 ; preds = %12<br>
%17 = bitcast i32* %4 to i8*<br>
%18 = call i32 @cudaSetupArgument(i8* %17, i64 ptrtoint (i32* getelementptr (i32* null, i32 1) to i64), i64 ptrtoint (i32* getelementptr ({ float*, float*, float*, i32 }* null, i64 0, i32 3) to i64))<br>
%19 = icmp eq i32 %18, 0<br>
br i1 %19, label %20, label %22</p>
<p dir="ltr">; <label>:20 ; preds = %16<br>
%21 = call i32 @cudaLaunch(i8* bitcast (void (float*, float*, float*, i32)* @_Z9vectorAddPKfS0_Pfi to i8*))<br>
br label %22</p>
<p dir="ltr">; <label>:22 ; preds = %20, %16, %12, %8, %0<br>
ret void<br>
}</p>
<p dir="ltr">//======================= code end =====================<br></p>
<p dir="ltr">From the emitted IR code, I can see NO addrspace declarations, which should be specific to<br>
NVVM IR. So this may be the problem.</p>
<p dir="ltr">Could anyone tell me how to fix this?</p>
<p dir="ltr">Any help is appreciated!</p>
<p dir="ltr"> suxing<br>
</p>