<table border="1" cellspacing="0" cellpadding="8">
<tr>
<th>Issue</th>
<td>
<a href=https://github.com/llvm/llvm-project/issues/55675>55675</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>
Twice as many kernel arguments as specified in SPIR
</td>
</tr>
<tr>
<th>Labels</th>
<td>
new issue
</td>
</tr>
<tr>
<th>Assignees</th>
<td>
</td>
</tr>
<tr>
<th>Reporter</th>
<td>
jonasbhjulstad
</td>
</tr>
</table>
<pre>
I have defined a simple kernel function that takes one global input argument, which is compiled from OpenCL for C++ (-cl-std=clc++2021) in clang.
```
__kernel void foo(global float* x0)
{
int a = 0;
}
```
The kernel (kernel.cpp, contained in same directory as test executable) is compiled to non-readable SPIR with this command:
```
clang -c -target spirv32 -cl-kernel-arg-info -Xclang -no-opaque-pointers -cl-std=clc++2021 kernel.cpp -o kernel.spv
```
It is fully possible to load the kernel into OpenCL (as in the example below), but the recognized number of kernel input arguments is wrong. Two (float*) arguments are recognized by OpenCL (using `clGetKernelInfo(kernel, CL_KERNEL_NUM_ARGS, ...)`) but I'm only able to assign one kernel argument using `clSetKernelArg(.)`. Is something missing in the SPIR translation process, or does the additional kernel argument serve another purpose?
Readable format is attained with the following command:
```
clang -c -target spirv32 -cl-kernel-arg-info -Xclang -no-opaque-pointers -cl-std=clc++2021 -emit-llvm kernel.cpp -o kernel.ll
llvm-spirv --spirv-text kernel.ll -o kernel.rspv
```
Which gives the following readable output:
```
119734787 65536 393230 19 0
2 Capability Addresses
2 Capability Linkage
2 Capability Kernel
5 ExtInstImport 1 "OpenCL.std"
3 MemoryModel 1 2
4 EntryPoint 6 15 "foo"
3 Source 3 300000
3 Name 6 "foo"
3 Name 16 "x0"
5 Decorate 6 LinkageAttributes "foo" Export
4 Decorate 7 Alignment 4
4 Decorate 10 Alignment 4
4 Decorate 13 Alignment 4
4 Decorate 16 Alignment 4
4 TypeInt 11 32 0
4 Constant 11 14 0
2 TypeVoid 2
3 TypeFloat 3 32
4 TypePointer 4 5 3
4 TypeFunction 5 2 4
4 TypePointer 9 7 4
4 TypePointer 12 7 11
5 Function 2 6 2 5
3 FunctionParameter 4 7
2 Label 8
4 Variable 9 10 7
4 Variable 12 13 7
5 Store 10 7 2 4
5 Store 13 14 2 4
1 Return
1 FunctionEnd
5 Function 2 15 2 5
3 FunctionParameter 4 16
2 Label 17
5 FunctionCall 2 18 6 16
1 Return
1 FunctionEnd
```
Minimal "working example"
```
#define CL_TARGET_OPENCL_VERSION 300
#include <CL/cl.h>
#include <iostream>
#include <string.h>
#include <stdio.h>
#include <stdlib.h>
#include <string>
#include <fstream>
#include <cassert>
#include <array>
#include <random>
#include <sstream>
#include <cmath>
#define SUCCESS 0
#define FAILURE 1
struct CLG_Instance
{
cl_uint numPlatforms;
cl_platform_id platform;
cl_uint numDevices = 0;
std::vector<cl_device_id> device_ids;
cl_context context;
cl_command_queue commandQueue;
cl_program program;
size_t global_mem_size;
size_t global_mem_cache_size;
size_t max_work_group_size;
};
/* convert the kernel file into a string */
std::string convertToString(const char *filename)
{
std::ifstream t(filename);
std::stringstream buffer;
buffer << t.rdbuf();
return buffer.str();
}
const char pwd[] = "/home/deb/Documents/OpenCL-Graph/test/";
int main(int argc, char *argv[])
{
CLG_Instance clInstance;
cl_uint numPlatforms; //the NO. of platforms
cl_int status = clGetPlatformIDs(0, NULL, &numPlatforms);
if (numPlatforms > 0)
{
cl_platform_id* platforms =
(cl_platform_id*)malloc(numPlatforms * sizeof(cl_platform_id));
status = clGetPlatformIDs(numPlatforms, platforms, NULL);
clInstance.platform = platforms[0];
free(platforms);
}
cl_uint deviceType = CL_DEVICE_TYPE_GPU;
status = clGetDeviceIDs(clInstance.platform, deviceType, 0, NULL, &clInstance.numDevices);
clInstance.device_ids.resize(clInstance.numDevices);
status = clGetDeviceIDs(clInstance.platform, deviceType, clInstance.numDevices, clInstance.device_ids.data(), NULL);
clInstance.context = clCreateContext(NULL, 1, clInstance.device_ids.data(), NULL, NULL, NULL);
clInstance.commandQueue = clCreateCommandQueueWithProperties(clInstance.context, clInstance.device_ids[0], NULL, NULL);
clGetDeviceInfo(clInstance.device_ids[0], CL_DEVICE_GLOBAL_MEM_CACHE_SIZE, sizeof(size_t), &clInstance.global_mem_cache_size, NULL);
clGetDeviceInfo(clInstance.device_ids[0], CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(size_t), &clInstance.global_mem_size, NULL);
clGetDeviceInfo(clInstance.device_ids[0], CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(size_t), &clInstance.max_work_group_size, NULL);
int err = 0;
std::string programBinary = convertToString((std::string(pwd) + "kernel.spv").c_str());
long unsigned int programSize = sizeof(char)*programBinary.length();
clInstance.program = clCreateProgramWithIL(clInstance.context, (const void*) programBinary.data(), sizeof(char)*programBinary.length(), &err);
assert(err == CL_SUCCESS);
status = clBuildProgram(clInstance.program, 1, clInstance.device_ids.data(), NULL, NULL, NULL);
if (status == CL_BUILD_PROGRAM_FAILURE)
{
size_t log_size;
clGetProgramBuildInfo(clInstance.program, clInstance.device_ids[0], CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
// Allocate memory for the log
char *log = (char *)malloc(log_size);
// Get the log
clGetProgramBuildInfo(clInstance.program, clInstance.device_ids[0], CL_PROGRAM_BUILD_LOG, log_size, log, NULL);
// Print the log
printf("%s\n", log);
}
float x0[3] = {1.0,2.0,3.0};
float x_res[3] = {-10,-10,-10};
size_t inputBufferSize = sizeof(float)*3;
// size_t outputBufferSize = sizeof(float)*3;
cl_mem inputBuffer = clCreateBuffer(clInstance.context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, inputBufferSize, (void*) x0, &err);
assert(err == CL_SUCCESS);
// cl_mem outputBuffer = clCreateBuffer(clInstance.context, CL_MEM_WRITE_ONLY, outputBufferSize, NULL, &err);
// assert(err == CL_SUCCESS);
cl_kernel kernel;
kernel = clCreateKernel(clInstance.program, "foo", &err);
assert(err == CL_SUCCESS);
cl_uint num_args = 0;
size_t numargSize = sizeof(cl_uint);
status = clGetKernelInfo(kernel, CL_KERNEL_NUM_ARGS, numargSize, &num_args, &numargSize);
std::cout << "number of arguments: " << num_args << std::endl;
//get kernel argument names
std::vector<std::string> arg_names;
std::string arg_name;
for (cl_uint i = 0; i < num_args; i++)
{
size_t arg_name_size = 0;
status = clGetKernelArgInfo(kernel, i, CL_KERNEL_ARG_NAME, 0, NULL, &arg_name_size);
arg_name.resize(arg_name_size);
status = clGetKernelArgInfo(kernel, i, CL_KERNEL_ARG_NAME, arg_name_size, (void*) arg_name.data(), NULL);
arg_names.push_back(std::string(arg_name));
}
for (cl_uint i = 0; i < num_args; i++)
{
std::cout << arg_names[i] << std::endl;
}
//get kernel argument type names
std::vector<std::string> arg_type_names;
std::string arg_type_name;
for (cl_uint i = 0; i < num_args; i++)
{
size_t arg_type_name_size = 0;
status = clGetKernelArgInfo(kernel, i, CL_KERNEL_ARG_TYPE_NAME, 0, NULL, &arg_type_name_size);
arg_type_name.resize(arg_type_name_size);
status = clGetKernelArgInfo(kernel, i, CL_KERNEL_ARG_TYPE_NAME, arg_type_name_size, (void*) arg_type_name.data(), NULL);
arg_type_names.push_back(std::string(arg_type_name));
}
//print arg_type_names
for (cl_uint i = 0; i < num_args; i++)
{
std::cout << arg_type_names[i] << std::endl;
}
// err = clSetKernelArg(kernel, 1, sizeof(float), &df);
assert(err == CL_SUCCESS);
// err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &outputBuffer);
// assert(err == CL_SUCCESS);
size_t globalWorkSize[1] = {8};
err = clEnqueueNDRangeKernel(clInstance.commandQueue, kernel, 1, NULL, globalWorkSize, NULL, 0, NULL, NULL);
assert(err == CL_SUCCESS);
// err = clEnqueueReadBuffer(clInstance.commandQueue, outputBuffer, CL_TRUE, 0, outputBufferSize, x_res, 0, NULL, NULL);
// assert(err == CL_SUCCESS);
status = clReleaseKernel(kernel); //Release kernel.
status = clReleaseProgram(clInstance.program); //Release the program object.
status = clReleaseMemObject(inputBuffer);
// status = clReleaseMemObject(outputBuffer);
status = clReleaseCommandQueue(clInstance.commandQueue); //Release Command queue.
status = clReleaseContext(clInstance.context); //Release context.
return SUCCESS;
}
```
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJzFGtty6jjya8iLCgrsEOAhD9ySpQ5JWEjOmdkXl7AFeI6xvLKdy379dkuyZRubJGcyM6kUYLXU9251S95y7-16QQ70mRGP7fyQeYSS2D9GASM_mQhZQHZp6CY-D0lyoAlJ6E8WEx4ysg_4lgbED6M0IVTs0yMLk5Y1JS8H3z0QPyYuP0Z-ACh3gh_JQ8TC6ZLsuCDTljWBf9Kyhm03aMeJ17JnbuCqYatr9VrWCDATN6DhvtPqzlrdceuqq__lo-No_p65DxQ4B2SapV3AKXAyJq9dwKNXDybqB4E_PwSOCdAk3ZY9ySbMaumoz8dDrg-go3513ChCeV0eJlSqDjiO6RFU6QvmJly8ERqThMUJYa_MTRO6DZiUrKCbhJOQh23BqIdgslkt1uTFTw6gbzXvSEPQz7iWO6kg0nZJOwETsITEkS-ebYugXhWbbQC0_XDHSfs3PT3kbR7R_6asHXHQBRMxIQ2GIEZW0ubZUxw917KzSFC2XRoEbyTiceyjRCAhGAQkNToEojxzCNAnaMkPJZy9Uul7WxbwFzQe6HcL_oUw0Cnfh_7_QGlhetwyQfjOICx6YYxcvAgOvkMeXziSyHwC1W-mUVFCu30rMJXGPugKRHODW5Z8k3QWoMbc_sjbdOl8m6_v50vn_unOGa9vNzja6XSQ9yt0P8n-omUNjhA1oBaqVUJBO_tQRpKWIWOLFChvMspjsQfCGm2HLGIS8yMDF4GZRz-WK7QOpQclgoZxQGXgRoK7LI6RMwg-j0MA4zzqeT7CIWKqHMRMQEagIYd5gkSpAGOyln1TDIl15rEQ0UcqDU8THQjafxEWgCGRuX_ckdvs6CftIHg-1vt0EChOcEZbUidt9d1O2GtiphXWiKZAUJ8_ZCbc-89a40YbebjzNAHPbdJKrzca2JeD4YBc9fv2FbFHtmV3SW9EukTNsMiURnTrB37yRsaeJ8DQQK0GuPTDn3TP6kDfdGaTkD6ZvyaLME4Wx4iLhPQgGCwVFh3Uq2XpiTa5Y0dIcnfcg8U9ko1fknmYiLcVWoRckV4fEcgEbVZueCpcRmxid_EvH7_H9HlVs0ACehKCWd3KmZ1B-Aqa4Cot4ThJhA9Rh2rI8cxfpTAZh_mqARkHEIfS6y9Pwb3uO3D7HfhVLfzxLWILGOn1CLh4Nx-fctA7VYDeZcHIuOA7bnRGIzh0g1kNtWiVUK9UMADBPgCLkJtsM-8DpsvaRSPQST2kZwEIOCu6eGaGHLEFhrAAfcZmBlhRASZUXA1KKCyypJDvyTAn-p0KX0bHCA0wOB0HTkDzg9wJNrDZMjXXiJUP26hLM94ja5akIiwx0csZnYceaZSt139XOHTSGul6htts1ZRCMgGcQ4ySq09zV5dx1OedH_pHiqWK9cLFT0w4emOFkfrVlq3qP9zRHmEfmz86D6v5PTx9n683i4d7jNQSJcv2QzdIPcgo9nS6bFk3btA5tOx5HdzncQI579gABiAw2bga0o7Pz0EDf3sGjLgbgLuzbLmwSTORNECpEPStAQa7r8cbpT1PFDbTkizaMJun6XS-2ZATi92MF8un9Zz0iuYBEqmbgDVvnYXMKS6rKYTdwEkxS0M9tYJiAXfyOK-I9YRIAxzIPtnv6pwMyYw9-y7m3VJpjZPkjjyG_2dZGaOcgePJ6YAYxCX5wwkHWGDjDqy_T8GysnCgDEhZVmf8Gx9ORBF8D4FK9Lep_bujGApAJ9FNjXNkRwdHzs9wqXtgTfOO9NXB4HP2gqdRZRb0GmaFsuYNNiwg4TM4XLFU3kGXoOplaMykKxNZxt5kdtZ61TCN4ZFvlNtbQxf3FOIeqMCFiC6EXFVujEaG-wyfr0ODJFg9F1YVRS0T1wu26W7HRHGeGkHnhn-SdIQHI4C2gk6oxKdmQ60hqlPyFk2VjEaw6MVr9Set_kx6HiY56-bAkd8bj23hc8ZdVfLDb1XLtG8FjQ7wiP2ZVL9VsQn69BHqWeBCNoxi78p2T2sSnp8V0dMm03hdMQDBBfNYLIjdEIREMnWDrnD_0MFeJ8qhhaW4EnAmqYo62axkaBYzkHfYRa7vn5ZL_G5ZVyU6JQtorP4O-5_iNIIBWmimwWGMqCPFSCFPoCtHhbUzYhRy8ocuWl0NlGD_Crh7wgdgxlDiu5plo4o_SQ89p5myJqYFBecaqyI0JuxksyV2s7Q_6aJPlJftBANnHEb1eh9VPLuYVVVaxDJM0oENeTb_vpjOncffV3PndvV0ar-q0CorK5Fr-EdhDRV8qnpMYZFJ8qeuo9jOp5qE3oGuBPNfifx7mP60FA20pg08ejShOuHUWj8PuXxxti8pFqeQ_BI21XuUNcz01_scyer3eyyY3a7ChwH8gG58JXgE-4LPKurL9tQmHjN3PsOXDOKCAxhbqdOS9_Aal75dPkzGS-dufudMx9N_zZ3N4j9znJKHvNpctcLKnlm_MTfq8Qt5_UUu_1r-7sa_OT8e1t-c2_XD0-pTPNaVLuf8sco2pi0mxLkqUFcruhCb-CEVb8p_T-oX5LW8DBPpC-Z7og6RrcKpJG79o47rmOqh4qkBB7ppiIdv8sQ2yZjYgJiSBbO_wE4vMYxLfHagFtpDiV4uTirJLys1izG5UmMYjotlYxjmJRsebOsTyzL9ctr4HLvK2GCdKu-61bGG2nJ6s9FdR3V2KTlPUj_wtHCV5JwNfmEaLKcbVasYdjTbk6fFcuas1g-36_GdozukvH6RKIp9kBRJ1e0B35eL9Qyu6getWZT4NCQL4n4gUDPuFK_Lh9u6rTdnpzZBqPoJy0QyxnIJj5yO8lBO3rNg7QgIKnLo-hUAuloeZkPFsuvDhEErDYT-FoUZPuXv8wm1wPZKYOzXMh4hSPUnkE36wMU0lD81hYovlmu3Kkl58YA3Uf2JnTcog0mvg5a25KcNn4WOsLDKgcKpsrDdwyXm86SVLPiyvBWZyHbqNLvpGxHMF3ZZIKUgjUMdUH8KSamIBXcs8lFKiGqoMRNO1fa6no9nzsP98ncQPx-cPqx-d_71sHl0Vo9rnFyRVWfSQg597X5F8jtRkxaxqKZfkPHHevE4l0LKW5qKzis5oUYAzcyvyAEC6KMGfa1VxJvdeRYE-qYvvxriOD9y_6qtptAeO9Bz1x4uKVeFGTChZh9XGM7uYZ-94jO0TFct2TOPObxKVhczLpg5OxMBfZk7zfx-EiYRedehJhU0IJ9zTCz0gob9UTkGXqhV7_jwPCc-d0BXrbrsOS521LpzVV02q5zRuCDGFsTPzSh_GuHkiLq2-9B-nRGTm0CNbzRbeiz2J8b2yxYHazv347t53c5cIlw1snRzPcF0v-8u-RJGK1SqeTBn63zDW5Ui7kRpfHC21P1ZW5HnRj8tuWtON77eG-qiyrhrf-KrPfS9wKnbzM_HUYJHM78aTLj4gxGVT_3bwiqn-JfFljzDOhdgZRaa_DOfVQ619xd_HfN19Ooiz7D68fAzTvKBGDRuciYQjUfLWrfqiv9AmBaof1Gs5kcQJ2_RGKv2Sg20KWal_3m7-oLpTxSLn2ZJVZeGp2JleJbSrzBZuuH6wcVPWb70J71C-zGsbziMYPNQ3sPdz9Y03NcWi6UbOhCrInuWASpsFCClRNHc7H2JnbQ4-LJTfSVflqVsH5kxHtdPJr3VVfaqy_u4XH_GwsV8t2YBo7ExUWYHXEhKf4qinp-9-3Qe6_njoJG538qwYieeHZrx7R-wb75D4Y4dH-Q8eT9XGxUFbb2HoSmwmqhPS4ZvdokaSYleS6RnvSOluU-o6x8rpirT0bM6p0nSkNRXrpnTNNy05i-LXHjXtjeyR_Qi8ZOAXT---C7Dl1xBnrdqZRQjII6Y6-989XIsvpl4kYrg-pAkEXY3uqryk0O6RcXBA75-p7_a4A_KPDd-HKcYJDf9_tWgf3G47vZtxq4ubdu1e_2B1Xcv-4PdgHnuJe12e8PdRYBv3cTXeEMLvRV7IRIF9qX92YV_bXUtq9u3Lnvw173qWMPBdkc92u-6owEd9luXXXakPr7u93zscLG_ENeSpW0Ku91lN_DjJDZA9TInY9fqQviCpsmBi-s_eEjj7eGPNADDeReSg2spwf8BdAR5Zw">